[Mlir-commits] [mlir] [mlir][SPIR-V] Restrict GroupNonUniform ops to Subgroup scope (PR #213087)
Arseniy Obolenskiy
llvmlistbot at llvm.org
Thu Jul 30 10:46:48 PDT 2026
https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/213087
spirv-val now limits execution scope for GroupNonUniform* ops to Subgroup, except OpGroupNonUniformRotateKHR which still allows Workgroup (see https://github.com/KhronosGroup/SPIRV-Tools/pull/6811). Tighten the ODS trait accordingly and stop lowering GPU non-uniform reductions to a Workgroup scope op
Follow-up to #212928
>From df2d30cb711f9566a4611907a435aba75be72506 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 30 Jul 2026 19:45:24 +0200
Subject: [PATCH] [mlir][SPIR-V] Restrict GroupNonUniform ops to Subgroup scope
spirv-val now limits execution scope for GroupNonUniform* ops to Subgroup, except OpGroupNonUniformRotateKHR which still allows Workgroup (see https://github.com/KhronosGroup/SPIRV-Tools/pull/6811). Tighten the ODS trait accordingly and stop lowering GPU non-uniform reductions to a Workgroup scope op
Follow-up to #212928
---
.../Dialect/SPIRV/IR/SPIRVNonUniformOps.td | 102 ++++++-------
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 14 +-
mlir/test/Conversion/ConvertToSPIRV/gpu.mlir | 8 +-
.../Conversion/GPUToSPIRV/reductions.mlir | 119 ++++++++++-----
mlir/test/Dialect/SPIRV/IR/availability.mlir | 2 +-
.../Dialect/SPIRV/IR/non-uniform-ops.mlir | 140 +++++++++---------
mlir/test/Dialect/SPIRV/IR/target-env.mlir | 2 +-
.../SPIRV/Transforms/vce-deduction.mlir | 2 +-
.../lib/Dialect/SPIRV/TestAvailability.cpp | 2 +-
9 files changed, 218 insertions(+), 173 deletions(-)
diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td
index 0a0e91a4c49fb..e473a8e0efb3c 100644
--- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td
+++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td
@@ -18,7 +18,7 @@
class SPIRV_GroupNonUniformArithmeticOp<string mnemonic, Type type,
list<Trait> traits = []> : SPIRV_Op<mnemonic, !listconcat([
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>
], traits)> {
let arguments = (ins
@@ -42,7 +42,7 @@ class SPIRV_GroupNonUniformArithmeticOp<string mnemonic, Type type,
// -----
def SPIRV_GroupNonUniformBallotOp : SPIRV_Op<"GroupNonUniformBallot",[
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Result is a bitfield value combining the Predicate value from all
@@ -61,7 +61,7 @@ def SPIRV_GroupNonUniformBallotOp : SPIRV_Op<"GroupNonUniformBallot",[
size of the group) is the higher bit number of the last bitmask needed
to represent all bits of the group invocations.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
Predicate must be a Boolean type.
@@ -98,7 +98,7 @@ def SPIRV_GroupNonUniformBallotOp : SPIRV_Op<"GroupNonUniformBallot",[
// -----
def SPIRV_GroupNonUniformBallotFindLSBOp : SPIRV_Op<"GroupNonUniformBallotFindLSB", [
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Find the least significant bit set to 1 in Value, considering only the
@@ -158,7 +158,7 @@ def SPIRV_GroupNonUniformBallotFindLSBOp : SPIRV_Op<"GroupNonUniformBallotFindLS
// -----
def SPIRV_GroupNonUniformBallotFindMSBOp : SPIRV_Op<"GroupNonUniformBallotFindMSB", [
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Find the most significant bit set to 1 in Value, considering only the
@@ -219,7 +219,7 @@ def SPIRV_GroupNonUniformBallotFindMSBOp : SPIRV_Op<"GroupNonUniformBallotFindMS
def SPIRV_GroupNonUniformBroadcastOp : SPIRV_Op<"GroupNonUniformBroadcast", [
Pure, AllTypesMatch<["value", "result"]>,
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Result is the Value of the invocation identified by the id Id to all
@@ -230,7 +230,7 @@ def SPIRV_GroupNonUniformBroadcastOp : SPIRV_Op<"GroupNonUniformBroadcast", [
Result Type must be a scalar or vector of floating-point type, integer
type, or Boolean type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The type of Value must be the same as Result Type.
@@ -249,7 +249,7 @@ def SPIRV_GroupNonUniformBroadcastOp : SPIRV_Op<"GroupNonUniformBroadcast", [
%vector_value = ... : vector<4xf32>
%id = ... : i32
%0 = spirv.GroupNonUniformBroadcast <Subgroup> %scalar_value, %id : f32, i32
- %1 = spirv.GroupNonUniformBroadcast <Workgroup> %vector_value, %id :
+ %1 = spirv.GroupNonUniformBroadcast <Subgroup> %vector_value, %id :
vector<4xf32>, i32
```
}];
@@ -335,7 +335,7 @@ def SPIRV_GroupNonUniformBroadcastFirstOp : SPIRV_Op<"GroupNonUniformBroadcastFi
// -----
def SPIRV_GroupNonUniformElectOp : SPIRV_Op<"GroupNonUniformElect", [
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Result is true only in the active invocation with the lowest id in the
@@ -345,12 +345,12 @@ def SPIRV_GroupNonUniformElectOp : SPIRV_Op<"GroupNonUniformElect", [
let description = [{
Result Type must be a Boolean type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
#### Example:
```mlir
- %0 = spirv.GroupNonUniformElect <Workgroup> : i1
+ %0 = spirv.GroupNonUniformElect <Subgroup> : i1
```
}];
@@ -385,7 +385,7 @@ def SPIRV_GroupNonUniformFAddOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
let description = [{
Result Type must be a scalar or vector of floating-point type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is 0. If Operation is ClusteredReduce,
ClusterSize must be specified.
@@ -408,7 +408,7 @@ def SPIRV_GroupNonUniformFAddOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
- %0 = spirv.GroupNonUniformFAdd <Workgroup> <Reduce> %scalar : f32 -> f32
+ %0 = spirv.GroupNonUniformFAdd <Subgroup> <Reduce> %scalar : f32 -> f32
%1 = spirv.GroupNonUniformFAdd <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xf32>, i32 -> vector<4xf32>
```
}];
@@ -432,7 +432,7 @@ def SPIRV_GroupNonUniformFMaxOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
let description = [{
Result Type must be a scalar or vector of floating-point type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is -INF. If Operation is ClusteredReduce,
ClusterSize must be specified.
@@ -458,7 +458,7 @@ def SPIRV_GroupNonUniformFMaxOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
- %0 = spirv.GroupNonUniformFMax <Workgroup> <Reduce> %scalar : f32 -> f32
+ %0 = spirv.GroupNonUniformFMax <Subgroup> <Reduce> %scalar : f32 -> f32
%1 = spirv.GroupNonUniformFMax <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xf32>, i32 -> vector<4xf32>
```
}];
@@ -482,7 +482,7 @@ def SPIRV_GroupNonUniformFMinOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
let description = [{
Result Type must be a scalar or vector of floating-point type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is +INF. If Operation is ClusteredReduce,
ClusterSize must be specified.
@@ -508,7 +508,7 @@ def SPIRV_GroupNonUniformFMinOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
- %0 = spirv.GroupNonUniformFMin <Workgroup> <Reduce> %scalar : f32 -> i32
+ %0 = spirv.GroupNonUniformFMin <Subgroup> <Reduce> %scalar : f32 -> i32
%1 = spirv.GroupNonUniformFMin <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xf32>, i32 -> vector<4xf32>
```
}];
@@ -532,7 +532,7 @@ def SPIRV_GroupNonUniformFMulOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
let description = [{
Result Type must be a scalar or vector of floating-point type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is 1. If Operation is ClusteredReduce,
ClusterSize must be specified.
@@ -555,7 +555,7 @@ def SPIRV_GroupNonUniformFMulOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
- %0 = spirv.GroupNonUniformFMul <Workgroup> <Reduce> %scalar : f32 -> f32
+ %0 = spirv.GroupNonUniformFMul <Subgroup> <Reduce> %scalar : f32 -> f32
%1 = spirv.GroupNonUniformFMul <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xf32>, i32 -> vector<4xf32>
```
}];
@@ -579,7 +579,7 @@ def SPIRV_GroupNonUniformIAddOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
let description = [{
Result Type must be a scalar or vector of integer type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is 0. If Operation is ClusteredReduce,
ClusterSize must be specified.
@@ -600,7 +600,7 @@ def SPIRV_GroupNonUniformIAddOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformIAdd <Workgroup> <Reduce> %scalar : i32 -> i32
+ %0 = spirv.GroupNonUniformIAdd <Subgroup> <Reduce> %scalar : i32 -> i32
%1 = spirv.GroupNonUniformIAdd <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
}];
@@ -624,7 +624,7 @@ def SPIRV_GroupNonUniformIMulOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
let description = [{
Result Type must be a scalar or vector of integer type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is 1. If Operation is ClusteredReduce,
ClusterSize must be specified.
@@ -645,7 +645,7 @@ def SPIRV_GroupNonUniformIMulOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformIMul <Workgroup> <Reduce> %scalar : i32 -> i32
+ %0 = spirv.GroupNonUniformIMul <Subgroup> <Reduce> %scalar : i32 -> i32
%1 = spirv.GroupNonUniformIMul <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
}];
@@ -671,7 +671,7 @@ def SPIRV_GroupNonUniformSMaxOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
let description = [{
Result Type must be a scalar or vector of integer type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is INT_MIN. If Operation is
ClusteredReduce, ClusterSize must be specified.
@@ -692,7 +692,7 @@ def SPIRV_GroupNonUniformSMaxOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformSMax <Workgroup> <Reduce> %scalar : i32
+ %0 = spirv.GroupNonUniformSMax <Subgroup> <Reduce> %scalar : i32
%1 = spirv.GroupNonUniformSMax <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
}];
@@ -718,7 +718,7 @@ def SPIRV_GroupNonUniformSMinOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
let description = [{
Result Type must be a scalar or vector of integer type.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is INT_MAX. If Operation is
ClusteredReduce, ClusterSize must be specified.
@@ -739,7 +739,7 @@ def SPIRV_GroupNonUniformSMinOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformSMin <Workgroup> <Reduce> %scalar : i32 -> i32
+ %0 = spirv.GroupNonUniformSMin <Subgroup> <Reduce> %scalar : i32 -> i32
%1 = spirv.GroupNonUniformSMin <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
}];
@@ -756,7 +756,7 @@ def SPIRV_GroupNonUniformSMinOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
def SPIRV_GroupNonUniformShuffleOp : SPIRV_Op<"GroupNonUniformShuffle", [
Pure, AllTypesMatch<["value", "result"]>,
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Result is the Value of the invocation identified by the id Id.
@@ -766,7 +766,7 @@ def SPIRV_GroupNonUniformShuffleOp : SPIRV_Op<"GroupNonUniformShuffle", [
Result Type must be a scalar or vector of floating-point type, integer
type, or Boolean type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The type of Value must be the same as Result Type.
@@ -810,7 +810,7 @@ def SPIRV_GroupNonUniformShuffleOp : SPIRV_Op<"GroupNonUniformShuffle", [
def SPIRV_GroupNonUniformShuffleDownOp : SPIRV_Op<"GroupNonUniformShuffleDown", [
Pure, AllTypesMatch<["value", "result"]>,
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Result is the Value of the invocation identified by the current
@@ -821,7 +821,7 @@ def SPIRV_GroupNonUniformShuffleDownOp : SPIRV_Op<"GroupNonUniformShuffleDown",
Result Type must be a scalar or vector of floating-point type, integer
type, or Boolean type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The type of Value must be the same as Result Type.
@@ -867,7 +867,7 @@ def SPIRV_GroupNonUniformShuffleDownOp : SPIRV_Op<"GroupNonUniformShuffleDown",
def SPIRV_GroupNonUniformShuffleUpOp : SPIRV_Op<"GroupNonUniformShuffleUp", [
Pure, AllTypesMatch<["value", "result"]>,
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Result is the Value of the invocation identified by the current
@@ -878,7 +878,7 @@ def SPIRV_GroupNonUniformShuffleUpOp : SPIRV_Op<"GroupNonUniformShuffleUp", [
Result Type must be a scalar or vector of floating-point type, integer
type, or Boolean type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The type of Value must be the same as Result Type.
@@ -923,7 +923,7 @@ def SPIRV_GroupNonUniformShuffleUpOp : SPIRV_Op<"GroupNonUniformShuffleUp", [
def SPIRV_GroupNonUniformShuffleXorOp : SPIRV_Op<"GroupNonUniformShuffleXor", [
Pure, AllTypesMatch<["value", "result"]>,
- SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Workgroup", "Subgroup"]>]> {
+ SPIRV_ExecutionScopeAttrIs<"execution_scope", ["Subgroup"]>]> {
let summary = [{
Result is the Value of the invocation identified by the current
@@ -934,7 +934,7 @@ def SPIRV_GroupNonUniformShuffleXorOp : SPIRV_Op<"GroupNonUniformShuffleXor", [
Result Type must be a scalar or vector of floating-point type, integer
type, or Boolean type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The type of Value must be the same as Result Type.
@@ -989,7 +989,7 @@ def SPIRV_GroupNonUniformUMaxOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
Result Type must be a scalar or vector of integer type, whose
Signedness operand is 0.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is 0. If Operation is ClusteredReduce,
ClusterSize must be specified.
@@ -1010,7 +1010,7 @@ def SPIRV_GroupNonUniformUMaxOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformUMax <Workgroup> <Reduce> %scalar : i32 -> i32
+ %0 = spirv.GroupNonUniformUMax <Subgroup> <Reduce> %scalar : i32 -> i32
%1 = spirv.GroupNonUniformUMax <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
}];
@@ -1037,7 +1037,7 @@ def SPIRV_GroupNonUniformUMinOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
Result Type must be a scalar or vector of integer type, whose
Signedness operand is 0.
- Execution must be Workgroup or Subgroup Scope.
+ Execution must be Subgroup Scope.
The identity I for Operation is UINT_MAX. If Operation is
ClusteredReduce, ClusterSize must be specified.
@@ -1058,7 +1058,7 @@ def SPIRV_GroupNonUniformUMinOp : SPIRV_GroupNonUniformArithmeticOp<"GroupNonUni
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformUMin <Workgroup> <Reduce> %scalar : i32 -> i32
+ %0 = spirv.GroupNonUniformUMin <Subgroup> <Reduce> %scalar : i32 -> i32
%1 = spirv.GroupNonUniformUMin <Subgroup> <ClusteredReduce> %vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
}];
@@ -1084,7 +1084,7 @@ def SPIRV_GroupNonUniformBitwiseAndOp :
let description = [{
Result Type must be a scalar or vector of integer type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The identity I for Operation is ~0. If Operation is ClusteredReduce,
ClusterSize must be present.
@@ -1105,7 +1105,7 @@ def SPIRV_GroupNonUniformBitwiseAndOp :
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformBitwiseAnd <Workgroup> <Reduce> %scalar : i32 -> i32
+ %0 = spirv.GroupNonUniformBitwiseAnd <Subgroup> <Reduce> %scalar : i32 -> i32
%1 = spirv.GroupNonUniformBitwiseAnd <Subgroup> <ClusteredReduce>
%vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
@@ -1134,7 +1134,7 @@ def SPIRV_GroupNonUniformBitwiseOrOp :
let description = [{
Result Type must be a scalar or vector of integer type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The identity I for Operation is 0. If Operation is ClusteredReduce,
ClusterSize must be present.
@@ -1155,7 +1155,7 @@ def SPIRV_GroupNonUniformBitwiseOrOp :
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformBitwiseOr <Workgroup> <Reduce> %scalar : i32 -> i32
+ %0 = spirv.GroupNonUniformBitwiseOr <Subgroup> <Reduce> %scalar : i32 -> i32
%1 = spirv.GroupNonUniformBitwiseOr <Subgroup> <ClusteredReduce>
%vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
@@ -1184,7 +1184,7 @@ def SPIRV_GroupNonUniformBitwiseXorOp :
let description = [{
Result Type must be a scalar or vector of integer type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The identity I for Operation is 0. If Operation is ClusteredReduce,
ClusterSize must be present.
@@ -1205,7 +1205,7 @@ def SPIRV_GroupNonUniformBitwiseXorOp :
%four = spirv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
- %0 = spirv.GroupNonUniformBitwiseXor <Workgroup> <Reduce> %scalar : i32 -> i32
+ %0 = spirv.GroupNonUniformBitwiseXor <Subgroup> <Reduce> %scalar : i32 -> i32
%1 = spirv.GroupNonUniformBitwiseXor <Subgroup> <ClusteredReduce>
%vector cluster_size(%four) : vector<4xi32>, i32 -> vector<4xi32>
```
@@ -1234,7 +1234,7 @@ def SPIRV_GroupNonUniformLogicalAndOp :
let description = [{
Result Type must be a scalar or vector of Boolean type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The identity I for Operation is ~0. If Operation is ClusteredReduce,
ClusterSize must be present.
@@ -1255,7 +1255,7 @@ def SPIRV_GroupNonUniformLogicalAndOp :
%four = spirv.Constant 4 : i32
%scalar = ... : i1
%vector = ... : vector<4xi1>
- %0 = spirv.GroupNonUniformLogicalAnd <Workgroup> <Reduce> %scalar : i1 -> i1
+ %0 = spirv.GroupNonUniformLogicalAnd <Subgroup> <Reduce> %scalar : i1 -> i1
%1 = spirv.GroupNonUniformLogicalAnd <Subgroup> <ClusteredReduce>
%vector cluster_size(%four) : vector<4xi1>, i32 -> vector<4xi1>
```
@@ -1284,7 +1284,7 @@ def SPIRV_GroupNonUniformLogicalOrOp :
let description = [{
Result Type must be a scalar or vector of Boolean type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The identity I for Operation is 0. If Operation is ClusteredReduce,
ClusterSize must be present.
@@ -1305,7 +1305,7 @@ def SPIRV_GroupNonUniformLogicalOrOp :
%four = spirv.Constant 4 : i32
%scalar = ... : i1
%vector = ... : vector<4xi1>
- %0 = spirv.GroupNonUniformLogicalOr <Workgroup> <Reduce> %scalar : i1 -> i1
+ %0 = spirv.GroupNonUniformLogicalOr <Subgroup> <Reduce> %scalar : i1 -> i1
%1 = spirv.GroupNonUniformLogicalOr <Subgroup> <ClusteredReduce>
%vector cluster_size(%four) : vector<4xi1>, i32 -> vector<4xi1>
```
@@ -1334,7 +1334,7 @@ def SPIRV_GroupNonUniformLogicalXorOp :
let description = [{
Result Type must be a scalar or vector of Boolean type.
- Execution is a Scope. It must be either Workgroup or Subgroup.
+ Execution is a Scope. It must be Subgroup.
The identity I for Operation is 0. If Operation is ClusteredReduce,
ClusterSize must be present.
@@ -1355,7 +1355,7 @@ def SPIRV_GroupNonUniformLogicalXorOp :
%four = spirv.Constant 4 : i32
%scalar = ... : i1
%vector = ... : vector<4xi1>
- %0 = spirv.GroupNonUniformLogicalXor <Workgroup> <Reduce> %scalar : i1 -> i1
+ %0 = spirv.GroupNonUniformLogicalXor <Subgroup> <Reduce> %scalar : i1 -> i1
%1 = spirv.GroupNonUniformLogicalXor <Subgroup> <ClusteredReduce>
%vector cluster_size(%four) : vector<4xi1>, i32 -> vector<4xi1>
```
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index 48623e2fac47c..609dbe57a9a46 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -694,16 +694,20 @@ template <typename UniformOp, typename NonUniformOp>
static Value createGroupReduceOpImpl(OpBuilder &builder, Location loc,
Value arg, bool isGroup, bool isUniform,
std::optional<uint32_t> clusterSize) {
+ spirv::Scope scope =
+ isGroup ? spirv::Scope::Workgroup : spirv::Scope::Subgroup;
+ // GroupNonUniform* ops only support Subgroup scope.
+ if (!isUniform && scope != spirv::Scope::Subgroup)
+ return Value();
+
Type type = arg.getType();
- auto scope = mlir::spirv::ScopeAttr::get(builder.getContext(),
- isGroup ? spirv::Scope::Workgroup
- : spirv::Scope::Subgroup);
+ auto scopeAttr = mlir::spirv::ScopeAttr::get(builder.getContext(), scope);
auto groupOp = spirv::GroupOperationAttr::get(
builder.getContext(), clusterSize.has_value()
? spirv::GroupOperation::ClusteredReduce
: spirv::GroupOperation::Reduce);
if (isUniform) {
- return UniformOp::create(builder, loc, type, scope, groupOp, arg)
+ return UniformOp::create(builder, loc, type, scopeAttr, groupOp, arg)
.getResult();
}
@@ -713,7 +717,7 @@ static Value createGroupReduceOpImpl(OpBuilder &builder, Location loc,
builder, loc, builder.getI32Type(),
builder.getIntegerAttr(builder.getI32Type(), *clusterSize));
- return NonUniformOp::create(builder, loc, type, scope, groupOp, arg,
+ return NonUniformOp::create(builder, loc, type, scopeAttr, groupOp, arg,
clusterSizeValue)
.getResult();
}
diff --git a/mlir/test/Conversion/ConvertToSPIRV/gpu.mlir b/mlir/test/Conversion/ConvertToSPIRV/gpu.mlir
index 6163f3464d1cc..2b5c82e5bd877 100644
--- a/mlir/test/Conversion/ConvertToSPIRV/gpu.mlir
+++ b/mlir/test/Conversion/ConvertToSPIRV/gpu.mlir
@@ -1,16 +1,16 @@
-// RUN: mlir-opt -test-convert-to-spirv -split-input-file %s | FileCheck %s
+// RUN: mlir-opt -test-convert-to-spirv -split-input-file -verify-diagnostics %s | FileCheck %s
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @all_reduce
- // CHECK-SAME: (%[[ARG0:.*]]: f32)
- // CHECK: %{{.*}} = spirv.GroupNonUniformFAdd <Workgroup> <Reduce> %[[ARG0]] : f32 -> f32
gpu.func @all_reduce(%arg0 : f32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ // expected-error @+1 {{functions in 'spirv.module' can only contain spirv.* ops}}
%reduced = gpu.all_reduce add %arg0 {} : (f32) -> (f32)
gpu.return
}
diff --git a/mlir/test/Conversion/GPUToSPIRV/reductions.mlir b/mlir/test/Conversion/GPUToSPIRV/reductions.mlir
index 99c9bcab78187..2e09a75704b40 100644
--- a/mlir/test/Conversion/GPUToSPIRV/reductions.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/reductions.mlir
@@ -20,18 +20,18 @@ gpu.module @kernels {
// -----
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
-
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @test
- // CHECK-SAME: (%[[ARG:.*]]: f32)
gpu.func @test(%arg : f32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %{{.*}} = spirv.GroupNonUniformFAdd <Workgroup> <Reduce> %[[ARG]] : f32 -> f32
- %reduced = gpu.all_reduce add %arg {} : (f32) -> (f32)
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
+ %r0 = gpu.all_reduce add %arg {} : (f32) -> (f32)
gpu.return
}
}
@@ -60,18 +60,18 @@ gpu.module @kernels {
// -----
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
-
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @test
- // CHECK-SAME: (%[[ARG:.*]]: i32)
gpu.func @test(%arg : i32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %{{.*}} = spirv.GroupNonUniformIAdd <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
- %reduced = gpu.all_reduce add %arg {} : (i32) -> (i32)
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
+ %r0 = gpu.all_reduce add %arg {} : (i32) -> (i32)
gpu.return
}
}
@@ -180,18 +180,18 @@ gpu.module @kernels {
// -----
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
-
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @test
- // CHECK-SAME: (%[[ARG:.*]]: f32)
gpu.func @test(%arg : f32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %{{.*}} = spirv.GroupNonUniformFMul <Workgroup> <Reduce> %[[ARG]] : f32 -> f32
- %reduced = gpu.all_reduce mul %arg {} : (f32) -> (f32)
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
+ %r0 = gpu.all_reduce mul %arg {} : (f32) -> (f32)
gpu.return
}
}
@@ -220,18 +220,18 @@ gpu.module @kernels {
// -----
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
-
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @test
- // CHECK-SAME: (%[[ARG:.*]]: i32)
gpu.func @test(%arg : i32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %{{.*}} = spirv.GroupNonUniformIMul <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
- %reduced = gpu.all_reduce mul %arg {} : (i32) -> (i32)
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
+ %r0 = gpu.all_reduce mul %arg {} : (i32) -> (i32)
gpu.return
}
}
@@ -340,18 +340,18 @@ gpu.module @kernels {
// -----
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
-
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @test
- // CHECK-SAME: (%[[ARG:.*]]: f32)
gpu.func @test(%arg : f32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %{{.*}} = spirv.GroupNonUniformFMin <Workgroup> <Reduce> %[[ARG]] : f32 -> f32
- %reduced = gpu.all_reduce minnumf %arg {} : (f32) -> (f32)
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
+ %r0 = gpu.all_reduce minnumf %arg {} : (f32) -> (f32)
gpu.return
}
}
@@ -382,18 +382,39 @@ gpu.module @kernels {
// -----
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @test
- // CHECK-SAME: (%[[ARG:.*]]: i32)
- gpu.func @test(%arg : i32) kernel
+ gpu.func @minsi(%arg : i32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %{{.*}} = spirv.GroupNonUniformUMin <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
%r0 = gpu.all_reduce minsi %arg {} : (i32) -> (i32)
+ gpu.return
+ }
+}
+
+}
+
+// -----
+
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
+} {
+
+gpu.module @kernels {
+ gpu.func @minui(%arg : i32) kernel
+ attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
%r1 = gpu.all_reduce minui %arg {} : (i32) -> (i32)
gpu.return
}
@@ -507,18 +528,18 @@ gpu.module @kernels {
// -----
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
-
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @test
- // CHECK-SAME: (%[[ARG:.*]]: f32)
gpu.func @test(%arg : f32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %{{.*}} = spirv.GroupNonUniformFMax <Workgroup> <Reduce> %[[ARG]] : f32 -> f32
- %reduced = gpu.all_reduce maxnumf %arg {} : (f32) -> (f32)
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
+ %r0 = gpu.all_reduce maxnumf %arg {} : (f32) -> (f32)
gpu.return
}
}
@@ -549,19 +570,39 @@ gpu.module @kernels {
// -----
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
module attributes {
gpu.container_module,
spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
} {
gpu.module @kernels {
- // CHECK-LABEL: spirv.func @test
- // CHECK-SAME: (%[[ARG:.*]]: i32)
- gpu.func @test(%arg : i32) kernel
+ gpu.func @maxsi(%arg : i32) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %{{.*}} = spirv.GroupNonUniformSMax <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
- // CHECK: %{{.*}} = spirv.GroupNonUniformUMax <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
%r0 = gpu.all_reduce maxsi %arg {} : (i32) -> (i32)
+ gpu.return
+ }
+}
+
+}
+
+// -----
+
+// GroupNonUniform ops only support Subgroup scope, so a Workgroup-scope
+// non-uniform gpu.all_reduce fails to legalize.
+
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
+} {
+
+gpu.module @kernels {
+ gpu.func @maxui(%arg : i32) kernel
+ attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ // expected-error @+1 {{failed to legalize operation 'gpu.all_reduce'}}
%r1 = gpu.all_reduce maxui %arg {} : (i32) -> (i32)
gpu.return
}
diff --git a/mlir/test/Dialect/SPIRV/IR/availability.mlir b/mlir/test/Dialect/SPIRV/IR/availability.mlir
index 4ef242bdc5b16..c0bf80b709ebd 100644
--- a/mlir/test/Dialect/SPIRV/IR/availability.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/availability.mlir
@@ -26,7 +26,7 @@ func.func @subgroup_ballot(%predicate: i1) -> vector<4xi32> {
// CHECK: max version: v1.6
// CHECK: extensions: [ ]
// CHECK: capabilities: [ [GroupNonUniformBallot] ]
- %0 = spirv.GroupNonUniformBallot <Workgroup> %predicate : vector<4xi32>
+ %0 = spirv.GroupNonUniformBallot <Subgroup> %predicate : vector<4xi32>
return %0: vector<4xi32>
}
diff --git a/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir b/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir
index 7c3ceaf10b56d..9c927331a08b4 100644
--- a/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir
@@ -5,15 +5,15 @@
//===----------------------------------------------------------------------===//
func.func @group_non_uniform_ballot(%predicate: i1) -> vector<4xi32> {
- // CHECK: %{{.*}} = spirv.GroupNonUniformBallot <Workgroup> %{{.*}}: vector<4xi32>
- %0 = spirv.GroupNonUniformBallot <Workgroup> %predicate : vector<4xi32>
+ // CHECK: %{{.*}} = spirv.GroupNonUniformBallot <Subgroup> %{{.*}}: vector<4xi32>
+ %0 = spirv.GroupNonUniformBallot <Subgroup> %predicate : vector<4xi32>
return %0: vector<4xi32>
}
// -----
func.func @group_non_uniform_ballot(%predicate: i1) -> vector<4xi32> {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformBallot <Device> %predicate : vector<4xi32>
return %0: vector<4xi32>
}
@@ -22,7 +22,7 @@ func.func @group_non_uniform_ballot(%predicate: i1) -> vector<4xi32> {
func.func @group_non_uniform_ballot(%predicate: i1) -> vector<4xsi32> {
// expected-error @+1 {{op result #0 must be vector of 8/16/32/64-bit signless/unsigned integer values of length 4 of ranks 1, but got 'vector<4xsi32>'}}
- %0 = spirv.GroupNonUniformBallot <Workgroup> %predicate : vector<4xsi32>
+ %0 = spirv.GroupNonUniformBallot <Subgroup> %predicate : vector<4xsi32>
return %0: vector<4xsi32>
}
@@ -41,7 +41,7 @@ func.func @group_non_uniform_ballot_find_lsb(%value : vector<4xi32>) -> i32 {
// -----
func.func @group_non_uniform_ballot_find_lsb(%value : vector<4xi32>) -> i32 {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformBallotFindLSB <Device> %value : vector<4xi32>, i32
return %0: i32
}
@@ -69,7 +69,7 @@ func.func @group_non_uniform_ballot_find_msb(%value : vector<4xi32>) -> i32 {
// -----
func.func @group_non_uniform_ballot_find_msb(%value : vector<4xi32>) -> i32 {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformBallotFindMSB <Device> %value : vector<4xi32>, i32
return %0: i32
}
@@ -90,8 +90,8 @@ func.func @group_non_uniform_ballot_find_msb(%value : vector<4xi32>) -> si32 {
func.func @group_non_uniform_broadcast_scalar(%value: f32) -> f32 {
%one = spirv.Constant 1 : i32
- // CHECK: spirv.GroupNonUniformBroadcast <Workgroup> %{{.*}}, %{{.*}} : f32, i32
- %0 = spirv.GroupNonUniformBroadcast <Workgroup> %value, %one : f32, i32
+ // CHECK: spirv.GroupNonUniformBroadcast <Subgroup> %{{.*}}, %{{.*}} : f32, i32
+ %0 = spirv.GroupNonUniformBroadcast <Subgroup> %value, %one : f32, i32
return %0: f32
}
@@ -108,7 +108,7 @@ func.func @group_non_uniform_broadcast_vector(%value: vector<4xf32>) -> vector<4
func.func @group_non_uniform_broadcast_negative_scope(%value: f32, %localid: i32 ) -> f32 {
%one = spirv.Constant 1 : i32
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformBroadcast <Device> %value, %one : f32, i32
return %0: f32
}
@@ -126,7 +126,7 @@ func.func @group_non_uniform_broadcast_negative_non_const(%value: f32, %localid:
func.func @group_non_uniform_broadcast_bf16(%value: bf16) -> bf16 {
%one = spirv.Constant 1 : i32
// expected-error @+1 {{op operand #0 must be 8/16/32/64-bit integer or 16/32/64-bit float or bool or vector of bool or 8/16/32/64-bit integer or 16/32/64-bit float values of length 2/3/4/8/16 of ranks 1, but got 'bf16'}}
- %0 = spirv.GroupNonUniformBroadcast <Workgroup> %value, %one : bf16, i32
+ %0 = spirv.GroupNonUniformBroadcast <Subgroup> %value, %one : bf16, i32
return %0: bf16
}
@@ -135,7 +135,7 @@ func.func @group_non_uniform_broadcast_bf16(%value: bf16) -> bf16 {
func.func @group_non_uniform_broadcast_float8(%value: f8E4M3FN) -> f8E4M3FN {
%one = spirv.Constant 1 : i32
// expected-error @+1 {{op operand #0 must be 8/16/32/64-bit integer or 16/32/64-bit float or bool or vector of bool or 8/16/32/64-bit integer or 16/32/64-bit float values of length 2/3/4/8/16 of ranks 1, but got 'f8E4M3FN'}}
- %0 = spirv.GroupNonUniformBroadcast <Workgroup> %value, %one : f8E4M3FN, i32
+ %0 = spirv.GroupNonUniformBroadcast <Subgroup> %value, %one : f8E4M3FN, i32
return %0: f8E4M3FN
}
@@ -194,15 +194,15 @@ func.func @group_non_uniform_broadcast_first_negative_type_mismatch(%value: f32)
// CHECK-LABEL: @group_non_uniform_elect
func.func @group_non_uniform_elect() -> i1 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformElect <Workgroup> : i1
- %0 = spirv.GroupNonUniformElect <Workgroup> : i1
+ // CHECK: %{{.+}} = spirv.GroupNonUniformElect <Subgroup> : i1
+ %0 = spirv.GroupNonUniformElect <Subgroup> : i1
return %0: i1
}
// -----
func.func @group_non_uniform_elect() -> i1 {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformElect <CrossDevice> : i1
return %0: i1
}
@@ -215,16 +215,16 @@ func.func @group_non_uniform_elect() -> i1 {
// CHECK-LABEL: @group_non_uniform_fadd_reduce
func.func @group_non_uniform_fadd_reduce(%val: f32) -> f32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformFAdd <Workgroup> <Reduce> %{{.+}} : f32 -> f32
- %0 = spirv.GroupNonUniformFAdd <Workgroup> <Reduce> %val : f32 -> f32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformFAdd <Subgroup> <Reduce> %{{.+}} : f32 -> f32
+ %0 = spirv.GroupNonUniformFAdd <Subgroup> <Reduce> %val : f32 -> f32
return %0: f32
}
// CHECK-LABEL: @group_non_uniform_fadd_clustered_reduce
func.func @group_non_uniform_fadd_clustered_reduce(%val: vector<2xf32>) -> vector<2xf32> {
%four = spirv.Constant 4 : i32
- // CHECK: %{{.+}} = spirv.GroupNonUniformFAdd <Workgroup> <ClusteredReduce> %{{.+}} cluster_size(%{{.+}}) : vector<2xf32>, i32 -> vector<2xf32>
- %0 = spirv.GroupNonUniformFAdd <Workgroup> <ClusteredReduce> %val cluster_size(%four) : vector<2xf32>, i32 -> vector<2xf32>
+ // CHECK: %{{.+}} = spirv.GroupNonUniformFAdd <Subgroup> <ClusteredReduce> %{{.+}} cluster_size(%{{.+}}) : vector<2xf32>, i32 -> vector<2xf32>
+ %0 = spirv.GroupNonUniformFAdd <Subgroup> <ClusteredReduce> %val cluster_size(%four) : vector<2xf32>, i32 -> vector<2xf32>
return %0: vector<2xf32>
}
@@ -234,16 +234,16 @@ func.func @group_non_uniform_fadd_clustered_reduce(%val: vector<2xf32>) -> vecto
// CHECK-LABEL: @group_non_uniform_fmul_reduce
func.func @group_non_uniform_fmul_reduce(%val: f32) -> f32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformFMul <Workgroup> <Reduce> %{{.+}} : f32 -> f32
- %0 = spirv.GroupNonUniformFMul <Workgroup> <Reduce> %val : f32 -> f32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformFMul <Subgroup> <Reduce> %{{.+}} : f32 -> f32
+ %0 = spirv.GroupNonUniformFMul <Subgroup> <Reduce> %val : f32 -> f32
return %0: f32
}
// CHECK-LABEL: @group_non_uniform_fmul_clustered_reduce
func.func @group_non_uniform_fmul_clustered_reduce(%val: vector<2xf32>) -> vector<2xf32> {
%four = spirv.Constant 4 : i32
- // CHECK: %{{.+}} = spirv.GroupNonUniformFMul <Workgroup> <ClusteredReduce> %{{.+}} cluster_size(%{{.+}}) : vector<2xf32>, i32 -> vector<2xf32>
- %0 = spirv.GroupNonUniformFMul <Workgroup> <ClusteredReduce> %val cluster_size(%four) : vector<2xf32>, i32 -> vector<2xf32>
+ // CHECK: %{{.+}} = spirv.GroupNonUniformFMul <Subgroup> <ClusteredReduce> %{{.+}} cluster_size(%{{.+}}) : vector<2xf32>, i32 -> vector<2xf32>
+ %0 = spirv.GroupNonUniformFMul <Subgroup> <ClusteredReduce> %val cluster_size(%four) : vector<2xf32>, i32 -> vector<2xf32>
return %0: vector<2xf32>
}
@@ -251,7 +251,7 @@ func.func @group_non_uniform_fmul_clustered_reduce(%val: vector<2xf32>) -> vecto
func.func @group_non_uniform_bf16_fmul_reduce(%val: bf16) -> bf16 {
// expected-error @+1 {{op operand #0 must be 16/32/64-bit float or fixed-length vector of 16/32/64-bit float values of length 2/3/4/8/16 of ranks 1, but got 'bf16'}}
- %0 = spirv.GroupNonUniformFMul <Workgroup> <Reduce> %val : bf16 -> bf16
+ %0 = spirv.GroupNonUniformFMul <Subgroup> <Reduce> %val : bf16 -> bf16
return %0: bf16
}
@@ -263,8 +263,8 @@ func.func @group_non_uniform_bf16_fmul_reduce(%val: bf16) -> bf16 {
// CHECK-LABEL: @group_non_uniform_fmax_reduce
func.func @group_non_uniform_fmax_reduce(%val: f32) -> f32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformFMax <Workgroup> <Reduce> %{{.+}} : f32 -> f32
- %0 = spirv.GroupNonUniformFMax <Workgroup> <Reduce> %val : f32 -> f32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformFMax <Subgroup> <Reduce> %{{.+}} : f32 -> f32
+ %0 = spirv.GroupNonUniformFMax <Subgroup> <Reduce> %val : f32 -> f32
return %0: f32
}
@@ -272,7 +272,7 @@ func.func @group_non_uniform_fmax_reduce(%val: f32) -> f32 {
func.func @group_non_uniform_bf16_fmax_reduce(%val: bf16) -> bf16 {
// expected-error @+1 {{op operand #0 must be 16/32/64-bit float or fixed-length vector of 16/32/64-bit float values of length 2/3/4/8/16 of ranks 1, but got 'bf16'}}
- %0 = spirv.GroupNonUniformFMax <Workgroup> <Reduce> %val : bf16 -> bf16
+ %0 = spirv.GroupNonUniformFMax <Subgroup> <Reduce> %val : bf16 -> bf16
return %0: bf16
}
@@ -284,8 +284,8 @@ func.func @group_non_uniform_bf16_fmax_reduce(%val: bf16) -> bf16 {
// CHECK-LABEL: @group_non_uniform_fmin_reduce
func.func @group_non_uniform_fmin_reduce(%val: f32) -> f32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformFMin <Workgroup> <Reduce> %{{.+}} : f32 -> f32
- %0 = spirv.GroupNonUniformFMin <Workgroup> <Reduce> %val : f32 -> f32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformFMin <Subgroup> <Reduce> %{{.+}} : f32 -> f32
+ %0 = spirv.GroupNonUniformFMin <Subgroup> <Reduce> %val : f32 -> f32
return %0: f32
}
@@ -297,23 +297,23 @@ func.func @group_non_uniform_fmin_reduce(%val: f32) -> f32 {
// CHECK-LABEL: @group_non_uniform_iadd_reduce
func.func @group_non_uniform_iadd_reduce(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformIAdd <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformIAdd <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformIAdd <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformIAdd <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
// CHECK-LABEL: @group_non_uniform_iadd_clustered_reduce
func.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
%four = spirv.Constant 4 : i32
- // CHECK: %{{.+}} = spirv.GroupNonUniformIAdd <Workgroup> <ClusteredReduce> %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>, i32 -> vector<2xi32>
- %0 = spirv.GroupNonUniformIAdd <Workgroup> <ClusteredReduce> %val cluster_size(%four) : vector<2xi32>, i32 -> vector<2xi32>
+ // CHECK: %{{.+}} = spirv.GroupNonUniformIAdd <Subgroup> <ClusteredReduce> %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>, i32 -> vector<2xi32>
+ %0 = spirv.GroupNonUniformIAdd <Subgroup> <ClusteredReduce> %val cluster_size(%four) : vector<2xi32>, i32 -> vector<2xi32>
return %0: vector<2xi32>
}
// -----
func.func @group_non_uniform_iadd_reduce(%val: i32) -> i32 {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformIAdd <Device> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -322,7 +322,7 @@ func.func @group_non_uniform_iadd_reduce(%val: i32) -> i32 {
func.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
// expected-error @+1 {{cluster size operand must be provided for 'ClusteredReduce' group operation}}
- %0 = spirv.GroupNonUniformIAdd <Workgroup> <ClusteredReduce> %val : vector<2xi32> -> vector<2xi32>
+ %0 = spirv.GroupNonUniformIAdd <Subgroup> <ClusteredReduce> %val : vector<2xi32> -> vector<2xi32>
return %0: vector<2xi32>
}
@@ -330,7 +330,7 @@ func.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vecto
func.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>, %size: i32) -> vector<2xi32> {
// expected-error @+1 {{cluster size operand must come from a constant op}}
- %0 = spirv.GroupNonUniformIAdd <Workgroup> <ClusteredReduce> %val cluster_size(%size) : vector<2xi32>, i32 -> vector<2xi32>
+ %0 = spirv.GroupNonUniformIAdd <Subgroup> <ClusteredReduce> %val cluster_size(%size) : vector<2xi32>, i32 -> vector<2xi32>
return %0: vector<2xi32>
}
@@ -339,7 +339,7 @@ func.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>, %size: i
func.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
%five = spirv.Constant 5 : i32
// expected-error @+1 {{cluster size operand must be a power of two}}
- %0 = spirv.GroupNonUniformIAdd <Workgroup> <ClusteredReduce> %val cluster_size(%five) : vector<2xi32>, i32 -> vector<2xi32>
+ %0 = spirv.GroupNonUniformIAdd <Subgroup> <ClusteredReduce> %val cluster_size(%five) : vector<2xi32>, i32 -> vector<2xi32>
return %0: vector<2xi32>
}
@@ -351,16 +351,16 @@ func.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vecto
// CHECK-LABEL: @group_non_uniform_imul_reduce
func.func @group_non_uniform_imul_reduce(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformIMul <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformIMul <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformIMul <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformIMul <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
// CHECK-LABEL: @group_non_uniform_imul_clustered_reduce
func.func @group_non_uniform_imul_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
%four = spirv.Constant 4 : i32
- // CHECK: %{{.+}} = spirv.GroupNonUniformIMul <Workgroup> <ClusteredReduce> %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>, i32 -> vector<2xi32>
- %0 = spirv.GroupNonUniformIMul <Workgroup> <ClusteredReduce> %val cluster_size(%four) : vector<2xi32>, i32 -> vector<2xi32>
+ // CHECK: %{{.+}} = spirv.GroupNonUniformIMul <Subgroup> <ClusteredReduce> %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>, i32 -> vector<2xi32>
+ %0 = spirv.GroupNonUniformIMul <Subgroup> <ClusteredReduce> %val cluster_size(%four) : vector<2xi32>, i32 -> vector<2xi32>
return %0: vector<2xi32>
}
@@ -372,8 +372,8 @@ func.func @group_non_uniform_imul_clustered_reduce(%val: vector<2xi32>) -> vecto
// CHECK-LABEL: @group_non_uniform_smax_reduce
func.func @group_non_uniform_smax_reduce(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformSMax <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformSMax <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformSMax <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformSMax <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -385,8 +385,8 @@ func.func @group_non_uniform_smax_reduce(%val: i32) -> i32 {
// CHECK-LABEL: @group_non_uniform_smin_reduce
func.func @group_non_uniform_smin_reduce(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformSMin <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformSMin <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformSMin <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformSMin <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -429,7 +429,7 @@ func.func @group_non_uniform_shuffle_float8(%val: f8E4M3FN, %id: i32) -> f8E4M3F
// -----
func.func @group_non_uniform_shuffle(%val: vector<2xf32>, %id: i32) -> vector<2xf32> {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformShuffle <Device> %val, %id : vector<2xf32>, i32
return %0: vector<2xf32>
}
@@ -481,7 +481,7 @@ func.func @group_non_uniform_shuffle_xor_float8(%val: f8E4M3FN, %id: i32) -> f8E
// -----
func.func @group_non_uniform_shuffle(%val: vector<2xf32>, %id: i32) -> vector<2xf32> {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformShuffleXor <Device> %val, %id : vector<2xf32>, i32
return %0: vector<2xf32>
}
@@ -533,7 +533,7 @@ func.func @group_non_uniform_shuffle_up_float8(%val: f8E4M3FN, %id: i32) -> f8E4
// -----
func.func @group_non_uniform_shuffle(%val: vector<2xf32>, %id: i32) -> vector<2xf32> {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformShuffleUp <Device> %val, %id : vector<2xf32>, i32
return %0: vector<2xf32>
}
@@ -585,7 +585,7 @@ func.func @group_non_uniform_shuffle_down_float8(%val: f8E4M3FN, %id: i32) -> f8
// -----
func.func @group_non_uniform_shuffle(%val: vector<2xf32>, %id: i32) -> vector<2xf32> {
- // expected-error @+1 {{execution_scope must be 'Workgroup' or 'Subgroup'}}
+ // expected-error @+1 {{execution_scope must be 'Subgroup'}}
%0 = spirv.GroupNonUniformShuffleDown <Device> %val, %id : vector<2xf32>, i32
return %0: vector<2xf32>
}
@@ -606,8 +606,8 @@ func.func @group_non_uniform_shuffle(%val: vector<2xf32>, %id: si32) -> vector<2
// CHECK-LABEL: @group_non_uniform_umax_reduce
func.func @group_non_uniform_umax_reduce(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformUMax <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformUMax <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformUMax <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformUMax <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -619,8 +619,8 @@ func.func @group_non_uniform_umax_reduce(%val: i32) -> i32 {
// CHECK-LABEL: @group_non_uniform_umin_reduce
func.func @group_non_uniform_umin_reduce(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformUMin <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformUMin <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformUMin <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformUMin <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -632,8 +632,8 @@ func.func @group_non_uniform_umin_reduce(%val: i32) -> i32 {
// CHECK-LABEL: @group_non_uniform_bitwise_and
func.func @group_non_uniform_bitwise_and(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformBitwiseAnd <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformBitwiseAnd <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformBitwiseAnd <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformBitwiseAnd <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -641,7 +641,7 @@ func.func @group_non_uniform_bitwise_and(%val: i32) -> i32 {
func.func @group_non_uniform_bitwise_and(%val: i1) -> i1 {
// expected-error @+1 {{operand #0 must be 8/16/32/64-bit integer or fixed-length vector of 8/16/32/64-bit integer values of length 2/3/4/8/16 of ranks 1, but got 'i1'}}
- %0 = spirv.GroupNonUniformBitwiseAnd <Workgroup> <Reduce> %val : i1 -> i1
+ %0 = spirv.GroupNonUniformBitwiseAnd <Subgroup> <Reduce> %val : i1 -> i1
return %0: i1
}
@@ -653,8 +653,8 @@ func.func @group_non_uniform_bitwise_and(%val: i1) -> i1 {
// CHECK-LABEL: @group_non_uniform_bitwise_or
func.func @group_non_uniform_bitwise_or(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformBitwiseOr <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformBitwiseOr <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformBitwiseOr <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformBitwiseOr <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -662,7 +662,7 @@ func.func @group_non_uniform_bitwise_or(%val: i32) -> i32 {
func.func @group_non_uniform_bitwise_or(%val: i1) -> i1 {
// expected-error @+1 {{operand #0 must be 8/16/32/64-bit integer or fixed-length vector of 8/16/32/64-bit integer values of length 2/3/4/8/16 of ranks 1, but got 'i1'}}
- %0 = spirv.GroupNonUniformBitwiseOr <Workgroup> <Reduce> %val : i1 -> i1
+ %0 = spirv.GroupNonUniformBitwiseOr <Subgroup> <Reduce> %val : i1 -> i1
return %0: i1
}
@@ -674,8 +674,8 @@ func.func @group_non_uniform_bitwise_or(%val: i1) -> i1 {
// CHECK-LABEL: @group_non_uniform_bitwise_xor
func.func @group_non_uniform_bitwise_xor(%val: i32) -> i32 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformBitwiseXor <Workgroup> <Reduce> %{{.+}} : i32 -> i32
- %0 = spirv.GroupNonUniformBitwiseXor <Workgroup> <Reduce> %val : i32 -> i32
+ // CHECK: %{{.+}} = spirv.GroupNonUniformBitwiseXor <Subgroup> <Reduce> %{{.+}} : i32 -> i32
+ %0 = spirv.GroupNonUniformBitwiseXor <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -683,7 +683,7 @@ func.func @group_non_uniform_bitwise_xor(%val: i32) -> i32 {
func.func @group_non_uniform_bitwise_xor(%val: i1) -> i1 {
// expected-error @+1 {{operand #0 must be 8/16/32/64-bit integer or fixed-length vector of 8/16/32/64-bit integer values of length 2/3/4/8/16 of ranks 1, but got 'i1'}}
- %0 = spirv.GroupNonUniformBitwiseXor <Workgroup> <Reduce> %val : i1 -> i1
+ %0 = spirv.GroupNonUniformBitwiseXor <Subgroup> <Reduce> %val : i1 -> i1
return %0: i1
}
@@ -695,8 +695,8 @@ func.func @group_non_uniform_bitwise_xor(%val: i1) -> i1 {
// CHECK-LABEL: @group_non_uniform_logical_and
func.func @group_non_uniform_logical_and(%val: i1) -> i1 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformLogicalAnd <Workgroup> <Reduce> %{{.+}} : i1 -> i1
- %0 = spirv.GroupNonUniformLogicalAnd <Workgroup> <Reduce> %val : i1 -> i1
+ // CHECK: %{{.+}} = spirv.GroupNonUniformLogicalAnd <Subgroup> <Reduce> %{{.+}} : i1 -> i1
+ %0 = spirv.GroupNonUniformLogicalAnd <Subgroup> <Reduce> %val : i1 -> i1
return %0: i1
}
@@ -704,7 +704,7 @@ func.func @group_non_uniform_logical_and(%val: i1) -> i1 {
func.func @group_non_uniform_logical_and(%val: i32) -> i32 {
// expected-error @+1 {{operand #0 must be bool or fixed-length vector of bool values of length 2/3/4/8/16 of ranks 1, but got 'i32'}}
- %0 = spirv.GroupNonUniformLogicalAnd <Workgroup> <Reduce> %val : i32 -> i32
+ %0 = spirv.GroupNonUniformLogicalAnd <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -716,8 +716,8 @@ func.func @group_non_uniform_logical_and(%val: i32) -> i32 {
// CHECK-LABEL: @group_non_uniform_logical_or
func.func @group_non_uniform_logical_or(%val: i1) -> i1 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformLogicalOr <Workgroup> <Reduce> %{{.+}} : i1 -> i1
- %0 = spirv.GroupNonUniformLogicalOr <Workgroup> <Reduce> %val : i1 -> i1
+ // CHECK: %{{.+}} = spirv.GroupNonUniformLogicalOr <Subgroup> <Reduce> %{{.+}} : i1 -> i1
+ %0 = spirv.GroupNonUniformLogicalOr <Subgroup> <Reduce> %val : i1 -> i1
return %0: i1
}
@@ -725,7 +725,7 @@ func.func @group_non_uniform_logical_or(%val: i1) -> i1 {
func.func @group_non_uniform_logical_or(%val: i32) -> i32 {
// expected-error @+1 {{operand #0 must be bool or fixed-length vector of bool values of length 2/3/4/8/16 of ranks 1, but got 'i32'}}
- %0 = spirv.GroupNonUniformLogicalOr <Workgroup> <Reduce> %val : i32 -> i32
+ %0 = spirv.GroupNonUniformLogicalOr <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
@@ -737,8 +737,8 @@ func.func @group_non_uniform_logical_or(%val: i32) -> i32 {
// CHECK-LABEL: @group_non_uniform_logical_xor
func.func @group_non_uniform_logical_xor(%val: i1) -> i1 {
- // CHECK: %{{.+}} = spirv.GroupNonUniformLogicalXor <Workgroup> <Reduce> %{{.+}} : i1 -> i1
- %0 = spirv.GroupNonUniformLogicalXor <Workgroup> <Reduce> %val : i1 -> i1
+ // CHECK: %{{.+}} = spirv.GroupNonUniformLogicalXor <Subgroup> <Reduce> %{{.+}} : i1 -> i1
+ %0 = spirv.GroupNonUniformLogicalXor <Subgroup> <Reduce> %val : i1 -> i1
return %0: i1
}
@@ -746,7 +746,7 @@ func.func @group_non_uniform_logical_xor(%val: i1) -> i1 {
func.func @group_non_uniform_logical_xor(%val: i32) -> i32 {
// expected-error @+1 {{operand #0 must be bool or fixed-length vector of bool values of length 2/3/4/8/16 of ranks 1, but got 'i32'}}
- %0 = spirv.GroupNonUniformLogicalXor <Workgroup> <Reduce> %val : i32 -> i32
+ %0 = spirv.GroupNonUniformLogicalXor <Subgroup> <Reduce> %val : i32 -> i32
return %0: i32
}
diff --git a/mlir/test/Dialect/SPIRV/IR/target-env.mlir b/mlir/test/Dialect/SPIRV/IR/target-env.mlir
index 273aa33e15eae..7d807bbbfbd4f 100644
--- a/mlir/test/Dialect/SPIRV/IR/target-env.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/target-env.mlir
@@ -71,7 +71,7 @@ func.func @cmp_exchange_weak_unsupported_version(%ptr: !spirv.ptr<i32, Workgroup
func.func @group_non_uniform_ballot_suitable_version(%predicate: i1) -> vector<4xi32> attributes {
spirv.target_env = #spirv.target_env<#spirv.vce<v1.4, [GroupNonUniformBallot], []>, #spirv.resource_limits<>>
} {
- // CHECK: spirv.GroupNonUniformBallot <Workgroup>
+ // CHECK: spirv.GroupNonUniformBallot <Subgroup>
%0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
return %0: vector<4xi32>
}
diff --git a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
index 5ac467d143807..ad9653e55bfe1 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir
@@ -27,7 +27,7 @@ spirv.module Logical GLSL450 attributes {
#spirv.vce<v1.5, [Shader, GroupNonUniformBallot], []>, #spirv.resource_limits<>>
} {
spirv.func @group_non_uniform_ballot(%predicate : i1) -> vector<4xi32> "None" {
- %0 = spirv.GroupNonUniformBallot <Workgroup> %predicate : vector<4xi32>
+ %0 = spirv.GroupNonUniformBallot <Subgroup> %predicate : vector<4xi32>
spirv.ReturnValue %0: vector<4xi32>
}
}
diff --git a/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp b/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp
index 5643a0ff5b91c..ae995a970bade 100644
--- a/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp
+++ b/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp
@@ -168,7 +168,7 @@ struct ConvertToGroupNonUniformBallot : RewritePattern {
PatternRewriter &rewriter) const override {
Value predicate = op->getOperand(0);
rewriter.replaceOpWithNewOp<spirv::GroupNonUniformBallotOp>(
- op, op->getResult(0).getType(), spirv::Scope::Workgroup, predicate);
+ op, op->getResult(0).getType(), spirv::Scope::Subgroup, predicate);
return success();
}
};
More information about the Mlir-commits
mailing list