Add support for the SPIR-V extension SPV_KHR_uniform_group_instructions #82064
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@llvm/pr-subscribers-backend-spir-v Author: Vyacheslav Levytskyy (VyacheslavLevytskyy) ChangesThis PR is to add support for the SPIR-V extension SPV_KHR_uniform_group_instructions that adds new instructions to SPIR-V to support additional group operations within uniform control flow. Patch is 21.86 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/82064.diff 7 Files Affected:
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 8721b900c8beee..b0ce2d33665acb 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -93,6 +93,15 @@ struct IntelSubgroupsBuiltin {
#define GET_IntelSubgroupsBuiltins_DECL
#define GET_IntelSubgroupsBuiltins_IMPL
+struct GroupUniformBuiltin {
+ StringRef Name;
+ uint32_t Opcode;
+ bool IsLogical;
+};
+
+#define GET_GroupUniformBuiltins_DECL
+#define GET_GroupUniformBuiltins_IMPL
+
struct GetBuiltin {
StringRef Name;
InstructionSet::InstructionSet Set;
@@ -974,6 +983,57 @@ static bool generateIntelSubgroupsInst(const SPIRV::IncomingCall *Call,
return true;
}
+static bool generateGroupUniformInst(const SPIRV::IncomingCall *Call,
+ MachineIRBuilder &MIRBuilder,
+ SPIRVGlobalRegistry *GR) {
+ const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
+ MachineFunction &MF = MIRBuilder.getMF();
+ const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
+ if (!ST->canUseExtension(
+ SPIRV::Extension::SPV_KHR_uniform_group_instructions)) {
+ std::string DiagMsg = std::string(Builtin->Name) +
+ ": the builtin requires the following SPIR-V "
+ "extension: SPV_KHR_uniform_group_instructions";
+ report_fatal_error(DiagMsg.c_str(), false);
+ }
+ const SPIRV::GroupUniformBuiltin *GroupUniform =
+ SPIRV::lookupGroupUniformBuiltin(Builtin->Name);
+ MachineRegisterInfo *MRI = MIRBuilder.getMRI();
+
+ Register GroupResultReg = Call->ReturnRegister;
+ MRI->setRegClass(GroupResultReg, &SPIRV::IDRegClass);
+
+ // Scope
+ Register ScopeReg = Call->Arguments[0];
+ MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass);
+
+ // Group Operation
+ Register ConstGroupOpReg = Call->Arguments[1];
+ const MachineInstr *Const = getDefInstrMaybeConstant(ConstGroupOpReg, MRI);
+ if (!Const || Const->getOpcode() != TargetOpcode::G_CONSTANT)
+ report_fatal_error(
+ "expect a constant group operation for a uniform group instruction",
+ false);
+ const MachineOperand &ConstOperand = Const->getOperand(1);
+ if (!ConstOperand.isCImm())
+ report_fatal_error("uniform group instructions: group operation must be an "
+ "integer constant",
+ false);
+
+ // Value
+ Register ValueReg = Call->Arguments[2];
+ MRI->setRegClass(ValueReg, &SPIRV::IDRegClass);
+
+ auto MIB = MIRBuilder.buildInstr(GroupUniform->Opcode)
+ .addDef(GroupResultReg)
+ .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+ .addUse(ScopeReg);
+ addNumImm(ConstOperand.getCImm()->getValue(), MIB);
+ MIB.addUse(ValueReg);
+
+ return true;
+}
+
// These queries ask for a single size_t result for a given dimension index, e.g
// size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
// these values are all vec3 types, so we need to extract the correct index or
@@ -2053,6 +2113,8 @@ std::optional<bool> lowerBuiltin(const StringRef DemangledCall,
return generateLoadStoreInst(Call.get(), MIRBuilder, GR);
case SPIRV::IntelSubgroups:
return generateIntelSubgroupsInst(Call.get(), MIRBuilder, GR);
+ case SPIRV::GroupUniform:
+ return generateGroupUniformInst(Call.get(), MIRBuilder, GR);
}
return false;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
index 4013dd22f4ab57..900183ed9a8b48 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
@@ -55,6 +55,7 @@ def AsyncCopy : BuiltinGroup;
def VectorLoadStore : BuiltinGroup;
def LoadStore : BuiltinGroup;
def IntelSubgroups : BuiltinGroup;
+def GroupUniform : BuiltinGroup;
//===----------------------------------------------------------------------===//
// Class defining a demangled builtin record. The information in the record
@@ -604,7 +605,10 @@ class GroupBuiltin<string name, Op operation> {
!eq(operation, OpGroupNonUniformBallotFindMSB));
bit IsLogical = !or(!eq(operation, OpGroupNonUniformLogicalAnd),
!eq(operation, OpGroupNonUniformLogicalOr),
- !eq(operation, OpGroupNonUniformLogicalXor));
+ !eq(operation, OpGroupNonUniformLogicalXor),
+ !eq(operation, OpGroupLogicalAndKHR),
+ !eq(operation, OpGroupLogicalOrKHR),
+ !eq(operation, OpGroupLogicalXorKHR));
bit NoGroupOperation = !or(IsElect, IsAllOrAny, IsAllEqual,
IsBallot, IsInverseBallot,
IsBallotBitExtract, IsBallotFindBit,
@@ -872,6 +876,51 @@ defm : DemangledGroupBuiltin<"group_non_uniform_scan_inclusive_logical_xors", Wo
defm : DemangledGroupBuiltin<"group_non_uniform_scan_exclusive_logical_xors", WorkOrSub, OpGroupNonUniformLogicalXor>;
defm : DemangledGroupBuiltin<"group_clustered_reduce_logical_xor", WorkOrSub, OpGroupNonUniformLogicalXor>;
+// cl_khr_work_group_uniform_arithmetic / SPV_KHR_uniform_group_instructions
+defm : DemangledGroupBuiltin<"group_reduce_imul", OnlyWork, OpGroupIMulKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_mulu", OnlyWork, OpGroupIMulKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_muls", OnlyWork, OpGroupIMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_imul", OnlyWork, OpGroupIMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_mulu", OnlyWork, OpGroupIMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_muls", OnlyWork, OpGroupIMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_exclusive_imul", OnlyWork, OpGroupIMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_exclusive_mulu", OnlyWork, OpGroupIMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_exclusive_muls", OnlyWork, OpGroupIMulKHR>;
+
+defm : DemangledGroupBuiltin<"group_reduce_mulf", OnlyWork, OpGroupFMulKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_mulh", OnlyWork, OpGroupFMulKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_muld", OnlyWork, OpGroupFMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_mulf", OnlyWork, OpGroupFMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_mulh", OnlyWork, OpGroupFMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_muld", OnlyWork, OpGroupFMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_exclusive_mulf", OnlyWork, OpGroupFMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_exclusive_mulh", OnlyWork, OpGroupFMulKHR>;
+defm : DemangledGroupBuiltin<"group_scan_exclusive_muld", OnlyWork, OpGroupFMulKHR>;
+
+defm : DemangledGroupBuiltin<"group_scan_exclusive_and", OnlyWork, OpGroupBitwiseAndKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_and", OnlyWork, OpGroupBitwiseAndKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_and", OnlyWork, OpGroupBitwiseAndKHR>;
+
+defm : DemangledGroupBuiltin<"group_scan_exclusive_or", OnlyWork, OpGroupBitwiseOrKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_or", OnlyWork, OpGroupBitwiseOrKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_or", OnlyWork, OpGroupBitwiseOrKHR>;
+
+defm : DemangledGroupBuiltin<"group_scan_exclusive_xor", OnlyWork, OpGroupBitwiseXorKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_xor", OnlyWork, OpGroupBitwiseXorKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_xor", OnlyWork, OpGroupBitwiseXorKHR>;
+
+defm : DemangledGroupBuiltin<"group_scan_exclusive_logical_and", OnlyWork, OpGroupLogicalAndKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_logical_and", OnlyWork, OpGroupLogicalAndKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_logical_and", OnlyWork, OpGroupLogicalAndKHR>;
+
+defm : DemangledGroupBuiltin<"group_scan_exclusive_logical_or", OnlyWork, OpGroupLogicalOrKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_logical_or", OnlyWork, OpGroupLogicalOrKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_logical_or", OnlyWork, OpGroupLogicalOrKHR>;
+
+defm : DemangledGroupBuiltin<"group_scan_exclusive_logical_xor", OnlyWork, OpGroupLogicalXorKHR>;
+defm : DemangledGroupBuiltin<"group_scan_inclusive_logical_xor", OnlyWork, OpGroupLogicalXorKHR>;
+defm : DemangledGroupBuiltin<"group_reduce_logical_xor", OnlyWork, OpGroupLogicalXorKHR>;
+
//===----------------------------------------------------------------------===//
// Class defining a sub group builtin that should be translated into a
// SPIR-V instruction using the SPV_INTEL_subgroups extension.
@@ -928,6 +977,52 @@ foreach i = ["", "2", "4", "8", "16"] in {
}
// OpSubgroupImageBlockReadINTEL and OpSubgroupImageBlockWriteINTEL are to be resolved later on (in code)
+//===----------------------------------------------------------------------===//
+// Class defining a builtin for group operations within uniform control flow.
+// It should be translated into a SPIR-V instruction using
+// the SPV_KHR_uniform_group_instructions extension.
+//
+// name is the demangled name of the given builtin.
+// opcode specifies the SPIR-V operation code of the generated instruction.
+//===----------------------------------------------------------------------===//
+class GroupUniformBuiltin<string name, Op operation> {
+ string Name = name;
+ Op Opcode = operation;
+ bit IsLogical = !or(!eq(operation, OpGroupLogicalAndKHR),
+ !eq(operation, OpGroupLogicalOrKHR),
+ !eq(operation, OpGroupLogicalXorKHR));
+}
+
+// Table gathering all the Intel sub group builtins.
+def GroupUniformBuiltins : GenericTable {
+ let FilterClass = "GroupUniformBuiltin";
+ let Fields = ["Name", "Opcode", "IsLogical"];
+}
+
+// Function to lookup group builtins by their name and set.
+def lookupGroupUniformBuiltin : SearchIndex {
+ let Table = GroupUniformBuiltins;
+ let Key = ["Name"];
+}
+
+// Multiclass used to define incoming builtin records for
+// the SPV_KHR_uniform_group_instructions extension
+// and corresponding work group builtin records.
+multiclass DemangledGroupUniformBuiltin<string name, bits<8> minNumArgs, bits<8> maxNumArgs, Op operation> {
+ def : DemangledBuiltin<!strconcat("__spirv_Group", name), OpenCL_std, GroupUniform, minNumArgs, maxNumArgs>;
+ def : GroupUniformBuiltin<!strconcat("__spirv_Group", name), operation>;
+}
+
+// cl_khr_work_group_uniform_arithmetic / SPV_KHR_uniform_group_instructions
+defm : DemangledGroupUniformBuiltin<"IMulKHR", 3, 3, OpGroupIMulKHR>;
+defm : DemangledGroupUniformBuiltin<"FMulKHR", 3, 3, OpGroupFMulKHR>;
+defm : DemangledGroupUniformBuiltin<"BitwiseAndKHR", 3, 3, OpGroupBitwiseAndKHR>;
+defm : DemangledGroupUniformBuiltin<"BitwiseOrKHR", 3, 3, OpGroupBitwiseOrKHR>;
+defm : DemangledGroupUniformBuiltin<"BitwiseXorKHR", 3, 3, OpGroupBitwiseXorKHR>;
+defm : DemangledGroupUniformBuiltin<"LogicalAndKHR", 3, 3, OpGroupLogicalAndKHR>;
+defm : DemangledGroupUniformBuiltin<"LogicalOrKHR", 3, 3, OpGroupLogicalOrKHR>;
+defm : DemangledGroupUniformBuiltin<"LogicalXorKHR", 3, 3, OpGroupLogicalXorKHR>;
+
//===----------------------------------------------------------------------===//
// Class defining a get builtin record used for lowering builtin calls such as
// "get_sub_group_eq_mask" or "get_global_id" to SPIR-V instructions.
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td b/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
index 904fef1d6c82f9..873a959e828fc5 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
+++ b/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
@@ -773,6 +773,8 @@ def OpConstantFunctionPointerINTEL: Op<5600, (outs ID:$res), (ins TYPE:$ty, ID:$
def OpFunctionPointerCallINTEL: Op<5601, (outs ID:$res), (ins TYPE:$ty, ID:$funPtr, variable_ops), "$res = OpFunctionPointerCallINTEL $ty $funPtr">;
// 3.49.21. Group and Subgroup Instructions
+
+// - SPV_INTEL_subgroups
def OpSubgroupShuffleINTEL: Op<5571, (outs ID:$res), (ins TYPE:$type, ID:$data, ID:$invocationId),
"$res = OpSubgroupShuffleINTEL $type $data $invocationId">;
def OpSubgroupShuffleDownINTEL: Op<5572, (outs ID:$res), (ins TYPE:$type, ID:$current, ID:$next, ID:$delta),
@@ -789,3 +791,21 @@ def OpSubgroupImageBlockReadINTEL: Op<5577, (outs ID:$res), (ins TYPE:$type, ID:
"$res = OpSubgroupImageBlockReadINTEL $type $image $coordinate">;
def OpSubgroupImageBlockWriteINTEL: Op<5578, (outs), (ins ID:$image, ID:$coordinate, ID:$data),
"OpSubgroupImageBlockWriteINTEL $image $coordinate $data">;
+
+// - SPV_KHR_uniform_group_instructions
+def OpGroupIMulKHR: Op<6401, (outs ID:$res), (ins TYPE:$type, ID:$scope, i32imm:$groupOp, ID:$value),
+ "$res = OpGroupIMulKHR $type $scope $groupOp $value">;
+def OpGroupFMulKHR: Op<6402, (outs ID:$res), (ins TYPE:$type, ID:$scope, i32imm:$groupOp, ID:$value),
+ "$res = OpGroupFMulKHR $type $scope $groupOp $value">;
+def OpGroupBitwiseAndKHR: Op<6403, (outs ID:$res), (ins TYPE:$type, ID:$scope, i32imm:$groupOp, ID:$value),
+ "$res = OpGroupBitwiseAndKHR $type $scope $groupOp $value">;
+def OpGroupBitwiseOrKHR: Op<6404, (outs ID:$res), (ins TYPE:$type, ID:$scope, i32imm:$groupOp, ID:$value),
+ "$res = OpGroupBitwiseOrKHR $type $scope $groupOp $value">;
+def OpGroupBitwiseXorKHR: Op<6405, (outs ID:$res), (ins TYPE:$type, ID:$scope, i32imm:$groupOp, ID:$value),
+ "$res = OpGroupBitwiseXorKHR $type $scope $groupOp $value">;
+def OpGroupLogicalAndKHR: Op<6406, (outs ID:$res), (ins TYPE:$type, ID:$scope, i32imm:$groupOp, ID:$value),
+ "$res = OpGroupLogicalAndKHR $type $scope $groupOp $value">;
+def OpGroupLogicalOrKHR: Op<6407, (outs ID:$res), (ins TYPE:$type, ID:$scope, i32imm:$groupOp, ID:$value),
+ "$res = OpGroupLogicalOrKHR $type $scope $groupOp $value">;
+def OpGroupLogicalXorKHR: Op<6408, (outs ID:$res), (ins TYPE:$type, ID:$scope, i32imm:$groupOp, ID:$value),
+ "$res = OpGroupLogicalXorKHR $type $scope $groupOp $value">;
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 688b98ffa67477..baca07fbbec41e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -1008,6 +1008,19 @@ void addInstrRequirements(const MachineInstr &MI,
Reqs.addCapability(SPIRV::Capability::FunctionPointersINTEL);
}
break;
+ case SPIRV::OpGroupIMulKHR:
+ case SPIRV::OpGroupFMulKHR:
+ case SPIRV::OpGroupBitwiseAndKHR:
+ case SPIRV::OpGroupBitwiseOrKHR:
+ case SPIRV::OpGroupBitwiseXorKHR:
+ case SPIRV::OpGroupLogicalAndKHR:
+ case SPIRV::OpGroupLogicalOrKHR:
+ case SPIRV::OpGroupLogicalXorKHR:
+ if (ST.canUseExtension(SPIRV::Extension::SPV_KHR_uniform_group_instructions)) {
+ Reqs.addExtension(SPIRV::Extension::SPV_KHR_uniform_group_instructions);
+ Reqs.addCapability(SPIRV::Capability::GroupUniformArithmeticKHR);
+ }
+ break;
case SPIRV::OpFunctionPointerCallINTEL:
if (ST.canUseExtension(SPIRV::Extension::SPV_INTEL_function_pointers)) {
Reqs.addExtension(SPIRV::Extension::SPV_INTEL_function_pointers);
diff --git a/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp b/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
index 354cd5d9b297e7..da62d7721b85d3 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
@@ -42,6 +42,10 @@ cl::list<SPIRV::Extension::Extension> Extensions(
"use of local memory and work group barriers, and to "
"utilize specialized hardware to load and store blocks of "
"data from images or buffers."),
+ clEnumValN(SPIRV::Extension::SPV_KHR_uniform_group_instructions,
+ "SPV_KHR_uniform_group_instructions",
+ "Allows support for additional group operations within "
+ "uniform control flow."),
clEnumValN(SPIRV::Extension::SPV_KHR_no_integer_wrap_decoration,
"SPV_KHR_no_integer_wrap_decoration",
"Adds decorations to indicate that a given instruction does "
diff --git a/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td b/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
index ed05013642ac21..b11166dfc6dcd9 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
+++ b/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
@@ -455,6 +455,7 @@ defm BitInstructions : CapabilityOperand<6025, 0, 0, [SPV_KHR_bit_instructions],
defm ExpectAssumeKHR : CapabilityOperand<5629, 0, 0, [SPV_KHR_expect_assume], []>;
defm FunctionPointersINTEL : CapabilityOperand<5603, 0, 0, [SPV_INTEL_function_pointers], []>;
defm IndirectReferencesINTEL : CapabilityOperand<5604, 0, 0, [SPV_INTEL_function_pointers], []>;
+defm GroupUniformArithmeticKHR : CapabilityOperand<6400, 0, 0, [SPV_KHR_uniform_group_instructions], []>;
//===----------------------------------------------------------------------===//
// Multiclass used to define SourceLanguage enum values and at the same time
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_uniform_group_instructions/uniform-group-instructions.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_uniform_group_instructions/uniform-group-instructions.ll
new file mode 100644
index 00000000000000..39bf63ddae4fe7
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_uniform_group_instructions/uniform-group-instructions.ll
@@ -0,0 +1,80 @@
+; RUN: not llc -O0 -mtriple=spirv32-unknown-unknown %s -o %t.spvt 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-extensions=SPV_KHR_uniform_group_instructions %s -o - | FileCheck %s
+
+; CHECK-ERROR: LLVM ERROR: __spirv_GroupBitwiseAndKHR: the builtin requires the following SPIR-V extension: SPV_KHR_uniform_group_instructions
+
+; CHECK: Capability GroupUniformArithmeticKHR
+; CHECK: Extension "SPV_KHR_uniform_group_instructions"
+; CHECK-DAG: %[[TyInt:[0-9]+]] = OpTypeInt 32 0
+; CHECK-DAG: %[[TyBool:[0-9]+]] = OpTypeBool
+; CHECK-DAG: %[[TyFloat:[0-9]+]] = OpTypeFloat 16
+; CHECK-DAG: %[[Scope:[0-9]+]] = OpConstant %[[TyInt]] 2
+; CHECK-DAG: %[[ConstInt:[0-9]+]] = OpConstant %[[TyInt]] 0
+; CHECK-DAG: %[[ConstFloat:[0-9]+]] = OpConstant %[[TyFloat]] 0
+; CHECK-DAG: %[[ConstBool:[0-9]+]] = OpConstantFalse %[[TyBool]]
+
+; CHECK: OpGroupBitwiseAndKHR %[[TyInt]] %[[Scope]] 0 %[[ConstInt]]
+; CHECK: OpGroupBitwiseOrKHR %[[TyInt]] %[[Scope]] 0 %[[ConstInt]]
+; CHECK: OpGroupBitwiseXorKHR %[[TyInt]] %[[Scope]] 0 %[[ConstInt]]
+; CHECK: OpGroupLogicalAndKHR %[[TyBool]] %[[Scope]] 0 %[[ConstBool]]
+; CHECK: OpGroupLogicalOrKHR %[[TyBool]] %[[Scope]] 0 %[[ConstBool]]
+; CHECK: OpGroupLogicalXorKHR %[[TyBool]] %[[Scope]] 0 %[[ConstBool]]
+; CHECK: OpGroupIMulKHR %[[TyInt]] %[[Scope]] 0 %[[ConstInt]]
+; CHECK: OpGroupFMulKHR %[[TyFloat]] %[[Scope]] 0 %[[ConstFloat]]
+
+; CHECK: OpGroupBitwiseAndKHR %[[TyInt]] %[[Scope]] 0 %[[ConstInt]]
+; CHECK: OpGroupBitwiseOrKHR %[[TyInt]] %[[Scope]] 0 %[[ConstInt]]
+; CHECK: OpGroupBitwiseXorKHR %[[TyInt]] %[[Scope]] 0 %[[ConstInt]]
+; CHECK: OpGroupLogicalAndKHR %[[TyBool]] %[[Scope]] 0 %[[ConstBool]]
+; CHECK: OpGroupLogicalOrKHR %[[TyBool]] %[[Scope]] 0 %[[ConstBool]]
+; CHECK: OpGroupLogicalXorKHR %[[TyBool]] %[[Scope]] 0 %[[ConstBool]]
+; CHECK: OpGroupIMulKHR %[[TyInt]] %[[Scope]] 0 %[[ConstInt]]
+; CHECK: OpGroupFMulKHR %[[TyFloat]] %[[Scope]] 0 %[[ConstFloat]]
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+define dso_local spir_func void @test1() {
+entry:
+ %res1 = tail call spir_func i32 @_Z26__spirv_GroupBitwiseAndKHR(i32 2, i32 0, i32 0)
+ %res2 = tail call spir_func i32 @_Z25__spirv_GroupBitwiseOrKHR(i32 2, i32 0, i32 0)
+ %res3 = tail call spir_func i32 @_Z26__spirv_GroupBitwiseXorKHR(i32 2, i32 0, i32 0)
+ %res4 = tail call spir_func i1 @_Z26__spirv_GroupLogicalAndKHR(i32 2, i32 0, i1 false)
+ %res5 = tail call spir_func i1 @_Z25__spirv_GroupLogicalOrKHR(i32 2, i32 0, i1 false)
+ %res6 = tail call spir_func i1 @_Z26__spirv_GroupLogicalXorKHR(i32 2, i32 0, i1 false)
+ %res7 = tail call spir_func i32 @_Z20__spirv_GroupIMulKHR(i32 2, i32 0, i32 0)
+ %res8 = tail call spir_func half @_Z20__spirv_GroupFMulKHR(...
[truncated]
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
michalpaszkowski
approved these changes
Feb 19, 2024
VyacheslavLevytskyy
force-pushed
the
add_SPV_KHR_uniform_group_instructions
branch
from
05b61f2
to
824e22a
Compare
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR is to add support for the SPIR-V extension SPV_KHR_uniform_group_instructions that adds new instructions to SPIR-V to support additional group operations within uniform control flow.