1// RUN: mlir-opt -split-input-file -spirv-lower-abi-attrs %s | FileCheck %s 2 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.module 8 // 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, 1 13 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.Return 18 } 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.module 29 // 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, 1 32 // CHECK: spirv.ExecutionMode [[FN]] "SubgroupSize", 64 33 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.Return 38 } 39 } 40} 41