110 lines · plain
1// RUN: mlir-opt -test-convert-to-spirv="convert-gpu-modules=true run-signature-conversion=false run-vector-unrolling=false" -split-input-file %s | FileCheck %s2 3module attributes {4 gpu.container_module,5 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], []>, #spirv.resource_limits<>>6} {7 // CHECK-LABEL: func.func @main8 // CHECK: %[[C1:.*]] = arith.constant 1 : index9 // CHECK: gpu.launch_func @[[$KERNELS_1:.*]]::@[[$BUILTIN_WG_ID_X:.*]] blocks in (%[[C1]], %[[C1]], %[[C1]]) threads in (%[[C1]], %[[C1]], %[[C1]])10 // CHECK: gpu.launch_func @[[$KERNELS_2:.*]]::@[[$BUILTIN_WG_ID_Y:.*]] blocks in (%[[C1]], %[[C1]], %[[C1]]) threads in (%[[C1]], %[[C1]], %[[C1]])11 func.func @main() {12 %c1 = arith.constant 1 : index13 gpu.launch_func @kernels_1::@builtin_workgroup_id_x14 blocks in (%c1, %c1, %c1) threads in (%c1, %c1, %c1)15 gpu.launch_func @KERNELS_2::@builtin_workgroup_id_y16 blocks in (%c1, %c1, %c1) threads in (%c1, %c1, %c1)17 return18 }19 20 // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL45021 // CHECK: spirv.func @[[$BUILTIN_WG_ID_X]]22 // CHECK: spirv.mlir.addressof23 // CHECK: spirv.Load "Input"24 // CHECK: spirv.CompositeExtract25 gpu.module @kernels_1 {26 gpu.func @builtin_workgroup_id_x() kernel27 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {28 %0 = gpu.block_id x29 gpu.return30 }31 }32 // CHECK: gpu.module @[[$KERNELS_1]]33 // CHECK: gpu.func @[[$BUILTIN_WG_ID_X]]34 // CHECK gpu.block_id x35 // CHECK: gpu.return36 37 // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL45038 // CHECK: spirv.func @[[$BUILTIN_WG_ID_Y]]39 // CHECK: spirv.mlir.addressof40 // CHECK: spirv.Load "Input"41 // CHECK: spirv.CompositeExtract42 gpu.module @KERNELS_2 {43 gpu.func @builtin_workgroup_id_y() kernel44 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {45 %0 = gpu.block_id y46 gpu.return47 }48 }49 // CHECK: gpu.module @[[$KERNELS_2]]50 // CHECK: gpu.func @[[$BUILTIN_WG_ID_Y]]51 // CHECK gpu.block_id y52 // CHECK: gpu.return53}54 55// -----56 57module attributes {58 gpu.container_module,59 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>60} {61 // CHECK-LABEL: func.func @main62 // CHECK-SAME: %[[ARG0:.*]]: memref<2xi32>, %[[ARG1:.*]]: memref<4xi32>63 // CHECK: %[[C1:.*]] = arith.constant 1 : index64 // CHECK: gpu.launch_func @[[$KERNEL_MODULE:.*]]::@[[$KERNEL_FUNC:.*]] blocks in (%[[C1]], %[[C1]], %[[C1]]) threads in (%[[C1]], %[[C1]], %[[C1]]) args(%[[ARG0]] : memref<2xi32>, %[[ARG1]] : memref<4xi32>)65 func.func @main(%arg0 : memref<2xi32>, %arg2 : memref<4xi32>) {66 %c1 = arith.constant 1 : index67 gpu.launch_func @kernels::@kernel_foo68 blocks in (%c1, %c1, %c1) threads in (%c1, %c1, %c1)69 args(%arg0 : memref<2xi32>, %arg2 : memref<4xi32>)70 return71 }72 73 // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL45074 // CHECK: spirv.func @[[$KERNEL_FUNC]]75 // CHECK-SAME: %{{.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<2 x i32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}76 // CHECK-SAME: %{{.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<4 x i32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}77 gpu.module @kernels {78 gpu.func @kernel_foo(%arg0 : memref<2xi32>, %arg1 : memref<4xi32>)79 kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {80 // CHECK: spirv.Constant81 // CHECK: spirv.Constant dense<0>82 %idx0 = arith.constant 0 : index83 %vec0 = arith.constant dense<[0, 0]> : vector<2xi32>84 // CHECK: spirv.AccessChain85 // CHECK: spirv.Load "StorageBuffer"86 %val = memref.load %arg0[%idx0] : memref<2xi32>87 // CHECK: spirv.CompositeInsert88 %vec = vector.insert %val, %vec0[%idx0] : i32 into vector<2xi32>89 // CHECK: spirv.VectorShuffle90 %shuffle = vector.shuffle %vec, %vec[3, 2, 1, 0] : vector<2xi32>, vector<2xi32>91 %res = vector.extract %shuffle[%idx0] : i32 from vector<4xi32>92 // CHECK: spirv.AccessChain93 // CHECK: spirv.Store "StorageBuffer"94 memref.store %res, %arg1[%idx0]: memref<4xi32>95 // CHECK: spirv.Return96 gpu.return97 }98 }99 // CHECK: gpu.module @[[$KERNEL_MODULE]]100 // CHECK: gpu.func @[[$KERNEL_FUNC]]101 // CHECK-SAME: %{{.*}}: memref<2xi32>, %{{.*}}: memref<4xi32>102 // CHECK: arith.constant103 // CHECK: memref.load104 // CHECK: vector.insert105 // CHECK: vector.shuffle106 // CHECK: vector.extract107 // CHECK: memref.store108 // CHECK: gpu.return109}110