41 lines · plain
1// RUN: mlir-opt --split-input-file --convert-gpu-to-spirv %s | FileCheck %s2 3module attributes {gpu.container_module} {4 // CHECK-LABEL: spirv.module @{{.*}} GLSL4505 gpu.module @kernels [#spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>] {6 // CHECK: spirv.func @load_kernel7 // 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 : index10 // CHECK: %[[PTR:.*]] = spirv.AccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}11 // CHECK-NEXT: {{%.*}} = spirv.Load "StorageBuffer" %[[PTR]] : f3212 %0 = memref.load %arg0[%c0, %c0] : memref<12x4xf32>13 // CHECK: spirv.Return14 gpu.return15 }16 }17}18 19// -----20// Checks that the `-convert-gpu-to-spirv` pass selects the first21// `spirv.target_env` from the `targets` array attribute attached to `gpu.module`.22module attributes {gpu.container_module} {23 // CHECK-LABEL: spirv.module @{{.*}} GLSL45024 // CHECK-SAME: #spirv.target_env<#spirv.vce<v1.4, [Shader], [SPV_KHR_storage_buffer_storage_class]>25 gpu.module @kernels [26 #spirv.target_env<#spirv.vce<v1.4, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>,27 #spirv.target_env<#spirv.vce<v1.0, [Kernel], []>, #spirv.resource_limits<>>,28 #spirv.target_env<#spirv.vce<v1.0, [Shader], []>, #spirv.resource_limits<>>] {29 // CHECK: spirv.func @load_kernel30 // 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)>})31 gpu.func @load_kernel(%arg0: memref<12x4xf32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {32 %c0 = arith.constant 0 : index33 // CHECK: %[[PTR:.*]] = spirv.AccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}34 // CHECK-NEXT: {{%.*}} = spirv.Load "StorageBuffer" %[[PTR]] : f3235 %0 = memref.load %arg0[%c0, %c0] : memref<12x4xf32>36 // CHECK: spirv.Return37 gpu.return38 }39 }40}41