142 lines · plain
1// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s2 3module attributes {4 gpu.container_module,5 spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniformShuffle], []>, #spirv.resource_limits<subgroup_size = 16>>6} {7 8gpu.module @kernels {9 // CHECK-LABEL: spirv.func @shuffle_xor()10 gpu.func @shuffle_xor() kernel11 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {12 %mask = arith.constant 8 : i3213 %width = arith.constant 16 : i3214 %val = arith.constant 42.0 : f3215 16 // CHECK: %[[MASK:.+]] = spirv.Constant 8 : i3217 // CHECK: %[[VAL:.+]] = spirv.Constant 4.200000e+01 : f3218 // CHECK: %{{.+}} = spirv.GroupNonUniformShuffleXor <Subgroup> %[[VAL]], %[[MASK]] : f32, i3219 // CHECK: %{{.+}} = spirv.Constant true20 %result, %valid = gpu.shuffle xor %val, %mask, %width : f3221 gpu.return22 }23}24 25}26 27// -----28 29module attributes {30 gpu.container_module,31 spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniformShuffle], []>, #spirv.resource_limits<subgroup_size = 32>>32} {33 34gpu.module @kernels {35 gpu.func @shuffle_xor() kernel36 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {37 %mask = arith.constant 8 : i3238 %width = arith.constant 16 : i3239 %val = arith.constant 42.0 : f3240 41 // Cannot convert due to shuffle width and target subgroup size mismatch42 // expected-error @+1 {{failed to legalize operation 'gpu.shuffle'}}43 %result, %valid = gpu.shuffle xor %val, %mask, %width : f3244 gpu.return45 }46}47 48}49 50// -----51 52module attributes {53 gpu.container_module,54 spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniformShuffle], []>, #spirv.resource_limits<subgroup_size = 16>>55} {56 57gpu.module @kernels {58 // CHECK-LABEL: spirv.func @shuffle_idx()59 gpu.func @shuffle_idx() kernel60 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {61 %mask = arith.constant 8 : i3262 %width = arith.constant 16 : i3263 %val = arith.constant 42.0 : f3264 65 // CHECK: %[[MASK:.+]] = spirv.Constant 8 : i3266 // CHECK: %[[VAL:.+]] = spirv.Constant 4.200000e+01 : f3267 // CHECK: %{{.+}} = spirv.GroupNonUniformShuffle <Subgroup> %[[VAL]], %[[MASK]] : f32, i3268 // CHECK: %{{.+}} = spirv.Constant true69 %result, %valid = gpu.shuffle idx %val, %mask, %width : f3270 gpu.return71 }72}73 74}75 76// -----77 78module attributes {79 gpu.container_module,80 spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], []>,81 #spirv.resource_limits<subgroup_size = 16>>82} {83 84gpu.module @kernels {85 // CHECK-LABEL: spirv.func @shuffle_down()86 gpu.func @shuffle_down() kernel87 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {88 %offset = arith.constant 4 : i3289 %width = arith.constant 16 : i3290 %val = arith.constant 42.0 : f3291 92 // CHECK: %[[OFFSET:.+]] = spirv.Constant 4 : i3293 // CHECK: %[[WIDTH:.+]] = spirv.Constant 16 : i3294 // CHECK: %[[VAL:.+]] = spirv.Constant 4.200000e+01 : f3295 // CHECK: %{{.+}} = spirv.GroupNonUniformShuffleDown <Subgroup> %[[VAL]], %[[OFFSET]] : f32, i3296 97 // CHECK: %[[INVOCATION_ID_ADDR:.+]] = spirv.mlir.addressof @__builtin__SubgroupLocalInvocationId__ : !spirv.ptr<i32, Input>98 // CHECK: %[[LANE_ID:.+]] = spirv.Load "Input" %[[INVOCATION_ID_ADDR]] : i3299 // CHECK: %[[VAL_LANE_ID:.+]] = spirv.IAdd %[[LANE_ID]], %[[OFFSET]] : i32100 // CHECK: %[[VALID:.+]] = spirv.ULessThan %[[VAL_LANE_ID]], %[[WIDTH]] : i32101 102 %result, %valid = gpu.shuffle down %val, %offset, %width : f32103 gpu.return104 }105}106 107}108 109// -----110 111module attributes {112 gpu.container_module,113 spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], []>,114 #spirv.resource_limits<subgroup_size = 16>>115} {116 117gpu.module @kernels {118 // CHECK-LABEL: spirv.func @shuffle_up()119 gpu.func @shuffle_up() kernel120 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {121 %offset = arith.constant 4 : i32122 %width = arith.constant 16 : i32123 %val = arith.constant 42.0 : f32124 125 // CHECK: %[[OFFSET:.+]] = spirv.Constant 4 : i32126 // CHECK: %[[WIDTH:.+]] = spirv.Constant 16 : i32127 // CHECK: %[[VAL:.+]] = spirv.Constant 4.200000e+01 : f32128 // CHECK: %{{.+}} = spirv.GroupNonUniformShuffleUp <Subgroup> %[[VAL]], %[[OFFSET]] : f32, i32129 130 // CHECK: %[[INVOCATION_ID_ADDR:.+]] = spirv.mlir.addressof @__builtin__SubgroupLocalInvocationId__ : !spirv.ptr<i32, Input>131 // CHECK: %[[LANE_ID:.+]] = spirv.Load "Input" %[[INVOCATION_ID_ADDR]] : i32132 // CHECK: %[[VAL_LANE_ID:.+]] = spirv.ISub %[[LANE_ID]], %[[OFFSET]] : i32133 // CHECK: %[[CST0:.+]] = spirv.Constant 0 : i32134 // CHECK: %[[VALID:.+]] = spirv.SGreaterThanEqual %[[VAL_LANE_ID]], %[[CST0]] : i32135 136 %result, %valid = gpu.shuffle up %val, %offset, %width : f32137 gpu.return138 }139}140 141}142