75 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, [Addresses, Kernel, 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 @{{.*}} Physical32 OpenCL16 // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>17 // INDEX64-LABEL: spirv.module @{{.*}} Physical64 OpenCL18 // INDEX64: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi64>, 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-NOT: spirv.UConvert26 %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, [Addresses, Kernel, Int64], []>, #spirv.resource_limits<>>37} {38 // INDEX32-LABEL: spirv.module @{{.*}} Physical32 OpenCL39 // INDEX32: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input>40 // INDEX64-LABEL: spirv.module @{{.*}} Physical64 OpenCL41 // INDEX64: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input>42 gpu.module @kernels {43 gpu.func @builtin_subgroup_size() kernel44 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 i6448 %0 = gpu.subgroup_size : index49 gpu.return50 }51 }52}53 54// -----55 56module attributes {57 gpu.container_module,58 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Addresses, Kernel, Int64], []>, #spirv.resource_limits<>>59} {60 // INDEX32-LABEL: spirv.module @{{.*}} Physical32 OpenCL61 // INDEX32: spirv.GlobalVariable [[LANEID:@.*]] built_in("SubgroupLocalInvocationId") : !spirv.ptr<i32, Input>62 // INDEX64-LABEL: spirv.module @{{.*}} Physical64 OpenCL63 // INDEX64: spirv.GlobalVariable [[LANEID:@.*]] built_in("SubgroupLocalInvocationId") : !spirv.ptr<i32, Input>64 gpu.module @kernels {65 gpu.func @builtin_laneid() kernel66 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {67 // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LANEID]]68 // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]69 // INDEX64: spirv.UConvert %{{.+}} : i32 to i6470 %0 = gpu.lane_id71 gpu.return72 }73 }74}75