diff --git a/llvm/docs/SPIRVUsage.rst b/llvm/docs/SPIRVUsage.rst index 589ee7646ce17..d27177a4541a4 100644 --- a/llvm/docs/SPIRVUsage.rst +++ b/llvm/docs/SPIRVUsage.rst @@ -298,7 +298,7 @@ SPIR-V backend, along with their descriptions and argument details. - `[Type, Type, Any Integer]` - Inserts an element into an aggregate type at a specified index. Allows for building and modifying arrays and vectors. * - `int_spv_const_composite` - - 32-bit Integer + - Type - `[Vararg]` - Constructs a composite type from given elements. Key for creating arrays, structs, and vectors from individual components. * - `int_spv_bitcast` 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>]>; 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(I) || isa(I); } -static bool isAggrToReplace(const Value *V) { - return isa(V) || isa(V) || +static bool isAggrConstForceInt32(const Value *V) { + return isa(V) || isa(V) || + isa(V) || (isa(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 Args, Value *Op, Instruction *I, - IRBuilder<> &B, std::queue &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(Op)) { - SmallVector Args(AggrC->op_begin(), AggrC->op_end()); - BuildCompositeIntrinsic(AggrC, Args, Op, I, B, Worklist, KeepInst, - *this); - } else if (auto *AggrC = dyn_cast(Op)) { + Constant *AggrConst = nullptr; + Type *ResTy = nullptr; + if (auto *COp = dyn_cast(Op)) { + AggrConst = cast(COp); + ResTy = COp->getType(); + } else if (auto *COp = dyn_cast(Op)) { + AggrConst = cast(COp); + ResTy = B.getInt32Ty(); + } else if (auto *COp = dyn_cast(Op)) { + AggrConst = cast(COp); + ResTy = B.getInt32Ty(); + } else if (auto *COp = dyn_cast(Op)) { + AggrConst = cast(COp); + ResTy = B.getInt32Ty(); + } else if (auto *COp = dyn_cast(Op)) { + if (!Op->getType()->isVectorTy()) { + AggrConst = cast(COp); + ResTy = B.getInt32Ty(); + } + } + if (AggrConst) { SmallVector 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(Op) && - !Op->getType()->isVectorTy()) { - auto *AggrC = cast(Op); - SmallVector Args(AggrC->op_begin(), AggrC->op_end()); - BuildCompositeIntrinsic(AggrC, Args, Op, I, B, Worklist, KeepInst, - *this); + if (auto *COp = dyn_cast(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 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, + call void @llvm.genx.svm.scatter.v8i1.v8i64.v8i32(<8 x i1> , i32 0, <8 x i64> %sum_r, <8 x i32> ) + ret void +} + +declare void @llvm.genx.svm.scatter.v8i1.v8i64.v8i32(<8 x i1>, i32, <8 x i64>, <8 x i32>)