From 6fba979ac7a6f5ff5a27bfef091c6762bd0ba18a Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Thu, 15 Apr 2021 19:17:53 +0300 Subject: [PATCH 01/13] [sycl-post-link] Add property with the default values of specialization constants This patch introduces a new property which contains all the defaults for spec constants. It is done by setting of metadata with the default values: ``` ; Compilation line: ; sycl-post-link -spec-const=default -S llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll -o property.table !5 = !{i32 42} !6 = !{%struct.ComposConst { i32 1, double 2.000000e+00, %struct.myConst { i32 13, float 0x4020666660000000 } }} ``` Property set looks this way: ``` [SYCL/specialization constants default values] all=2|gAAAAAAAAAgKAAAA ``` --- llvm/include/llvm/Support/PropertySetIO.h | 2 + llvm/lib/Support/PropertySetIO.cpp | 1 + llvm/tools/sycl-post-link/SpecConstants.cpp | 44 +++++++++++++++++++- llvm/tools/sycl-post-link/SpecConstants.h | 2 +- llvm/tools/sycl-post-link/sycl-post-link.cpp | 16 +++++++ sycl/include/CL/sycl/detail/pi.h | 2 + 6 files changed, 64 insertions(+), 3 deletions(-) diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index e571a3679c1f7..cb6c70cc63c72 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -185,6 +185,8 @@ class PropertySetRegistry { "SYCL/specialization constants"; static constexpr char SYCL_COMPOSITE_SPECIALIZATION_CONSTANTS[] = "SYCL/composite specialization constants"; + static constexpr char SYCL_SPEC_CONSTANTS_DEFAULT_VALUES[] = + "SYCL/specialization constants default values"; static constexpr char SYCL_DEVICELIB_REQ_MASK[] = "SYCL/devicelib req mask"; static constexpr char SYCL_KERNEL_PARAM_OPT_INFO[] = "SYCL/kernel param opt"; static constexpr char SYCL_MISC_PROP[] = "SYCL/misc properties"; diff --git a/llvm/lib/Support/PropertySetIO.cpp b/llvm/lib/Support/PropertySetIO.cpp index f583c91088762..c7acf2c1c6ab8 100644 --- a/llvm/lib/Support/PropertySetIO.cpp +++ b/llvm/lib/Support/PropertySetIO.cpp @@ -195,6 +195,7 @@ void PropertyValue::copy(const PropertyValue &P) { constexpr char PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS[]; constexpr char PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK[]; +constexpr char PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES[]; constexpr char PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO[]; constexpr char PropertySetRegistry::SYCL_COMPOSITE_SPECIALIZATION_CONSTANTS[]; constexpr char PropertySetRegistry::SYCL_MISC_PROP[]; diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index ac1a505dd095e..47def197ce72c 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -48,6 +48,9 @@ constexpr char SPIRV_GET_SPEC_CONST_COMPOSITE[] = // associated information) encountered in the module constexpr char SPEC_CONST_MD_STRING[] = "sycl.specialization-constants"; +constexpr char SPEC_CONST_DEFAULT_VAL_MD_STRING[] = + "SYCL_SPEC_CONST_DEFAULT_VAL"; + void AssertRelease(bool Cond, const char *Msg) { if (!Cond) report_fatal_error((Twine("SpecConstants.cpp: ") + Msg).str().c_str()); @@ -215,6 +218,17 @@ std::string mangleFuncItanium(StringRef BaseName, const FunctionType *FT) { return Res; } +void setSpecConstDefaultValueMetadata(Instruction *I, StringRef SymID, + Value *Default) { + LLVMContext &Ctx = I->getContext(); + SmallVector MDOperands; + MDOperands.push_back(MDString::get(Ctx, SymID)); + MDOperands.push_back(ConstantAsMetadata::get(cast(Default))); + + MDNode *Entry = MDNode::get(Ctx, MDOperands); + I->setMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING, Entry); +} + /// Recursively iterates over a composite type in order to collect information /// about its scalar elements. void collectCompositeElementsInfoRecursive( @@ -543,7 +557,6 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, // specialization constants is passed as a 3rd argument of intrinsic. Value *RTBuffer = IsComposite ? CI->getArgOperand(3) : CI->getArgOperand(2); - // Add the string literal to a "spec const string literal ID" -> // "offset" map, uniquing the integer offsets if this is new // literal. @@ -583,6 +596,13 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, GEP, PointerType::get(SCTy, GEP->getAddressSpace()), "bc", CI); Replacement = new LoadInst(SCTy, BitCast, "load", CI); + + Value *GlobVar = (IsComposite ? CI->getOperand(2) : CI->getOperand(1)) + ->stripPointerCasts(); + Value *DefaultValue = + cast(GlobVar)->getInitializer()->getOperand(0); + + setSpecConstDefaultValueMetadata(Load, SymID, DefaultValue); } else { // Replace the intrinsic with default C++ value for the spec constant // type. @@ -624,7 +644,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, } bool SpecConstantsPass::collectSpecConstantMetadata(Module &M, - SpecIDMapTy &IDMap) { + SpecIDMapTy &IDMap, std::vector *v) { NamedMDNode *MD = M.getOrInsertNamedMetadata(SPEC_CONST_MD_STRING); if (!MD) return false; @@ -636,6 +656,26 @@ bool SpecConstantsPass::collectSpecConstantMetadata(Module &M, return static_cast(C->getUniqueInteger().getZExtValue()); }; + // const MDNode *N = I.getMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING); + // if (N) { + // const auto *MDSym = cast(N->getOperand(0)); + // const auto *MDInt = cast(N->getOperand(1)); + + // const auto *Constant = MDInt->getValue(); + // if (auto IntConst = dyn_cast(Constant)) { + // auto Val = IntConst->getValue().getZExtValue(); + + // char *a_begin = reinterpret_cast(&Val); + // copy_n(a_begin, IntConst->getType()->getScalarSizeInBits() / CHAR_BIT, + // back_inserter(*v)); + + // } else if (auto FPConst = dyn_cast(Constant)) { + // auto Val = FPConst->getValue(); + // // How to get correct float/double value? + // } + // // else if Composite type OR use the handling above in recursive + // // collection of struct scalars + for (const auto *Node : MD->operands()) { StringRef ID = cast(Node->getOperand(0).get())->getString(); assert((Node->getNumOperands() - 1) % 3 == 0 && diff --git a/llvm/tools/sycl-post-link/SpecConstants.h b/llvm/tools/sycl-post-link/SpecConstants.h index 505b37c3c97a7..abe1faa69a0db 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.h +++ b/llvm/tools/sycl-post-link/SpecConstants.h @@ -60,7 +60,7 @@ class SpecConstantsPass : public llvm::PassInfoMixin { // Searches given module for occurrences of specialization constant-specific // metadata and builds "spec constant name" -> vector<"spec constant int ID"> // map - static bool collectSpecConstantMetadata(llvm::Module &M, SpecIDMapTy &IDMap); + static bool collectSpecConstantMetadata(llvm::Module &M, SpecIDMapTy &IDMap, std::vector *v = nullptr); private: bool SetValAtRT; diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index ce3f85185a7c9..7993ef4dca924 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -421,6 +421,22 @@ static string_vector saveDeviceImageProperty( TmpSpecIDMap); } } + if (ImgPSInfo.DoSpecConst) { //} && ImgPSInfo.SetSpecConstAtRT) { + if (ImgPSInfo.SpecConstsMet) { + // extract spec constant maps per each module + SpecIDMapTy TmpSpecIDMap; + std::vector vec; + SpecConstantsPass::collectSpecConstantMetadata(*ResultModules[I].get(), + TmpSpecIDMap, &vec); + using SpecDefaultMapTy = std::map>; + SpecDefaultMapTy TmpSpecDefaultMap; + + TmpSpecDefaultMap["all"] = vec; + PropSet.add( + llvm::util::PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES, + TmpSpecDefaultMap); + } + } if (ImgPSInfo.EmitKernelParamInfo) { // extract kernel parameter optimization info per module ModuleAnalysisManager MAM; diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index a6cdadf664310..8c6a480ef28d0 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -673,6 +673,8 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4; /// PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS defined in /// PropertySetIO.h #define __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP "SYCL/specialization constants" +#define __SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP \ + "SYCL/specialization constants default values" /// PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK defined in PropertySetIO.h #define __SYCL_PI_PROPERTY_SET_DEVICELIB_REQ_MASK "SYCL/devicelib req mask" /// PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO defined in PropertySetIO.h From 45db4f145c1356b2769d60ec520258988a8e7706 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Fri, 30 Apr 2021 19:31:41 +0300 Subject: [PATCH 02/13] Support floats and composite types --- llvm/tools/sycl-post-link/SpecConstants.cpp | 133 ++++++++++++++----- llvm/tools/sycl-post-link/SpecConstants.h | 8 +- llvm/tools/sycl-post-link/sycl-post-link.cpp | 17 +-- 3 files changed, 112 insertions(+), 46 deletions(-) diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index 47def197ce72c..d2e7e4942b2cd 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -47,9 +47,10 @@ constexpr char SPIRV_GET_SPEC_CONST_COMPOSITE[] = // Name of the metadata which holds a list of all specialization constants (with // associated information) encountered in the module constexpr char SPEC_CONST_MD_STRING[] = "sycl.specialization-constants"; - +// Name of the metadata which holds a default value list of all specialization +// constants encountered in the module constexpr char SPEC_CONST_DEFAULT_VAL_MD_STRING[] = - "SYCL_SPEC_CONST_DEFAULT_VAL"; + "sycl.specialization-constants-default-values"; void AssertRelease(bool Cond, const char *Msg) { if (!Cond) @@ -218,15 +219,16 @@ std::string mangleFuncItanium(StringRef BaseName, const FunctionType *FT) { return Res; } -void setSpecConstDefaultValueMetadata(Instruction *I, StringRef SymID, - Value *Default) { +MDNode *generateSpecConstDefaultValueMetadata(Instruction *I, StringRef SymID, + Value *Default) { LLVMContext &Ctx = I->getContext(); SmallVector MDOperands; - MDOperands.push_back(MDString::get(Ctx, SymID)); + // MDOperands.push_back(MDString::get(Ctx, SymID)); MDOperands.push_back(ConstantAsMetadata::get(cast(Default))); - MDNode *Entry = MDNode::get(Ctx, MDOperands); - I->setMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING, Entry); + // MDNode *Entry = MDNode::get(Ctx, MDOperands); + return MDNode::get(Ctx, MDOperands); + // I->setMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING, Entry); } /// Recursively iterates over a composite type in order to collect information @@ -279,6 +281,76 @@ void collectCompositeElementsInfoRecursive( } } +/// Recursively iterates over a composite type in order to collect information +/// about its scalar elements. +void collectCompositeElementsDefaultValuesRecursive( + const Module &M, Constant *C, unsigned &Index, unsigned &Offset, + std::vector &DefaultValues) { + Type *Ty = C->getType(); + if (auto *ArrTy = dyn_cast(Ty)) { + for (size_t I = 0; I < ArrTy->getNumElements(); ++I) { + Constant *El = cast(C->getOperand(I)); + // TODO: this is a spot for potential optimization: for arrays we could + // just make a single recursive call here and use it to populate + // DefaultValues in a loop. + collectCompositeElementsDefaultValuesRecursive(M, El, Index, Offset, + DefaultValues); + } + } else if (auto *StructTy = dyn_cast(Ty)) { + const StructLayout *SL = M.getDataLayout().getStructLayout(StructTy); + for (size_t I = 0, E = StructTy->getNumElements(); I < E; ++I) { + Constant *El = cast(C->getOperand(I)); + // When handling elements of a structure, we do not use manually + // calculated offsets (which are sum of sizes of all previously + // encountered elements), but instead rely on data provided for us by + // DataLayout, because the structure can be unpacked, i.e. padded in + // order to ensure particular alignment of its elements. + unsigned LocalOffset = Offset + SL->getElementOffset(I); + + // If there was some alignment, fill the data between values with zeros. + while (LocalOffset != DefaultValues.size()) + DefaultValues.push_back(0); + + collectCompositeElementsDefaultValuesRecursive(M, El, Index, LocalOffset, + DefaultValues); + } + // Update "global" offset according to the total size of a handled struct + // type. + Offset += SL->getSizeInBytes(); + } else if (auto *VecTy = dyn_cast(Ty)) { + for (size_t I = 0; I < VecTy->getNumElements(); ++I) { + Constant *El = cast(C->getOperand(I)); + // TODO: this is a spot for potential optimization: for vectors we could + // just make a single recursive call here and use it to populate + // DefaultValues in a loop. + collectCompositeElementsDefaultValuesRecursive(M, El, Index, Offset, + DefaultValues); + } + } else { // Assume that we encountered some scalar element + int NumBytes = Ty->getScalarSizeInBits() / CHAR_BIT; + char *CharPtr; + + if (auto IntConst = dyn_cast(C)) { + auto Val = IntConst->getValue().getZExtValue(); + CharPtr = reinterpret_cast(&Val); + } else if (auto FPConst = dyn_cast(C)) { + auto Val = FPConst->getValue(); + + if (NumBytes == 4) { + float v = Val.convertToFloat(); + CharPtr = reinterpret_cast(&v); + } else if (NumBytes == 8) { + double v = Val.convertToDouble(); + CharPtr = reinterpret_cast(&v); + } + } + std::copy_n(CharPtr, NumBytes, std::back_inserter(DefaultValues)); + Index++; + Offset += Ty->getPrimitiveSizeInBits() / 8 + + (Ty->getPrimitiveSizeInBits() % 8 != 0); + } +} + MDNode *generateSpecConstantMetadata(const Module &M, StringRef SymbolicID, Type *SCTy, ArrayRef IDs, bool IsNativeSpecConstant) { @@ -449,6 +521,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, StringMap> IDMap; StringMap OffsetMap; MapVector SCMetadata; + SmallVector DefaultsMetadata; // Iterate through all declarations of instances of function template // template T __sycl_get*SpecConstantValue(const char *ID) @@ -595,14 +668,17 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, BitCastInst *BitCast = new BitCastInst( GEP, PointerType::get(SCTy, GEP->getAddressSpace()), "bc", CI); - Replacement = new LoadInst(SCTy, BitCast, "load", CI); + Instruction *Inst = new LoadInst(SCTy, BitCast, "load", CI); + Replacement = Inst; Value *GlobVar = (IsComposite ? CI->getOperand(2) : CI->getOperand(1)) ->stripPointerCasts(); Value *DefaultValue = cast(GlobVar)->getInitializer()->getOperand(0); - setSpecConstDefaultValueMetadata(Load, SymID, DefaultValue); + if (IsNewSpecConstant) + DefaultsMetadata.push_back(generateSpecConstDefaultValueMetadata( + Inst, SymID, DefaultValue)); } else { // Replace the intrinsic with default C++ value for the spec constant // type. @@ -640,15 +716,24 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, for (const auto &P : SCMetadata) MD->addOperand(P.second); + NamedMDNode *MDDefaults = + M.getOrInsertNamedMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING); + for (const auto &P : DefaultsMetadata) + MDDefaults->addOperand(P); + return IRModified ? PreservedAnalyses::none() : PreservedAnalyses::all(); } -bool SpecConstantsPass::collectSpecConstantMetadata(Module &M, - SpecIDMapTy &IDMap, std::vector *v) { +bool SpecConstantsPass::collectSpecConstantMetadata( + Module &M, SpecIDMapTy &IDMap, std::vector &DefaultValues) { NamedMDNode *MD = M.getOrInsertNamedMetadata(SPEC_CONST_MD_STRING); if (!MD) return false; + NamedMDNode *N = M.getOrInsertNamedMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING); + if (!N) + return false; + auto ExtractIntegerFromMDNodeOperand = [=](const MDNode *N, unsigned OpNo) -> unsigned { Constant *C = @@ -656,26 +741,6 @@ bool SpecConstantsPass::collectSpecConstantMetadata(Module &M, return static_cast(C->getUniqueInteger().getZExtValue()); }; - // const MDNode *N = I.getMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING); - // if (N) { - // const auto *MDSym = cast(N->getOperand(0)); - // const auto *MDInt = cast(N->getOperand(1)); - - // const auto *Constant = MDInt->getValue(); - // if (auto IntConst = dyn_cast(Constant)) { - // auto Val = IntConst->getValue().getZExtValue(); - - // char *a_begin = reinterpret_cast(&Val); - // copy_n(a_begin, IntConst->getType()->getScalarSizeInBits() / CHAR_BIT, - // back_inserter(*v)); - - // } else if (auto FPConst = dyn_cast(Constant)) { - // auto Val = FPConst->getValue(); - // // How to get correct float/double value? - // } - // // else if Composite type OR use the handling above in recursive - // // collection of struct scalars - for (const auto *Node : MD->operands()) { StringRef ID = cast(Node->getOperand(0).get())->getString(); assert((Node->getNumOperands() - 1) % 3 == 0 && @@ -689,6 +754,12 @@ bool SpecConstantsPass::collectSpecConstantMetadata(Module &M, IDMap[ID] = Descs; } + unsigned Index = 0, Offset = 0; + for (const auto *Node : N->operands()) { + auto *Constant = cast(Node->getOperand(0))->getValue(); + collectCompositeElementsDefaultValuesRecursive(M, Constant, Index, Offset, + DefaultValues); + } return true; } diff --git a/llvm/tools/sycl-post-link/SpecConstants.h b/llvm/tools/sycl-post-link/SpecConstants.h index abe1faa69a0db..29dc8356e39a9 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.h +++ b/llvm/tools/sycl-post-link/SpecConstants.h @@ -58,9 +58,11 @@ class SpecConstantsPass : public llvm::PassInfoMixin { llvm::ModuleAnalysisManager &MAM); // Searches given module for occurrences of specialization constant-specific - // metadata and builds "spec constant name" -> vector<"spec constant int ID"> - // map - static bool collectSpecConstantMetadata(llvm::Module &M, SpecIDMapTy &IDMap, std::vector *v = nullptr); + // metadata and builds: + // 1. "spec constant name" -> vector<"spec constant int ID"> map + // 2. vector of default values for every spec constant + static bool collectSpecConstantMetadata(llvm::Module &M, SpecIDMapTy &IDMap, + std::vector &DefaultValues); private: bool SetValAtRT; diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 7993ef4dca924..cf6d84da8eb5a 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -414,24 +414,17 @@ static string_vector saveDeviceImageProperty( if (ImgPSInfo.SpecConstsMet) { // extract spec constant maps per each module SpecIDMapTy TmpSpecIDMap; - SpecConstantsPass::collectSpecConstantMetadata(*ResultModules[I].get(), - TmpSpecIDMap); + std::vector DefaultValues; + SpecConstantsPass::collectSpecConstantMetadata( + *ResultModules[I].get(), TmpSpecIDMap, DefaultValues); PropSet.add( llvm::util::PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS, TmpSpecIDMap); - } - } - if (ImgPSInfo.DoSpecConst) { //} && ImgPSInfo.SetSpecConstAtRT) { - if (ImgPSInfo.SpecConstsMet) { - // extract spec constant maps per each module - SpecIDMapTy TmpSpecIDMap; - std::vector vec; - SpecConstantsPass::collectSpecConstantMetadata(*ResultModules[I].get(), - TmpSpecIDMap, &vec); + using SpecDefaultMapTy = std::map>; SpecDefaultMapTy TmpSpecDefaultMap; - TmpSpecDefaultMap["all"] = vec; + TmpSpecDefaultMap["all"] = DefaultValues; PropSet.add( llvm::util::PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES, TmpSpecDefaultMap); From eb594cf1ff8605bffcfa1419933533e32a76e9ac Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Wed, 5 May 2021 18:18:18 +0300 Subject: [PATCH 03/13] Apply suggestions from code review --- llvm/include/llvm/Support/PropertySetIO.h | 7 ++ .../spec-constants/SYCL-2020.ll | 8 ++ llvm/tools/sycl-post-link/SpecConstants.cpp | 104 ++++++++---------- llvm/tools/sycl-post-link/SpecConstants.h | 13 ++- llvm/tools/sycl-post-link/sycl-post-link.cpp | 22 ++-- .../basic_tests/SYCL-2020-spec-constants.cpp | 8 +- 6 files changed, 89 insertions(+), 73 deletions(-) diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index cb6c70cc63c72..546531479c61b 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -203,6 +203,13 @@ class PropertySetRegistry { PropSet.insert(std::make_pair(Prop.first, PropertyValue(Prop.second))); } + // Function to add a property to a given category (property set name). + template + void add(StringRef Category, StringRef PropName, T &PropVal) { + auto &PropSet = PropSetMap[Category]; + PropSet.insert(std::make_pair(PropName, PropertyValue(PropVal))); + } + // Parses and creates a property set registry. static Expected> read(const MemoryBuffer *Buf); diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll index 6661b7f0b861c..48fe9349170d3 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll @@ -132,6 +132,9 @@ attributes #3 = { nounwind } ; CHECK: !sycl.specialization-constants = !{![[#ID0:]], ![[#ID1:]], ![[#ID2:]], ![[#ID3:]]} ; +; CHECK-DEF: !sycl.specialization-constants-default-values = !{![[#ID4:]], ![[#ID5:]], ![[#ID6:]], ![[#ID7:]]} +; CHECK-RT-NOT: !sycl.specialization-constants-default-values +; ; CHECK: ![[#ID0]] = !{!"_ZTS14name_generatorIL_Z9id_doubleEE", i32 0, i32 0, i32 8} ; CHECK: ![[#ID1]] = !{!"_ZTS14name_generatorIL_Z6id_intEE", i32 1, i32 0, i32 4} ; @@ -148,3 +151,8 @@ attributes #3 = { nounwind } ; CHECK-RT-SAME: i32 [[#SCID7]], i32 4, i32 4, ; CHECK-RT-SAME: i32 [[#SCID8]], i32 8, i32 4, ; CHECK-RT-SAME: i32 [[#SCID9]], i32 16, i32 8} +; +; CHECK-DEF: ![[#ID4]] = !{double 3.140000e+00} +; CHECK-DEF: ![[#ID5]] = !{i32 42} +; CHECK-DEF: ![[#ID6]] = !{%struct.ComposConst { i32 1, double 2.000000e+00, %struct.myConst { i32 13, float 0x4020666660000000 } }} +; CHECK-DEF: ![[#ID7]] = !{%struct.ComposConst2 { i8 1, %struct.myConst { i32 52, float 0x40479999A0000000 }, double 2.000000e+00 }} diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index d2e7e4942b2cd..c5badb25b0281 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -222,13 +222,7 @@ std::string mangleFuncItanium(StringRef BaseName, const FunctionType *FT) { MDNode *generateSpecConstDefaultValueMetadata(Instruction *I, StringRef SymID, Value *Default) { LLVMContext &Ctx = I->getContext(); - SmallVector MDOperands; - // MDOperands.push_back(MDString::get(Ctx, SymID)); - MDOperands.push_back(ConstantAsMetadata::get(cast(Default))); - - // MDNode *Entry = MDNode::get(Ctx, MDOperands); - return MDNode::get(Ctx, MDOperands); - // I->setMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING, Entry); + return MDNode::get(Ctx, ConstantAsMetadata::get(cast(Default))); } /// Recursively iterates over a composite type in order to collect information @@ -282,18 +276,15 @@ void collectCompositeElementsInfoRecursive( } /// Recursively iterates over a composite type in order to collect information -/// about its scalar elements. +/// about default values of its scalar elements. void collectCompositeElementsDefaultValuesRecursive( - const Module &M, Constant *C, unsigned &Index, unsigned &Offset, + const Module &M, Constant *C, unsigned &Offset, std::vector &DefaultValues) { Type *Ty = C->getType(); if (auto *ArrTy = dyn_cast(Ty)) { for (size_t I = 0; I < ArrTy->getNumElements(); ++I) { Constant *El = cast(C->getOperand(I)); - // TODO: this is a spot for potential optimization: for arrays we could - // just make a single recursive call here and use it to populate - // DefaultValues in a loop. - collectCompositeElementsDefaultValuesRecursive(M, El, Index, Offset, + collectCompositeElementsDefaultValuesRecursive(M, El, Offset, DefaultValues); } } else if (auto *StructTy = dyn_cast(Ty)) { @@ -311,7 +302,7 @@ void collectCompositeElementsDefaultValuesRecursive( while (LocalOffset != DefaultValues.size()) DefaultValues.push_back(0); - collectCompositeElementsDefaultValuesRecursive(M, El, Index, LocalOffset, + collectCompositeElementsDefaultValuesRecursive(M, El, LocalOffset, DefaultValues); } // Update "global" offset according to the total size of a handled struct @@ -320,14 +311,12 @@ void collectCompositeElementsDefaultValuesRecursive( } else if (auto *VecTy = dyn_cast(Ty)) { for (size_t I = 0; I < VecTy->getNumElements(); ++I) { Constant *El = cast(C->getOperand(I)); - // TODO: this is a spot for potential optimization: for vectors we could - // just make a single recursive call here and use it to populate - // DefaultValues in a loop. - collectCompositeElementsDefaultValuesRecursive(M, El, Index, Offset, + collectCompositeElementsDefaultValuesRecursive(M, El, Offset, DefaultValues); } } else { // Assume that we encountered some scalar element - int NumBytes = Ty->getScalarSizeInBits() / CHAR_BIT; + int NumBytes = Ty->getScalarSizeInBits() / CHAR_BIT + + (Ty->getScalarSizeInBits() % 8 != 0); char *CharPtr; if (auto IntConst = dyn_cast(C)) { @@ -345,9 +334,7 @@ void collectCompositeElementsDefaultValuesRecursive( } } std::copy_n(CharPtr, NumBytes, std::back_inserter(DefaultValues)); - Index++; - Offset += Ty->getPrimitiveSizeInBits() / 8 + - (Ty->getPrimitiveSizeInBits() % 8 != 0); + Offset += NumBytes; } } @@ -572,6 +559,22 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, } StringRef SymID = getStringLiteralArg(CI, NameArgNo, DelInsts); Value *Replacement = nullptr; + Constant *DefaultValue = nullptr; + if (Is2020Intrinsic) { + // For SYCL 2020, there is a mechanism to specify the default value. + // It is stored as an initializer of a global variable referenced by + // the second argument of the intrinsic. + auto *GV = dyn_cast( + CI->getArgOperand(NameArgNo + 1)->stripPointerCasts()); + if (GV) { + auto *Initializer = GV->getInitializer(); + assert(isa(Initializer) && + "expected specialization_id instance"); + // specialization_id structure contains a single field which is the + // default value of corresponding specialization constant. + DefaultValue = Initializer->getAggregateElement(0u); + } + } if (SetValAtRT) { // 2. Spec constant value will be set at run time - then add the literal @@ -587,23 +590,6 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, IDs.push_back(NextID); } - Constant *DefaultValue = nullptr; - if (Is2020Intrinsic) { - // For SYCL 2020, there is a mechanism to specify the default value. - // It is stored as an initializer of a global variable referenced by - // the second argument of the intrinsic. - auto *GV = dyn_cast( - CI->getArgOperand(NameArgNo + 1)->stripPointerCasts()); - if (GV) { - auto *Initializer = GV->getInitializer(); - assert(isa(Initializer) && - "expected specialization_id instance"); - // specialization_id structure contains a single field which is the - // default value of corresponding specialization constant. - DefaultValue = Initializer->getAggregateElement(0u); - } - } - // 3. Transform to spirv intrinsic _Z*__spirv_SpecConstant* or // _Z*__spirv_SpecConstantComposite Replacement = emitSpecConstantRecursive(SCTy, CI, IDs, DefaultValue); @@ -671,11 +657,6 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, Instruction *Inst = new LoadInst(SCTy, BitCast, "load", CI); Replacement = Inst; - Value *GlobVar = (IsComposite ? CI->getOperand(2) : CI->getOperand(1)) - ->stripPointerCasts(); - Value *DefaultValue = - cast(GlobVar)->getInitializer()->getOperand(0); - if (IsNewSpecConstant) DefaultsMetadata.push_back(generateSpecConstDefaultValueMetadata( Inst, SymID, DefaultValue)); @@ -716,24 +697,23 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, for (const auto &P : SCMetadata) MD->addOperand(P.second); - NamedMDNode *MDDefaults = - M.getOrInsertNamedMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING); - for (const auto &P : DefaultsMetadata) - MDDefaults->addOperand(P); + // Emit default values metadata only in native (default) spec constants mode. + if (!SetValAtRT) { + NamedMDNode *MDDefaults = + M.getOrInsertNamedMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING); + for (const auto &P : DefaultsMetadata) + MDDefaults->addOperand(P); + } return IRModified ? PreservedAnalyses::none() : PreservedAnalyses::all(); } -bool SpecConstantsPass::collectSpecConstantMetadata( - Module &M, SpecIDMapTy &IDMap, std::vector &DefaultValues) { - NamedMDNode *MD = M.getOrInsertNamedMetadata(SPEC_CONST_MD_STRING); +bool SpecConstantsPass::collectSpecConstantMetadata(Module &M, + SpecIDMapTy &IDMap) { + NamedMDNode *MD = M.getNamedMetadata(SPEC_CONST_MD_STRING); if (!MD) return false; - NamedMDNode *N = M.getOrInsertNamedMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING); - if (!N) - return false; - auto ExtractIntegerFromMDNodeOperand = [=](const MDNode *N, unsigned OpNo) -> unsigned { Constant *C = @@ -754,10 +734,20 @@ bool SpecConstantsPass::collectSpecConstantMetadata( IDMap[ID] = Descs; } - unsigned Index = 0, Offset = 0; + + return true; +} + +bool SpecConstantsPass::collectSpecConstantDefaultValuesMetadata( + Module &M, std::vector &DefaultValues) { + NamedMDNode *N = M.getNamedMetadata(SPEC_CONST_DEFAULT_VAL_MD_STRING); + if (!N) + return false; + + unsigned Offset = 0; for (const auto *Node : N->operands()) { auto *Constant = cast(Node->getOperand(0))->getValue(); - collectCompositeElementsDefaultValuesRecursive(M, Constant, Index, Offset, + collectCompositeElementsDefaultValuesRecursive(M, Constant, Offset, DefaultValues); } diff --git a/llvm/tools/sycl-post-link/SpecConstants.h b/llvm/tools/sycl-post-link/SpecConstants.h index 29dc8356e39a9..0481b9d9985e9 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.h +++ b/llvm/tools/sycl-post-link/SpecConstants.h @@ -58,11 +58,14 @@ class SpecConstantsPass : public llvm::PassInfoMixin { llvm::ModuleAnalysisManager &MAM); // Searches given module for occurrences of specialization constant-specific - // metadata and builds: - // 1. "spec constant name" -> vector<"spec constant int ID"> map - // 2. vector of default values for every spec constant - static bool collectSpecConstantMetadata(llvm::Module &M, SpecIDMapTy &IDMap, - std::vector &DefaultValues); + // metadata and builds "spec constant name" -> vector<"spec constant int ID"> + // map + static bool collectSpecConstantMetadata(llvm::Module &M, SpecIDMapTy &IDMap); + // Searches given module for occurrences of specialization constant-specific + // metadata and builds vector of default values for every spec constant. + static bool + collectSpecConstantDefaultValuesMetadata(llvm::Module &M, + std::vector &DefaultValues); private: bool SetValAtRT; diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index cf6d84da8eb5a..ac3697ee59f2c 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -414,20 +414,24 @@ static string_vector saveDeviceImageProperty( if (ImgPSInfo.SpecConstsMet) { // extract spec constant maps per each module SpecIDMapTy TmpSpecIDMap; - std::vector DefaultValues; - SpecConstantsPass::collectSpecConstantMetadata( - *ResultModules[I].get(), TmpSpecIDMap, DefaultValues); + + SpecConstantsPass::collectSpecConstantMetadata(*ResultModules[I].get(), + TmpSpecIDMap); PropSet.add( llvm::util::PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS, TmpSpecIDMap); - using SpecDefaultMapTy = std::map>; - SpecDefaultMapTy TmpSpecDefaultMap; + // Add property with the default values of spec constants only in native + // (default) mode. + if (!ImgPSInfo.SetSpecConstAtRT) { + std::vector DefaultValues; + SpecConstantsPass::collectSpecConstantDefaultValuesMetadata( + *ResultModules[I].get(), DefaultValues); - TmpSpecDefaultMap["all"] = DefaultValues; - PropSet.add( - llvm::util::PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES, - TmpSpecDefaultMap); + PropSet.add(llvm::util::PropertySetRegistry:: + SYCL_SPEC_CONSTANTS_DEFAULT_VALUES, + "all", DefaultValues); + } } } if (ImgPSInfo.EmitKernelParamInfo) { diff --git a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp index 8bfb21a2f11d3..b74fd0d8d21eb 100644 --- a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp +++ b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp @@ -1,8 +1,8 @@ // RUN: %clangxx -fsycl -fsycl-device-only -c -o %t.bc %s // RUN: sycl-post-link %t.bc -spec-const=rt -o %t-split1.txt -// RUN: cat %t-split1_0.prop | FileCheck %s +// RUN: cat %t-split1_0.prop | FileCheck %s -check-prefixes=CHECK,CHECK-RT // RUN: sycl-post-link %t.bc -spec-const=default -o %t-split2.txt -// RUN: cat %t-split2_0.prop | FileCheck %s +// RUN: cat %t-split2_0.prop | FileCheck %s -check-prefixes=CHECK,CHECK-DEF // RUN: llvm-spirv -o %t-split1_0.spv -spirv-max-version=1.1 -spirv-ext=+all %t-split1_0.bc // RUN: llvm-spirv -o %t-split2_0.spv -spirv-max-version=1.1 -spirv-ext=+all %t-split2_0.bc // @@ -100,3 +100,7 @@ int main() { // CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL9uint32_idEEE=2| // CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL9uint64_idEEE=2| // FIXME: check line for half constant + +// CHECK-RT-NOT: [SYCL/specialization constants default values] +// CHECK-DEF: [SYCL/specialization constants default values] +// CHECK-DEF: all=2| \ No newline at end of file From c632c64e91ccee057b1c36b1cf647eee5edc0ee7 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Wed, 5 May 2021 18:23:50 +0300 Subject: [PATCH 04/13] Minor line changes --- llvm/tools/sycl-post-link/SpecConstants.cpp | 1 + llvm/tools/sycl-post-link/sycl-post-link.cpp | 2 -- sycl/test/basic_tests/SYCL-2020-spec-constants.cpp | 2 +- 3 files changed, 2 insertions(+), 3 deletions(-) diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index c5badb25b0281..e48370f5e0abe 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -616,6 +616,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, // specialization constants is passed as a 3rd argument of intrinsic. Value *RTBuffer = IsComposite ? CI->getArgOperand(3) : CI->getArgOperand(2); + // Add the string literal to a "spec const string literal ID" -> // "offset" map, uniquing the integer offsets if this is new // literal. diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index ac3697ee59f2c..1d4fd914bd082 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -414,7 +414,6 @@ static string_vector saveDeviceImageProperty( if (ImgPSInfo.SpecConstsMet) { // extract spec constant maps per each module SpecIDMapTy TmpSpecIDMap; - SpecConstantsPass::collectSpecConstantMetadata(*ResultModules[I].get(), TmpSpecIDMap); PropSet.add( @@ -427,7 +426,6 @@ static string_vector saveDeviceImageProperty( std::vector DefaultValues; SpecConstantsPass::collectSpecConstantDefaultValuesMetadata( *ResultModules[I].get(), DefaultValues); - PropSet.add(llvm::util::PropertySetRegistry:: SYCL_SPEC_CONSTANTS_DEFAULT_VALUES, "all", DefaultValues); diff --git a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp index b74fd0d8d21eb..1db6b8f2587da 100644 --- a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp +++ b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp @@ -103,4 +103,4 @@ int main() { // CHECK-RT-NOT: [SYCL/specialization constants default values] // CHECK-DEF: [SYCL/specialization constants default values] -// CHECK-DEF: all=2| \ No newline at end of file +// CHECK-DEF: all=2| From 4ca0b7647dc7de72d691c72a6be14073f53444ee Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Wed, 5 May 2021 19:39:25 +0300 Subject: [PATCH 05/13] Do no generate metadata if default value ptr is null --- llvm/tools/sycl-post-link/SpecConstants.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index e8ff4053d373b..8a7f3abc07212 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -661,7 +661,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, Instruction *Inst = new LoadInst(SCTy, BitCast, "load", CI); Replacement = Inst; - if (IsNewSpecConstant) + if (IsNewSpecConstant && DefaultValue) DefaultsMetadata.push_back(generateSpecConstDefaultValueMetadata( Inst, SymID, DefaultValue)); } else { From d6fa1f963480fd50c23bcd28de4231772e28c092 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Fri, 7 May 2021 19:59:59 +0300 Subject: [PATCH 06/13] Minor changes after CR --- llvm/include/llvm/Support/PropertySetIO.h | 2 +- llvm/tools/sycl-post-link/SpecConstants.cpp | 4 ++++ 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index 546531479c61b..23095f9e87ac0 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -205,7 +205,7 @@ class PropertySetRegistry { // Function to add a property to a given category (property set name). template - void add(StringRef Category, StringRef PropName, T &PropVal) { + void add(StringRef Category, StringRef PropName, const T &PropVal) { auto &PropSet = PropSetMap[Category]; PropSet.insert(std::make_pair(PropName, PropertyValue(PropVal))); } diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index 8a7f3abc07212..29a12a5248693 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -277,6 +277,9 @@ void collectCompositeElementsInfoRecursive( /// Recursively iterates over a composite type in order to collect information /// about default values of its scalar elements. +/// TODO: processing of composite spec constants here is similar to +/// collectCompositeElementsInfoRecursive. Possible place for improvement - +/// factor out the common code, e.g. using visitor pattern. void collectCompositeElementsDefaultValuesRecursive( const Module &M, Constant *C, unsigned &Offset, std::vector &DefaultValues) { @@ -567,6 +570,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, // the second argument of the intrinsic. auto *GV = dyn_cast( CI->getArgOperand(NameArgNo + 1)->stripPointerCasts()); + // Go through global variable if the second argument was not null. if (GV) { assert(GV->hasInitializer() && "expected initializer"); auto *Initializer = GV->getInitializer(); From 08de0d4c903f0c0407ef10695639b25ad8d4f69c Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Tue, 11 May 2021 13:08:43 +0300 Subject: [PATCH 07/13] Apply comments --- llvm/include/llvm/Support/PropertySetIO.h | 2 -- llvm/lib/Support/PropertySetIO.cpp | 1 - llvm/tools/sycl-post-link/SpecConstants.cpp | 14 ++++++-------- 3 files changed, 6 insertions(+), 11 deletions(-) diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index 23095f9e87ac0..2a6e1fbacc914 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -183,8 +183,6 @@ class PropertySetRegistry { // Specific property category names used by tools. static constexpr char SYCL_SPECIALIZATION_CONSTANTS[] = "SYCL/specialization constants"; - static constexpr char SYCL_COMPOSITE_SPECIALIZATION_CONSTANTS[] = - "SYCL/composite specialization constants"; static constexpr char SYCL_SPEC_CONSTANTS_DEFAULT_VALUES[] = "SYCL/specialization constants default values"; static constexpr char SYCL_DEVICELIB_REQ_MASK[] = "SYCL/devicelib req mask"; diff --git a/llvm/lib/Support/PropertySetIO.cpp b/llvm/lib/Support/PropertySetIO.cpp index c7acf2c1c6ab8..97016364e0d0a 100644 --- a/llvm/lib/Support/PropertySetIO.cpp +++ b/llvm/lib/Support/PropertySetIO.cpp @@ -197,7 +197,6 @@ constexpr char PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS[]; constexpr char PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK[]; constexpr char PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES[]; constexpr char PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO[]; -constexpr char PropertySetRegistry::SYCL_COMPOSITE_SPECIALIZATION_CONSTANTS[]; constexpr char PropertySetRegistry::SYCL_MISC_PROP[]; } // namespace util diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index 29a12a5248693..9544ef6fbd47c 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -219,9 +219,8 @@ std::string mangleFuncItanium(StringRef BaseName, const FunctionType *FT) { return Res; } -MDNode *generateSpecConstDefaultValueMetadata(Instruction *I, StringRef SymID, - Value *Default) { - LLVMContext &Ctx = I->getContext(); +MDNode *generateSpecConstDefaultValueMetadata(StringRef SymID, Value *Default) { + LLVMContext &Ctx = Default->getContext(); return MDNode::get(Ctx, ConstantAsMetadata::get(cast(Default))); } @@ -511,7 +510,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, StringMap> IDMap; StringMap OffsetMap; MapVector SCMetadata; - SmallVector DefaultsMetadata; + SmallVector DefaultsMetadata; // Iterate through all declarations of instances of function template // template T __sycl_get*SpecConstantValue(const char *ID) @@ -662,12 +661,11 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, BitCastInst *BitCast = new BitCastInst( GEP, PointerType::get(SCTy, GEP->getAddressSpace()), "bc", CI); - Instruction *Inst = new LoadInst(SCTy, BitCast, "load", CI); - Replacement = Inst; + Replacement = new LoadInst(SCTy, BitCast, "load", CI); if (IsNewSpecConstant && DefaultValue) - DefaultsMetadata.push_back(generateSpecConstDefaultValueMetadata( - Inst, SymID, DefaultValue)); + DefaultsMetadata.push_back( + generateSpecConstDefaultValueMetadata(SymID, DefaultValue)); } else { // Replace the intrinsic with default C++ value for the spec constant // type. From 6d258433c9cbb463e9e5632a167a76679af6b7fd Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Tue, 11 May 2021 13:55:19 +0300 Subject: [PATCH 08/13] remove changes from SYCL header pi.h --- sycl/include/CL/sycl/detail/pi.h | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 8c6a480ef28d0..a6cdadf664310 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -673,8 +673,6 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4; /// PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS defined in /// PropertySetIO.h #define __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP "SYCL/specialization constants" -#define __SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP \ - "SYCL/specialization constants default values" /// PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK defined in PropertySetIO.h #define __SYCL_PI_PROPERTY_SET_DEVICELIB_REQ_MASK "SYCL/devicelib req mask" /// PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO defined in PropertySetIO.h From 774051bf903024440348a650241a1f108e4c6b19 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Fri, 21 May 2021 15:03:40 +0300 Subject: [PATCH 09/13] Fix test failure for non-native spec consts --- sycl/source/detail/device_image_impl.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index bbfc7f24ca822..b68edae9ebe1a 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -256,7 +256,7 @@ class device_image_impl { // supposed to be called from c'tor. MSpecConstSymMap[std::string{SCName}].push_back( SpecConstDescT{/*ID*/ It[0], /*CompositeOffset*/ It[1], - /*Size*/ It[2], BlobOffset, HasDefaultValues}); + /*Size*/ It[2], BlobOffset}); BlobOffset += /*Size*/ It[2]; It += NumElements; } From 02293fbbe2104b1669968e82e15419a0eb70b772 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Mon, 24 May 2021 14:18:10 +0300 Subject: [PATCH 10/13] Temporarily disable spec constants unit tests until they support AOT binaries DefaultValues.DefaultValuesAreSet and DefaultValues.DefaultValuesAreOverriden tests only support JIT which is not the case for the default values. Once AOT binaries support is added, the tests can be enabled back. --- sycl/unittests/spec_constants/DefaultValues.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/unittests/spec_constants/DefaultValues.cpp b/sycl/unittests/spec_constants/DefaultValues.cpp index 655547f83d31d..a5605124a533a 100644 --- a/sycl/unittests/spec_constants/DefaultValues.cpp +++ b/sycl/unittests/spec_constants/DefaultValues.cpp @@ -205,7 +205,7 @@ static sycl::unittest::PiImage generateDefaultImage() { sycl::unittest::PiImage Img = generateDefaultImage(); sycl::unittest::PiImageArray ImgArray{Img}; -TEST(DefaultValues, DefaultValuesAreSet) { +TEST(DefaultValues, DISABLED_DefaultValuesAreSet) { sycl::platform Plt{sycl::default_selector()}; if (Plt.is_host()) { std::cerr << "Test is not supported on host, skipping\n"; @@ -238,7 +238,7 @@ TEST(DefaultValues, DefaultValuesAreSet) { EXPECT_EQ(SpecConstVal1, 8); } -TEST(DefaultValues, DefaultValuesAreOverriden) { +TEST(DefaultValues, DISABLED_DefaultValuesAreOverriden) { sycl::platform Plt{sycl::default_selector()}; if (Plt.is_host()) { std::cerr << "Test is not supported on host, skipping\n"; From c4e538a4eddb9837e9f87ea6c1b03e1dc515bbb3 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Tue, 25 May 2021 18:47:25 +0300 Subject: [PATCH 11/13] Minor improvement --- llvm/include/llvm/Support/PropertySetIO.h | 17 ++++++----------- 1 file changed, 6 insertions(+), 11 deletions(-) diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index 709d5e5b3effa..6ca3dc8aecdbc 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -43,6 +43,7 @@ #include #include #include +#include namespace llvm { namespace util { @@ -191,18 +192,12 @@ class PropertySetRegistry { // Function for bulk addition of an entire property set under given category // (property set name). - template - void add(StringRef Category, const std::map &Props) { - assert(PropSetMap.find(Category) == PropSetMap.end() && - "category already added"); - auto &PropSet = PropSetMap[Category]; + template void add(StringRef Category, const MapTy &Props) { + using KeyTy = typename MapTy::value_type::first_type; + static_assert(std::is_same::type, + llvm::StringRef>::value, + "wrong key type"); - for (const auto &Prop : Props) - PropSet.insert(std::make_pair(Prop.first, PropertyValue(Prop.second))); - } - - template - void add(StringRef Category, const MapVector &Props) { assert(PropSetMap.find(Category) == PropSetMap.end() && "category already added"); auto &PropSet = PropSetMap[Category]; From 9b0305bd34d1ec7fe6aa75e25189ee833a0fd663 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Tue, 25 May 2021 18:54:54 +0300 Subject: [PATCH 12/13] Remove RT changes from this PR Moving to https://github.com/intel/llvm/pull/3813 --- sycl/source/detail/device_image_impl.hpp | 2 +- sycl/unittests/spec_constants/DefaultValues.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index b68edae9ebe1a..bbfc7f24ca822 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -256,7 +256,7 @@ class device_image_impl { // supposed to be called from c'tor. MSpecConstSymMap[std::string{SCName}].push_back( SpecConstDescT{/*ID*/ It[0], /*CompositeOffset*/ It[1], - /*Size*/ It[2], BlobOffset}); + /*Size*/ It[2], BlobOffset, HasDefaultValues}); BlobOffset += /*Size*/ It[2]; It += NumElements; } diff --git a/sycl/unittests/spec_constants/DefaultValues.cpp b/sycl/unittests/spec_constants/DefaultValues.cpp index a5605124a533a..655547f83d31d 100644 --- a/sycl/unittests/spec_constants/DefaultValues.cpp +++ b/sycl/unittests/spec_constants/DefaultValues.cpp @@ -205,7 +205,7 @@ static sycl::unittest::PiImage generateDefaultImage() { sycl::unittest::PiImage Img = generateDefaultImage(); sycl::unittest::PiImageArray ImgArray{Img}; -TEST(DefaultValues, DISABLED_DefaultValuesAreSet) { +TEST(DefaultValues, DefaultValuesAreSet) { sycl::platform Plt{sycl::default_selector()}; if (Plt.is_host()) { std::cerr << "Test is not supported on host, skipping\n"; @@ -238,7 +238,7 @@ TEST(DefaultValues, DISABLED_DefaultValuesAreSet) { EXPECT_EQ(SpecConstVal1, 8); } -TEST(DefaultValues, DISABLED_DefaultValuesAreOverriden) { +TEST(DefaultValues, DefaultValuesAreOverriden) { sycl::platform Plt{sycl::default_selector()}; if (Plt.is_host()) { std::cerr << "Test is not supported on host, skipping\n"; From c3c798970df293911a6ba1d78bb7fbe2a1a82305 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Wed, 26 May 2021 12:15:02 +0300 Subject: [PATCH 13/13] Fix typo --- llvm/include/llvm/Support/PropertySetIO.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index 6ca3dc8aecdbc..cd7688d10f2db 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -194,7 +194,7 @@ class PropertySetRegistry { // (property set name). template void add(StringRef Category, const MapTy &Props) { using KeyTy = typename MapTy::value_type::first_type; - static_assert(std::is_same::type, + static_assert(std::is_same::type, llvm::StringRef>::value, "wrong key type");