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, [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 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 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} 53