[SPIR-V] Fix types of internal intrinsic functions and add a test case for __builtin_alloca() #92265
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 @llvm/pr-subscribers-llvm-ir Author: Vyacheslav Levytskyy (VyacheslavLevytskyy) ChangesThis PR generation of argument types of internal intrinsic functions The added test case serves two purposes: it is to check the above mentioned fix and to demonstrate that a call to __builtin_alloca() maps to instructions from SPV_INTEL_variable_length_array when this extension is available. Full diff: https://github.com/llvm/llvm-project/pull/92265.diff 3 Files Affected:
diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
index 931786ab96479..cc84decc43407 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -27,7 +27,7 @@ let TargetPrefix = "spv" in {
def int_spv_insertv : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_any_ty, llvm_vararg_ty]>;
def int_spv_extractelt : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_anyint_ty]>;
def int_spv_insertelt : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_any_ty, llvm_anyint_ty]>;
- def int_spv_const_composite : Intrinsic<[llvm_i32_ty], [llvm_vararg_ty]>;
+ def int_spv_const_composite : Intrinsic<[llvm_any_ty], [llvm_vararg_ty]>;
def int_spv_bitcast : Intrinsic<[llvm_any_ty], [llvm_any_ty]>;
def int_spv_ptrcast : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_metadata_ty, llvm_i32_ty], [ImmArg<ArgIndex<2>>]>;
def int_spv_switch : Intrinsic<[], [llvm_any_ty, llvm_vararg_ty]>;
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index c00066f5dca62..8964ad0fd559d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -167,8 +167,9 @@ static bool isMemInstrToReplace(Instruction *I) {
isa<ExtractValueInst>(I) || isa<AtomicCmpXchgInst>(I);
}
-static bool isAggrToReplace(const Value *V) {
- return isa<ConstantAggregate>(V) || isa<ConstantDataArray>(V) ||
+static bool isAggrConstForceInt32(const Value *V) {
+ return isa<ConstantArray>(V) || isa<ConstantStruct>(V) ||
+ isa<ConstantDataArray>(V) ||
(isa<ConstantAggregateZero>(V) && !V->getType()->isVectorTy());
}
@@ -576,36 +577,42 @@ void SPIRVEmitIntrinsics::preprocessCompositeConstants(IRBuilder<> &B) {
assert(I);
bool KeepInst = false;
for (const auto &Op : I->operands()) {
- auto BuildCompositeIntrinsic =
- [](Constant *AggrC, ArrayRef<Value *> Args, Value *Op, Instruction *I,
- IRBuilder<> &B, std::queue<Instruction *> &Worklist,
- bool &KeepInst, SPIRVEmitIntrinsics &SEI) {
- B.SetInsertPoint(I);
- auto *CCI =
- B.CreateIntrinsic(Intrinsic::spv_const_composite, {}, {Args});
- Worklist.push(CCI);
- I->replaceUsesOfWith(Op, CCI);
- KeepInst = true;
- SEI.AggrConsts[CCI] = AggrC;
- SEI.AggrConstTypes[CCI] = SEI.deduceNestedTypeHelper(AggrC);
- };
-
- if (auto *AggrC = dyn_cast<ConstantAggregate>(Op)) {
- SmallVector<Value *> Args(AggrC->op_begin(), AggrC->op_end());
- BuildCompositeIntrinsic(AggrC, Args, Op, I, B, Worklist, KeepInst,
- *this);
- } else if (auto *AggrC = dyn_cast<ConstantDataArray>(Op)) {
+ Constant *AggrConst = nullptr;
+ Type *ResTy = nullptr;
+ if (auto *COp = dyn_cast<ConstantVector>(Op)) {
+ AggrConst = cast<Constant>(COp);
+ ResTy = COp->getType();
+ } else if (auto *COp = dyn_cast<ConstantArray>(Op)) {
+ AggrConst = cast<Constant>(COp);
+ ResTy = B.getInt32Ty();
+ } else if (auto *COp = dyn_cast<ConstantStruct>(Op)) {
+ AggrConst = cast<Constant>(COp);
+ ResTy = B.getInt32Ty();
+ } else if (auto *COp = dyn_cast<ConstantDataArray>(Op)) {
+ AggrConst = cast<Constant>(COp);
+ ResTy = B.getInt32Ty();
+ } else if (auto *COp = dyn_cast<ConstantAggregateZero>(Op)) {
+ if (!Op->getType()->isVectorTy()) {
+ AggrConst = cast<Constant>(COp);
+ ResTy = B.getInt32Ty();
+ }
+ }
+ if (AggrConst) {
SmallVector<Value *> Args;
- for (unsigned i = 0; i < AggrC->getNumElements(); ++i)
- Args.push_back(AggrC->getElementAsConstant(i));
- BuildCompositeIntrinsic(AggrC, Args, Op, I, B, Worklist, KeepInst,
- *this);
- } else if (isa<ConstantAggregateZero>(Op) &&
- !Op->getType()->isVectorTy()) {
- auto *AggrC = cast<ConstantAggregateZero>(Op);
- SmallVector<Value *> Args(AggrC->op_begin(), AggrC->op_end());
- BuildCompositeIntrinsic(AggrC, Args, Op, I, B, Worklist, KeepInst,
- *this);
+ if (auto *COp = dyn_cast<ConstantDataSequential>(Op))
+ for (unsigned i = 0; i < COp->getNumElements(); ++i)
+ Args.push_back(COp->getElementAsConstant(i));
+ else
+ for (auto &COp : AggrConst->operands())
+ Args.push_back(COp);
+ B.SetInsertPoint(I);
+ auto *CI =
+ B.CreateIntrinsic(Intrinsic::spv_const_composite, {ResTy}, {Args});
+ Worklist.push(CI);
+ I->replaceUsesOfWith(Op, CI);
+ KeepInst = true;
+ AggrConsts[CI] = AggrConst;
+ AggrConstTypes[CI] = deduceNestedTypeHelper(AggrConst);
}
}
if (!KeepInst)
@@ -1054,8 +1061,8 @@ void SPIRVEmitIntrinsics::processGlobalValue(GlobalVariable &GV,
// by llvm IR general logic.
deduceElementTypeHelper(&GV);
Constant *Init = GV.getInitializer();
- Type *Ty = isAggrToReplace(Init) ? B.getInt32Ty() : Init->getType();
- Constant *Const = isAggrToReplace(Init) ? B.getInt32(1) : Init;
+ Type *Ty = isAggrConstForceInt32(Init) ? B.getInt32Ty() : Init->getType();
+ Constant *Const = isAggrConstForceInt32(Init) ? B.getInt32(1) : Init;
auto *InitInst = B.CreateIntrinsic(Intrinsic::spv_init_global,
{GV.getType(), Ty}, {&GV, Const});
InitInst->setArgOperand(1, Init);
@@ -1132,11 +1139,11 @@ void SPIRVEmitIntrinsics::processInstrAfterVisit(Instruction *I,
if (II && II->getIntrinsicID() == Intrinsic::spv_const_composite &&
TrackConstants) {
B.SetInsertPoint(I->getNextNode());
- Type *Ty = B.getInt32Ty();
auto t = AggrConsts.find(I);
assert(t != AggrConsts.end());
- auto *NewOp = buildIntrWithMD(Intrinsic::spv_track_constant, {Ty, Ty},
- t->second, I, {}, B);
+ auto *NewOp =
+ buildIntrWithMD(Intrinsic::spv_track_constant,
+ {II->getType(), II->getType()}, t->second, I, {}, B);
I->replaceAllUsesWith(NewOp);
NewOp->setArgOperand(0, I);
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/builtin_alloca.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/builtin_alloca.ll
new file mode 100644
index 0000000000000..4d6173e5b7232
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/builtin_alloca.ll
@@ -0,0 +1,48 @@
+; The goal of the test is to:
+; 1) check that composite constants of ConstantVector type preserve their
+; type and can be successfully used further in LLVM intrinsic functions;
+; 2) demonstrate that a call to __builtin_alloca() maps to instructions
+; from SPV_INTEL_variable_length_array when this extension is available.
+
+; Test LLVM IR is an artificial example, but it's similar to what can be
+; generated by DPC++ compiler from the code snippet:
+; ...
+; size_t Sz = ...;
+; queue Q;
+; Q.submit([&](sycl::handler &CGH) {
+; ...
+; CGH.single_task([=](sycl::kernel_handler KH) SYCL_ESIMD_KERNEL {
+; int *PrivateArray = (int *)__builtin_alloca(sizeof(int) * Sz);
+; ...
+; simd<int, 8> InitVec(100, 10);
+; InitVec.copy_to(PrivateArray);
+; ...
+; });
+; }).wait();
+; ...
+
+; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown %s -o %t.spvt 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR
+
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_variable_length_array %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_variable_length_array %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-ERROR: LLVM ERROR: array allocation: this instruction requires the following SPIR-V extension: SPV_INTEL_variable_length_array
+
+; CHECK-SPIRV: Capability VariableLengthArrayINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_variable_length_array"
+; CHECK-SPIRV: OpVariableLengthArrayINTEL %[[#]] %[[#]]
+
+define spir_kernel void @foo(i64 %_arg_sz) {
+entry:
+ %sz = shl i64 %_arg_sz, 2
+ %p1 = alloca i8, i64 %sz, align 8
+ %p4 = addrspacecast ptr %p1 to ptr addrspace(4)
+ %i = ptrtoint ptr addrspace(4) %p4 to i64
+ %splat_ins = insertelement <8 x i64> poison, i64 %i, i64 0
+ %splat_v = shufflevector <8 x i64> %splat_ins, <8 x i64> poison, <8 x i32> zeroinitializer
+ %sum_r = add <8 x i64> %splat_v, <i64 0, i64 4, i64 8, i64 12, i64 16, i64 20, i64 24, i64 28>
+ call void @llvm.genx.svm.scatter.v8i1.v8i64.v8i32(<8 x i1> <i1 true, i1 true, i1 true, i1 true, i1 true, i1 true, i1 true, i1 true>, i32 0, <8 x i64> %sum_r, <8 x i32> <i32 100, i32 110, i32 120, i32 130, i32 140, i32 150, i32 160, i32 170>)
+ ret void
+}
+
+declare void @llvm.genx.svm.scatter.v8i1.v8i64.v8i32(<8 x i1>, i32, <8 x i64>, <8 x i32>)
|
| @@ -27,7 +27,7 @@ let TargetPrefix = "spv" in { | |||
| def int_spv_insertv : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_any_ty, llvm_vararg_ty]>; | |||
| def int_spv_extractelt : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_anyint_ty]>; | |||
| def int_spv_insertelt : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_any_ty, llvm_anyint_ty]>; | |||
| def int_spv_const_composite : Intrinsic<[llvm_i32_ty], [llvm_vararg_ty]>; | |||
| def int_spv_const_composite : Intrinsic<[llvm_any_ty], [llvm_vararg_ty]>; | |||
There was a problem hiding this comment.
Please also change the description in SPIRVUsage file. It might make sense to add a comment somewhere in IntrinsicsSPIRV.td which will signal that any modifications/new intrinsics need to be also documented in SPIRVUsage.
There was a problem hiding this comment.
Indeed, thank you!
michalpaszkowski
approved these changes
May 17, 2024
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 generation of argument types of internal intrinsic functions
spv_const_compositeandspv_track_constant, so that composite constants of ConstantVector type preserve their correct type in transformation passes and can be successfully used further by LLVM intrinsic functions.The added test case serves two purposes: it is to check the above mentioned fix and to demonstrate that a call to __builtin_alloca() maps to instructions from SPV_INTEL_variable_length_array when this extension is available.