1// RUN: mlir-opt -allow-unregistered-dialect -convert-gpu-to-spirv="use-64bit-index=true" -verify-diagnostics -split-input-file %s -o - | FileCheck %s 2 3module attributes { 4 gpu.container_module, 5 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>> 6} { 7 gpu.module @kernels { 8 // CHECK-LABEL: spirv.module @{{.*}} Physical64 OpenCL 9 // CHECK: spirv.func 10 // CHECK-SAME: {{%.*}}: f32 11 // CHECK-NOT: spirv.interface_var_abi 12 // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.array<12 x f32>, CrossWorkgroup> 13 // CHECK-NOT: spirv.interface_var_abi 14 // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]> 15 // CHECK-LABEL: func.func @basic_module_structure 16 // CHECK-SAME: attributes {gpu.kernel} 17 gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) kernel 18 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 19 gpu.return 20 } 21 } 22 23 func.func @main() { 24 %0 = "op"() : () -> (f32) 25 %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) 26 %cst = arith.constant 1 : index 27 gpu.launch_func @kernels::@basic_module_structure 28 blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst) 29 args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) 30 return 31 } 32} 33 34// ----- 35 36module attributes { 37 gpu.container_module 38} { 39 gpu.module @kernels attributes { 40 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>> 41 } { 42 // CHECK-LABEL: spirv.module @{{.*}} Physical64 OpenCL 43 // CHECK-SAME: spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>> 44 // CHECK: spirv.func 45 // CHECK-SAME: {{%.*}}: f32 46 // CHECK-NOT: spirv.interface_var_abi 47 // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.array<12 x f32>, CrossWorkgroup> 48 // CHECK-NOT: spirv.interface_var_abi 49 // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]> 50 // CHECK-LABEL: func.func @basic_module_structure 51 // CHECK-SAME: attributes {gpu.kernel} 52 gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) kernel 53 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 54 gpu.return 55 } 56 } 57 58 func.func @main() { 59 %0 = "op"() : () -> (f32) 60 %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) 61 %cst = arith.constant 1 : index 62 gpu.launch_func @kernels::@basic_module_structure 63 blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst) 64 args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) 65 return 66 } 67} 68