442 lines · plain
1// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv="use-64bit-index=false" %s -o - | FileCheck %s --check-prefix=INDEX322// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv="use-64bit-index=true" %s -o - | FileCheck %s --check-prefix=INDEX643 4module attributes {5 gpu.container_module,6 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>7} {8 func.func @builtin() {9 %c0 = arith.constant 1 : index10 gpu.launch_func @kernels::@builtin_workgroup_id_x11 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)12 return13 }14 15 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL45016 // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>17 // INDEX64-LABEL: spirv.module @{{.*}} Logical GLSL45018 // INDEX64: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>19 gpu.module @kernels {20 gpu.func @builtin_workgroup_id_x() kernel21 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {22 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]23 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]24 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}25 // INDEX64: spirv.UConvert %{{.+}} : i32 to i6426 %0 = gpu.block_id x27 gpu.return28 }29 }30}31 32// -----33 34module attributes {35 gpu.container_module,36 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>37} {38 func.func @builtin() {39 %c0 = arith.constant 1 : index40 %c256 = arith.constant 256 : i3241 gpu.launch_func @kernels::@builtin_workgroup_id_y42 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)43 dynamic_shared_memory_size %c25644 return45 }46 47 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL45048 // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>49 gpu.module @kernels {50 gpu.func @builtin_workgroup_id_y() kernel51 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {52 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]53 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]54 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}55 %0 = gpu.block_id y56 gpu.return57 }58 }59}60 61// -----62 63module attributes {64 gpu.container_module,65 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>66} {67 func.func @builtin() {68 %c0 = arith.constant 1 : index69 gpu.launch_func @kernels::@builtin_workgroup_id_z70 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)71 return72 }73 74 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL45075 // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>76 gpu.module @kernels {77 gpu.func @builtin_workgroup_id_z() kernel78 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {79 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]80 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]81 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}82 %0 = gpu.block_id z83 gpu.return84 }85 }86}87 88// -----89 90module attributes {91 gpu.container_module,92 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>93} {94 func.func @builtin() {95 %c0 = arith.constant 1 : index96 gpu.launch_func @kernels::@builtin_workgroup_size_x97 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)98 return99 }100 101 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450102 gpu.module @kernels {103 gpu.func @builtin_workgroup_size_x() kernel104 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {105 // The constant value is obtained from the spirv.entry_point_abi.106 // Note that this ignores the workgroup size specification in gpu.launch.107 // We may want to define gpu.workgroup_size and convert it to the entry108 // point ABI we want here.109 // INDEX32: spirv.Constant 32 : i32110 %0 = gpu.block_dim x111 gpu.return112 }113 }114}115 116// -----117 118module attributes {119 gpu.container_module,120 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>121} {122 func.func @builtin() {123 %c0 = arith.constant 1 : index124 gpu.launch_func @kernels::@builtin_workgroup_size_y125 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)126 return127 }128 129 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450130 gpu.module @kernels {131 gpu.func @builtin_workgroup_size_y() kernel132 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {133 // The constant value is obtained from the spirv.entry_point_abi.134 // INDEX32: spirv.Constant 4 : i32135 %0 = gpu.block_dim y136 gpu.return137 }138 }139}140 141// -----142 143module attributes {144 gpu.container_module,145 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>146} {147 func.func @builtin() {148 %c0 = arith.constant 1 : index149 gpu.launch_func @kernels::@builtin_workgroup_size_z150 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)151 return152 }153 154 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450155 gpu.module @kernels {156 gpu.func @builtin_workgroup_size_z() kernel157 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {158 // The constant value is obtained from the spirv.entry_point_abi.159 // INDEX32: spirv.Constant 1 : i32160 %0 = gpu.block_dim z161 gpu.return162 }163 }164}165 166// -----167 168module attributes {169 gpu.container_module,170 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>171} {172 func.func @builtin() {173 %c0 = arith.constant 1 : index174 gpu.launch_func @kernels::@builtin_local_id_x175 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)176 return177 }178 179 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450180 // INDEX32: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") : !spirv.ptr<vector<3xi32>, Input>181 gpu.module @kernels {182 gpu.func @builtin_local_id_x() kernel183 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {184 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LOCALINVOCATIONID]]185 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]186 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}187 %0 = gpu.thread_id x188 gpu.return189 }190 }191}192 193// -----194 195module attributes {196 gpu.container_module,197 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>198} {199 func.func @builtin() {200 %c0 = arith.constant 1 : index201 gpu.launch_func @kernels::@builtin_num_workgroups_x202 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)203 return204 }205 206 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450207 // INDEX32: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") : !spirv.ptr<vector<3xi32>, Input>208 gpu.module @kernels {209 gpu.func @builtin_num_workgroups_x() kernel210 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {211 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMWORKGROUPS]]212 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]213 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}214 %0 = gpu.grid_dim x215 gpu.return216 }217 }218}219 220// -----221 222module attributes {223 gpu.container_module,224 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>225} {226 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450227 // INDEX32: spirv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId") : !spirv.ptr<i32, Input>228 gpu.module @kernels {229 gpu.func @builtin_subgroup_id() kernel230 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {231 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPID]]232 // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]233 %0 = gpu.subgroup_id : index234 gpu.return235 }236 }237}238 239// -----240 241module attributes {242 gpu.container_module,243 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>244} {245 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450246 // INDEX32: spirv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups") : !spirv.ptr<i32, Input>247 gpu.module @kernels {248 gpu.func @builtin_num_subgroups() kernel249 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {250 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMSUBGROUPS]]251 // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]252 %0 = gpu.num_subgroups : index253 gpu.return254 }255 }256}257 258// -----259 260module attributes {261 gpu.container_module,262 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>263} {264 func.func @builtin() {265 %c0 = arith.constant 1 : index266 gpu.launch_func @kernels::@builtin_workgroup_size_x267 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)268 return269 }270 271 // INDEX32-LABEL: spirv.module @{{.*}}272 // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>273 gpu.module @kernels {274 gpu.func @builtin_workgroup_size_x() kernel275 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {276 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]277 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]278 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}279 %0 = gpu.block_dim x280 gpu.return281 }282 }283}284 285// -----286 287module attributes {288 gpu.container_module,289 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>290} {291 func.func @builtin() {292 %c0 = arith.constant 1 : index293 gpu.launch_func @kernels::@builtin_workgroup_size_y294 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)295 return296 }297 298 // INDEX32-LABEL: spirv.module @{{.*}}299 // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>300 gpu.module @kernels {301 gpu.func @builtin_workgroup_size_y() kernel302 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {303 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]304 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]305 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}306 %0 = gpu.block_dim y307 gpu.return308 }309 }310}311 312// -----313 314module attributes {315 gpu.container_module,316 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>317} {318 func.func @builtin() {319 %c0 = arith.constant 1 : index320 gpu.launch_func @kernels::@builtin_workgroup_size_z321 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)322 return323 }324 325 // INDEX32-LABEL: spirv.module @{{.*}}326 // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>327 gpu.module @kernels {328 gpu.func @builtin_workgroup_size_z() kernel329 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {330 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]331 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]332 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}333 %0 = gpu.block_dim z334 gpu.return335 }336 }337}338 339// -----340 341module attributes {342 gpu.container_module,343 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>344} {345 func.func @builtin() {346 %c0 = arith.constant 1 : index347 gpu.launch_func @kernels::@builtin_global_id_x348 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)349 return350 }351 352 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450353 // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>354 gpu.module @kernels {355 gpu.func @builtin_global_id_x() kernel356 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {357 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]358 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]359 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}360 %0 = gpu.global_id x361 gpu.return362 }363 }364}365 366// -----367 368module attributes {369 gpu.container_module,370 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>371} {372 func.func @builtin() {373 %c0 = arith.constant 1 : index374 gpu.launch_func @kernels::@builtin_global_id_y375 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)376 return377 }378 379 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450380 // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>381 gpu.module @kernels {382 gpu.func @builtin_global_id_y() kernel383 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {384 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]385 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]386 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}387 %0 = gpu.global_id y388 gpu.return389 }390 }391}392 393// -----394 395module attributes {396 gpu.container_module,397 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>398} {399 func.func @builtin() {400 %c0 = arith.constant 1 : index401 gpu.launch_func @kernels::@builtin_global_id_z402 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)403 return404 }405 406 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450407 // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>408 gpu.module @kernels {409 gpu.func @builtin_global_id_z() kernel410 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {411 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]412 // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]413 // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}414 %0 = gpu.global_id z415 gpu.return416 }417 }418}419 420 421// -----422 423module attributes {424 gpu.container_module,425 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>426} {427 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450428 // INDEX32: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input>429 // INDEX64-LABEL: spirv.module @{{.*}} Logical GLSL450430 // INDEX64: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input>431 gpu.module @kernels {432 gpu.func @builtin_subgroup_size() kernel433 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {434 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]]435 // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]436 // INDEX64: spirv.UConvert %{{.+}} : i32 to i64437 %0 = gpu.subgroup_size : index438 gpu.return439 }440 }441}442