mlir/test/Conversion/GPUToSPIRV/builtins-opencl.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, [Addresses, Kernel, 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 @{{.*}} Physical32 OpenCL
  16   // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
  17   // INDEX64-LABEL:  spirv.module @{{.*}} Physical64 OpenCL
  18   // INDEX64: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi64>, 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-NOT: spirv.UConvert
  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, [Addresses, Kernel, Int64], []>, #spirv.resource_limits<>>
  37 } {
  38   // INDEX32-LABEL:  spirv.module @{{.*}} Physical32 OpenCL
  39   // INDEX32: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input>
  40   // INDEX64-LABEL:  spirv.module @{{.*}} Physical64 OpenCL
  41   // INDEX64: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input>
  42   gpu.module @kernels {
  43     gpu.func @builtin_subgroup_size() kernel
  44       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
  45       // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]]
  46       // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
  47       // INDEX64: spirv.UConvert %{{.+}} : i32 to i64
  48       %0 = gpu.subgroup_size : index
  49       gpu.return
  50     }
  51   }
  52 }