diff --git a/llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll b/llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll index b88f06e7a5dc7..f73c5ffeb7242 100644 --- a/llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll +++ b/llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll @@ -1,6 +1,10 @@ ; RUN: sycl-post-link -spec-const=default --ir-output-only %s -S -o - \ ; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue ; +; This test checks that composite specialization constants can be correctly +; initialized by sycl-post-link tool for AOT use-case (default initialization +; should be used according to the type of constant) +; ; TODO: consider adding a test case with vector type: the pass itself already ; supports this, but at the moment, sycl::vec type is not a POD type, which ; means we can't have it within a spec constant, i.e. we can't generate LLVM IR diff --git a/llvm/test/tools/sycl-post-link/composite-spec-constant.ll b/llvm/test/tools/sycl-post-link/composite-spec-constant.ll index 59e7fb4cee9c1..5e0d8b1508439 100644 --- a/llvm/test/tools/sycl-post-link/composite-spec-constant.ll +++ b/llvm/test/tools/sycl-post-link/composite-spec-constant.ll @@ -83,7 +83,7 @@ attributes #4 = { convergent } !0 = !{i32 1, !"wchar_size", i32 4} !1 = !{i32 1, i32 2} !2 = !{i32 4, i32 100000} -!3 = !{!"clang version 12.0.0 (/data/github.com/intel/llvm/clang 56ee5b054b5a1f2f703722fc414fcb05af18b40a)"} +!3 = !{!"clang version 12.0.0 "} !4 = !{i32 -1, i32 -1, i32 -1, i32 -1} !5 = !{i64 0, i64 16, !6, i64 16, i64 8, !6} !6 = !{!7, !7, i64 0} diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index b2d9b7e3d71d5..e27481d64a635 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -109,8 +109,6 @@ StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo, return Res; } -// TODO support spec constant types other than integer or -// floating-point. Value *getDefaultCPPValue(Type *T) { if (T->isIntegerTy()) return Constant::getIntegerValue(T, APInt(T->getScalarSizeInBits(), 0)); @@ -167,16 +165,13 @@ std::string manglePrimitiveType(const Type *T) { // llvm-spirv doesn't care about the mangling and the only intent here is to // make sure that we won't encounter redefinition error when we proceed two // spec constants with different types. - if (T->isStructTy()) { + if (T->isStructTy()) return T->getStructName().str(); - } - if (T->isArrayTy()) { + if (T->isArrayTy()) return "A" + manglePrimitiveType(T->getArrayElementType()); - } - if (auto *VecTy = dyn_cast(T)) { + if (auto *VecTy = dyn_cast(T)) return "Dv" + std::to_string(VecTy->getNumElements()) + "_" + manglePrimitiveType(VecTy->getElementType()); - } llvm_unreachable("unsupported spec const type"); return ""; } @@ -196,10 +191,9 @@ void setSpecConstSymIDMetadata(Instruction *I, StringRef SymID, LLVMContext &Ctx = I->getContext(); SmallVector MDOperands; MDOperands.push_back(MDString::get(Ctx, SymID)); - for (unsigned ID : IntIDs) { + for (unsigned ID : IntIDs) MDOperands.push_back( ConstantAsMetadata::get(ConstantInt::get(Ctx, APInt(32, ID)))); - } MDNode *Entry = MDNode::get(Ctx, MDOperands); I->setMetadata(SPEC_CONST_SYM_ID_MD_STRING, Entry); } @@ -218,7 +212,7 @@ getScalarSpecConstMetadata(const Instruction *I) { void collectCompositeElementsInfoRecursive( const Type *Ty, unsigned &Index, unsigned &Offset, - std::vector &Result) { + std::vector &Result) { if (auto *ArrTy = dyn_cast(Ty)) { for (size_t I = 0; I < ArrTy->getNumElements(); ++I) { // TODO: this is a spot for potential optimization: for arrays we could @@ -233,14 +227,14 @@ void collectCompositeElementsInfoRecursive( } } else if (auto *VecTy = dyn_cast(Ty)) { for (size_t I = 0; I < VecTy->getNumElements(); ++I) { - // TODO: this is a spot for potential optimization: for arrays we could + // TODO: this is a spot for potential optimization: for vectors we could // just make a single recursive call here and use it to populate Result // in a loop. collectCompositeElementsInfoRecursive(VecTy->getElementType(), Index, Offset, Result); } } else { // Assume that we encountered some scalar element - CompositeSpecConstDescriptor Desc; + CompositeSpecConstElementDescriptor Desc; Desc.ID = 0; // To be filled later Desc.Offset = Offset; Desc.Size = Ty->getPrimitiveSizeInBits() / 8; @@ -249,14 +243,16 @@ void collectCompositeElementsInfoRecursive( } } -std::pair> +std::pair> getCompositeSpecConstMetadata(const Instruction *I) { const MDNode *N = I->getMetadata(SPEC_CONST_SYM_ID_MD_STRING); if (!N) - return std::make_pair("", std::vector{}); + return std::make_pair("", + std::vector{}); const auto *MDSym = cast(N->getOperand(0)); - std::vector Result(N->getNumOperands() - 1); + std::vector Result(N->getNumOperands() - + 1); unsigned Index = 0, Offset = 0; collectCompositeElementsInfoRecursive(I->getType(), Index, Offset, Result); @@ -269,6 +265,21 @@ getCompositeSpecConstMetadata(const Instruction *I) { return std::make_pair(MDSym->getString(), Result); } +Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName, + ArrayRef Args, Instruction *InsertBefore) { + SmallVector ArgTys(Args.size()); + for (unsigned I = 0; I < Args.size(); ++I) { + ArgTys[I] = Args[I]->getType(); + } + auto *FT = FunctionType::get(RetTy, ArgTys, false /*isVarArg*/); + std::string FunctionName = mangleFuncItanium(BaseFunctionName, FT); + Module *M = InsertBefore->getFunction()->getParent(); + FunctionCallee FC = M->getOrInsertFunction(FunctionName, FT); + assert(FC.getCallee() && "SPIRV intrinsic creation failed"); + auto *Call = CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore); + return Call; +} + Instruction *emitSpecConstant(unsigned NumericID, Type *Ty, Instruction *InsertBefore) { Function *F = InsertBefore->getFunction(); @@ -279,37 +290,17 @@ Instruction *emitSpecConstant(unsigned NumericID, Type *Ty, Value *Def = getDefaultCPPValue(Ty); // ... Now replace the call with SPIRV intrinsic version. Value *Args[] = {ID, Def}; - constexpr size_t NArgs = sizeof(Args) / sizeof(Args[0]); - Type *ArgTys[NArgs] = {nullptr}; - for (unsigned int I = 0; I < NArgs; ++I) - ArgTys[I] = Args[I]->getType(); - FunctionType *FT = FunctionType::get(Ty, ArgTys, false /*isVarArg*/); - Module *M = F->getParent(); - std::string SPIRVName = mangleFuncItanium(SPIRV_GET_SPEC_CONST_VAL, FT); - FunctionCallee FC = M->getOrInsertFunction(SPIRVName, FT); - assert(FC.getCallee() && "SPIRV intrinsic creation failed"); - CallInst *SpecConstant = - CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore); - return SpecConstant; + return emitCall(Ty, SPIRV_GET_SPEC_CONST_VAL, Args, InsertBefore); } Instruction *emitSpecConstantComposite(Type *Ty, ArrayRef Elements, Instruction *InsertBefore) { - SmallVector ArgTys(Elements.size()); SmallVector Args(Elements.size()); for (unsigned I = 0; I < Elements.size(); ++I) { - ArgTys[I] = Elements[I]->getType(); Args[I] = cast(Elements[I]); } - auto *FT = FunctionType::get(Ty, ArgTys, false /*isVarArg*/); - Module *M = InsertBefore->getFunction()->getParent(); - std::string SPIRVName = mangleFuncItanium(SPIRV_GET_SPEC_CONST_COMPOSITE, FT); - FunctionCallee FC = M->getOrInsertFunction(SPIRVName, FT); - assert(FC.getCallee() && "SPIRV intrinsic creation failed"); - CallInst *SpecConstant = - CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore); - return SpecConstant; + return emitCall(Ty, SPIRV_GET_SPEC_CONST_COMPOSITE, Args, InsertBefore); } Instruction * @@ -416,13 +407,14 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, NextID += GeneratedIDs.size(); } - if (IsComposite) + if (IsComposite) { // __sycl_getCompositeSpecConstant returns through argument, so, the // only thing we need to do here is to store into a memory pointed by // that argument new StoreInst(SPIRVCall, CI->getArgOperand(0), CI); - else + } else { CI->replaceAllUsesWith(SPIRVCall); + } // Mark the instruction with list for later // recollection by collectSpecConstantMetadata method. @@ -442,13 +434,14 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, // 2a. Spec constant must be resolved at compile time - just replace // the intrinsic with default C++ value for the spec constant type. Value *Default = getDefaultCPPValue(SCTy); - if (IsComposite) + if (IsComposite) { // __sycl_getCompositeSpecConstant returns through argument, so, the // only thing we need to do here is to store into a memory pointed by // that argument new StoreInst(Default, CI->getArgOperand(0), CI); - else + } else { CI->replaceAllUsesWith(Default); + } } for (auto *I : DelInsts) { diff --git a/llvm/tools/sycl-post-link/SpecConstants.h b/llvm/tools/sycl-post-link/SpecConstants.h index 8b7ceeb7d30a9..dbb071b2cbd4f 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.h +++ b/llvm/tools/sycl-post-link/SpecConstants.h @@ -25,13 +25,23 @@ using namespace llvm; using ScalarSpecIDMapTy = std::map; -struct CompositeSpecConstDescriptor { +// Represents an element of a composite speciailization constant - at SYCL RT +// level composite specialization constants are being represented as a single +// byte-array, while at SPIR-V level they are represented by a number of scalar +// specialization constants. +struct CompositeSpecConstElementDescriptor { + // Encodes ID of a scalar specialization constants which is a leaf of some + // composite specialization constant. unsigned ID; + // Encodes offset from the beginning of composite, where scalar resides, i.e. + // location of the scalar value within a byte-array containing the whole + // composite specialization constant. unsigned Offset; + // Encodes size of scalar specialization constant. unsigned Size; }; using CompositeSpecIDMapTy = - std::map>; + std::map>; class SpecConstantsPass : public PassInfoMixin { public: