mlir/test/Conversion/GPUToSPIRV/builtins-vulkan.mlir

   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
   4 module 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
  34 module 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
  63 module 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
  90 module 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
 118 module 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
 143 module 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
 168 module 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
 195 module 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
 222 module 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
 241 module 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
 260 module 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
 287 module 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
 314 module 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
 341 module 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
 368 module 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
 395 module 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
 423 module 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 }