[llvm] [SPIR-V] Fix types of internal intrinsic functions and add a test case for __builtin_alloca() (PR #92265)
via llvm-commits
llvm-commits at lists.llvm.org
Wed May 15 07:02:54 PDT 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-spir-v
@llvm/pr-subscribers-llvm-ir
Author: Vyacheslav Levytskyy (VyacheslavLevytskyy)
<details>
<summary>Changes</summary>
This PR generation of argument types of internal intrinsic functions `spv_const_composite` and `spv_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.
---
Full diff: https://github.com/llvm/llvm-project/pull/92265.diff
3 Files Affected:
- (modified) llvm/include/llvm/IR/IntrinsicsSPIRV.td (+1-1)
- (modified) llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp (+43-36)
- (added) llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/builtin_alloca.ll (+48)
``````````diff
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>)
``````````
</details>
https://github.com/llvm/llvm-project/pull/92265
More information about the llvm-commits
mailing list