41 lines · plain
1// RUN: mlir-opt -split-input-file -spirv-lower-abi-attrs %s | FileCheck %s2 3module attributes {4 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>>5} {6 spirv.module Physical64 OpenCL {7 // CHECK-LABEL: spirv.module8 // CHECK: spirv.func [[FN:@.*]]({{%.*}}: f32, {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>9 // We cannot generate SubgroupSize execution mode without necessary capability -- leave it alone.10 // CHECK-SAME: #spirv.entry_point_abi<subgroup_size = 64>11 // CHECK: spirv.EntryPoint "Kernel" [[FN]]12 // CHECK: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 113 spirv.func @kernel(14 %arg0: f32,15 %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None"16 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} {17 spirv.Return18 }19 }20}21 22// -----23 24module attributes {25 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, SubgroupDispatch], []>, #spirv.resource_limits<>>26} {27 spirv.module Physical64 OpenCL {28 // CHECK-LABEL: spirv.module29 // CHECK: spirv.func [[FN:@.*]]({{%.*}}: f32, {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>30 // CHECK: spirv.EntryPoint "Kernel" [[FN]]31 // CHECK: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 132 // CHECK: spirv.ExecutionMode [[FN]] "SubgroupSize", 6433 spirv.func @kernel(34 %arg0: f32,35 %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None"36 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} {37 spirv.Return38 }39 }40}41