[llvm] Add support for the SPIR-V extension SPV_KHR_uniform_group_instructions (PR #82064)
Vyacheslav Levytskyy via llvm-commits
llvm-commits at lists.llvm.org
Fri Feb 16 15:06:24 PST 2024
https://github.com/VyacheslavLevytskyy created https://github.com/llvm/llvm-project/pull/82064
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.
>From d59cc1a9aab10141a4d3bc901d8f5011708fefaf Mon Sep 17 00:00:00 2001
From: "Levytskyy, Vyacheslav" <vyacheslav.levytskyy at intel.com>
Date: Fri, 16 Feb 2024 15:03:58 -0800
Subject: [PATCH] add support for the SPIR-V extension
SPV_KHR_uniform_group_instructions
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 62 ++++++++++++
llvm/lib/Target/SPIRV/SPIRVBuiltins.td | 97 ++++++++++++++++++-
llvm/lib/Target/SPIRV/SPIRVInstrInfo.td | 20 ++++
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 13 +++
llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp | 4 +
.../lib/Target/SPIRV/SPIRVSymbolicOperands.td | 1 +
.../uniform-group-instructions.ll | 80 +++++++++++++++
7 files changed, 276 insertions(+), 1 deletion(-)
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_uniform_group_instructions/uniform-group-instructions.ll
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(i32 2, i32 0, half 0xH0000)
+ ret void
+}
+
+define dso_local spir_func void @test2() {
+entry:
+ %res1 = tail call spir_func i32 @_Z21work_group_reduce_andi(i32 0)
+ %res2 = tail call spir_func i32 @_Z20work_group_reduce_ori(i32 0)
+ %res3 = tail call spir_func i32 @_Z21work_group_reduce_xori(i32 0)
+ %res4 = tail call spir_func i32 @_Z29work_group_reduce_logical_andi(i32 0)
+ %res5 = tail call spir_func i32 @_Z28work_group_reduce_logical_ori(i32 0)
+ %res6 = tail call spir_func i32 @_Z29work_group_reduce_logical_xori(i32 0)
+ %res7 = tail call spir_func i32 @_Z21work_group_reduce_muli(i32 0)
+ %res8 = tail call spir_func half @_Z21work_group_reduce_mulDh(half 0xH0000)
+ ret void
+}
+
+declare dso_local spir_func i32 @_Z26__spirv_GroupBitwiseAndKHR(i32, i32, i32)
+declare dso_local spir_func i32 @_Z25__spirv_GroupBitwiseOrKHR(i32, i32, i32)
+declare dso_local spir_func i32 @_Z26__spirv_GroupBitwiseXorKHR(i32, i32, i32)
+declare dso_local spir_func i1 @_Z26__spirv_GroupLogicalAndKHR(i32, i32, i1)
+declare dso_local spir_func i1 @_Z25__spirv_GroupLogicalOrKHR(i32, i32, i1)
+declare dso_local spir_func i1 @_Z26__spirv_GroupLogicalXorKHR(i32, i32, i1)
+declare dso_local spir_func i32 @_Z20__spirv_GroupIMulKHR(i32, i32, i32)
+declare dso_local spir_func half @_Z20__spirv_GroupFMulKHR(i32, i32, half)
+
+declare dso_local spir_func i32 @_Z21work_group_reduce_andi(i32)
+declare dso_local spir_func i32 @_Z20work_group_reduce_ori(i32)
+declare dso_local spir_func i32 @_Z21work_group_reduce_xori(i32)
+declare dso_local spir_func i32 @_Z29work_group_reduce_logical_andi(i32)
+declare dso_local spir_func i32 @_Z28work_group_reduce_logical_ori(i32)
+declare dso_local spir_func i32 @_Z29work_group_reduce_logical_xori(i32)
+declare dso_local spir_func i32 @_Z21work_group_reduce_muli(i32)
+declare dso_local spir_func half @_Z21work_group_reduce_mulDh(half)
More information about the llvm-commits
mailing list