1*52ca1499SLei Zhang// RUN: mlir-opt -split-input-file -spirv-lower-abi-attrs %s | FileCheck %s 29414a71aSKonrad Dobros 39414a71aSKonrad Dobrosmodule attributes { 45ab6ef75SJakub Kuderski spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>> 59414a71aSKonrad Dobros} { 65ab6ef75SJakub Kuderski spirv.module Physical64 OpenCL { 75ab6ef75SJakub Kuderski // CHECK-LABEL: spirv.module 85ab6ef75SJakub Kuderski // CHECK: spirv.func [[FN:@.*]]({{%.*}}: f32, {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup> 9*52ca1499SLei Zhang // We cannot generate SubgroupSize execution mode without necessary capability -- leave it alone. 10*52ca1499SLei Zhang // CHECK-SAME: #spirv.entry_point_abi<subgroup_size = 64> 115ab6ef75SJakub Kuderski // CHECK: spirv.EntryPoint "Kernel" [[FN]] 125ab6ef75SJakub Kuderski // CHECK: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 135ab6ef75SJakub Kuderski spirv.func @kernel( 149414a71aSKonrad Dobros %arg0: f32, 155ab6ef75SJakub Kuderski %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None" 16*52ca1499SLei Zhang attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} { 17*52ca1499SLei Zhang spirv.Return 18*52ca1499SLei Zhang } 19*52ca1499SLei Zhang } 20*52ca1499SLei Zhang} 21*52ca1499SLei Zhang 22*52ca1499SLei Zhang// ----- 23*52ca1499SLei Zhang 24*52ca1499SLei Zhangmodule attributes { 25*52ca1499SLei Zhang spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, SubgroupDispatch], []>, #spirv.resource_limits<>> 26*52ca1499SLei Zhang} { 27*52ca1499SLei Zhang spirv.module Physical64 OpenCL { 28*52ca1499SLei Zhang // CHECK-LABEL: spirv.module 29*52ca1499SLei Zhang // CHECK: spirv.func [[FN:@.*]]({{%.*}}: f32, {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup> 30*52ca1499SLei Zhang // CHECK: spirv.EntryPoint "Kernel" [[FN]] 31*52ca1499SLei Zhang // CHECK: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 32*52ca1499SLei Zhang // CHECK: spirv.ExecutionMode [[FN]] "SubgroupSize", 64 33*52ca1499SLei Zhang spirv.func @kernel( 34*52ca1499SLei Zhang %arg0: f32, 35*52ca1499SLei Zhang %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None" 36*52ca1499SLei Zhang attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} { 375ab6ef75SJakub Kuderski spirv.Return 389414a71aSKonrad Dobros } 399414a71aSKonrad Dobros } 409414a71aSKonrad Dobros} 41