[mlir][spirv] Add correct handling of Kernel and Addresses capabilities

[kdobros]

Summary

This change adds initial support needed to generate OpenCL compliant SPIRV.
If Kernel capability is declared then memory model becomes OpenCL.
If Addresses capability is declared then addressing model becomes Physical64.
Additionally for Kernel capability interface variable ABI attributes are not
generated as entry point function is expected to have normal arguments.

[mravishankar]

THanks! This is very welcome addition. Just a few minor comments below.

I am wondering if the Addresses capability should be done here. In general we have assumed that addressing mode is 32-bit, i.e. indexType is lowered to i32 in the SPIRVTypeConverter. We should make it 64 bits if the capability exists. It should then use 64-bit indexing out of the box, but that is untested. So seems like there is some work to be done on that aspect. The rest of the change seems unrelated.

Broader question though. RIght now there is an mlir-vulkan-runner that takes the generated SPIR-V binary and runs it to make sure the generated binary is correct. This also means that it is validated. There is no direct way of validating SPIR-V binary in MLIR tree, but wondering if you have used such a validation tool on the generated SPIR-V binary with Kernel capabilities?

mlir/test/Conversion/GPUToSPIRV/test_opencl_spirv.mlir

30 ↗

(On Diff #282871)

Nit: Please address missing newline at end of file.

[antiagainst]

Awesome, great to see that you are interested in bring up the OpenCL conversion path! :)

In D85196#2193857, @mravishankar wrote:

I am wondering if the Addresses capability should be done here.

+1. Overall I'd suggest to have a central place to host this target execution environment related information. For example, add a new ExecutionEnvironment.{h|cpp} or something to have utility functions returning the addressing/memory/execution/etc. model. This way it avoids scattering such information in all the patterns and it's a better separation of concerns I think.

In general we have assumed that addressing mode is 32-bit, i.e. indexType is lowered to i32 in the SPIRVTypeConverter. We should make it 64 bits if the capability exists. It should then use 64-bit indexing out of the box, but that is untested. So seems like there is some work to be done on that aspect. The rest of the change seems unrelated.

+1. Good point Mahesh. I think the type converter should have access to the target env, so it can change to emit a different bitwidth for index if that's desired. It shouldn't be too hard but I guess there are places we are assuming 32-bit as pointed by Mahesh and that needs to be fixed. It can happen as a next step though.

mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp

245

This seems a bit reversed. I think we should look into avoiding generating the ABI attributes in the first place if they are not useful?

[kdobros]

Separated functions to query execution environment related information
into ExecutionEnvironment{.h,.cpp}.
Removed generation of spv.interface_var_abi if it's not needed.
Fixed missing newline.

[kdobros]

@mravishankar @antiagainst
Thank you for the reviews, I have made some changes addressing your comments.

In D85196#2193857, @mravishankar wrote:

Broader question though. RIght now there is an mlir-vulkan-runner that takes the generated SPIR-V binary and runs it to make sure the generated binary is correct. This also means that it is validated. There is no direct way of validating SPIR-V binary in MLIR tree, but wondering if you have used such a validation tool on the generated SPIR-V binary with Kernel capabilities?

This is a part of upstreaming PlaidML required changes from https://github.com/plaidml/plaidml/pull/1237. I am currently working on adding OpenCL+SPIRV path there and I have managed to get full pipeline, including kernel compilation, working.
There are still some issues with extended set of instructions that I'm planning to look at later.
Do you think it may be useful to consider adding similar mlir-opencl-runner in the future?

In D85196#2193857, @mravishankar wrote:

I am wondering if the Addresses capability should be done here. In general we have assumed that addressing mode is 32-bit, i.e. indexType is lowered to i32 in the SPIRVTypeConverter. We should make it 64 bits if the capability exists. It should then use 64-bit indexing out of the box, but that is untested. So seems like there is some work to be done on that aspect. The rest of the change seems unrelated.

I was looking at both Addresses and Kernel capabilities together, because they are both required to make the binary compile. Do you think I should split Addresses into separate change?

In D85196#2194730, @antiagainst wrote:

+1. Good point Mahesh. I think the type converter should have access to the target env, so it can change to emit a different bitwidth for index if that's desired. It shouldn't be too hard but I guess there are places we are assuming 32-bit as pointed by Mahesh and that needs to be fixed. It can happen as a next step though.

Agree, but if it's not a problem, I'd like to do this in separate change. It's currently not required to compile generated SPIR-V.
There is also the issue of selecting between Physical32 vs Physical64 and I have no idea how to select between those. Seems to me that spv.target_env may need some additional field, but I was little worried about breaking compatibility with other projects.

[mravishankar]

In D85196#2196051, @kdobros wrote:

@mravishankar @antiagainst
Thank you for the reviews, I have made some changes addressing your comments.

In D85196#2193857, @mravishankar wrote:

Broader question though. RIght now there is an mlir-vulkan-runner that takes the generated SPIR-V binary and runs it to make sure the generated binary is correct. This also means that it is validated. There is no direct way of validating SPIR-V binary in MLIR tree, but wondering if you have used such a validation tool on the generated SPIR-V binary with Kernel capabilities?

This is a part of upstreaming PlaidML required changes from https://github.com/plaidml/plaidml/pull/1237. I am currently working on adding OpenCL+SPIRV path there and I have managed to get full pipeline, including kernel compilation, working.
There are still some issues with extended set of instructions that I'm planning to look at later.
Do you think it may be useful to consider adding similar mlir-opencl-runner in the future?

I think that would be very useful and a good way to make sure as we add OpenCL features we dont regress on Vulkan support. We heavily use Vulkan within IREE, so if something breaks they will show up there (cause of more extensive testing there). There are some mlir-vulkan-runner tests that will gives basic functionality testing. But having an OpenCL runner will provide the same stability for Compute capabilities, i.e. make sure that the OpenCL support doesnt regress with changes that are meant for Vulkan. Since there is no "in-tree" verifier, having a runner would be the best way to hook that up.

In D85196#2193857, @mravishankar wrote:

I am wondering if the Addresses capability should be done here. In general we have assumed that addressing mode is 32-bit, i.e. indexType is lowered to i32 in the SPIRVTypeConverter. We should make it 64 bits if the capability exists. It should then use 64-bit indexing out of the box, but that is untested. So seems like there is some work to be done on that aspect. The rest of the change seems unrelated.

I was looking at both Addresses and Kernel capabilities together, because they are both required to make the binary compile. Do you think I should split Addresses into separate change?

I think now the change looks better to me. So I just have a few more comments, and it should be fine to land.

In D85196#2194730, @antiagainst wrote:

+1. Good point Mahesh. I think the type converter should have access to the target env, so it can change to emit a different bitwidth for index if that's desired. It shouldn't be too hard but I guess there are places we are assuming 32-bit as pointed by Mahesh and that needs to be fixed. It can happen as a next step though.

Agree, but if it's not a problem, I'd like to do this in separate change. It's currently not required to compile generated SPIR-V.

There is also the issue of selecting between Physical32 vs Physical64 and I have no idea how to select between those. Seems to me that spv.target_env may need some additional field, but I was little worried about breaking compatibility with other projects.
spv.target_env definitely needs more attributes. With the mlir-vulkan-runner we get some basic level of testing. Completely fine with looking into this in a separate change.

mlir/include/mlir/Dialect/SPIRV/ExecutionEnvironment.h

22 ↗

(On Diff #283195)

Mostly nit: If you just include SPIRVTypes.h you will get the enum definitions. THis is fine as well, so do not have strong opinion on this.

31 ↗

(On Diff #283195)

I think we just use Optional<...> in MLIR AFAIK.

38 ↗

(On Diff #283195)

I think most of these functions can go into TargetAndABI.h ?

mlir/lib/Dialect/SPIRV/ExecutionEnvironment.cpp

20 ↗

(On Diff #283195)

I think most of these can go into TargetAndABI.cpp?

23 ↗

(On Diff #283195)

Maybe add a TODO here to say that the Physical64 is being hard-wired here, but it should really come from the TargetEnvAttr itself.

[kdobros]

mlir/include/mlir/Dialect/SPIRV/ExecutionEnvironment.h

31 ↗

(On Diff #283195)

I'm not sure about this, FailureOr<...> is defined inside MLIR, so if it's not used shouldn't it be removed?
I personally like the implication that if value is not present then something is wrong and function failed, and that's exactly the case here.
Just to be clear, I'm fine with changing this to Optional, but to me this seems a little confusing.

38 ↗

(On Diff #283195)

Sure, looks like a fitting place to host these functions.
Just to clarify, is it okay to move all of them there?

[mravishankar]

mlir/include/mlir/Dialect/SPIRV/ExecutionEnvironment.h

31 ↗

(On Diff #283195)

I got confused with something else (which I though was from core LLVM). I am not totally opposed to it, but its not used at all in the SPIR-V dialect. I think Optional is more consistent.

38 ↗

(On Diff #283195)

I think it is OK to move it there.

[kdobros]

Addressed review comments.
Minor style fixes - added explicit types if they don't hurt readability.

[rriddle]

mlir/include/mlir/Dialect/SPIRV/ExecutionEnvironment.h

31 ↗

(On Diff #283195)

FailureOr is more recent, and should be used when it is more readable than Optional. This is a similar situation to why LogicalResult is preferred over bool.

[mravishankar]

THanks! Looks good to me!

[antiagainst]

Cool, thanks for the contribution!

mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h

67

... whether the given SPIR-V target ..

mlir/test/Conversion/GPUToSPIRV/test_opencl_spirv.mlir

1 ↗

(On Diff #283678)

This should be named as model-structure-opencl.mlir or something to follow the convention.

[mravishankar]

mlir/test/Conversion/GPUToSPIRV/test_opencl_spirv.mlir

11 ↗

(On Diff #283678)

Forgot to mention this, but could you refactor the CHECK* lines a bit. It is more readable IMO if you group all the CHECK statements together and insert spaces to

• align the ':' after CHECK* Statement
• indent the checked code appropriately.

So

// CHECK-LABEL: func @foo
//  CHECK-SAME:  %[[ARG0:.*]]
//       CHECK:  %[[T0:.*]] =

[kdobros]

Addressed comments.
Refactored not needed CHECKs in .mlir tests.
Brought back FailureOr, as per comment it is preferred over Optional.

[kdobros]

@mravishankar @antiagainst
I think I addressed all the comments.
If you think the change is still good as is, could you commit on my behalf? I don't have commit access.
Name and email: Konrad Dobros <[e-mail]>

[mravishankar]

Will commit this. THanks!

mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h

 StringRef getInterfaceVarABIAttrName();
 
 /// Gets the InterfaceVarABIAttr given its fields.
 InterfaceVarABIAttr getInterfaceVarABIAttr(unsigned descriptorSet,
                                            unsigned binding,
                                            Optional<StorageClass> storageClass,
                                            MLIRContext *context);
 
+/// Returns whether the given SPIR-V target (described by TargetEnvAttr) needs
+/// ABI attributes for interface variables (spv.interface_var_abi).
+bool needsInterfaceVarABIAttrs(TargetEnvAttr targetAttr);
+
 /// Returns the attribute name for specifying entry point information.
 StringRef getEntryPointABIAttrName();
 
 /// Gets the EntryPointABIAttr given its fields.
 EntryPointABIAttr getEntryPointABIAttr(ArrayRef<int32_t> localSize,
                                        MLIRContext *context);
 
 /// Queries the entry point ABI on the nearest function-like op containing the

 /// containing the given `op`.
 TargetEnvAttr lookupTargetEnv(Operation *op);
 
 /// Queries the target environment recursively from enclosing symbol table ops
 /// containing the given `op` or returns the default target environment as
 /// returned by getDefaultTargetEnv() if not provided.
 TargetEnvAttr lookupTargetEnvOrDefault(Operation *op);
 
+/// Returns addressing model selected based on target environment.
+AddressingModel getAddressingModel(TargetEnvAttr targetAttr);
+
+/// Returns execution model selected based on target environment.
+/// Returns failure if it cannot be selected.
+FailureOr<ExecutionModel> getExecutionModel(TargetEnvAttr targetAttr);
+
+/// Returns memory model selected based on target environment.
+/// Returns failure if it cannot be selected.
+FailureOr<MemoryModel> getMemoryModel(TargetEnvAttr targetAttr);
+
 } // namespace spirv
 } // namespace mlir
 
 #endif // MLIR_DIALECT_SPIRV_TARGETANDABI_H

mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp

 //===- ConvertGPUToSPIRV.cpp - Convert GPU ops to SPIR-V dialect ----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the conversion patterns from GPU ops to SPIR-V dialect.
 //
 //===----------------------------------------------------------------------===//
 #include "mlir/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVLowering.h"
 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
+#include "mlir/Dialect/SPIRV/TargetAndABI.h"
 #include "mlir/IR/Module.h"
 
 using namespace mlir;
 
 namespace {
 /// Pattern lowering GPU block/thread size/id to loading SPIR-V invocation
 /// builtin variables.
 template <typename SourceOp, spirv::BuiltIn builtin>
lowerAsEntryFunction(gpu::GPUFuncOp funcOp, SPIRVTypeConverter &typeConverter,
                      spirv::EntryPointABIAttr entryPointInfo,
                      ArrayRef<spirv::InterfaceVarABIAttr> argABIInfo) {
   auto fnType = funcOp.getType();
   if (fnType.getNumResults()) {
     funcOp.emitError("SPIR-V lowering only supports entry functions"
                      "with no return values right now");
     return nullptr;
   }
-  if (fnType.getNumInputs() != argABIInfo.size()) {
+  if (!argABIInfo.empty() && fnType.getNumInputs() != argABIInfo.size()) {
     funcOp.emitError(
-        "lowering as entry functions requires ABI info for all arguments");
+        "lowering as entry functions requires ABI info for all arguments "
+        "or none of them");
     return nullptr;
   }
   // Update the signature to valid SPIR-V types and add the ABI
   // attributes. These will be "materialized" by using the
   // LowerABIAttributesPass.
   TypeConverter::SignatureConversion signatureConverter(fnType.getNumInputs());
   {
     for (auto argType : enumerate(funcOp.getType().getInputs())) {

 }
 
 /// Populates `argABI` with spv.interface_var_abi attributes for lowering
 /// gpu.func to spv.func if no arguments have the attributes set
 /// already. Returns failure if any argument has the ABI attribute set already.
 static LogicalResult
 getDefaultABIAttrs(MLIRContext *context, gpu::GPUFuncOp funcOp,
                    SmallVectorImpl<spirv::InterfaceVarABIAttr> &argABI) {
+  spirv::TargetEnvAttr targetEnv = spirv::lookupTargetEnvOrDefault(funcOp);
+  if (!spirv::needsInterfaceVarABIAttrs(targetEnv))
+    return success();
+
   for (auto argIndex : llvm::seq<unsigned>(0, funcOp.getNumArguments())) {
     if (funcOp.getArgAttrOfType<spirv::InterfaceVarABIAttr>(
             argIndex, spirv::getInterfaceVarABIAttrName()))
       return failure();
     // Vulkan's interface variable requirements needs scalars to be wrapped in a
     // struct. The struct held in storage buffer.
     Optional<spirv::StorageClass> sc;
     if (funcOp.getArgument(argIndex).getType().isIntOrIndexOrFloat())

 
 //===----------------------------------------------------------------------===//
 // ModuleOp with gpu.module.
 //===----------------------------------------------------------------------===//
 
 LogicalResult GPUModuleConversion::matchAndRewrite(
     gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
+  spirv::TargetEnvAttr targetEnv = spirv::lookupTargetEnvOrDefault(moduleOp);
+  spirv::AddressingModel addressingModel = spirv::getAddressingModel(targetEnv);
+  FailureOr<spirv::MemoryModel> memoryModel = spirv::getMemoryModel(targetEnv);
+  if (failed(memoryModel))
+    return moduleOp.emitRemark("match failure: could not selected memory model "
+                               "based on 'spv.target_env'");
+
   auto spvModule = rewriter.create<spirv::ModuleOp>(
-      moduleOp.getLoc(), spirv::AddressingModel::Logical,
-      spirv::MemoryModel::GLSL450);
+      moduleOp.getLoc(), addressingModel, memoryModel.getValue());
 
   // Move the region from the module op into the SPIR-V module.
   Region &spvModuleRegion = spvModule.body();
   rewriter.inlineRegionBefore(moduleOp.body(), spvModuleRegion,
                               spvModuleRegion.begin());
   // The spv.module build method adds a block with a terminator. Remove that
   // block. The terminator of the module op in the remaining block will be
   // legalized later.

mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp

 //===----------------------------------------------------------------------===//
 
 LogicalResult
 mlir::spirv::setABIAttrs(spirv::FuncOp funcOp,
                          spirv::EntryPointABIAttr entryPointInfo,
                          ArrayRef<spirv::InterfaceVarABIAttr> argABIInfo) {
   // Set the attributes for argument and the function.
   StringRef argABIAttrName = spirv::getInterfaceVarABIAttrName();
-  for (auto argIndex : llvm::seq<unsigned>(0, funcOp.getNumArguments())) {
+  for (auto argIndex : llvm::seq<unsigned>(0, argABIInfo.size())) {
     funcOp.setArgAttr(argIndex, argABIAttrName, argABIInfo[argIndex]);
   }
   funcOp.setAttr(spirv::getEntryPointABIAttrName(), entryPointInfo);
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // SPIR-V ConversionTarget

mlir/lib/Dialect/SPIRV/TargetAndABI.cpp

 spirv::InterfaceVarABIAttr
 spirv::getInterfaceVarABIAttr(unsigned descriptorSet, unsigned binding,
                               Optional<spirv::StorageClass> storageClass,
                               MLIRContext *context) {
   return spirv::InterfaceVarABIAttr::get(descriptorSet, binding, storageClass,
                                          context);
 }
 
+bool spirv::needsInterfaceVarABIAttrs(spirv::TargetEnvAttr targetAttr) {
+  for (spirv::Capability cap : targetAttr.getCapabilities()) {
+    if (cap == spirv::Capability::Kernel)
+      return false;
+    if (cap == spirv::Capability::Shader)
+      return true;
+  }
+  return false;
+}
+
 StringRef spirv::getEntryPointABIAttrName() { return "spv.entry_point_abi"; }
 
 spirv::EntryPointABIAttr
 spirv::getEntryPointABIAttr(ArrayRef<int32_t> localSize, MLIRContext *context) {
   assert(localSize.size() == 3);
   return spirv::EntryPointABIAttr::get(
       DenseElementsAttr::get<int32_t>(
           VectorType::get(3, IntegerType::get(32, context)), localSize)

 }
 
 spirv::TargetEnvAttr spirv::lookupTargetEnvOrDefault(Operation *op) {
   if (spirv::TargetEnvAttr attr = spirv::lookupTargetEnv(op))
     return attr;
 
   return getDefaultTargetEnv(op->getContext());
 }
+
+spirv::AddressingModel
+spirv::getAddressingModel(spirv::TargetEnvAttr targetAttr) {
+  for (spirv::Capability cap : targetAttr.getCapabilities()) {
+    // TODO: Physical64 is hard-coded here, but some information should come
+    // from TargetEnvAttr to selected between Physical32 and Physical64.
+    if (cap == Capability::Kernel)
+      return spirv::AddressingModel::Physical64;
+  }
+  // Logical addressing doesn't need any capabilities so return it as default.
+  return spirv::AddressingModel::Logical;
+}
+
+FailureOr<spirv::ExecutionModel>
+spirv::getExecutionModel(spirv::TargetEnvAttr targetAttr) {
+  for (spirv::Capability cap : targetAttr.getCapabilities()) {
+    if (cap == spirv::Capability::Kernel)
+      return spirv::ExecutionModel::Kernel;
+    if (cap == spirv::Capability::Shader)
+      return spirv::ExecutionModel::GLCompute;
+  }
+  return failure();
+}
+
+FailureOr<spirv::MemoryModel>
+spirv::getMemoryModel(spirv::TargetEnvAttr targetAttr) {
+  for (spirv::Capability cap : targetAttr.getCapabilities()) {
+    if (cap == spirv::Capability::Addresses)
+      return spirv::MemoryModel::OpenCL;
+    if (cap == spirv::Capability::Shader)
+      return spirv::MemoryModel::GLSL450;
+  }
+  return failure();
+}

mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp

static LogicalResult lowerEntryPointABIAttr(spirv::FuncOp funcOp,
   builder.setInsertionPoint(spirvModule.body().front().getTerminator());
 
   // Adds the spv.EntryPointOp after collecting all the interface variables
   // needed.
   SmallVector<Attribute, 1> interfaceVars;
   if (failed(getInterfaceVariables(funcOp, interfaceVars))) {
     return failure();
   }
+
+  spirv::TargetEnvAttr targetEnv = spirv::lookupTargetEnv(funcOp);
+  FailureOr<spirv::ExecutionModel> executionModel =
+      spirv::getExecutionModel(targetEnv);
+  if (failed(executionModel))
+    return funcOp.emitRemark("lower entry point failure: could not select "
+                             "execution model based on 'spv.target_env'");
+
   builder.create<spirv::EntryPointOp>(
-      funcOp.getLoc(), spirv::ExecutionModel::GLCompute, funcOp, interfaceVars);
+      funcOp.getLoc(), executionModel.getValue(), funcOp, interfaceVars);
+
   // Specifies the spv.ExecutionModeOp.
   auto localSizeAttr = entryPointAttr.local_size();
   SmallVector<int32_t, 3> localSize(localSizeAttr.getValues<int32_t>());
   builder.create<spirv::ExecutionModeOp>(
       funcOp.getLoc(), funcOp, spirv::ExecutionMode::LocalSize, localSize);
   funcOp.removeAttr(entryPointAttrName);
   return success();
 }
typeConverter.addSourceMaterialization([](OpBuilder &builder,
                                             spirv::PointerType type,
                                             ValueRange inputs, Location loc) {
     if (inputs.size() != 1 || !inputs[0].getType().isa<spirv::PointerType>())
       return Value();
     return builder.create<spirv::BitcastOp>(loc, type, inputs[0]).getResult();
   });
 
   OwningRewritePatternList patterns;
   patterns.insert<ProcessInterfaceVarABI>(context, typeConverter);
 
   ConversionTarget target(*context);
   // "Legal" function ops should have no interface variable ABI attributes.
   target.addDynamicallyLegalOp<spirv::FuncOp>([&](spirv::FuncOp op) {
     StringRef attrName = spirv::getInterfaceVarABIAttrName();
     for (unsigned i = 0, e = op.getNumArguments(); i < e; ++i)
       if (op.getArgAttr(i, attrName))
         return false;

mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir

+// RUN: mlir-opt -allow-unregistered-dialect -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
+
+module attributes {
+  gpu.container_module,
+  spv.target_env = #spv.target_env<
+    #spv.vce<v1.0, [Kernel, Addresses], []>,
+    {max_compute_workgroup_invocations = 128 : i32,
+     max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  gpu.module @kernels {
+    // CHECK-LABEL: spv.module Physical64 OpenCL
+    //       CHECK:   spv.func
+    //  CHECK-SAME:     {{%.*}}: f32
+    //   CHECK-NOT:     spv.interface_var_abi
+    //  CHECK-SAME:     {{%.*}}: !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=4> [0]>, CrossWorkgroup>
+    //   CHECK-NOT:     spv.interface_var_abi
+    //  CHECK-SAME:     spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
+    gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, 11>) kernel
+        attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
+      gpu.return
+    }
+  }
+
+  func @main() {
+    %0 = "op"() : () -> (f32)
+    %1 = "op"() : () -> (memref<12xf32, 11>)
+    %cst = constant 1 : index
+    "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@basic_module_structure }
+        : (index, index, index, index, index, index, f32, memref<12xf32, 11>) -> ()
+    return
+  }
+}

mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir

+// RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
+
+module attributes {
+  spv.target_env = #spv.target_env<
+    #spv.vce<v1.0, [Kernel, Addresses], []>,
+    {max_compute_workgroup_invocations = 128 : i32,
+     max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  spv.module Physical64 OpenCL {
+    // CHECK-LABEL: spv.module
+    //       CHECK:   spv.func [[FN:@.*]](
+    //  CHECK-SAME:     {{%.*}}: f32
+    //  CHECK-SAME:     {{%.*}}: !spv.ptr<!spv.struct<!spv.array<12 x f32>>, CrossWorkgroup>
+    //       CHECK:   spv.EntryPoint "Kernel" [[FN]]
+    //       CHECK:   spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
+    spv.func @kernel(
+      %arg0: f32,
+      %arg1: !spv.ptr<!spv.struct<!spv.array<12 x f32>>, CrossWorkgroup>) "None"
+    attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
+      spv.Return
+    }
+  }
+}