1// RUN: mlir-opt -allow-unregistered-dialect -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s 2 3module attributes {gpu.container_module} { 4 gpu.module @kernels { 5 // CHECK: spirv.module @{{.*}} Logical GLSL450 { 6 // CHECK-LABEL: spirv.func @basic_module_structure 7 // CHECK-SAME: {{%.*}}: f32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0), StorageBuffer>} 8 // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>} 9 // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]> 10 gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel 11 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 12 // CHECK: spirv.Return 13 gpu.return 14 } 15 } 16 17 func.func @main() { 18 %0 = "op"() : () -> (f32) 19 %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<StorageBuffer>>) 20 %cst = arith.constant 1 : index 21 gpu.launch_func @kernels::@basic_module_structure 22 blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst) 23 args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) 24 return 25 } 26} 27 28// ----- 29 30module attributes {gpu.container_module} { 31 gpu.module @kernels { 32 // CHECK: spirv.module @{{.*}} Logical GLSL450 { 33 // CHECK-LABEL: spirv.func @basic_module_structure_preset_ABI 34 // CHECK-SAME: {{%[a-zA-Z0-9_]*}}: f32 35 // CHECK-SAME: spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer> 36 // CHECK-SAME: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer> 37 // CHECK-SAME: spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)> 38 // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]> 39 gpu.func @basic_module_structure_preset_ABI( 40 %arg0 : f32 41 {spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer>}, 42 %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>> 43 {spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)>}) kernel 44 attributes 45 {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 46 // CHECK: spirv.Return 47 gpu.return 48 } 49 } 50} 51 52// ----- 53 54module attributes {gpu.container_module} { 55 gpu.module @kernels { 56 // expected-error @below {{failed to legalize operation 'gpu.func'}} 57 // expected-remark @below {{match failure: missing 'spirv.entry_point_abi' attribute}} 58 gpu.func @missing_entry_point_abi(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel { 59 gpu.return 60 } 61 } 62 63 func.func @main() { 64 %0 = "op"() : () -> (f32) 65 %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<StorageBuffer>>) 66 %cst = arith.constant 1 : index 67 gpu.launch_func @kernels::@missing_entry_point_abi 68 blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst) 69 args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) 70 return 71 } 72} 73 74// ----- 75 76module attributes {gpu.container_module} { 77 gpu.module @kernels { 78 // expected-error @below {{failed to legalize operation 'gpu.func'}} 79 // expected-remark @below {{match failure: missing 'spirv.interface_var_abi' attribute at argument 1}} 80 gpu.func @missing_entry_point_abi( 81 %arg0 : f32 82 {spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer>}, 83 %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel 84 attributes 85 {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 86 gpu.return 87 } 88 } 89} 90 91// ----- 92 93module attributes {gpu.container_module} { 94 gpu.module @kernels { 95 // expected-error @below {{failed to legalize operation 'gpu.func'}} 96 // expected-remark @below {{match failure: missing 'spirv.interface_var_abi' attribute at argument 0}} 97 gpu.func @missing_entry_point_abi( 98 %arg0 : f32, 99 %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>> 100 {spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)>}) kernel 101 attributes 102 {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 103 gpu.return 104 } 105 } 106} 107 108// ----- 109 110module attributes {gpu.container_module} { 111 gpu.module @kernels { 112 // CHECK-LABEL: spirv.func @barrier 113 gpu.func @barrier(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel 114 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 115 // CHECK: spirv.ControlBarrier <Workgroup>, <Workgroup>, <AcquireRelease|WorkgroupMemory> 116 gpu.barrier 117 gpu.return 118 } 119 } 120 121 func.func @main() { 122 %0 = "op"() : () -> (f32) 123 %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<StorageBuffer>>) 124 %cst = arith.constant 1 : index 125 gpu.launch_func @kernels::@barrier 126 blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst) 127 args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) 128 return 129 } 130} 131