Skip to content

Commit 71d3d92

Browse files
committed
Add test
1 parent 5d05b68 commit 71d3d92

File tree

1 file changed

+17
-0
lines changed

1 file changed

+17
-0
lines changed
Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
// RUN: mlir-opt --spirv-attach-target='caps=Shader exts=SPV_KHR_storage_buffer_storage_class' --convert-gpu-to-spirv %s -o - | FileCheck %s
2+
3+
module attributes {gpu.container_module} {
4+
// CHECK-LABEL: spirv.module @{{.*}} GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
5+
gpu.module @kernels {
6+
// CHECK: spirv.func @load_kernel
7+
// 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)>})
8+
gpu.func @load_kernel(%arg0: memref<12x4xf32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
9+
%c0 = arith.constant 0 : index
10+
// CHECK: %[[PTR:.*]] = spirv.AccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
11+
// CHECK-NEXT: {{%.*}} = spirv.Load "StorageBuffer" %[[PTR]] : f32
12+
%0 = memref.load %arg0[%c0, %c0] : memref<12x4xf32>
13+
// CHECK: spirv.Return
14+
gpu.return
15+
}
16+
}
17+
}

0 commit comments

Comments
 (0)