[mlir][spirv] Allow controlling subgroup size

[antiagainst]

Summary

This commit extends the ResourceLimitsAttr to support specifying
a minimal and maximal subgroup size, and extends EntryPointABIAttr
to support specifying the requested subgroup size. This is possible
now in Vulkan with the VK_EXT_subgroup_size_control extension.
For OpenCL it's possible to use the SubgroupSize execution mode
directly.

mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td

 
 class SPIRV_Attr<string attrName, string attrMnemonic>
     : AttrDef<SPIRV_Dialect, attrName> {
   let mnemonic = attrMnemonic;
 }
 
 // For entry functions, this attribute specifies information related to entry
 // points in the generated SPIR-V module:
-// 1) WorkGroup Size.
+// 1) [optional] Requested workgroup size.
+// 2) [optional] Requested subgroup size.
 def SPIRV_EntryPointABIAttr : SPIRV_Attr<"EntryPointABI", "entry_point_abi"> {
-  let parameters = (ins OptionalParameter<"DenseIntElementsAttr">:$local_size);
+  let parameters = (ins
+    OptionalParameter<"DenseI32ArrayAttr">:$workgroup_size,
+    OptionalParameter<"llvm::Optional<int>">:$subgroup_size
+  );
   let assemblyFormat = "`<` struct(params) `>`";
 }
 
 def SPIRV_ExtensionArrayAttr : TypedArrayAttrBase<
     SPIRV_ExtensionAttr, "SPIR-V extension array attribute">;
 
 def SPIRV_CapabilityArrayAttr : TypedArrayAttrBase<
     SPIRV_CapabilityAttr, "SPIR-V capability array attribute">;
let parameters = (ins
     DefaultValuedParameter<
       "ArrayAttr",
       "$_builder.getI32ArrayAttr({128, 128, 64})"
     >:$max_compute_workgroup_size,
 
     // The default number of invocations in each subgroup.
     DefaultValuedParameter<"int", "32">:$subgroup_size,
 
+    // The minimum supported size if the subgroup size is controllable.
+    OptionalParameter<"mlir::Optional<int>">:$min_subgroup_size,
+    // The maximum supported size if the subgroup size is controllable.
+    OptionalParameter<"mlir::Optional<int>">:$max_subgroup_size,
+
     // The configurations of cooperative matrix operations
     // supported. Default is an empty list.
     DefaultValuedParameter<
       "ArrayAttr",
       "nullptr"
     >:$cooperative_matrix_properties_nv
   );
   let assemblyFormat = "`<` struct(params) `>`";
 }
 
 #endif // MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI

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

 /// Returns whether the given SPIR-V target (described by TargetEnvAttr) needs
 /// ABI attributes for interface variables (spirv.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);
+EntryPointABIAttr getEntryPointABIAttr(MLIRContext *context,
+                                       ArrayRef<int32_t> workgroupSize = {},
+                                       llvm::Optional<int> subgroupSize = {});
 
 /// Queries the entry point ABI on the nearest function-like op containing the
 /// given `op`. Returns null attribute if not found.
 EntryPointABIAttr lookupEntryPointABI(Operation *op);
 
 /// Queries the local workgroup size from entry point ABI on the nearest
 /// function-like op containing the given `op`. Returns null attribute if not
 /// found.
-DenseIntElementsAttr lookupLocalWorkGroupSize(Operation *op);
+DenseI32ArrayAttr lookupLocalWorkGroupSize(Operation *op);
 
 /// Returns a default resource limits attribute that uses numbers from
 /// "Table 46. Required Limits" of the Vulkan spec.
 ResourceLimitsAttr getDefaultResourceLimits(MLIRContext *context);
 
 /// Returns the attribute name for specifying SPIR-V target environment.
 StringRef getTargetEnvAttrName();
 

mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp

auto spirvBuiltin =
       spirv::getBuiltinVariableValue(op, builtin, indexType, rewriter);
   rewriter.replaceOp(op, spirvBuiltin);
   return success();
 }
 
 LogicalResult WorkGroupSizeConversion::matchAndRewrite(
     gpu::BlockDimOp op, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
-  auto workGroupSizeAttr = spirv::lookupLocalWorkGroupSize(op);
+  DenseI32ArrayAttr workGroupSizeAttr = spirv::lookupLocalWorkGroupSize(op);
   if (!workGroupSizeAttr)
     return failure();
 
-  auto val = workGroupSizeAttr
-                 .getValues<int32_t>()[static_cast<int32_t>(op.getDimension())];
+  int val =
+      workGroupSizeAttr.asArrayRef()[static_cast<int32_t>(op.getDimension())];
   auto convertedType =
       getTypeConverter()->convertType(op.getResult().getType());
   if (!convertedType)
     return failure();
   rewriter.replaceOpWithNewOp<spirv::ConstantOp>(
       op, convertedType, IntegerAttr::get(convertedType, val));
   return success();
 }

mlir/lib/Conversion/LinalgToSPIRV/LinalgToSPIRV.cpp

LogicalResult SingleWorkgroupReduction::matchAndRewrite(
   auto originalOutputType = op->getOperand(1).getType().cast<MemRefType>();
 
   auto binaryOpKind = matchAsPerformingReduction(genericOp);
   if (!binaryOpKind)
     return failure();
 
   // Query the shader interface for local workgroup size to make sure the
   // invocation configuration fits with the input memref's shape.
-  DenseIntElementsAttr localSize = spirv::lookupLocalWorkGroupSize(genericOp);
-  if (!localSize)
+  DenseI32ArrayAttr workgroupSize = spirv::lookupLocalWorkGroupSize(genericOp);
+  if (!workgroupSize)
     return failure();
 
-  if ((*localSize.begin()).getSExtValue() != originalInputType.getDimSize(0))
+  if (workgroupSize.asArrayRef()[0] != originalInputType.getDimSize(0))
     return failure();
-  if (llvm::any_of(llvm::drop_begin(localSize.getValues<APInt>(), 1),
-                   [](const APInt &size) { return !size.isOneValue(); }))
+  if (llvm::any_of(workgroupSize.asArrayRef().drop_front(),
+                   [](int size) { return size != 1; }))
     return failure();
 
   // TODO: Query the target environment to make sure the current
   // workload fits in a local workgroup.
 
   Value convertedInput = adaptor.getOperands()[0];
   Value convertedOutput = adaptor.getOperands()[1];
   Location loc = genericOp.getLoc();

mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp

if (cap == spirv::Capability::Shader)
       return true;
   }
   return false;
 }
 
 StringRef spirv::getEntryPointABIAttrName() { return "spirv.entry_point_abi"; }
 
 spirv::EntryPointABIAttr
-spirv::getEntryPointABIAttr(ArrayRef<int32_t> localSize, MLIRContext *context) {
-  if (localSize.empty())
-    return spirv::EntryPointABIAttr::get(context, nullptr);
-
-  assert(localSize.size() == 3);
-  return spirv::EntryPointABIAttr::get(
-      context, DenseElementsAttr::get<int32_t>(
-                   VectorType::get(3, IntegerType::get(context, 32)), localSize)
-                   .cast<DenseIntElementsAttr>());
+spirv::getEntryPointABIAttr(MLIRContext *context,
+                            ArrayRef<int32_t> workgroupSize,
+                            llvm::Optional<int> subgroupSize) {
+  DenseI32ArrayAttr workgroupSizeAttr;
+  if (!workgroupSize.empty()) {
+    assert(workgroupSize.size() == 3);
+    workgroupSizeAttr = DenseI32ArrayAttr::get(context, workgroupSize);
+  }
+  return spirv::EntryPointABIAttr::get(context, workgroupSizeAttr,
+                                       /*subgroupSize=*/llvm::None);
 }
 
 spirv::EntryPointABIAttr spirv::lookupEntryPointABI(Operation *op) {
   while (op && !isa<FunctionOpInterface>(op))
     op = op->getParentOp();
   if (!op)
     return {};
 
   if (auto attr = op->getAttrOfType<spirv::EntryPointABIAttr>(
           spirv::getEntryPointABIAttrName()))
     return attr;
 
   return {};
 }
 
-DenseIntElementsAttr spirv::lookupLocalWorkGroupSize(Operation *op) {
+DenseI32ArrayAttr spirv::lookupLocalWorkGroupSize(Operation *op) {
   if (auto entryPoint = spirv::lookupEntryPointABI(op))
-    return entryPoint.getLocalSize();
+    return entryPoint.getWorkgroupSize();
 
   return {};
 }
 
 spirv::ResourceLimitsAttr
 spirv::getDefaultResourceLimits(MLIRContext *context) {
   // All the fields have default values. Here we just provide a nicer way to
   // construct a default resource limit attribute.
   Builder b(context);
   return spirv::ResourceLimitsAttr::get(
       context,
       /*max_compute_shared_memory_size=*/16384,
       /*max_compute_workgroup_invocations=*/128,
       /*max_compute_workgroup_size=*/b.getI32ArrayAttr({128, 128, 64}),
       /*subgroup_size=*/32,
+      /*min_subgroup_size=*/llvm::None,
+      /*max_subgroup_size=*/llvm::None,
       /*cooperative_matrix_properties_nv=*/ArrayAttr());
 }
 
 StringRef spirv::getTargetEnvAttrName() { return "spirv.target_env"; }
 
 spirv::TargetEnvAttr spirv::getDefaultTargetEnv(MLIRContext *context) {
   auto triple = spirv::VerCapExtAttr::get(spirv::Version::V_1_0,
                                           {spirv::Capability::Shader},

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

 //===- LowerABIAttributesPass.cpp - Decorate composite type ---------------===//
 //
 // 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 a pass to lower attributes that specify the shader ABI
 // for the functions in the generated SPIR-V module.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/Transforms/Passes.h"
 
+#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h"
 #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h"
 #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
+#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h"
 #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h"
 #include "mlir/Dialect/SPIRV/Utils/LayoutUtils.h"
+#include "mlir/IR/BuiltinAttributes.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "llvm/ADT/SetVector.h"
 
 namespace mlir {
 namespace spirv {
 #define GEN_PASS_DEF_SPIRVLOWERABIATTRIBUTES
 #include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc"
 } // namespace spirv
static LogicalResult lowerEntryPointABIAttr(spirv::FuncOp funcOp,
 
   // Adds the spirv.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);
+  spirv::TargetEnvAttr targetEnvAttr = spirv::lookupTargetEnv(funcOp);
+  spirv::TargetEnv targetEnv(targetEnvAttr);
   FailureOr<spirv::ExecutionModel> executionModel =
-      spirv::getExecutionModel(targetEnv);
+      spirv::getExecutionModel(targetEnvAttr);
   if (failed(executionModel))
     return funcOp.emitRemark("lower entry point failure: could not select "
                              "execution model based on 'spirv.target_env'");
 
   builder.create<spirv::EntryPointOp>(funcOp.getLoc(), executionModel.value(),
                                       funcOp, interfaceVars);
 
   // Specifies the spirv.ExecutionModeOp.
-  auto localSizeAttr = entryPointAttr.getLocalSize();
-  if (localSizeAttr) {
-    auto values = localSizeAttr.getValues<int32_t>();
-    SmallVector<int32_t, 3> localSize(values);
-    builder.create<spirv::ExecutionModeOp>(
-        funcOp.getLoc(), funcOp, spirv::ExecutionMode::LocalSize, localSize);
+  if (DenseI32ArrayAttr workgroupSizeAttr = entryPointAttr.getWorkgroupSize()) {
+    Optional<ArrayRef<spirv::Capability>> caps =
+        spirv::getCapabilities(spirv::ExecutionMode::LocalSize);
+    if (!caps || targetEnv.allows(*caps)) {
+      builder.create<spirv::ExecutionModeOp>(funcOp.getLoc(), funcOp,
+                                             spirv::ExecutionMode::LocalSize,
+                                             workgroupSizeAttr.asArrayRef());
+      // Erase workgroup size.
+      entryPointAttr = spirv::EntryPointABIAttr::get(
+          entryPointAttr.getContext(), DenseI32ArrayAttr(),
+          entryPointAttr.getSubgroupSize());
+    }
+  }
+  if (Optional<int> subgroupSize = entryPointAttr.getSubgroupSize()) {
+    Optional<ArrayRef<spirv::Capability>> caps =
+        spirv::getCapabilities(spirv::ExecutionMode::SubgroupSize);
+    if (!caps || targetEnv.allows(*caps)) {
+      builder.create<spirv::ExecutionModeOp>(funcOp.getLoc(), funcOp,
+                                             spirv::ExecutionMode::SubgroupSize,
+                                             *subgroupSize);
+      // Erase subgroup size.
+      entryPointAttr = spirv::EntryPointABIAttr::get(
+          entryPointAttr.getContext(), entryPointAttr.getWorkgroupSize(),
+          llvm::None);
+    }
+  }
+  if (entryPointAttr.getWorkgroupSize() || entryPointAttr.getSubgroupSize())
+    funcOp->setAttr(entryPointAttrName, entryPointAttr);
+  else
     funcOp->removeAttr(entryPointAttrName);
-  }
   return success();
 }
 
 namespace {
 /// A pattern to convert function signature according to interface variable ABI
 /// attributes.
 ///
 /// Specifically, this pattern creates global variables according to interface

mlir/test/Conversion/GPUToSPIRV/builtins.mlir

 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_x() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.block_id x
       gpu.return
     }
   }
 }
gpu.launch_func @kernels::@builtin_workgroup_id_y
         dynamic_shared_memory_size %c256
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_y() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = gpu.block_id y
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_z() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = gpu.block_id z
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_x() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {
       // The constant value is obtained from the spirv.entry_point_abi.
       // Note that this ignores the workgroup size specification in gpu.launch.
       // We may want to define gpu.workgroup_size and convert it to the entry
       // point ABI we want here.
       // CHECK: spirv.Constant 32 : i32
       %0 = gpu.block_dim x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_y
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_y() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // The constant value is obtained from the spirv.entry_point_abi.
       // CHECK: spirv.Constant 4 : i32
       %0 = gpu.block_dim y
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_z() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // The constant value is obtained from the spirv.entry_point_abi.
       // CHECK: spirv.Constant 1 : i32
       %0 = gpu.block_dim z
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_local_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
   gpu.module @kernels {
     gpu.func @builtin_local_id_x() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LOCALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.thread_id x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_num_workgroups_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
   gpu.module @kernels {
     gpu.func @builtin_num_workgroups_x() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMWORKGROUPS]]
       // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.grid_dim x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId")
   gpu.module @kernels {
     gpu.func @builtin_subgroup_id() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPID]]
       // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_id : index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups")
   gpu.module @kernels {
     gpu.func @builtin_num_subgroups() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMSUBGROUPS]]
       // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
       %0 = gpu.num_subgroups : index
       gpu.return
     }
   }
 }
 
gpu.launch_func @kernels::@builtin_global_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
   gpu.module @kernels {
     gpu.func @builtin_global_id_x() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.global_id x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_global_id_y
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
   gpu.module @kernels {
     gpu.func @builtin_global_id_y() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = gpu.global_id y
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_global_id_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
   gpu.module @kernels {
     gpu.func @builtin_global_id_z() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = gpu.global_id z
       gpu.return
     }
   }
 }
 
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
   // CHECK: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize")
   gpu.module @kernels {
     gpu.func @builtin_subgroup_size() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]]
       // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_size : index
       gpu.return
     }
   }
 }

mlir/test/Conversion/GPUToSPIRV/entry-point.mlir

 // RUN: mlir-opt -test-spirv-entry-point-abi %s | FileCheck %s -check-prefix=DEFAULT
 // RUN: mlir-opt -test-spirv-entry-point-abi="workgroup-size=32" %s | FileCheck %s -check-prefix=WG32
 
 //      DEFAULT: gpu.func @foo()
-// DEFAULT-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<1> : vector<3xi32>>
+// DEFAULT-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>
 
 //      WG32: gpu.func @foo()
-// WG32-SAME:  spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>
+// WG32-SAME:  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>
 
 gpu.module @kernels {
   gpu.func @foo() kernel {
     gpu.return
   }
 }

mlir/test/Conversion/GPUToSPIRV/gpu-to-spirv.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spirv.func @basic_module_structure
     // CHECK-SAME: {{%.*}}: f32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0), StorageBuffer>}
     // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}
-    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]> : vector<3xi32>>
+    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
     gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // CHECK: spirv.Return
       gpu.return
     }
   }
 
   func.func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<StorageBuffer>>)

 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spirv.func @basic_module_structure_preset_ABI
     // CHECK-SAME: {{%[a-zA-Z0-9_]*}}: f32
     // CHECK-SAME: spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer>
     // CHECK-SAME: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
     // CHECK-SAME: spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)>
-    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]> : vector<3xi32>>
+    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
     gpu.func @basic_module_structure_preset_ABI(
       %arg0 : f32
         {spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer>},
       %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>
         {spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)>}) kernel
       attributes
-        {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+        {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // CHECK: spirv.Return
       gpu.return
     }
   }
 }
 
 // -----
 
module attributes {gpu.container_module} {
   gpu.module @kernels {
     // expected-error @below {{failed to legalize operation 'gpu.func'}}
     // expected-remark @below {{match failure: missing 'spirv.interface_var_abi' attribute at argument 1}}
     gpu.func @missing_entry_point_abi(
       %arg0 : f32
         {spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer>},
       %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel
     attributes
-      {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // expected-error @below {{failed to legalize operation 'gpu.func'}}
     // expected-remark @below {{match failure: missing 'spirv.interface_var_abi' attribute at argument 0}}
     gpu.func @missing_entry_point_abi(
       %arg0 : f32,
       %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>
         {spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)>}) kernel
     attributes
-      {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // CHECK-LABEL: spirv.func @barrier
     gpu.func @barrier(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // CHECK: spirv.ControlBarrier <Workgroup>, <Workgroup>, <AcquireRelease|WorkgroupMemory>
       gpu.barrier
       gpu.return
     }
   }
 
   func.func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<StorageBuffer>>)
     %cst = arith.constant 1 : index
     gpu.launch_func @kernels::@barrier
         blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
         args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>)
     return
   }
 }

mlir/test/Conversion/GPUToSPIRV/load-store.mlir

gpu.module @kernels {
     // CHECK-SAME: %[[ARG0:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}
     // CHECK-SAME: %[[ARG1:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}
     // CHECK-SAME: %[[ARG2:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 2)>}
     // CHECK-SAME: %[[ARG3:.*]]: i32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 3), StorageBuffer>}
     // CHECK-SAME: %[[ARG4:.*]]: i32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 4), StorageBuffer>}
     // CHECK-SAME: %[[ARG5:.*]]: i32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 5), StorageBuffer>}
     // CHECK-SAME: %[[ARG6:.*]]: i32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 6), StorageBuffer>}
     gpu.func @load_store_kernel(%arg0: memref<12x4xf32, #spirv.storage_class<StorageBuffer>>, %arg1: memref<12x4xf32, #spirv.storage_class<StorageBuffer>>, %arg2: memref<12x4xf32, #spirv.storage_class<StorageBuffer>>, %arg3: index, %arg4: index, %arg5: index, %arg6: index) kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
       // CHECK: %[[ADDRESSWORKGROUPID:.*]] = spirv.mlir.addressof @[[$WORKGROUPIDVAR]]
       // CHECK: %[[WORKGROUPID:.*]] = spirv.Load "Input" %[[ADDRESSWORKGROUPID]]
       // CHECK: %[[WORKGROUPIDX:.*]] = spirv.CompositeExtract %[[WORKGROUPID]]{{\[}}0 : i32{{\]}}
       // CHECK: %[[ADDRESSLOCALINVOCATIONID:.*]] = spirv.mlir.addressof @[[$LOCALINVOCATIONIDVAR]]
       // CHECK: %[[LOCALINVOCATIONID:.*]] = spirv.Load "Input" %[[ADDRESSLOCALINVOCATIONID]]
       // CHECK: %[[LOCALINVOCATIONIDX:.*]] = spirv.CompositeExtract %[[LOCALINVOCATIONID]]{{\[}}0 : i32{{\]}}
       %0 = gpu.block_id x
       %1 = gpu.block_id y

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

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

mlir/test/Conversion/GPUToSPIRV/shuffle.mlir

 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniformShuffle], []>, #spirv.resource_limits<subgroup_size = 16>>
 } {
 
 gpu.module @kernels {
   // CHECK-LABEL:  spirv.func @shuffle_xor()
   gpu.func @shuffle_xor() kernel
-    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
     %mask = arith.constant 8 : i32
     %width = arith.constant 16 : i32
     %val = arith.constant 42.0 : f32
 
     // CHECK: %[[MASK:.+]] = spirv.Constant 8 : i32
     // CHECK: %[[VAL:.+]] = spirv.Constant 4.200000e+01 : f32
     // CHECK: %{{.+}} = spirv.Constant true
     // CHECK: %{{.+}} = spirv.GroupNonUniformShuffleXor <Subgroup> %[[VAL]], %[[MASK]] : f32, i32
     %result, %valid = gpu.shuffle xor %val, %mask, %width : f32
     gpu.return
   }
 }
 
 }
 
 // -----
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniformShuffle], []>, #spirv.resource_limits<subgroup_size = 32>>
 } {
 
 gpu.module @kernels {
   gpu.func @shuffle_xor() kernel
-    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
     %mask = arith.constant 8 : i32
     %width = arith.constant 16 : i32
     %val = arith.constant 42.0 : f32
 
     // Cannot convert due to shuffle width and target subgroup size mismatch
     // expected-error @+1 {{failed to legalize operation 'gpu.shuffle'}}
     %result, %valid = gpu.shuffle xor %val, %mask, %width : f32
     gpu.return
   }
 }
 
 }
 
 // -----
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniformShuffle], []>, #spirv.resource_limits<subgroup_size = 16>>
 } {
 
 gpu.module @kernels {
   // CHECK-LABEL:  spirv.func @shuffle_idx()
   gpu.func @shuffle_idx() kernel
-    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
+    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
     %mask = arith.constant 8 : i32
     %width = arith.constant 16 : i32
     %val = arith.constant 42.0 : f32
 
     // CHECK: %[[MASK:.+]] = spirv.Constant 8 : i32
     // CHECK: %[[VAL:.+]] = spirv.Constant 4.200000e+01 : f32
     // CHECK: %{{.+}} = spirv.Constant true
     // CHECK: %{{.+}} = spirv.GroupNonUniformShuffle <Subgroup> %[[VAL]], %[[MASK]] : f32, i32
     %result, %valid = gpu.shuffle idx %val, %mask, %width : f32
     gpu.return
   }
 }
 
 }

mlir/test/Conversion/GPUToSPIRV/wmma-ops-to-spirv.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -convert-gpu-to-spirv -split-input-file -verify-diagnostics %s | FileCheck %s
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, CooperativeMatrixNV, Float16], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix]>, #spirv.resource_limits<>>} {
   gpu.module @kernels {
     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spirv.func @gpu_wmma_load_op
     // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<512 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}
-    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]> : vector<3xi32>>
+    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
     gpu.func @gpu_wmma_load_op(%arg0 : memref<32x32xf16, #spirv.storage_class<StorageBuffer>>) kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       %i = arith.constant 16 : index
       %j = arith.constant 16 : index
       // CHECK: {{%.*}} = spirv.NV.CooperativeMatrixLoad {{%.*}}, {{%.*}}, {{%.*}} :  !spirv.ptr<f32, StorageBuffer> as !spirv.coopmatrix<16x16xf16, Subgroup>
       %0 = gpu.subgroup_mma_load_matrix %arg0[%i, %j] {leadDimension = 32 : index} : memref<32x32xf16, #spirv.storage_class<StorageBuffer>> -> !gpu.mma_matrix<16x16xf16, "COp">
       // CHECK: spirv.Return
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, CooperativeMatrixNV, Float16], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix]>, #spirv.resource_limits<>>} {
   gpu.module @kernels {
     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spirv.func @gpu_wmma_store_op
     // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<512 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}
     // CHECK-SAME: {{%.*}}: !spirv.coopmatrix<16x16xf16, Subgroup> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>})
-    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]> : vector<3xi32>>
+    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
     gpu.func @gpu_wmma_store_op(%arg0 : memref<32x32xf16, #spirv.storage_class<StorageBuffer>>, %arg1 : !gpu.mma_matrix<16x16xf16, "COp">) kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       %i = arith.constant 16 : index
       %j = arith.constant 16 : index
       //  CHECK: spirv.NV.CooperativeMatrixStore {{%.*}}, {{%.*}}, {{%.*}}, {{%.*}} : !spirv.ptr<f32, StorageBuffer>, !spirv.coopmatrix<16x16xf16, Subgroup>
       gpu.subgroup_mma_store_matrix %arg1, %arg0[%i,%j] {leadDimension= 32 : index} : !gpu.mma_matrix<16x16xf16, "COp">, memref<32x32xf16,  #spirv.storage_class<StorageBuffer>>
       // CHECK: spirv.Return
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, CooperativeMatrixNV, Float16], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix]>, #spirv.resource_limits<>>} {
   gpu.module @kernels {
     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spirv.func @gpu_wmma_mma_op
     // CHECK-SAME: {{%.*}}: !spirv.coopmatrix<16x16xf16, Subgroup> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}
     // CHECK-SAME: {{%.*}}: !spirv.coopmatrix<16x16xf16, Subgroup> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}
     // CHECK-SAME: {{%.*}}: !spirv.coopmatrix<16x16xf16, Subgroup> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 2)>})
-    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]> : vector<3xi32>>
+    // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
     gpu.func @gpu_wmma_mma_op(%A : !gpu.mma_matrix<16x16xf16, "AOp">, %B : !gpu.mma_matrix<16x16xf16, "BOp">, %C : !gpu.mma_matrix<16x16xf16, "COp">) kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // CHECK: {{%.*}} = spirv.NV.CooperativeMatrixMulAdd {{%.*}}, {{%.*}}, {{%.*}} : !spirv.coopmatrix<16x16xf16, Subgroup>, !spirv.coopmatrix<16x16xf16, Subgroup> -> !spirv.coopmatrix<16x16xf16, Subgroup>
       %D = gpu.subgroup_mma_compute %A, %B, %C : !gpu.mma_matrix<16x16xf16, "AOp">, !gpu.mma_matrix<16x16xf16, "BOp"> -> !gpu.mma_matrix<16x16xf16, "COp">
       // CHECK: spirv.Return
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, CooperativeMatrixNV, Float16], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix]>, #spirv.resource_limits<>>} {
   gpu.module @kernels {
     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spirv.func @gpu_wmma_constant_op
     gpu.func @gpu_wmma_constant_op() kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // CHECK: {{%.*}} = spirv.Constant
       %cst = arith.constant 1.0 : f16
       // CHECK: {{%.*}} = spirv.CompositeConstruct {{%.*}} : (f16) -> !spirv.coopmatrix<16x16xf16, Subgroup>
       %C = gpu.subgroup_mma_constant_matrix %cst : !gpu.mma_matrix<16x16xf16, "COp">
       // CHECK: spirv.Return
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [Shader, CooperativeMatrixNV, Float16], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix]>, #spirv.resource_limits<>>} {
   gpu.module @kernels {
     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spirv.func @gpu_wmma_elementwise_op
     // CHECK-SAME: {{%.*}}: !spirv.coopmatrix<16x16xf16, Subgroup> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}
     // CHECK-SAME: {{%.*}}: !spirv.coopmatrix<16x16xf16, Subgroup> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>})
     gpu.func @gpu_wmma_elementwise_op(%A : !gpu.mma_matrix<16x16xf16, "COp">, %B : !gpu.mma_matrix<16x16xf16, "COp">) kernel
-      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // CHECK:  {{%.*}} = spirv.FAdd {{%.*}}, {{%.*}} : !spirv.coopmatrix<16x16xf16, Subgroup>
       %C = gpu.subgroup_mma_elementwise addf %A, %B : (!gpu.mma_matrix<16x16xf16, "COp">, !gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf16, "COp">
       // CHECK:  {{%.*}} = spirv.FNegate {{%.*}} : !spirv.coopmatrix<16x16xf16, Subgroup>
       %D = gpu.subgroup_mma_elementwise negatef %C : (!gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf16, "COp">
       // CHECK:  {{%.*}} = spirv.FDiv {{%.*}}, {{%.*}} : !spirv.coopmatrix<16x16xf16, Subgroup>
       %E = gpu.subgroup_mma_elementwise divf %D, %A : (!gpu.mma_matrix<16x16xf16, "COp">, !gpu.mma_matrix<16x16xf16, "COp">) -> !gpu.mma_matrix<16x16xf16, "COp">
       // CHECK: spirv.Return
       gpu.return
     }
   }
 }
 No newline at end of file

mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir

 // CHECK:          spirv.AtomicIAdd "Device" "AcquireRelease" %[[OUTPTR]], %[[ADD]]
 // CHECK:          spirv.Branch ^bb2
 // CHECK:        ^bb2:
 // CHECK:          spirv.mlir.merge
 // CHECK:        }
 // CHECK:        spirv.Return
 
 func.func @single_workgroup_reduction(%input: memref<16xi32, #spirv.storage_class<StorageBuffer>>, %output: memref<1xi32, #spirv.storage_class<StorageBuffer>>) attributes {
-  spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 1, 1]>: vector<3xi32>>
+  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>
 } {
   linalg.generic #single_workgroup_reduction_trait
       ins(%input : memref<16xi32, #spirv.storage_class<StorageBuffer>>)
      outs(%output : memref<1xi32, #spirv.storage_class<StorageBuffer>>) {
     ^bb(%in: i32, %out: i32):
       %sum = arith.addi %in, %out : i32
       linalg.yield %sum : i32
   }
#single_workgroup_reduction_trait = {
   ]
 }
 
 module attributes {
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spirv.resource_limits<>>
 } {
 func.func @single_workgroup_reduction(%input: memref<16xi32, #spirv.storage_class<StorageBuffer>>, %output: memref<1xi32, #spirv.storage_class<StorageBuffer>>) attributes {
-  spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 1, 1]>: vector<3xi32>>
+  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>
 } {
   // expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
   linalg.generic #single_workgroup_reduction_trait
       ins(%input : memref<16xi32, #spirv.storage_class<StorageBuffer>>)
      outs(%output : memref<1xi32, #spirv.storage_class<StorageBuffer>>) {
     ^bb(%in: i32, %out: i32):
       %sum = arith.addi %in, %out : i32
       linalg.yield %sum : i32
#single_workgroup_reduction_trait = {
   ]
 }
 
 module attributes {
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spirv.resource_limits<>>
 } {
 func.func @single_workgroup_reduction(%input: memref<16x8xi32, #spirv.storage_class<StorageBuffer>>, %output: memref<16xi32, #spirv.storage_class<StorageBuffer>>) attributes {
-  spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[16, 8, 1]>: vector<3xi32>>
+  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 8, 1]>
 } {
   // expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
   linalg.generic #single_workgroup_reduction_trait
       ins(%input : memref<16x8xi32, #spirv.storage_class<StorageBuffer>>)
      outs(%output : memref<16xi32, #spirv.storage_class<StorageBuffer>>) {
     ^bb(%in: i32, %out: i32):
       %sum = arith.addi %in, %out : i32
       linalg.yield %sum : i32
   }
   spirv.Return
 }
 }

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

spirv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spirv.mlir.addressof @bar_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
       spirv.Return
     }
     spirv.EntryPoint "GLCompute" @bar
     spirv.ExecutionMode @bar "LocalSize", 1, 1, 1
   }
 
   gpu.module @foo {
-    gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<1> : vector<3xi32>>} {
+    gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
       gpu.return
     }
   }
 
   func.func @main() {
     %buffer = memref.alloc() : memref<6xi32>
     %one = arith.constant 1 : index
     gpu.launch_func @foo::@bar blocks in (%one, %one, %one)
         threads in (%one, %one, %one) args(%buffer : memref<6xi32>)
     return
   }
 }

mlir/test/Dialect/SPIRV/IR/target-and-abi.mlir

 // expected-error @+1 {{'spirv.entry_point_abi' attribute must be an entry point ABI attribute}}
 func.func @spv_entry_point() attributes {
   spirv.entry_point_abi = 64
 } { return }
 
 // -----
 
 func.func @spv_entry_point() attributes {
-  // expected-error @+2 {{failed to parse SPIRV_EntryPointABIAttr parameter 'local_size' which is to be a `DenseIntElementsAttr`}}
-  // expected-error @+1 {{invalid kind of attribute specified}}
-  spirv.entry_point_abi = #spirv.entry_point_abi<local_size = 64>
+  // expected-error @+2 {{failed to parse SPIRV_EntryPointABIAttr parameter 'workgroup_size' which is to be a `DenseI32ArrayAttr`}}
+  // expected-error @+1 {{expected '['}}
+  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = 64>
 } { return }
 
 // -----
 
 func.func @spv_entry_point() attributes {
-  // CHECK: {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[64, 1, 1]> : vector<3xi32>>}
-  spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[64, 1, 1]>: vector<3xi32>>
+  // CHECK: {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [64, 1, 1]>}
+  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [64, 1, 1]>
 } { return }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spirv.interface_var_abi
 //===----------------------------------------------------------------------===//
 

 // expected-error @+1 {{'spirv.interface_var_abi' attribute cannot specify storage class when attaching to a non-scalar value}}
 func.func @interface_var(
   %arg0 : memref<4xf32> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1), Uniform>}
 ) { return }
 
 // -----
 
 //===----------------------------------------------------------------------===//
+// spirv.resource_limits
+//===----------------------------------------------------------------------===//
+
+// CHECK-LABEL: func @resource_limits_all_default()
+func.func @resource_limits_all_default() attributes {
+  // CHECK-SAME: #spirv.resource_limits<>
+  limits = #spirv.resource_limits<>
+} { return }
+
+// -----
+
+// CHECK-LABEL: func @resource_limits_min_max_subgroup_size()
+func.func @resource_limits_min_max_subgroup_size() attributes {
+  // CHECK-SAME: #spirv.resource_limits<min_subgroup_size = 32, max_subgroup_size = 64>
+  limits = #spirv.resource_limits<min_subgroup_size = 32, max_subgroup_size=64>
+} { return }
+
+// -----
+
+//===----------------------------------------------------------------------===//
 // spirv.target_env
 //===----------------------------------------------------------------------===//
 
 func.func @target_env() attributes {
   // CHECK:      spirv.target_env = #spirv.target_env<
   // CHECK-SAME:   #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
   // CHECK-SAME:   #spirv.resource_limits<max_compute_workgroup_size = [128, 64, 64]>>
   spirv.target_env = #spirv.target_env<

mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir

spirv.module Logical GLSL450 {
 
   spirv.func @baz_no_return_another_control(%arg0: i32) -> () "Inline|Pure" {
     spirv.Return
   }
 
   spirv.func @kernel(
     %arg0: f32,
     %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None"
-  attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
+  attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {
     spirv.Return
   }
 
   spirv.func @kernel_different_attr(
     %arg0: f32,
     %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None"
-  attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[64, 1, 1]> : vector<3xi32>>} {
+  attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [64, 1, 1]>} {
     spirv.Return
   }
 }
 }

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

-// RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
+// RUN: mlir-opt -split-input-file -spirv-lower-abi-attrs %s | FileCheck %s
 
 module attributes {
   spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>>
 } {
   spirv.module Physical64 OpenCL {
     // CHECK-LABEL: spirv.module
     //       CHECK:   spirv.func [[FN:@.*]]({{%.*}}: f32, {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>
+    // We cannot generate SubgroupSize execution mode without necessary capability -- leave it alone.
+    //  CHECK-SAME:      #spirv.entry_point_abi<subgroup_size = 64>
     //       CHECK:   spirv.EntryPoint "Kernel" [[FN]]
     //       CHECK:   spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
     spirv.func @kernel(
       %arg0: f32,
       %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None"
-    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
+    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} {
+      spirv.Return
+    }
+  }
+}
+
+// -----
+
+module attributes {
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, SubgroupDispatch], []>, #spirv.resource_limits<>>
+} {
+  spirv.module Physical64 OpenCL {
+    // CHECK-LABEL: spirv.module
+    //       CHECK:   spirv.func [[FN:@.*]]({{%.*}}: f32, {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>
+    //       CHECK:   spirv.EntryPoint "Kernel" [[FN]]
+    //       CHECK:   spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
+    //       CHECK:   spirv.ExecutionMode [[FN]] "SubgroupSize", 64
+    spirv.func @kernel(
+      %arg0: f32,
+      %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None"
+    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} {
       spirv.Return
     }
   }
 }

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

-// RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
+// RUN: mlir-opt -split-input-file -spirv-lower-abi-attrs %s | FileCheck %s
 
 module attributes {
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spirv.module
 spirv.module Logical GLSL450 {
-  // CHECK-DAG:    spirv.GlobalVariable [[VAR0:@.*]] bind(0, 0) : !spirv.ptr<!spirv.struct<(f32 [0])>, StorageBuffer>
-  // CHECK-DAG:    spirv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
-  // CHECK:    spirv.func [[FN:@.*]]()
+  //  CHECK-DAG:    spirv.GlobalVariable [[VAR0:@.*]] bind(0, 0) : !spirv.ptr<!spirv.struct<(f32 [0])>, StorageBuffer>
+  //  CHECK-DAG:    spirv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+  //      CHECK:    spirv.func [[FN:@.*]]()
+  // We cannot generate SubgroupSize execution mode for Shader capability -- leave it alone.
+  // CHECK-SAME:      #spirv.entry_point_abi<subgroup_size = 64>
   spirv.func @kernel(
     %arg0: f32
            {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0), StorageBuffer>},
     %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, StorageBuffer>
            {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}) "None"
-  attributes {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
+  attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1], subgroup_size = 64>} {
     // CHECK: [[ARG1:%.*]] = spirv.mlir.addressof [[VAR1]]
     // CHECK: [[ADDRESSARG0:%.*]] = spirv.mlir.addressof [[VAR0]]
     // CHECK: [[CONST0:%.*]] = spirv.Constant 0 : i32
     // CHECK: [[ARG0PTR:%.*]] = spirv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]]
     // CHECK: [[ARG0:%.*]] = spirv.Load "StorageBuffer" [[ARG0PTR]]
     // CHECK: spirv.Return
     spirv.Return
   }
   // CHECK: spirv.EntryPoint "GLCompute" [[FN]]
   // CHECK: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
 } // end spirv.module
 
 } // end module

mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir

spirv.func @load_store_kernel(
     %arg3: i32
     {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 3), StorageBuffer>},
     %arg4: i32
     {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 4), StorageBuffer>},
     %arg5: i32
     {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 5), StorageBuffer>},
     %arg6: i32
     {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 6), StorageBuffer>}) "None"
-  attributes  {spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
+  attributes  {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {
     // CHECK: [[ADDRESSARG6:%.*]] = spirv.mlir.addressof [[VAR6]]
     // CHECK: [[CONST6:%.*]] = spirv.Constant 0 : i32
     // CHECK: [[ARG6PTR:%.*]] = spirv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]
     // CHECK: {{%.*}} = spirv.Load "StorageBuffer" [[ARG6PTR]]
     // CHECK: [[ADDRESSARG5:%.*]] = spirv.mlir.addressof [[VAR5]]
     // CHECK: [[CONST5:%.*]] = spirv.Constant 0 : i32
     // CHECK: [[ARG5PTR:%.*]] = spirv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]]
     // CHECK: {{%.*}} = spirv.Load "StorageBuffer" [[ARG5PTR]]

mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp

void TestSpirvEntryPointABIPass::runOnOperation() {
   StringRef attrName = spirv::getEntryPointABIAttrName();
   for (gpu::GPUFuncOp gpuFunc : gpuModule.getOps<gpu::GPUFuncOp>()) {
     if (!gpu::GPUDialect::isKernel(gpuFunc) || gpuFunc->getAttr(attrName))
       continue;
     SmallVector<int32_t, 3> workgroupSizeVec(workgroupSize.begin(),
                                              workgroupSize.end());
     workgroupSizeVec.resize(3, 1);
     gpuFunc->setAttr(attrName,
-                     spirv::getEntryPointABIAttr(workgroupSizeVec, context));
+                     spirv::getEntryPointABIAttr(context, workgroupSizeVec));
   }
 }
 
 namespace mlir {
 void registerTestSpirvEntryPointABIPass() {
   PassRegistration<TestSpirvEntryPointABIPass>();
 }
 } // namespace mlir

mlir/test/mlir-spirv-cpu-runner/double.mlir

 // RUN: mlir-spirv-cpu-runner %s -e main --entry-point-result=void --shared-libs=%mlir_lib_dir/libmlir_runner_utils%shlibext,%mlir_lib_dir/libmlir_test_spirv_cpu_runner_c_wrappers%shlibext
 
 // CHECK: [8,  8,  8,  8,  8,  8]
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>,
     #spirv.resource_limits<
      max_compute_workgroup_invocations = 128,
      max_compute_workgroup_size = [128, 128, 64]>>
 } {
   gpu.module @kernels {
     gpu.func @double(%arg0 : memref<6xi32>, %arg1 : memref<6xi32>)
-      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
+      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
       %factor = arith.constant 2 : i32
 
       %i0 = arith.constant 0 : index
       %i1 = arith.constant 1 : index
       %i2 = arith.constant 2 : index
       %i3 = arith.constant 3 : index
       %i4 = arith.constant 4 : index
       %i5 = arith.constant 5 : index

mlir/test/mlir-spirv-cpu-runner/simple_add.mlir

 // RUN: mlir-spirv-cpu-runner %s -e main --entry-point-result=void --shared-libs=%mlir_lib_dir/libmlir_runner_utils%shlibext,%mlir_lib_dir/libmlir_test_spirv_cpu_runner_c_wrappers%shlibext
 
 // CHECK: [[[7.7,    0,    0], [7.7,    0,    0], [7.7,    0,    0]], [[0,    7.7,    0], [0,    7.7,    0], [0,    7.7,    0]], [[0,    0,    7.7], [0,    0,    7.7], [0,    0,    7.7]]]
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_8bit_storage]>,
     #spirv.resource_limits<
      max_compute_workgroup_invocations = 128,
      max_compute_workgroup_size = [128, 128, 64]>>
 } {
   gpu.module @kernels {
     gpu.func @sum(%arg0 : memref<3xf32>, %arg1 : memref<3x3xf32>, %arg2 :  memref<3x3x3xf32>)
-      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
+      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
       %i0 = arith.constant 0 : index
       %i1 = arith.constant 1 : index
       %i2 = arith.constant 2 : index
 
       %x = memref.load %arg0[%i0] : memref<3xf32>
       %y = memref.load %arg1[%i0, %i0] : memref<3x3xf32>
       %sum = arith.addf %x, %y : f32
 

mlir/test/mlir-vulkan-runner/addf.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%mlir_lib_dir/libvulkan-runtime-wrappers%shlibext,%mlir_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK: [3.3,  3.3,  3.3,  3.3,  3.3,  3.3,  3.3,  3.3]
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
 } {
   gpu.module @kernels {
     gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
-      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
+      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
       %0 = gpu.block_id x
       %1 = memref.load %arg0[%0] : memref<8xf32>
       %2 = memref.load %arg1[%0] : memref<8xf32>
       %3 = arith.addf %1, %2 : f32
       memref.store %3, %arg2[%0] : memref<8xf32>
       gpu.return
     }
   }

mlir/test/mlir-vulkan-runner/addi.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%mlir_lib_dir/libvulkan-runtime-wrappers%shlibext,%mlir_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK-COUNT-64: [3, 3, 3, 3, 3, 3, 3, 3]
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
 } {
   gpu.module @kernels {
     gpu.func @kernel_addi(%arg0 : memref<8xi32>, %arg1 : memref<8x8xi32>, %arg2 : memref<8x8x8xi32>)
-      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
+      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
       %x = gpu.block_id x
       %y = gpu.block_id y
       %z = gpu.block_id z
       %0 = memref.load %arg0[%x] : memref<8xi32>
       %1 = memref.load %arg1[%y, %x] : memref<8x8xi32>
       %2 = arith.addi %0, %1 : i32
       memref.store %2, %arg2[%z, %y, %x] : memref<8x8x8xi32>
       gpu.return

mlir/test/mlir-vulkan-runner/addi8.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%mlir_lib_dir/libvulkan-runtime-wrappers%shlibext,%mlir_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK-COUNT-64: [3, 3, 3, 3, 3, 3, 3, 3]
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_8bit_storage]>, #spirv.resource_limits<>>
 } {
   gpu.module @kernels {
     gpu.func @kernel_addi(%arg0 : memref<8xi8>, %arg1 : memref<8x8xi8>, %arg2 : memref<8x8x8xi32>)
-      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
+      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
       %x = gpu.block_id x
       %y = gpu.block_id y
       %z = gpu.block_id z
       %0 = memref.load %arg0[%x] : memref<8xi8>
       %1 = memref.load %arg1[%y, %x] : memref<8x8xi8>
       %2 = arith.addi %0, %1 : i8
       %3 = arith.extui %2 : i8 to i32
       memref.store %3, %arg2[%z, %y, %x] : memref<8x8x8xi32>

mlir/test/mlir-vulkan-runner/mulf.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%mlir_lib_dir/libvulkan-runtime-wrappers%shlibext,%mlir_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK-COUNT-4: [6, 6, 6, 6]
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
 } {
   gpu.module @kernels {
     gpu.func @kernel_mul(%arg0 : memref<4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<4x4xf32>)
-      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
+      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
       %x = gpu.block_id x
       %y = gpu.block_id y
       %1 = memref.load %arg0[%x, %y] : memref<4x4xf32>
       %2 = memref.load %arg1[%x, %y] : memref<4x4xf32>
       %3 = arith.mulf %1, %2 : f32
       memref.store %3, %arg2[%x, %y] : memref<4x4xf32>
       gpu.return
     }

mlir/test/mlir-vulkan-runner/subf.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%mlir_lib_dir/libvulkan-runtime-wrappers%shlibext,%mlir_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK-COUNT-32: [2.2, 2.2, 2.2, 2.2]
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
     #spirv.resource_limits<>>
 } {
   gpu.module @kernels {
     gpu.func @kernel_sub(%arg0 : memref<8x4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<8x4x4xf32>)
-      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
+      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
       %x = gpu.block_id x
       %y = gpu.block_id y
       %z = gpu.block_id z
       %1 = memref.load %arg0[%x, %y, %z] : memref<8x4x4xf32>
       %2 = memref.load %arg1[%y, %z] : memref<4x4xf32>
       %3 = arith.subf %1, %2 : f32
       memref.store %3, %arg2[%x, %y, %z] : memref<8x4x4xf32>
       gpu.return

mlir/test/mlir-vulkan-runner/time.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%mlir_lib_dir/libvulkan-runtime-wrappers%shlibext,%mlir_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK: Compute shader execution time
 // CHECK: Command buffer submit time
 // CHECK: Wait idle time
 
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<
     #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
 } {
   gpu.module @kernels {
     gpu.func @kernel_add(%arg0 : memref<16384xf32>, %arg1 : memref<16384xf32>, %arg2 : memref<16384xf32>)
-      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<local_size = dense<[128, 1, 1]>: vector<3xi32>>} {
+      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [128, 1, 1]>} {
       %bid = gpu.block_id x
       %tid = gpu.thread_id x
       %cst = arith.constant 128 : index
       %b = arith.muli %bid, %cst : index
       %0 = arith.addi %b, %tid : index
       %1 = memref.load %arg0[%0] : memref<16384xf32>
       %2 = memref.load %arg1[%0] : memref<16384xf32>
       %3 = arith.addf %1, %2 : f32