1// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv="use-64bit-index=false" %s -o - | FileCheck %s --check-prefix=INDEX32 2// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv="use-64bit-index=true" %s -o - | FileCheck %s --check-prefix=INDEX64 3 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 : index 10 gpu.launch_func @kernels::@builtin_workgroup_id_x 11 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 12 return 13 } 14 15 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 16 // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input> 17 // INDEX64-LABEL: spirv.module @{{.*}} Logical GLSL450 18 // INDEX64: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input> 19 gpu.module @kernels { 20 gpu.func @builtin_workgroup_id_x() kernel 21 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 i64 26 %0 = gpu.block_id x 27 gpu.return 28 } 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 : index 40 %c256 = arith.constant 256 : i32 41 gpu.launch_func @kernels::@builtin_workgroup_id_y 42 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 43 dynamic_shared_memory_size %c256 44 return 45 } 46 47 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 48 // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input> 49 gpu.module @kernels { 50 gpu.func @builtin_workgroup_id_y() kernel 51 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 y 56 gpu.return 57 } 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 : index 69 gpu.launch_func @kernels::@builtin_workgroup_id_z 70 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 71 return 72 } 73 74 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 75 // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input> 76 gpu.module @kernels { 77 gpu.func @builtin_workgroup_id_z() kernel 78 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 z 83 gpu.return 84 } 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 : index 96 gpu.launch_func @kernels::@builtin_workgroup_size_x 97 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 98 return 99 } 100 101 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 102 gpu.module @kernels { 103 gpu.func @builtin_workgroup_size_x() kernel 104 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 entry 108 // point ABI we want here. 109 // INDEX32: spirv.Constant 32 : i32 110 %0 = gpu.block_dim x 111 gpu.return 112 } 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 : index 124 gpu.launch_func @kernels::@builtin_workgroup_size_y 125 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 126 return 127 } 128 129 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 130 gpu.module @kernels { 131 gpu.func @builtin_workgroup_size_y() kernel 132 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 : i32 135 %0 = gpu.block_dim y 136 gpu.return 137 } 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 : index 149 gpu.launch_func @kernels::@builtin_workgroup_size_z 150 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 151 return 152 } 153 154 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 155 gpu.module @kernels { 156 gpu.func @builtin_workgroup_size_z() kernel 157 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 : i32 160 %0 = gpu.block_dim z 161 gpu.return 162 } 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 : index 174 gpu.launch_func @kernels::@builtin_local_id_x 175 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 176 return 177 } 178 179 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 180 // INDEX32: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") : !spirv.ptr<vector<3xi32>, Input> 181 gpu.module @kernels { 182 gpu.func @builtin_local_id_x() kernel 183 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 x 188 gpu.return 189 } 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 : index 201 gpu.launch_func @kernels::@builtin_num_workgroups_x 202 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 203 return 204 } 205 206 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 207 // INDEX32: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") : !spirv.ptr<vector<3xi32>, Input> 208 gpu.module @kernels { 209 gpu.func @builtin_num_workgroups_x() kernel 210 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 x 215 gpu.return 216 } 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 GLSL450 227 // INDEX32: spirv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId") : !spirv.ptr<i32, Input> 228 gpu.module @kernels { 229 gpu.func @builtin_subgroup_id() kernel 230 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 : index 234 gpu.return 235 } 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 GLSL450 246 // INDEX32: spirv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups") : !spirv.ptr<i32, Input> 247 gpu.module @kernels { 248 gpu.func @builtin_num_subgroups() kernel 249 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 : index 253 gpu.return 254 } 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 : index 266 gpu.launch_func @kernels::@builtin_workgroup_size_x 267 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 268 return 269 } 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() kernel 275 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 x 280 gpu.return 281 } 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 : index 293 gpu.launch_func @kernels::@builtin_workgroup_size_y 294 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 295 return 296 } 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() kernel 302 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 y 307 gpu.return 308 } 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 : index 320 gpu.launch_func @kernels::@builtin_workgroup_size_z 321 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 322 return 323 } 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() kernel 329 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 z 334 gpu.return 335 } 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 : index 347 gpu.launch_func @kernels::@builtin_global_id_x 348 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 349 return 350 } 351 352 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 353 // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input> 354 gpu.module @kernels { 355 gpu.func @builtin_global_id_x() kernel 356 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 x 361 gpu.return 362 } 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 : index 374 gpu.launch_func @kernels::@builtin_global_id_y 375 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 376 return 377 } 378 379 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 380 // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input> 381 gpu.module @kernels { 382 gpu.func @builtin_global_id_y() kernel 383 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 y 388 gpu.return 389 } 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 : index 401 gpu.launch_func @kernels::@builtin_global_id_z 402 blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) 403 return 404 } 405 406 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450 407 // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input> 408 gpu.module @kernels { 409 gpu.func @builtin_global_id_z() kernel 410 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 z 415 gpu.return 416 } 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 GLSL450 428 // INDEX32: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input> 429 // INDEX64-LABEL: spirv.module @{{.*}} Logical GLSL450 430 // INDEX64: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input> 431 gpu.module @kernels { 432 gpu.func @builtin_subgroup_size() kernel 433 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 i64 437 %0 = gpu.subgroup_size : index 438 gpu.return 439 } 440 } 441} 442