1466aa585SJakub Kuderski// RUN: mlir-opt --split-input-file --spirv-lower-abi-attrs --verify-diagnostics %s \ 2466aa585SJakub Kuderski// RUN: | FileCheck %s 33b35f9d8SLei Zhang 458df5e6dSLei Zhangmodule attributes { 55ab6ef75SJakub Kuderski spirv.target_env = #spirv.target_env< 6ce82530cSJakub Kuderski #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>> 758df5e6dSLei Zhang} { 858df5e6dSLei Zhang 95ab6ef75SJakub Kuderski// CHECK-LABEL: spirv.module 105ab6ef75SJakub Kuderskispirv.module Logical GLSL450 { 115ab6ef75SJakub Kuderski // CHECK-DAG: spirv.GlobalVariable [[VAR0:@.*]] bind(0, 0) : !spirv.ptr<!spirv.struct<(f32 [0])>, StorageBuffer> 125ab6ef75SJakub Kuderski // CHECK-DAG: spirv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer> 135ab6ef75SJakub Kuderski // CHECK: spirv.func [[FN:@.*]]() 1452ca1499SLei Zhang // We cannot generate SubgroupSize execution mode for Shader capability -- leave it alone. 1552ca1499SLei Zhang // CHECK-SAME: #spirv.entry_point_abi<subgroup_size = 64> 165ab6ef75SJakub Kuderski spirv.func @kernel( 173b35f9d8SLei Zhang %arg0: f32 185ab6ef75SJakub Kuderski {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0), StorageBuffer>}, 195ab6ef75SJakub Kuderski %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, StorageBuffer> 205ab6ef75SJakub Kuderski {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}) "None" 2152ca1499SLei Zhang attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} { 225ab6ef75SJakub Kuderski // CHECK: [[ADDRESSARG0:%.*]] = spirv.mlir.addressof [[VAR0]] 235ab6ef75SJakub Kuderski // CHECK: [[CONST0:%.*]] = spirv.Constant 0 : i32 245ab6ef75SJakub Kuderski // CHECK: [[ARG0PTR:%.*]] = spirv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]] 255ab6ef75SJakub Kuderski // CHECK: [[ARG0:%.*]] = spirv.Load "StorageBuffer" [[ARG0PTR]] 26*b95dfa39Sfabrizio-indirli // CHECK: [[ARG1:%.*]] = spirv.mlir.addressof [[VAR1]] 275ab6ef75SJakub Kuderski // CHECK: spirv.Return 285ab6ef75SJakub Kuderski spirv.Return 293b35f9d8SLei Zhang } 305ab6ef75SJakub Kuderski // CHECK: spirv.EntryPoint "GLCompute" [[FN]] 315ab6ef75SJakub Kuderski // CHECK: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 325ab6ef75SJakub Kuderski} // end spirv.module 3358df5e6dSLei Zhang 3458df5e6dSLei Zhang} // end module 35466aa585SJakub Kuderski 36466aa585SJakub Kuderski// ----- 37466aa585SJakub Kuderski 38466aa585SJakub Kuderskimodule { 39466aa585SJakub Kuderski// expected-error@+1 {{'spirv.module' op missing SPIR-V target env attribute}} 40466aa585SJakub Kuderskispirv.module Logical GLSL450 {} 41466aa585SJakub Kuderski} // end module 42*b95dfa39Sfabrizio-indirli 43*b95dfa39Sfabrizio-indirli// ----- 44*b95dfa39Sfabrizio-indirli 45*b95dfa39Sfabrizio-indirli// CHECK-LABEL: spirv.module 46*b95dfa39Sfabrizio-indirli// Test case with SPIRV version 1.4: all the interface's storage variables are passed to OpEntryPoint 47*b95dfa39Sfabrizio-indirlispirv.module Logical GLSL450 attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>} { 48*b95dfa39Sfabrizio-indirli // CHECK-DAG: spirv.GlobalVariable [[VAR0:@.*]] bind(0, 0) : !spirv.ptr<!spirv.struct<(f32 [0])>, StorageBuffer> 49*b95dfa39Sfabrizio-indirli // CHECK-DAG: spirv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer> 50*b95dfa39Sfabrizio-indirli // CHECK: spirv.func [[FN:@.*]]() 51*b95dfa39Sfabrizio-indirli // CHECK-SAME: #spirv.entry_point_abi<subgroup_size = 64> 52*b95dfa39Sfabrizio-indirli spirv.func @kernel( 53*b95dfa39Sfabrizio-indirli %arg0: f32 54*b95dfa39Sfabrizio-indirli {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0), StorageBuffer>}, 55*b95dfa39Sfabrizio-indirli %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, StorageBuffer> 56*b95dfa39Sfabrizio-indirli {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}) "None" 57*b95dfa39Sfabrizio-indirli attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} { 58*b95dfa39Sfabrizio-indirli // CHECK: [[ADDRESSARG0:%.*]] = spirv.mlir.addressof [[VAR0]] 59*b95dfa39Sfabrizio-indirli // CHECK: [[CONST0:%.*]] = spirv.Constant 0 : i32 60*b95dfa39Sfabrizio-indirli // CHECK: [[ARG0PTR:%.*]] = spirv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]] 61*b95dfa39Sfabrizio-indirli // CHECK: [[ARG0:%.*]] = spirv.Load "StorageBuffer" [[ARG0PTR]] 62*b95dfa39Sfabrizio-indirli // CHECK: [[ARG1:%.*]] = spirv.mlir.addressof [[VAR1]] 63*b95dfa39Sfabrizio-indirli // CHECK: spirv.Return 64*b95dfa39Sfabrizio-indirli spirv.Return 65*b95dfa39Sfabrizio-indirli } 66*b95dfa39Sfabrizio-indirli // CHECK: spirv.EntryPoint "GLCompute" [[FN]], [[VAR0]], [[VAR1]] 67*b95dfa39Sfabrizio-indirli // CHECK: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 68*b95dfa39Sfabrizio-indirli} // end spirv.module 69