From 49eda0e4fa020458e4194ac9d9f56c3be3f75b20 Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Thu, 26 Sep 2024 12:01:20 -0700 Subject: [PATCH 1/3] [SYCL] Fix ESIMD split detection in module properties computation Signed-off-by: Sarnie, Nick --- .../include/llvm/SYCLLowerIR/ModuleSplitter.h | 4 ++ .../SYCLLowerIR/ComputeModuleRuntimeInfo.cpp | 32 +++++++++----- llvm/lib/SYCLLowerIR/ModuleSplitter.cpp | 19 ++++++++ .../tools/sycl-post-link/sycl-esimd/assert.ll | 43 +++++++++++++++++++ .../sycl-esimd/basic-esimd-lower.ll | 2 +- llvm/tools/sycl-post-link/sycl-post-link.cpp | 7 +-- 6 files changed, 90 insertions(+), 17 deletions(-) create mode 100644 llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll diff --git a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h index 0da3706ad3626..c692adb940a57 100644 --- a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h @@ -37,6 +37,8 @@ class OptionCategory; namespace module_split { +constexpr char SYCL_ESIMD_SPLIT_MD_NAME[] = "sycl-split-status"; + extern cl::OptionCategory &getModuleSplitCategory(); enum IRSplitMode { @@ -221,6 +223,8 @@ class ModuleDesc { return *Reqs; } + void saveSplitInformationAsMetadata(); + #ifndef NDEBUG void verifyESIMDProperty() const; void dump() const; diff --git a/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp b/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp index e7d6877b89a9a..d0c22be9d3eae 100644 --- a/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp +++ b/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp @@ -28,6 +28,21 @@ constexpr int DebugModuleProps = 0; #endif namespace llvm::sycl { +namespace { +module_split::SyclEsimdSplitStatus +getSYCLESIMDSplitStatusFromMetadata(const Module &M) { + auto *SplitMD = M.getNamedMetadata(module_split::SYCL_ESIMD_SPLIT_MD_NAME); + assert(SplitMD && "Unexpected metadata"); + auto *MDOp = SplitMD->getOperand(0); + assert(MDOp && "Unexpected metadata operand"); + const auto &MDConst = MDOp->getOperand(0); + auto *MDVal = mdconst::dyn_extract_or_null(MDConst); + uint8_t Val = MDVal->getZExtValue(); + assert(Val < 3 && "Unexpected value for split metadata"); + auto AsEnum = static_cast(Val); + return AsEnum; +} +} // namespace bool isModuleUsingAsan(const Module &M) { for (const auto &F : M) { @@ -305,16 +320,11 @@ PropSetRegTy computeModuleProperties(const Module &M, GV.getName()); } } - bool SeenESIMDFunction = false; - bool SeenSYCLFunction = false; - for (const auto &F : M) { - if (llvm::module_split::isESIMDFunction(F)) - SeenESIMDFunction = true; - else if (utils::isSYCLExternalFunction(&F) && - !F.getName().starts_with("__itt")) - SeenSYCLFunction = true; - } - if (SeenESIMDFunction && !SeenSYCLFunction) + + module_split::SyclEsimdSplitStatus SplitType = + getSYCLESIMDSplitStatusFromMetadata(M); + + if (SplitType == module_split::SyclEsimdSplitStatus::ESIMD_ONLY) PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "isEsimdImage", true); { StringRef RegAllocModeAttr = "sycl-register-alloc-mode"; @@ -359,7 +369,7 @@ PropSetRegTy computeModuleProperties(const Module &M, // 'if' below essentially preserves the behavior (presumably mistakenly) // implemented in intel/llvm#8763: ignore 'optLevel' property for images which // were produced my merge after ESIMD split - if (!SeenESIMDFunction || !SeenSYCLFunction) { + if (SplitType != module_split::SyclEsimdSplitStatus::SYCL_AND_ESIMD) { // Handle sycl-optlevel property int OptLevel = -1; for (const Function *F : EntryPoints) { diff --git a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp index 6068ce58f414f..904424f93dae6 100644 --- a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp +++ b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp @@ -26,6 +26,7 @@ #include "llvm/SYCLLowerIR/DeviceGlobals.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" #include "llvm/SYCLLowerIR/SYCLUtils.h" +#include "llvm/SYCLLowerIR/SpecConstants.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Error.h" #include "llvm/Support/FileSystem.h" @@ -798,6 +799,23 @@ void ModuleDesc::dump() const { } #endif // NDEBUG +void ModuleDesc::saveSplitInformationAsMetadata() { + // Add metadata to the module so we can identify what kind of SYCL/ESIMD split + // later. + auto *SplitMD = M->getOrInsertNamedMetadata(SYCL_ESIMD_SPLIT_MD_NAME); + auto *SplitMDOp = MDNode::get( + M->getContext(), ConstantAsMetadata::get(ConstantInt::get( + Type::getInt8Ty(M->getContext()), + static_cast(EntryPoints.Props.HasESIMD)))); + SplitMD->addOperand(SplitMDOp); + + // Add metadata to the module so we can identify it as the default value spec + // constants split later. + if (isSpecConstantDefault()) + M->getOrInsertNamedMetadata( + SpecConstantsPass::SPEC_CONST_DEFAULT_VAL_MODULE_MD_STRING); +} + void EntryPointGroup::saveNames(std::vector &Dest) const { Dest.reserve(Dest.size() + Functions.size()); std::transform(Functions.begin(), Functions.end(), @@ -1291,6 +1309,7 @@ static Expected saveModuleDesc(ModuleDesc &MD, std::string Prefix, bool OutputAssembly) { SplitModule SM; Prefix += OutputAssembly ? ".ll" : ".bc"; + MD.saveSplitInformationAsMetadata(); Error E = saveModuleIRInFile(MD.getModule(), Prefix, OutputAssembly); if (E) return E; diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll new file mode 100644 index 0000000000000..8a391f8e7c83c --- /dev/null +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll @@ -0,0 +1,43 @@ +; RUN: sycl-post-link -properties -split-esimd -S < %s -o %t.table +; RUN: FileCheck %s -input-file=%t_esimd_0.prop + +; Verify we mark a image with an ESIMD kernel with the isEsimdImage property + +; CHECK: isEsimdImage=1|1 + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" +target triple = "spir64-unknown-unknown" + +%"struct.sycl::_V1::detail::AssertHappened" = type { i32, [257 x i8], [257 x i8], [129 x i8], i32, i64, i64, i64, i64, i64, i64 } +%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" } +%"class.sycl::_V1::detail::array" = type { [1 x i64] } + +@.str = private unnamed_addr addrspace(1) constant [10 x i8] c"Id != 400\00", align 1 +@.str.1 = private unnamed_addr addrspace(1) constant [8 x i8] c"foo.cpp\00", align 1 +@__PRETTY_FUNCTION__ = private unnamed_addr addrspace(1) constant [56 x i8] c"auto main()::(anonymous class)::operator()(id<1>) const\00", align 1 +@SPIR_AssertHappenedMem = linkonce_odr dso_local addrspace(1) global %"struct.sycl::_V1::detail::AssertHappened" zeroinitializer, align 8 + +declare void @llvm.assume(i1 noundef) #2 + +define weak_odr dso_local spir_kernel void @esimd_kernel() local_unnamed_addr #0 !sycl_explicit_simd !0 { +entry: + tail call spir_func void @__assert_fail(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)), ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)), i32 noundef 13, ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @__PRETTY_FUNCTION__ to ptr addrspace(4))) #12 + ret void +} + +define weak dso_local spir_func void @__assert_fail(ptr addrspace(4) noundef %expr, ptr addrspace(4) noundef %file, i32 noundef %line, ptr addrspace(4) noundef %func) #1 { +entry: + tail call spir_func void @__devicelib_assert_fail(ptr addrspace(4) noundef %expr, ptr addrspace(4) noundef %file, i32 noundef %line, ptr addrspace(4) noundef %func) #1 + ret void +} + +define weak dso_local spir_func void @__devicelib_assert_fail(ptr addrspace(4) noundef %expr, ptr addrspace(4) noundef %file, i32 noundef %line, ptr addrspace(4) noundef %func) #2 { +entry: + ret void +} + +attributes #0 = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="foo.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" } +attributes #1 = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="/localdisk2/nsarnie/llvm_assert2/libdevice/fallback-cassert.cpp" "sycl-optlevel"="2" } +attributes #2 = { convergent nounwind } + +!0 = !{} diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll index 3773fd1048ba2..23e09d6bad4ed 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll @@ -55,7 +55,7 @@ attributes #0 = { "sycl-module-id"="a.cpp" } ; CHECK-NO-LOWERING: } ; With -O0, we only lower ESIMD code, but no other optimizations -; CHECK-O0: define dso_local spir_kernel void @ESIMD_kernel() #{{[0-9]}} !sycl_explicit_simd !3 !intel_reqd_sub_group_size !4 { +; CHECK-O0: define dso_local spir_kernel void @ESIMD_kernel() #{{[0-9]}} !sycl_explicit_simd !{{[0-9]}} !intel_reqd_sub_group_size !{{[0-9]}} { ; CHECK-O0: entry: ; CHECK-O0: %0 = load <3 x i64>, {{.*}} addrspacecast {{.*}} @__spirv_BuiltInGlobalInvocationId ; CHECK-O0: %1 = extractelement <3 x i64> %0, i64 0 diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 3ea65b2492a5c..3800c5875e44f 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -419,6 +419,7 @@ void saveModule(std::vector> &OutTables, module_split::ModuleDesc &MD, int I, StringRef IRFilename) { IrPropSymFilenameTriple BaseTriple; StringRef Suffix = getModuleSuffix(MD); + MD.saveSplitInformationAsMetadata(); if (!IRFilename.empty()) { // don't save IR, just record the filename BaseTriple.Ir = IRFilename.str(); @@ -509,10 +510,6 @@ processSpecConstantsWithDefaultValues(const module_split::ModuleDesc &MD) { assert(NewModuleDesc->Props.SpecConstsMet && "This property should be true since the presence of SpecConsts " "has been checked before the run of the pass"); - // Add metadata to the module so we can identify it as the default value split - // later. - NewModuleDesc->getModule().getOrInsertNamedMetadata( - SpecConstantsPass::SPEC_CONST_DEFAULT_VAL_MODULE_MD_STRING); NewModuleDesc->rebuildEntryPoints(); return NewModuleDesc; } @@ -791,7 +788,7 @@ processInputModule(std::unique_ptr M) { // to keep the optimizer from wrongfully removing them. llvm.compiler.used // symbols are usually removed at backend lowering, but this is handled here // for SPIR-V since SYCL compilation uses llvm-spirv, not the SPIR-V backend. - if (auto Triple = M->getTargetTriple().find("spir") != std::string::npos) + if (M->getTargetTriple().find("spir") != std::string::npos) Modified |= removeDeviceGlobalFromCompilerUsed(*M.get()); // Instrument each image scope device globals if the module has been From 1dc7bd6e50172e54dfacd649c4ba422e9417a0f1 Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Fri, 27 Sep 2024 08:37:03 -0700 Subject: [PATCH 2/3] rename md based on feedback Signed-off-by: Sarnie, Nick --- llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h index c692adb940a57..e622db50dd364 100644 --- a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h @@ -37,7 +37,7 @@ class OptionCategory; namespace module_split { -constexpr char SYCL_ESIMD_SPLIT_MD_NAME[] = "sycl-split-status"; +constexpr char SYCL_ESIMD_SPLIT_MD_NAME[] = "sycl-esimd-split-status"; extern cl::OptionCategory &getModuleSplitCategory(); From 082bd44c827a05e5037b096abbf694328d3c7920 Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Fri, 27 Sep 2024 11:21:26 -0700 Subject: [PATCH 3/3] clean up lit test attrib Signed-off-by: Sarnie, Nick --- llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll index 8a391f8e7c83c..59c159a544476 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/assert.ll @@ -37,7 +37,7 @@ entry: } attributes #0 = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="foo.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" } -attributes #1 = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="/localdisk2/nsarnie/llvm_assert2/libdevice/fallback-cassert.cpp" "sycl-optlevel"="2" } +attributes #1 = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="bar.cpp" "sycl-optlevel"="2" } attributes #2 = { convergent nounwind } !0 = !{}