1- // RUN: mlir-opt --convert-gpu-to-spirv %s | FileCheck %s
1+ // RUN: mlir-opt --split-input-file -- convert-gpu-to-spirv %s | FileCheck %s
22
33module attributes {gpu.container_module } {
44 // CHECK-LABEL: spirv.module @{{.*}} GLSL450
@@ -15,3 +15,25 @@ module attributes {gpu.container_module} {
1515 }
1616 }
1717}
18+
19+ // -----
20+
21+ module attributes {gpu.container_module } {
22+ // CHECK-LABEL: spirv.module @{{.*}} GLSL450
23+ // CHECK-SAME: #spirv.target_env<#spirv.vce<v1.4, [Shader], [SPV_KHR_storage_buffer_storage_class]>
24+ gpu.module @kernels [
25+ #spirv.target_env <#spirv.vce <v1.4 , [Shader ], [SPV_KHR_storage_buffer_storage_class ]>, #spirv.resource_limits <>>,
26+ #spirv.target_env <#spirv.vce <v1.0 , [Kernel ], []>, #spirv.resource_limits <>>,
27+ #spirv.target_env <#spirv.vce <v1.0 , [Shader ], []>, #spirv.resource_limits <>>] {
28+ // CHECK: spirv.func @load_kernel
29+ // CHECK-SAME: %[[ARG:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>})
30+ gpu.func @load_kernel (%arg0: memref <12 x4 xf32 >) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi <workgroup_size = [16 , 1 , 1 ]>} {
31+ %c0 = arith.constant 0 : index
32+ // CHECK: %[[PTR:.*]] = spirv.AccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
33+ // CHECK-NEXT: {{%.*}} = spirv.Load "StorageBuffer" %[[PTR]] : f32
34+ %0 = memref.load %arg0 [%c0 , %c0 ] : memref <12 x4 xf32 >
35+ // CHECK: spirv.Return
36+ gpu.return
37+ }
38+ }
39+ }
0 commit comments