From fb3bebd4df1fd23f6cc7c8bd25774eedee711146 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Mon, 20 Mar 2023 15:21:32 -0700 Subject: [PATCH 01/41] Add frontend changes Signed-off-by: Arvind Sudarsanam --- clang/lib/CodeGen/CodeGenFunction.cpp | 12 +++++++++--- .../CodeGenSYCL/sycl-add-opt-level-attrib.cpp | 18 ++++++++++++++++++ 2 files changed, 27 insertions(+), 3 deletions(-) create mode 100644 clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 82424462f47e2..8251834b3651a 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -589,11 +589,17 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD, && !FD->hasAttr()) return; - // TODO Module identifier is not reliable for this purpose since two modules - // can have the same ID, needs improvement - if (getLangOpts().SYCLIsDevice) + if (getLangOpts().SYCLIsDevice) { + // TODO Module identifier is not reliable for this purpose since two modules + // can have the same ID, needs improvement Fn->addFnAttr("sycl-module-id", Fn->getParent()->getModuleIdentifier()); + // Here, we add a function attribute 'sycl-optlevel' to store the + // optimization level. + int SYCLOptLevel = CGM.getCodeGenOpts().OptimizationLevel; + assert(SYCLOptLevel >= 0 && "Invalid optimization level!"); + Fn->addFnAttr("sycl-optlevel", std::to_string(SYCLOptLevel)); + } llvm::LLVMContext &Context = getLLVMContext(); if (FD->hasAttr() || FD->hasAttr()) diff --git a/clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp b/clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp new file mode 100644 index 0000000000000..e41a5a71ef990 --- /dev/null +++ b/clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp @@ -0,0 +1,18 @@ +// RUN: %clangxx %s -O0 -S -o %t.ll -fsycl-device-only +// RUN: FileCheck %s --input-file %t.ll -check-prefixes=CHECK-IR +// CHECK-IR: define weak_odr dso_local spir_kernel void @{{.*}}main{{.*}}sycl{{.*}}handler{{.*}}() #[[ATTR:[0-9]+]] +// CHECK-IR: attributes #[[ATTR]] = { {{.*}} "sycl-optlevel"="0" {{.*}}} + +// This test checks adding of the attribute 'sycl-optlevel' +// by the clang front-end + +#include + +int main() { + sycl::queue q; + q.submit([&](sycl::handler &h) { + h.single_task([=]() {}); + }); + return 0; +} + From d3a9e24e100ddc5864d91b497711f47b3c10177e Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 23 Mar 2023 12:36:35 -0700 Subject: [PATCH 02/41] Adding support to propagate frontend compile time options to backend Signed-off-by: Arvind Sudarsanam --- .../tools/sycl-post-link/sycl-opt-level.ll | 50 +++++++++++++++++++ llvm/tools/sycl-post-link/ModuleSplitter.cpp | 19 ++++++- llvm/tools/sycl-post-link/ModuleSplitter.h | 7 +++ llvm/tools/sycl-post-link/sycl-post-link.cpp | 2 + .../program_manager/program_manager.cpp | 26 ++++++++++ 5 files changed, 102 insertions(+), 2 deletions(-) create mode 100644 llvm/test/tools/sycl-post-link/sycl-opt-level.ll diff --git a/llvm/test/tools/sycl-post-link/sycl-opt-level.ll b/llvm/test/tools/sycl-post-link/sycl-opt-level.ll new file mode 100644 index 0000000000000..948c54332f17d --- /dev/null +++ b/llvm/test/tools/sycl-post-link/sycl-opt-level.ll @@ -0,0 +1,50 @@ +; This test checks parsing of the attribute 'sycl-optlevel' +; by the sycl-post-link-tool: +; In addition to splitting requested by user, the kernels are also split based +; on their optimization levels. +; sycl-post-link adds 'optLevel' property to the device binary + +; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: FileCheck %s -input-file=%t.table +; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-OPT-LEVEL-PROP-0 +; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-OPT-LEVEL-PROP-1 + +; CHECK: [Code|Properties|Symbols] +; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym +; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym + +; CHECK-OPT-LEVEL-PROP-0: optLevel=1|0 +; CHECK-OPT-LEVEL-PROP-1: optLevel=1|2 + +; ModuleID = 'final.bc' +source_filename = "llvm-link" +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" +target triple = "spir64-unknown-unknown" + +; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none) +define dso_local spir_func noundef i32 @_Z3fooii(i32 noundef %a, i32 noundef %b) local_unnamed_addr #0 { +entry: + %sub = sub nsw i32 %a, %b + ret i32 %sub +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define dso_local spir_func noundef i32 @_Z3booii(i32 noundef %a, i32 noundef %b) #1 { +entry: + %retval = alloca i32, align 4 + %a.addr = alloca i32, align 4 + %b.addr = alloca i32, align 4 + %retval.ascast = addrspacecast i32* %retval to i32 addrspace(4)* + %a.addr.ascast = addrspacecast i32* %a.addr to i32 addrspace(4)* + %b.addr.ascast = addrspacecast i32* %b.addr to i32 addrspace(4)* + store i32 %a, i32 addrspace(4)* %a.addr.ascast, align 4 + store i32 %b, i32 addrspace(4)* %b.addr.ascast, align 4 + %0 = load i32, i32 addrspace(4)* %a.addr.ascast, align 4 + %1 = load i32, i32 addrspace(4)* %b.addr.ascast, align 4 + %add = add nsw i32 %0, %1 + ret i32 %add +} + +attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test3.cpp" "sycl-optlevel"="2" } +attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test2.cpp" "sycl-optlevel"="0" } + diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index c571d4c25f0c4..4d4aa70dfb4c8 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -42,6 +42,7 @@ constexpr char GLOBAL_SCOPE_NAME[] = ""; constexpr char SYCL_SCOPE_NAME[] = ""; constexpr char ESIMD_SCOPE_NAME[] = ""; constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd"; +constexpr char ATTR_OPT_LEVEL[] = "sycl-optlevel"; bool hasIndirectFunctionsOrCalls(const Module &M) { for (const auto &F : M.functions()) { @@ -674,7 +675,8 @@ void ModuleDesc::dump() const { llvm::errs() << " ESIMD:" << toString(EntryPoints.Props.HasESIMD) << ", SpecConstMet:" << (Props.SpecConstsMet ? "YES" : "NO") << ", LargeGRF:" - << (EntryPoints.Props.UsesLargeGRF ? "YES" : "NO") << "\n"; + << (EntryPoints.Props.UsesLargeGRF ? "YES" : "NO") + << ", OptLevel:" << EntryPoints.getOptLevel() << "\n"; dumpEntryPoints(entries(), EntryPoints.GroupId.c_str(), 1); llvm::errs() << "}\n"; } @@ -713,6 +715,7 @@ namespace { struct UsedOptionalFeatures { SmallVector Aspects; bool UsesLargeGRF = false; + int OptLevel = -1; SmallVector ReqdWorkGroupSize; // TODO: extend this further with reqd-sub-group-size and other properties @@ -735,6 +738,11 @@ struct UsedOptionalFeatures { if (F->hasFnAttribute(::sycl::kernel_props::ATTR_LARGE_GRF)) UsesLargeGRF = true; + if (F->hasFnAttribute(ATTR_OPT_LEVEL)) + if (F->getFnAttribute(ATTR_OPT_LEVEL).getValueAsString() + .getAsInteger(10, OptLevel)) + OptLevel = -1; + if (const MDNode *MDN = F->getMetadata("reqd_work_group_size")) { size_t NumOperands = MDN->getNumOperands(); assert(NumOperands >= 1 && NumOperands <= 3 && @@ -748,10 +756,12 @@ struct UsedOptionalFeatures { llvm::hash_code AspectsHash = llvm::hash_combine_range(Aspects.begin(), Aspects.end()); llvm::hash_code LargeGRFHash = llvm::hash_value(UsesLargeGRF); + llvm::hash_code OptLevelHash = llvm::hash_value(OptLevel); llvm::hash_code ReqdWorkGroupSizeHash = llvm::hash_combine_range( ReqdWorkGroupSize.begin(), ReqdWorkGroupSize.end()); Hash = static_cast( - llvm::hash_combine(AspectsHash, LargeGRFHash, ReqdWorkGroupSizeHash)); + llvm::hash_combine(AspectsHash, LargeGRFHash, OptLevelHash, + ReqdWorkGroupSizeHash)); } std::string generateModuleName(StringRef BaseName) const { @@ -773,6 +783,9 @@ struct UsedOptionalFeatures { if (UsesLargeGRF) Ret += "-large-grf"; + if (OptLevel != -1) + Ret += "-O" + std::to_string(OptLevel); + return Ret; } @@ -869,6 +882,8 @@ getSplitterByOptionalFeatures(ModuleDesc &&MD, // Propagate LargeGRF flag to entry points group if (Features.UsesLargeGRF) MDProps.UsesLargeGRF = true; + if (Features.OptLevel != -1) + MDProps.OptLevel = Features.OptLevel; Groups.emplace_back( Features.generateModuleName(MD.getEntryPointGroup().GroupId), std::move(EntryPoints), MDProps); diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index ffc29c163e2b6..09b983fe9e5b3 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -57,6 +57,8 @@ struct EntryPointGroup { SyclEsimdSplitStatus HasESIMD = SyclEsimdSplitStatus::SYCL_AND_ESIMD; // Whether any of the EPs use large GRF mode. bool UsesLargeGRF = false; + // front-end opt level for kernel compilation + int OptLevel = -1; // Scope represented by EPs in a group EntryPointsGroupScope Scope = Scope_Global; @@ -67,6 +69,7 @@ struct EntryPointGroup { : SyclEsimdSplitStatus::SYCL_AND_ESIMD; Res.UsesLargeGRF = UsesLargeGRF || Other.UsesLargeGRF; // Scope remains global + // OptLevel is expected to be the same for both merging EPGs return Res; } }; @@ -93,6 +96,9 @@ struct EntryPointGroup { // Tells if some entry points use large GRF mode. bool isLargeGRF() const { return Props.UsesLargeGRF; } + // Returns opt level + int getOptLevel() const { return Props.OptLevel; } + void saveNames(std::vector &Dest) const; void rebuildFromNames(const std::vector &Names, const Module &M); void rebuild(const Module &M); @@ -147,6 +153,7 @@ class ModuleDesc { bool isESIMD() const { return EntryPoints.isEsimd(); } bool isSYCL() const { return EntryPoints.isSycl(); } bool isLargeGRF() const { return EntryPoints.isLargeGRF(); } + int getOptLevel() const { return EntryPoints.getOptLevel(); } const EntryPointSet &entries() const { return EntryPoints.Functions; } const EntryPointGroup &getEntryPointGroup() const { return EntryPoints; } diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index f9110752bc331..1c3f7bacec289 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -453,6 +453,8 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, } if (MD.isLargeGRF()) PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true}); + if (MD.getOptLevel() != -1) + PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"optLevel", MD.getOptLevel()}); { std::vector FuncNames = getKernelNamesUsingAssert(M); for (const StringRef &FName : FuncNames) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index a197193d35432..e60d0ed051925 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -359,6 +359,15 @@ static bool getUint32PropAsBool(const RTDeviceBinaryImage &Img, return Prop && (DeviceBinaryProperty(Prop).asUint32() != 0); } +static int getUint32PropAsInt(const RTDeviceBinaryImage &Img, + const char *PropName) { + pi_device_binary_property Prop = Img.getProperty(PropName); + if (!Prop) + return -1; + return (int)(DeviceBinaryProperty(Prop).asUint32()); +} + + static void appendCompileOptionsFromImage(std::string &CompileOpts, const RTDeviceBinaryImage &Img, const std::vector &Devs, @@ -381,6 +390,7 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, // TODO: Remove isDoubleGRF check in next ABI break bool isLargeGRF = getUint32PropAsBool(Img, "isLargeGRF") || getUint32PropAsBool(Img, "isDoubleGRF"); + int optLevel = getUint32PropAsInt(Img, "optLevel"); // The -vc-codegen option is always preserved for ESIMD kernels, regardless // of the contents SYCL_PROGRAM_COMPILE_OPTIONS environment variable. if (isEsimdImage) { @@ -400,6 +410,22 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, // metadata. CompileOpts += isEsimdImage ? "-doubleGRF" : "-ze-opt-large-register-file"; } + // Add optimization flags + if (Plugin.getBackend() == backend::ext_oneapi_level_zero) { + if (!CompileOpts.empty()) + CompileOpts += " "; + switch (optLevel) { + case 0: CompileOpts += "-ze-opt-disable"; break; + case 1: + case 2: CompileOpts += "-ze-opt-level=1"; break; + case 3: CompileOpts += "-ze-opt-level=2"; break; + } + } else if (Plugin.getBackend() == backend::opencl) { + if (!CompileOpts.empty()) + CompileOpts += " "; + if (optLevel == 0) + CompileOpts += "-cl-opt-disable"; + } if ((Plugin.getBackend() == backend::ext_oneapi_level_zero || Plugin.getBackend() == backend::opencl) && std::all_of(Devs.begin(), Devs.end(), From 395d550004f88662227c7d516af2fdfadeaa1538 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 23 Mar 2023 20:07:43 -0700 Subject: [PATCH 03/41] Minor changes Signed-off-by: Arvind Sudarsanam --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 9 +++++---- sycl/source/detail/program_manager/program_manager.cpp | 4 ++-- 2 files changed, 7 insertions(+), 6 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 4d4aa70dfb4c8..05f3171505fba 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -756,12 +756,12 @@ struct UsedOptionalFeatures { llvm::hash_code AspectsHash = llvm::hash_combine_range(Aspects.begin(), Aspects.end()); llvm::hash_code LargeGRFHash = llvm::hash_value(UsesLargeGRF); - llvm::hash_code OptLevelHash = llvm::hash_value(OptLevel); llvm::hash_code ReqdWorkGroupSizeHash = llvm::hash_combine_range( ReqdWorkGroupSize.begin(), ReqdWorkGroupSize.end()); + llvm::hash_code OptLevelHash = llvm::hash_value(OptLevel); Hash = static_cast( - llvm::hash_combine(AspectsHash, LargeGRFHash, OptLevelHash, - ReqdWorkGroupSizeHash)); + llvm::hash_combine(AspectsHash, LargeGRFHash, ReqdWorkGroupSizeHash, + OptLevelHash)); } std::string generateModuleName(StringRef BaseName) const { @@ -821,7 +821,8 @@ struct UsedOptionalFeatures { return false; } - return IsEmpty == Other.IsEmpty && UsesLargeGRF == Other.UsesLargeGRF; + return IsEmpty == Other.IsEmpty && UsesLargeGRF == Other.UsesLargeGRF && + OptLevel == Other.OptLevel; } unsigned hash() const { return static_cast(Hash); } diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index e60d0ed051925..414fd9fc52d61 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -412,7 +412,7 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, } // Add optimization flags if (Plugin.getBackend() == backend::ext_oneapi_level_zero) { - if (!CompileOpts.empty()) + if (!CompileOpts.empty() && (optLevel != -1)) CompileOpts += " "; switch (optLevel) { case 0: CompileOpts += "-ze-opt-disable"; break; @@ -421,7 +421,7 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, case 3: CompileOpts += "-ze-opt-level=2"; break; } } else if (Plugin.getBackend() == backend::opencl) { - if (!CompileOpts.empty()) + if (!CompileOpts.empty() && (optLevel == 0)) CompileOpts += " "; if (optLevel == 0) CompileOpts += "-cl-opt-disable"; From 0467da0c87b5efac562a40fc0bd3bca07723e3e4 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 23 Mar 2023 21:17:38 -0700 Subject: [PATCH 04/41] Modify order of split modules in the sycl-post-link table Signed-off-by: Arvind Sudarsanam --- .../device-code-split/per-aspect-split-1.ll | 16 ++++++++-------- .../per-reqd-wg-size-split-2.ll | 6 +++--- .../tools/sycl-post-link/sycl-esimd-large-grf.ll | 16 ++++++++-------- llvm/test/tools/sycl-post-link/sycl-large-grf.ll | 12 ++++++------ 4 files changed, 25 insertions(+), 25 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll index faec71a602ffd..687e50e0f6ffd 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll @@ -10,29 +10,29 @@ ; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ +; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M1-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \ +; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \ +; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M1-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \ +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ +; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M1-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \ +; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \ +; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M1-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \ +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll index c724ca3284909..533de091641ca 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll @@ -4,14 +4,14 @@ ; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; -; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \ +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M0-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 \ ; RUN: --implicit-check-not kernel2 ; -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M2-SYMS \ +; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel3 ; -; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M1-SYMS \ +; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M1-SYMS \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 \ ; RUN: --implicit-check-not kernel3 diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll b/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll index 413380ae70432..223c262af914f 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll @@ -9,16 +9,16 @@ ; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table -; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR -; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM -; RUN: FileCheck %s -input-file=%t_esimd_1.sym --check-prefixes CHECK-ESIMD-SYM -; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM +; RUN: FileCheck %s -input-file=%t_esimd_large_grf_1.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR +; RUN: FileCheck %s -input-file=%t_esimd_large_grf_1.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP +; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM +; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-SYM +; RUN: FileCheck %s -input-file=%t_esimd_large_grf_1.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM ; CHECK: [Code|Properties|Symbols] -; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym -; CHECK: {{.*}}esimd-large-grf.ll.tmp_1.ll|{{.*}}esimd-large-grf.ll.tmp_1.prop|{{.*}}esimd-large-grf.ll.tmp_1.sym -; CHECK: {{.*}}esimd-large-grf.ll.tmp_esimd_1.ll|{{.*}}esimd-large-grf.ll.tmp_esimd_1.prop|{{.*}}esimd-large-grf.ll.tmp_esimd_1.sym +; CHECK: {{.*}}esimd-large-grf.ll.tmp_0.ll|{{.*}}esimd-large-grf.ll.tmp_0.prop|{{.*}}esimd-large-grf.ll.tmp_0.sym +; CHECK: {{.*}}esimd-large-grf.ll.tmp_esimd_0.ll|{{.*}}esimd-large-grf.ll.tmp_esimd_0.prop|{{.*}}esimd-large-grf.ll.tmp_esimd_0.sym +; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym ; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1 ; CHECK-ESIMD-LargeGRF-PROP: isLargeGRF=1|1 diff --git a/llvm/test/tools/sycl-post-link/sycl-large-grf.ll b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll index 26340ddad59e9..cb76430e8c4a2 100644 --- a/llvm/test/tools/sycl-post-link/sycl-large-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll @@ -9,14 +9,14 @@ ; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table -; RUN: FileCheck %s -input-file=%t_large_grf_0.ll --check-prefixes CHECK-LARGE-GRF-IR -; RUN: FileCheck %s -input-file=%t_large_grf_0.prop --check-prefixes CHECK-LARGE-GRF-PROP -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM -; RUN: FileCheck %s -input-file=%t_large_grf_0.sym --check-prefixes CHECK-LARGE-GRF-SYM +; RUN: FileCheck %s -input-file=%t_large_grf_1.ll --check-prefixes CHECK-LARGE-GRF-IR +; RUN: FileCheck %s -input-file=%t_large_grf_1.prop --check-prefixes CHECK-LARGE-GRF-PROP +; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM +; RUN: FileCheck %s -input-file=%t_large_grf_1.sym --check-prefixes CHECK-LARGE-GRF-SYM ; CHECK: [Code|Properties|Symbols] -; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym -; CHECK: {{.*}}-large-grf.ll.tmp_1.ll|{{.*}}-large-grf.ll.tmp_1.prop|{{.*}}-large-grf.ll.tmp_1.sym +; CHECK: {{.*}}-large-grf.ll.tmp_0.ll|{{.*}}-large-grf.ll.tmp_0.prop|{{.*}}-large-grf.ll.tmp_0.sym +; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym ; CHECK-LARGE-GRF-PROP: isLargeGRF=1|1 From 1cc4d14a04f9a2ffac910d720a9ba859245c6744 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 23 Mar 2023 22:36:15 -0700 Subject: [PATCH 05/41] Add documentation Signed-off-by: Arvind Sudarsanam --- .../design/PropagateCompilerFlagsToRuntime.md | 66 +++++++++++++++++++ sycl/doc/index.rst | 1 + 2 files changed, 67 insertions(+) create mode 100644 sycl/doc/design/PropagateCompilerFlagsToRuntime.md diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md new file mode 100644 index 0000000000000..f839e85f1486d --- /dev/null +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -0,0 +1,66 @@ +# Propagation of optimization levels used by front-end compiler to backend compiler + +In order to ease the process of debugging, there is a user requirement to compile different modules with different levels of optimization. This document proposes a compiler flow that will enable propagation of compiler options specified for front-end to the runtimes and eventually to the backend. Currently, only O0/O1/O2/O3 options are handled. + +**NOTE**: This is not a final version. The document is still in progress. + +## Background + +When building an application with several source and object files, it is possible to specify the optimization parameters individually for each source file/object file (for each invocation of the DPCPP compiler). SYCL runtime should pass the original optimization options (e.g. -O0 or -O2) used when building an object file to the device backend compiler. This will improve the debugging experience by selectively disabling/enabling optimizations for each source file, and therefore achieving better debuggability and better performance as needed. + +The current behavior, is that the device backend optimization options are set by setting an environment variable (SYCL_PROGRAM_COMPILE_OPTIONS). If the -O0 option is specified, the runtime will pass -cl-opt-disable option to the backend device compiler for {*}all kernels{*}, essentially disabling optimizations globally. Otherwise, if the -O0 option is not specified for linker, it will not pass -cl-opt-disable option at all, therefore making the kernels mostly undebuggable, regardless of the original front-end compiler options. + +Here is an example that demonstrates this pain point: + +``` +clang++ -c test_host.cpp -o test_host.o +clang++ -c -fsycl test_device_1.cpp -o test_device_1.o +clang++ -c -fsycl -O0 test_device_2.cpp -o test_device_2.o +clang++ -fsycl -o test test_host.o test_device_1.o test_device_2.o +``` + +In this scenario, the fat binary is 'test' and there are no compilation flags sent across to the backend compiler. Though the user wanted to have full debuggability with test_device_2.c module, some of the debuggability is lost. + +Another scenario is shown below: + +``` +clang++ -c -O0 -fsycl -g test.cpp -o test.o +clang++ -fsycl test.o -o test +``` + +In this scenario, the fat binary is 'test' and there are no compilation flags sent across to the backend compiler. Though the user wanted to have full debuggability with test.cpp module, some of the debuggability is lost. The user was not able to set a breakpoint inside device code. + +## Requirements + +In order to support module-level debuggability, the user will compile different module files with different levels of optimization. These optimization levels must be preserved and made use of during every stage of compilation. Following are the requirements for this feature. +- If the user specifies '-Ox' as a front-end compile option for a particular module, this option must be preserved during compilation, linking, AOT compilation as well as JIT compilation. +- If the user specifies '-Ox' option using the environment variable, this option will override any front-end compile option and the new option will be preserved during AOT and JIT compilation. +- If the user specifies '-O0' option, SYCL runtime needs to pass the appropriate backend option to AOT and JIT compilation stages. + +The following table specifies the appropriate backend options for level-zero and OpenCL backends. + +| Front-end option | L0 backend option | OpenCL backend option | +| ---------------- | ----------------- | --------------------- | +| -O0 | -ze-opt-disable | -cl-opt-disable | +| -O1 | -ze-opt-level=1 | /* no option */ | +| -O2 | -ze-opt-level=1 | /* no option */ | +| -O3 | -ze-opt-level=2 | /* no option */ | + + +## Proposed design + +This chapter discusses changes required in various stages of the compilation pipeline. + + +### Changes to the DPC++ front-end + +For each SYCL kernel, we add a new function attribute that is named 'sycl-optlevel'. Value of this attribute is set to the optimization level which was used to compile the overlying module. + +### Changes to the sycl-post-link tool + +During sycl-post-link stage, a set of optional kernel features are combined to form a hash value and this hash value is used as a key to split a module into multiple sub-modules. Current list of optional kernel features include: (1) SYCL aspects (2) large-grf mode (3) reqd-work-group-size. In this design, we add the optimization level associated with the kernel into this list. This helps us to split the kernels based on their optimization level. +The sycl-post-link tool also adds a new property into the "SYCL/misc properties" property set for each device code module. This entry will be used to store the optimization level. Name of this property is 'optLevel' and the value is stored as a 32-bit integer. If there is a module where user did not specify an optimization module, there is no new entry in the property set. + +### Changes to the SYCL runtime + +In SYCL runtime, the device image properties can be accessed to extract the associated optLevel. Once the optimization level is available, a query is made to identify if the backend is a level-zero backend or OpenCL backend. The table provided in the 'Requirements' section is used as a guide to identify the appropriate backend option. This backend option is added to the existing list of compiler options and is sent to the backend. An improvement to this design will be to add a new plugin API which can backend specific optimization option for a given front-end option. \ No newline at end of file diff --git a/sycl/doc/index.rst b/sycl/doc/index.rst index 3ff96f2139716..57188fa144e1b 100644 --- a/sycl/doc/index.rst +++ b/sycl/doc/index.rst @@ -49,6 +49,7 @@ Design Documents for the oneAPI DPC++ Compiler design/KernelFusionJIT design/NonRelocatableDeviceCode design/DeviceAspectTraitDesign + design/PropagateCompilerFlagsToRuntime New OpenCL Extensions New SPIR-V Extensions From 570d7209cabd3cdf2be38c2bd4158cc7db9bca6b Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 23 Mar 2023 23:11:50 -0700 Subject: [PATCH 06/41] Fix formatting issues Signed-off-by: Arvind Sudarsanam --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 10 +++++----- llvm/tools/sycl-post-link/sycl-post-link.cpp | 3 ++- .../detail/program_manager/program_manager.cpp | 13 +++++++++---- 3 files changed, 16 insertions(+), 10 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 05f3171505fba..75546fe42b166 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -739,8 +739,9 @@ struct UsedOptionalFeatures { UsesLargeGRF = true; if (F->hasFnAttribute(ATTR_OPT_LEVEL)) - if (F->getFnAttribute(ATTR_OPT_LEVEL).getValueAsString() - .getAsInteger(10, OptLevel)) + if (F->getFnAttribute(ATTR_OPT_LEVEL) + .getValueAsString() + .getAsInteger(10, OptLevel)) OptLevel = -1; if (const MDNode *MDN = F->getMetadata("reqd_work_group_size")) { @@ -759,9 +760,8 @@ struct UsedOptionalFeatures { llvm::hash_code ReqdWorkGroupSizeHash = llvm::hash_combine_range( ReqdWorkGroupSize.begin(), ReqdWorkGroupSize.end()); llvm::hash_code OptLevelHash = llvm::hash_value(OptLevel); - Hash = static_cast( - llvm::hash_combine(AspectsHash, LargeGRFHash, ReqdWorkGroupSizeHash, - OptLevelHash)); + Hash = static_cast(llvm::hash_combine( + AspectsHash, LargeGRFHash, ReqdWorkGroupSizeHash, OptLevelHash)); } std::string generateModuleName(StringRef BaseName) const { diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 1c3f7bacec289..4ce4658a602ec 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -454,7 +454,8 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, if (MD.isLargeGRF()) PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true}); if (MD.getOptLevel() != -1) - PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"optLevel", MD.getOptLevel()}); + PropSet[PropSetRegTy::SYCL_MISC_PROP].insert( + {"optLevel", MD.getOptLevel()}); { std::vector FuncNames = getKernelNamesUsingAssert(M); for (const StringRef &FName : FuncNames) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 78acac03cc81c..e8d33865b2b7d 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -367,7 +367,6 @@ static int getUint32PropAsInt(const RTDeviceBinaryImage &Img, return (int)(DeviceBinaryProperty(Prop).asUint32()); } - static void appendCompileOptionsFromImage(std::string &CompileOpts, const RTDeviceBinaryImage &Img, const std::vector &Devs, @@ -415,10 +414,16 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, if (!CompileOpts.empty() && (optLevel != -1)) CompileOpts += " "; switch (optLevel) { - case 0: CompileOpts += "-ze-opt-disable"; break; + case 0: + CompileOpts += "-ze-opt-disable"; + break; case 1: - case 2: CompileOpts += "-ze-opt-level=1"; break; - case 3: CompileOpts += "-ze-opt-level=2"; break; + case 2: + CompileOpts += "-ze-opt-level=1"; + break; + case 3: + CompileOpts += "-ze-opt-level=2"; + break; } } else if (Plugin.getBackend() == backend::opencl) { if (!CompileOpts.empty() && (optLevel == 0)) From 992614ed8d632e6222dab2a17b1660184b5243cb Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Mon, 27 Mar 2023 09:22:55 -0700 Subject: [PATCH 07/41] Add plugin API to get backend option Signed-off-by: Arvind Sudarsanam --- sycl/include/sycl/detail/pi.def | 1 + sycl/include/sycl/detail/pi.h | 10 ++++++ sycl/plugins/cuda/pi_cuda.cpp | 14 +++++++++ .../esimd_emulator/pi_esimd_emulator.cpp | 12 +++++++ sycl/plugins/hip/pi_hip.cpp | 14 +++++++++ sycl/plugins/level_zero/pi_level_zero.cpp | 30 ++++++++++++++++++ sycl/plugins/opencl/pi_opencl.cpp | 16 ++++++++++ sycl/source/detail/plugin.hpp | 9 ++++++ .../program_manager/program_manager.cpp | 31 +++++++------------ sycl/test/abi/pi_level_zero_symbol_check.dump | 1 + sycl/test/abi/pi_opencl_symbol_check.dump | 1 + sycl/unittests/helpers/PiMockPlugin.hpp | 5 +++ 12 files changed, 125 insertions(+), 19 deletions(-) diff --git a/sycl/include/sycl/detail/pi.def b/sycl/include/sycl/detail/pi.def index eda09035c883e..4b80a83998b35 100644 --- a/sycl/include/sycl/detail/pi.def +++ b/sycl/include/sycl/detail/pi.def @@ -152,4 +152,5 @@ _PI_API(piGetDeviceAndHostTimer) _PI_API(piextEnqueueDeviceGlobalVariableWrite) _PI_API(piextEnqueueDeviceGlobalVariableRead) +_PI_API(piPluginGetBackendOptimizationOption) #undef _PI_API diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 8f6a20ea3771b..f00af76229f78 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -1954,6 +1954,16 @@ __SYCL_EXPORT pi_result piTearDown(void *PluginParameter); /// timestamp __SYCL_EXPORT pi_result piPluginGetLastError(char **message); +/// API to get backend specific optimization option. +/// \param opt_level is an integer that contains frontend optimization level. +/// \param backend_option is used to return the backend optimization option +/// corresponding to frontend optimization level. +/// +/// \return PI_SUCCESS is returned always. If a valid option is not +/// available, an empty string is returned. +__SYCL_EXPORT pi_result +piPluginGetBackendOptimizationOption(int opt_level, char **backend_option); + /// Queries device for it's global timestamp in nanoseconds, and updates /// HostTime with the value of the host timer at the closest possible point in /// time to that at which DeviceTime was returned. diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 2d88978d87780..ca7e0d800c21b 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -80,6 +80,18 @@ pi_result cuda_piPluginGetLastError(char **message) { return ErrorMessageCode; } +// Optimization strings +char EmptyStr[2] = ""; + +// Returns plugin specific backend optimization option. +// Return empty string for cuda. +// TODO: Determine correct string to be passed. +pi_result cuda_piPluginGetBackendOptimizationOption(int opt_level, + char **backend_option) { + *backend_option = &EmptyStr[0]; + return PI_SUCCESS; +} + // Iterates over the event wait list, returns correct pi_result error codes. // Invokes the callback for the latest event of each queue in the wait list. // The callback must take a single pi_event argument and return a pi_result. @@ -5743,6 +5755,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piPluginGetLastError, cuda_piPluginGetLastError) _PI_CL(piTearDown, cuda_piTearDown) _PI_CL(piGetDeviceAndHostTimer, cuda_piGetDeviceAndHostTimer) + _PI_CL(piPluginGetBackendOptimizationOption, + cuda_piPluginGetBackendOptimizationOption) #undef _PI_CL diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 5bb0ce881e79f..1c7ae537c1525 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -167,6 +167,18 @@ pi_result piPluginGetLastError(char **message) { return ErrorMessageCode; } +// Optimization strings +char EmptyStr[2] = ""; + +// Returns plugin specific backend optimization option. +// Return empty string for esimd emulator. +// TODO: Determine correct string to be passed. +pi_result piPluginGetBackendOptimizationOption(int opt_level, + char **backend_option) { + *backend_option = &EmptyStr[0]; + return PI_SUCCESS; +} + using IDBuilder = sycl::detail::Builder; template diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index b3b2276fdfa5a..6d7283c9f7b55 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -132,6 +132,18 @@ pi_result hip_piPluginGetLastError(char **message) { return ErrorMessageCode; } +// Optimization strings +char EmptyStr[2] = ""; + +// Returns plugin specific backend optimization option. +// Return empty string for hip. +// TODO: Determine correct string to be passed. +pi_result hip_piPluginGetBackendOptimizationOption(int opt_level, + char **backend_option) { + *backend_option = EmptyStr; + return PI_SUCCESS; +} + // Iterates over the event wait list, returns correct pi_result error codes. // Invokes the callback for the latest event of each queue in the wait list. // The callback must take a single pi_event argument and return a pi_result. @@ -5535,6 +5547,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piPluginGetLastError, hip_piPluginGetLastError) _PI_CL(piTearDown, hip_piTearDown) _PI_CL(piGetDeviceAndHostTimer, hip_piGetDeviceAndHostTimer) + _PI_CL(piPluginGetBackendOptimizationOption, + hip_piPluginGetBackendOptimizationOption) #undef _PI_CL diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 7fbffe6d804d4..54fb08a67dd26 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2126,6 +2126,36 @@ pi_result piPluginGetLastError(char **message) { return pi2ur::piPluginGetLastError(message); } +// Optimization strings +char EmptyStr[2] = ""; +char NoOptStr[16] = "-ze-opt-disable"; +char O1OptStr[16] = "-ze-opt-level=1"; +char O2OptStr[16] = "-ze-opt-level=2"; + +// Returns plugin specific backend optimization option. +// Return '-ze-opt-disable' for opt_level = 0. +// Return '-ze-opt-level=1' for opt_level = 1/2. +// Return '-ze-opt-level=2' for opt_level = 3. +pi_result piPluginGetBackendOptimizationOption(int opt_level, + char **backend_option) { + switch (opt_level) { + case 0: + *backend_option = &NoOptStr[0]; + break; + case 1: + case 2: + *backend_option = &O1OptStr[0]; + break; + case 3: + *backend_option = &O2OptStr[0]; + break; + default: + *backend_option = &EmptyStr[0]; + break; + } + return PI_SUCCESS; +} + pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, pi_uint32 NumEntries, pi_device *Devices, pi_uint32 *NumDevices) { diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index f2ea816b023f1..99a23b10015e1 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -93,6 +93,20 @@ pi_result piPluginGetLastError(char **message) { return ErrorMessageCode; } +// Optimization strings +char EmptyStr[2] = ""; +char NoOptStr[16] = "-cl-opt-disable"; + +// Returns plugin specific backend optimization option. +// Return '-cl-opt-disable' for opt_level = 0 and '' for others. +pi_result piPluginGetBackendOptimizationOption(int opt_level, + char **backend_option) { + *backend_option = EmptyStr; + if (opt_level == 0) + *backend_option = NoOptStr; + return PI_SUCCESS; +} + static cl_int getPlatformVersion(cl_platform_id plat, OCLV::OpenCLVersion &version) { cl_int ret_err = CL_INVALID_VALUE; @@ -2095,6 +2109,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piPluginGetLastError, piPluginGetLastError) _PI_CL(piTearDown, piTearDown) _PI_CL(piGetDeviceAndHostTimer, piGetDeviceAndHostTimer) + _PI_CL(piPluginGetBackendOptimizationOption, + piPluginGetBackendOptimizationOption) #undef _PI_CL diff --git a/sycl/source/detail/plugin.hpp b/sycl/source/detail/plugin.hpp index b2ed64d591218..4fe3366cc95dc 100644 --- a/sycl/source/detail/plugin.hpp +++ b/sycl/source/detail/plugin.hpp @@ -231,6 +231,15 @@ class plugin { void *getLibraryHandle() { return MLibraryHandle; } int unload() { return RT::unloadPlugin(MLibraryHandle); } + // Get backend optimization option + void getBackendOptimizationOption(int opt_level, + char **backend_option) const { + [[maybe_unused]] auto pi_result = + call_nocheck( + opt_level, backend_option); + return; + } + // return the index of PiPlatforms. // If not found, add it and return its index. // The function is expected to be called in a thread safe manner. diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index e8d33865b2b7d..08efadffe436b 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -410,26 +410,19 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, CompileOpts += isEsimdImage ? "-doubleGRF" : "-ze-opt-large-register-file"; } // Add optimization flags - if (Plugin.getBackend() == backend::ext_oneapi_level_zero) { - if (!CompileOpts.empty() && (optLevel != -1)) - CompileOpts += " "; - switch (optLevel) { - case 0: - CompileOpts += "-ze-opt-disable"; - break; - case 1: - case 2: - CompileOpts += "-ze-opt-level=1"; - break; - case 3: - CompileOpts += "-ze-opt-level=2"; - break; - } - } else if (Plugin.getBackend() == backend::opencl) { - if (!CompileOpts.empty() && (optLevel == 0)) + // Assigning space for 16 characters. + char *backend_option = nullptr; + // Empty string is returned in backend_option when no appropriate backend + // option is available for a given opt level. + Plugin.getBackendOptimizationOption(optLevel, &backend_option); + if (backend_option && backend_option[0] != '\0') { + if (!CompileOpts.empty()) CompileOpts += " "; - if (optLevel == 0) - CompileOpts += "-cl-opt-disable"; + CompileOpts += std::string(backend_option); + } else { + // emit warning + if (SYCLConfig::get() >= 2) + std::clog << "Optimization level not propagated to backend"; } if ((Plugin.getBackend() == backend::ext_oneapi_level_zero || Plugin.getBackend() == backend::opencl) && diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index fbefe601f3675..4d08a5250077a 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -59,6 +59,7 @@ piMemRelease piMemRetain piPlatformGetInfo piPlatformsGet +piPluginGetBackendOptimizationOption piPluginGetLastError piPluginInit piProgramBuild diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index 7925dfcbc6b53..2a2d49d5e1aeb 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -22,6 +22,7 @@ piMemBufferCreate piMemBufferPartition piMemImageCreate piPlatformsGet +piPluginGetBackendOptimizationOption piPluginGetLastError piPluginInit piProgramCreate diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index c2ac5e6863b8b..17d51f4ebf4e4 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -1121,6 +1121,11 @@ inline pi_result mock_piPluginGetLastError(char **message) { return PI_SUCCESS; } +inline pi_result mock_piPluginGetBackendOptimizationOption(int opt_level, + char **option) { + return PI_SUCCESS; +} + // Returns the wall-clock timestamp of host for deviceTime and hostTime inline pi_result mock_piGetDeviceAndHostTimer(pi_device device, uint64_t *deviceTime, From e2e678ab5b5b9446ccd636338c1763476f9c9ac5 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Mon, 27 Mar 2023 09:38:06 -0700 Subject: [PATCH 08/41] Address review comments on tests Signed-off-by: Arvind Sudarsanam --- clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp | 2 +- llvm/test/tools/sycl-post-link/sycl-opt-level.ll | 8 ++------ 2 files changed, 3 insertions(+), 7 deletions(-) diff --git a/clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp b/clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp index e41a5a71ef990..f5f7b8a452c04 100644 --- a/clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp +++ b/clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp @@ -1,6 +1,6 @@ // RUN: %clangxx %s -O0 -S -o %t.ll -fsycl-device-only // RUN: FileCheck %s --input-file %t.ll -check-prefixes=CHECK-IR -// CHECK-IR: define weak_odr dso_local spir_kernel void @{{.*}}main{{.*}}sycl{{.*}}handler{{.*}}() #[[ATTR:[0-9]+]] +// CHECK-IR: define {{.*}} spir_kernel void @{{.*}}main{{.*}}sycl{{.*}}handler{{.*}}() #[[ATTR:[0-9]+]] // CHECK-IR: attributes #[[ATTR]] = { {{.*}} "sycl-optlevel"="0" {{.*}}} // This test checks adding of the attribute 'sycl-optlevel' diff --git a/llvm/test/tools/sycl-post-link/sycl-opt-level.ll b/llvm/test/tools/sycl-post-link/sycl-opt-level.ll index 948c54332f17d..349e6e445027d 100644 --- a/llvm/test/tools/sycl-post-link/sycl-opt-level.ll +++ b/llvm/test/tools/sycl-post-link/sycl-opt-level.ll @@ -16,19 +16,15 @@ ; CHECK-OPT-LEVEL-PROP-0: optLevel=1|0 ; CHECK-OPT-LEVEL-PROP-1: optLevel=1|2 -; ModuleID = 'final.bc' -source_filename = "llvm-link" 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" target triple = "spir64-unknown-unknown" -; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none) define dso_local spir_func noundef i32 @_Z3fooii(i32 noundef %a, i32 noundef %b) local_unnamed_addr #0 { entry: %sub = sub nsw i32 %a, %b ret i32 %sub } -; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone define dso_local spir_func noundef i32 @_Z3booii(i32 noundef %a, i32 noundef %b) #1 { entry: %retval = alloca i32, align 4 @@ -45,6 +41,6 @@ entry: ret i32 %add } -attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test3.cpp" "sycl-optlevel"="2" } -attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test2.cpp" "sycl-optlevel"="0" } +attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "sycl-module-id"="test3.cpp" "sycl-optlevel"="2" } +attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "sycl-module-id"="test2.cpp" "sycl-optlevel"="0" } From d89738e5b963226f73efae765f616ebc5637115c Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Mon, 27 Mar 2023 10:34:53 -0700 Subject: [PATCH 09/41] Address review comments and other changes to documentation Signed-off-by: Arvind Sudarsanam --- .../design/PropagateCompilerFlagsToRuntime.md | 22 +++++++++---------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index f839e85f1486d..8adb3e9f2a4a9 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -1,14 +1,14 @@ # Propagation of optimization levels used by front-end compiler to backend compiler -In order to ease the process of debugging, there is a user requirement to compile different modules with different levels of optimization. This document proposes a compiler flow that will enable propagation of compiler options specified for front-end to the runtimes and eventually to the backend. Currently, only O0/O1/O2/O3 options are handled. +In order to ease the process of debugging, there is a user requirement to compile different modules with different levels of optimization. This document proposes a compiler flow that will enable propagation of compiler options specified for front-end to the runtimes and eventually to the backend. Currently, only `O0`/`O1`/`O2`/`O3` options are handled. **NOTE**: This is not a final version. The document is still in progress. ## Background -When building an application with several source and object files, it is possible to specify the optimization parameters individually for each source file/object file (for each invocation of the DPCPP compiler). SYCL runtime should pass the original optimization options (e.g. -O0 or -O2) used when building an object file to the device backend compiler. This will improve the debugging experience by selectively disabling/enabling optimizations for each source file, and therefore achieving better debuggability and better performance as needed. +When building an application with several source and object files, it is possible to specify the optimization parameters individually for each source file/object file (for each invocation of the DPCPP compiler). The SYCL runtime should pass the original optimization options (e.g. `-O0` or `-O2`) used when building an object file to the device backend compiler. This will improve the debugging experience by selectively disabling/enabling optimizations for each source file, and therefore achieving better debuggability and better performance as needed. -The current behavior, is that the device backend optimization options are set by setting an environment variable (SYCL_PROGRAM_COMPILE_OPTIONS). If the -O0 option is specified, the runtime will pass -cl-opt-disable option to the backend device compiler for {*}all kernels{*}, essentially disabling optimizations globally. Otherwise, if the -O0 option is not specified for linker, it will not pass -cl-opt-disable option at all, therefore making the kernels mostly undebuggable, regardless of the original front-end compiler options. +The current behavior is that the device backend optimization options can be propagated to the backend by setting the environment variable `SYCL_PROGRAM_COMPILE_OPTIONS`. For example, If `-O0` option is specified when using the OpenCL backend, the SYCL runtime will pass `-cl-opt-disable` option to the backend device compiler for {*}all modules{*} essentially disabling optimizations globally. Otherwise, if the `-O0` option is not specified for linker, it will not pass `-cl-opt-disable` option at all, therefore making the kernels mostly undebuggable, regardless of the original front-end compiler options. Here is an example that demonstrates this pain point: @@ -32,10 +32,10 @@ In this scenario, the fat binary is 'test' and there are no compilation flags se ## Requirements -In order to support module-level debuggability, the user will compile different module files with different levels of optimization. These optimization levels must be preserved and made use of during every stage of compilation. Following are the requirements for this feature. -- If the user specifies '-Ox' as a front-end compile option for a particular module, this option must be preserved during compilation, linking, AOT compilation as well as JIT compilation. -- If the user specifies '-Ox' option using the environment variable, this option will override any front-end compile option and the new option will be preserved during AOT and JIT compilation. -- If the user specifies '-O0' option, SYCL runtime needs to pass the appropriate backend option to AOT and JIT compilation stages. +In order to support module-level debuggability, the user will compile different module files with different levels of optimization. These optimization levels must be preserved and made use of during the backend compilation. Following are the requirements for this feature. +- If the user specifies `-Ox` as a front-end compile option for a particular module, this option must be preserved during backend JIT compilation. +- If the user specifies `-Ox` option using the environment variable, this option will override any front-end compile option and the new option will be preserved during JIT compilation. +- If the user specifies `-O0` option, SYCL runtime needs to pass the appropriate backend option to JIT compilation stages. The following table specifies the appropriate backend options for level-zero and OpenCL backends. @@ -54,13 +54,13 @@ This chapter discusses changes required in various stages of the compilation pip ### Changes to the DPC++ front-end -For each SYCL kernel, we add a new function attribute that is named 'sycl-optlevel'. Value of this attribute is set to the optimization level which was used to compile the overlying module. +For each SYCL kernel, we add a new function attribute that is named `sycl-optlevel`. Value of this attribute is set to the optimization level which was used to compile the overlying module. ### Changes to the sycl-post-link tool -During sycl-post-link stage, a set of optional kernel features are combined to form a hash value and this hash value is used as a key to split a module into multiple sub-modules. Current list of optional kernel features include: (1) SYCL aspects (2) large-grf mode (3) reqd-work-group-size. In this design, we add the optimization level associated with the kernel into this list. This helps us to split the kernels based on their optimization level. -The sycl-post-link tool also adds a new property into the "SYCL/misc properties" property set for each device code module. This entry will be used to store the optimization level. Name of this property is 'optLevel' and the value is stored as a 32-bit integer. If there is a module where user did not specify an optimization module, there is no new entry in the property set. +During `sycl-post-link` stage, a set of optional kernel features are combined to form a hash value and this hash value is used as a key to split a module into multiple sub-modules. Current list of optional kernel features include: (1) SYCL aspects (2) `large-grf` mode (3) `reqd-work-group-size`. In this design, we add the optimization level associated with the kernel into this list. This helps us to split the kernels based on their optimization level. +The `sycl-post-link` tool also adds a new property into the `SYCL/misc properties` property set for each device code module. This entry will be used to store the optimization level. Name of this property is 'optLevel' and the value is stored as a 32-bit integer. If there is a module where user did not specify an optimization module, there is no new entry in the property set. ### Changes to the SYCL runtime -In SYCL runtime, the device image properties can be accessed to extract the associated optLevel. Once the optimization level is available, a query is made to identify if the backend is a level-zero backend or OpenCL backend. The table provided in the 'Requirements' section is used as a guide to identify the appropriate backend option. This backend option is added to the existing list of compiler options and is sent to the backend. An improvement to this design will be to add a new plugin API which can backend specific optimization option for a given front-end option. \ No newline at end of file +In the SYCL runtime, the device image properties can be accessed to extract the associated optimization level. Once the optimization level is available, a query is made to identify the correct backend option. The table provided in the 'Requirements' section is used as a guide to identify the appropriate backend option. This backend option is added to the existing list of compiler options and is sent to the backend. \ No newline at end of file From 67ac752cd4b7b16e36da0bab0e7aa36e54decda5 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Mon, 27 Mar 2023 15:26:55 -0700 Subject: [PATCH 10/41] Minor correction Signed-off-by: Arvind Sudarsanam --- sycl/doc/design/PropagateCompilerFlagsToRuntime.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 8adb3e9f2a4a9..b08c9017b293d 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -52,7 +52,7 @@ The following table specifies the appropriate backend options for level-zero and This chapter discusses changes required in various stages of the compilation pipeline. -### Changes to the DPC++ front-end +### Changes to the clang front-end For each SYCL kernel, we add a new function attribute that is named `sycl-optlevel`. Value of this attribute is set to the optimization level which was used to compile the overlying module. From 87a6f147185037fa4595c90a973f980ac2442124 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Tue, 28 Mar 2023 18:40:27 -0700 Subject: [PATCH 11/41] Move attribute addition to a separate pass called in backend utils pass pipeline Signed-off-by: Arvind Sudarsanam --- clang/lib/CodeGen/BackendUtil.cpp | 4 +++ clang/lib/CodeGen/CodeGenFunction.cpp | 12 ++----- .../SYCLLowerIR/SYCLAddOptLevelAttribute.h | 33 +++++++++++++++++++ llvm/lib/Passes/PassBuilder.cpp | 1 + llvm/lib/Passes/PassRegistry.def | 1 + llvm/lib/SYCLLowerIR/CMakeLists.txt | 1 + .../SYCLLowerIR/SYCLAddOptLevelAttribute.cpp | 28 ++++++++++++++++ 7 files changed, 71 insertions(+), 9 deletions(-) create mode 100644 llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h create mode 100644 llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index 47b9d4e39f1a6..62848681b5e21 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -48,6 +48,7 @@ #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h" #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h" #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h" +#include "llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h" #include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h" #include "llvm/Support/BuryPointer.h" #include "llvm/Support/CommandLine.h" @@ -1031,6 +1032,9 @@ void EmitAssemblyHelper::RunOptimizationPipeline( MPM.addPass(SYCLPropagateAspectsUsagePass(/*ExcludeAspects=*/{}, /*ValidateAspects=*/false)); + // Add attribute corresponding to optimization level. + MPM.addPass(SYCLAddOptLevelAttributePass(CodeGenOpts.OptimizationLevel)); + // Add SPIRITTAnnotations pass to the pass manager if // -fsycl-instrument-device-code option was passed. This option can be // used only with spir triple. diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 8251834b3651a..9d372f3660409 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -589,17 +589,11 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD, && !FD->hasAttr()) return; - if (getLangOpts().SYCLIsDevice) { - // TODO Module identifier is not reliable for this purpose since two modules - // can have the same ID, needs improvement + // TODO Module identifier is not reliable for this purpose since two modules + // can have the same ID, needs improvement + if (getLangOpts().SYCLIsDevice) Fn->addFnAttr("sycl-module-id", Fn->getParent()->getModuleIdentifier()); - // Here, we add a function attribute 'sycl-optlevel' to store the - // optimization level. - int SYCLOptLevel = CGM.getCodeGenOpts().OptimizationLevel; - assert(SYCLOptLevel >= 0 && "Invalid optimization level!"); - Fn->addFnAttr("sycl-optlevel", std::to_string(SYCLOptLevel)); - } llvm::LLVMContext &Context = getLLVMContext(); if (FD->hasAttr() || FD->hasAttr()) diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h b/llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h new file mode 100644 index 0000000000000..d7d450715b973 --- /dev/null +++ b/llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h @@ -0,0 +1,33 @@ +//===---- SYCLAddOptLevelAttribute.cpp - SYCLAddOptLevelAttribute Pass --===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Pass adds 'sycl-optlevel' function attribute based on optimization level +// passed in. +// +//===----------------------------------------------------------------------===// +// +#ifndef LLVM_SYCL_ADD_OPT_LEVEL_ATTRIBUTE_H +#define LLVM_SYCL_ADD_OPT_LEVEL_ATTRIBUTE_H + +#include "llvm/IR/PassManager.h" + +namespace llvm { + +class SYCLAddOptLevelAttributePass + : public PassInfoMixin { +public: + SYCLAddOptLevelAttributePass(int OptLevel = -1) : OptLevel{OptLevel} {}; + PreservedAnalyses run(Module &M, ModuleAnalysisManager &); + +private: + int OptLevel; +}; + +} // namespace llvm + +#endif // LLVM_SYCL_ADD_OPT_LEVEL_ATTRIBUTE_H \ No newline at end of file diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp index 03b0a836354d4..91e3b0e526563 100644 --- a/llvm/lib/Passes/PassBuilder.cpp +++ b/llvm/lib/Passes/PassBuilder.cpp @@ -90,6 +90,7 @@ #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h" #include "llvm/SYCLLowerIR/LowerWGScope.h" #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h" +#include "llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h" #include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def index 1a715623eb277..e6e6bf8d20b82 100644 --- a/llvm/lib/Passes/PassRegistry.def +++ b/llvm/lib/Passes/PassRegistry.def @@ -139,6 +139,7 @@ MODULE_PASS("deadargelim-sycl", DeadArgumentEliminationSYCLPass()) MODULE_PASS("sycllowerwglocalmemory", SYCLLowerWGLocalMemoryPass()) MODULE_PASS("lower-esimd-kernel-attrs", SYCLFixupESIMDKernelWrapperMDPass()) MODULE_PASS("sycl-propagate-aspects-usage", SYCLPropagateAspectsUsagePass()) +MODULE_PASS("sycl-add-opt-level-attribute", SYCLAddOptLevelAttributePass()) MODULE_PASS("compile-time-properties", CompileTimePropertiesPass()) #undef MODULE_PASS diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index b8f6ab50d15b5..90e881a5a2a54 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -62,6 +62,7 @@ add_llvm_component_library(LLVMSYCLLowerIR LowerWGLocalMemory.cpp LowerWGScope.cpp MutatePrintfAddrspace.cpp + SYCLAddOptLevelAttribute.cpp SYCLPropagateAspectsUsage.cpp SYCLUtils.cpp diff --git a/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp b/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp new file mode 100644 index 0000000000000..afe76f2c9876b --- /dev/null +++ b/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp @@ -0,0 +1,28 @@ +//===---- SYCLAddOptLevelAttribute.cpp - SYCLAddOptLevelAttribute Pass ---===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===---------------------------------------------------------------------===// +// +// Pass adds 'sycl-optlevel' function attribute based on optimization level +// passed in. +//===---------------------------------------------------------------------===// + +#include "llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h" + +#include "llvm/IR/Module.h" + +using namespace llvm; + +PreservedAnalyses +SYCLAddOptLevelAttributePass::run(Module &M, ModuleAnalysisManager &MAM) { + for (Function &F : M.functions()) { + // Here, we add a function attribute 'sycl-optlevel' to store the + // optimization level. + assert(OptLevel >= 0 && "Invalid optimization level!"); + F.addFnAttr("sycl-optlevel", std::to_string(OptLevel)); + } + return PreservedAnalyses::all(); +} \ No newline at end of file From 97fdae740fbb137b2774e2ba7af73d3ee4df72f3 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Tue, 28 Mar 2023 19:23:41 -0700 Subject: [PATCH 12/41] Address review comments on documentation Signed-off-by: Arvind Sudarsanam --- .../design/PropagateCompilerFlagsToRuntime.md | 96 +++++++++++++++---- 1 file changed, 75 insertions(+), 21 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index b08c9017b293d..0bb4c93822a47 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -1,43 +1,75 @@ -# Propagation of optimization levels used by front-end compiler to backend compiler +# Propagation of optimization levels used by front-end compiler to backend -In order to ease the process of debugging, there is a user requirement to compile different modules with different levels of optimization. This document proposes a compiler flow that will enable propagation of compiler options specified for front-end to the runtimes and eventually to the backend. Currently, only `O0`/`O1`/`O2`/`O3` options are handled. +In order to ease the process of debugging, there is a user requirement to +compile different modules with different levels of optimization. This document +proposes a compiler flow that will enable propagation of compiler options +specified for front-end to the runtimes and eventually to the backend. +Currently, only `O0`/`O1`/`O2`/`O3` options are handled. **NOTE**: This is not a final version. The document is still in progress. ## Background -When building an application with several source and object files, it is possible to specify the optimization parameters individually for each source file/object file (for each invocation of the DPCPP compiler). The SYCL runtime should pass the original optimization options (e.g. `-O0` or `-O2`) used when building an object file to the device backend compiler. This will improve the debugging experience by selectively disabling/enabling optimizations for each source file, and therefore achieving better debuggability and better performance as needed. - -The current behavior is that the device backend optimization options can be propagated to the backend by setting the environment variable `SYCL_PROGRAM_COMPILE_OPTIONS`. For example, If `-O0` option is specified when using the OpenCL backend, the SYCL runtime will pass `-cl-opt-disable` option to the backend device compiler for {*}all modules{*} essentially disabling optimizations globally. Otherwise, if the `-O0` option is not specified for linker, it will not pass `-cl-opt-disable` option at all, therefore making the kernels mostly undebuggable, regardless of the original front-end compiler options. +When building an application with several source and object files, it is +possible to specify the optimization parameters individually for each source +file/object file (for each invocation of the DPCPP compiler). The SYCL runtime +should pass the original optimization options (e.g. `-O0` or `-O2`) used when +building an object file to the device backend compiler. This will improve the +debugging experience by selectively disabling/enabling optimizations for each +source file, and therefore achieving better debuggability and better performance +as needed. + +The current behavior is that the optimization level option is captured at link +time and converted into its backend-specific equivalent. This option is +propagated to the backend. For example, If `-O0` option is specified during +link-time when using the OpenCL backend, the SYCL runtime will pass +`-cl-opt-disable` option to the backend device compiler for {*}all modules{*} +essentially disabling optimizations for all modules. Otherwise, if the `-O0` +option is not specified for linker, it will not pass `-cl-opt-disable` option at +all, therefore making the kernels mostly undebuggable, regardless of the +original front-end compiler options. Link-time capturing of optimization option +is the essence of the current implementation and this leads to loss of +information about the compile-time options. Proposed design aims to rectify this +behavior. Here is an example that demonstrates this pain point: ``` clang++ -c test_host.cpp -o test_host.o clang++ -c -fsycl test_device_1.cpp -o test_device_1.o -clang++ -c -fsycl -O0 test_device_2.cpp -o test_device_2.o -clang++ -fsycl -o test test_host.o test_device_1.o test_device_2.o +clang++ -c -fsycl -g -O0 test_device_2.cpp -o test_device_2.o +clang++ -g -fsycl -o test test_host.o test_device_1.o test_device_2.o ``` -In this scenario, the fat binary is 'test' and there are no compilation flags sent across to the backend compiler. Though the user wanted to have full debuggability with test_device_2.c module, some of the debuggability is lost. +In this scenario, the fat binary is 'test' and there are no compilation flags +sent across to the backend compiler. Though the user wanted to have full +debuggability with test_device_2.cpp module, some of the debuggability is lost. Another scenario is shown below: ``` -clang++ -c -O0 -fsycl -g test.cpp -o test.o -clang++ -fsycl test.o -o test +clang++ -c -g -O0 -fsycl test.cpp -o test.o +clang++ -g -fsycl test.o -o test ``` -In this scenario, the fat binary is 'test' and there are no compilation flags sent across to the backend compiler. Though the user wanted to have full debuggability with test.cpp module, some of the debuggability is lost. The user was not able to set a breakpoint inside device code. +In this scenario, the fat binary is 'test' and there are no compilation flags +sent across to the backend compiler. Though the user wanted to have full +debuggability with test.cpp module, some of the debuggability is lost. The user +was not able to set a breakpoint inside device code. ## Requirements -In order to support module-level debuggability, the user will compile different module files with different levels of optimization. These optimization levels must be preserved and made use of during the backend compilation. Following are the requirements for this feature. -- If the user specifies `-Ox` as a front-end compile option for a particular module, this option must be preserved during backend JIT compilation. -- If the user specifies `-Ox` option using the environment variable, this option will override any front-end compile option and the new option will be preserved during JIT compilation. -- If the user specifies `-O0` option, SYCL runtime needs to pass the appropriate backend option to JIT compilation stages. +In order to support module-level debuggability, the user will compile different +module files with different levels of optimization. These optimization levels +must be preserved and made use of during the backend compilation. Following are +the requirements for this feature. +- If the user specifies `-Ox` as a front-end compile option for a particular +module, this option must be preserved during backend JIT compilation. +- If the user specifies `-O0` option, SYCL runtime needs to pass the appropriate +backend option to JIT compilation stages. -The following table specifies the appropriate backend options for level-zero and OpenCL backends. +The following table specifies the appropriate backend options for level-zero and +OpenCL backends. | Front-end option | L0 backend option | OpenCL backend option | | ---------------- | ----------------- | --------------------- | @@ -49,18 +81,40 @@ The following table specifies the appropriate backend options for level-zero and ## Proposed design -This chapter discusses changes required in various stages of the compilation pipeline. +This chapter discusses changes required in various stages of the compilation +pipeline. ### Changes to the clang front-end -For each SYCL kernel, we add a new function attribute that is named `sycl-optlevel`. Value of this attribute is set to the optimization level which was used to compile the overlying module. +For each SYCL kernel, we add a new function attribute that is named +`sycl-optlevel`. Value of this attribute is set to the optimization level which +was used to compile the overlying module. ### Changes to the sycl-post-link tool -During `sycl-post-link` stage, a set of optional kernel features are combined to form a hash value and this hash value is used as a key to split a module into multiple sub-modules. Current list of optional kernel features include: (1) SYCL aspects (2) `large-grf` mode (3) `reqd-work-group-size`. In this design, we add the optimization level associated with the kernel into this list. This helps us to split the kernels based on their optimization level. -The `sycl-post-link` tool also adds a new property into the `SYCL/misc properties` property set for each device code module. This entry will be used to store the optimization level. Name of this property is 'optLevel' and the value is stored as a 32-bit integer. If there is a module where user did not specify an optimization module, there is no new entry in the property set. +During device code split performed in the `sycl-post-link` tool, optimization +level attribute `sycl-optlevel` is treated as an optional feature, +i.e. device code split algorithm ensures that no kernels with different values +of sycl-optlevel are bundled into the same device image. See also optional +kernel features [design document](https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OptionalDeviceFeatures.md#changes-to-the-post-link-tool). +The `sycl-post-link` tool also adds a new property into the +`SYCL/misc properties` property set for each device code module. This entry will +be used to store the optimization level. Name of this property is 'optLevel' and +the value is stored as a 32-bit integer. If there is a module where user did not +specify an optimization module, there is no new entry in the property set. ### Changes to the SYCL runtime -In the SYCL runtime, the device image properties can be accessed to extract the associated optimization level. Once the optimization level is available, a query is made to identify the correct backend option. The table provided in the 'Requirements' section is used as a guide to identify the appropriate backend option. This backend option is added to the existing list of compiler options and is sent to the backend. \ No newline at end of file +In the SYCL runtime, the device image properties can be accessed to extract the +associated optimization level. Once the optimization level is available, a query +is made to the plugin to identify the correct backend option. This backend +option is added to the existing list of compiler options and is sent to the +backend. If the plugin returns empty string, then the user is informed of this +via a warning message. + +### Changes to the plugin + +In the level-zero and OpenCL plugins, the table provided in the 'Requirements' +section is used as a guide to identify the appropriate backend option. For other +plugins (HIP, cuda, and ESIMD emulator), empty string is returned. \ No newline at end of file From 1f2015e2a8188940fe2ed3411fe4c7641b3ff43a Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Tue, 28 Mar 2023 19:25:24 -0700 Subject: [PATCH 13/41] minor typo Signed-off-by: Arvind Sudarsanam --- sycl/doc/design/PropagateCompilerFlagsToRuntime.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 0bb4c93822a47..2b0842ebacacf 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -23,8 +23,8 @@ The current behavior is that the optimization level option is captured at link time and converted into its backend-specific equivalent. This option is propagated to the backend. For example, If `-O0` option is specified during link-time when using the OpenCL backend, the SYCL runtime will pass -`-cl-opt-disable` option to the backend device compiler for {*}all modules{*} -essentially disabling optimizations for all modules. Otherwise, if the `-O0` +`-cl-opt-disable` option to the backend device compiler for all modules +essentially disabling optimizations globally. Otherwise, if the `-O0` option is not specified for linker, it will not pass `-cl-opt-disable` option at all, therefore making the kernels mostly undebuggable, regardless of the original front-end compiler options. Link-time capturing of optimization option From 2cc0bcac2ed349d36fe05ced03e3475c5bc01e79 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 29 Mar 2023 06:41:56 -0700 Subject: [PATCH 14/41] Removed check for a specific attribute id Signed-off-by: Arvind Sudarsanam --- sycl/test/check_device_code/group_barrier.cpp | 48 +++++++++---------- 1 file changed, 24 insertions(+), 24 deletions(-) diff --git a/sycl/test/check_device_code/group_barrier.cpp b/sycl/test/check_device_code/group_barrier.cpp index 1b94e6300619b..bf9069eea8558 100644 --- a/sycl/test/check_device_code/group_barrier.cpp +++ b/sycl/test/check_device_code/group_barrier.cpp @@ -38,29 +38,29 @@ int main() { }); return 0; } -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 3, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 4, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 3, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 2, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 1, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 0, i32 912) #2 +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 3, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 4, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 3, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 2, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 1, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 0, i32 912) -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912) #2 +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912) -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912) #2 -// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912) #2 +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912) +// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912) From b77fc0fb42062d6236321f23f7e648d0be3738c9 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 29 Mar 2023 06:47:51 -0700 Subject: [PATCH 15/41] Move test from llvm-test-suite Signed-off-by: Arvind Sudarsanam --- .../sycl-opt-level.cpp | 50 +++++++++++++++++++ 1 file changed, 50 insertions(+) create mode 100644 sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level.cpp diff --git a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level.cpp b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level.cpp new file mode 100644 index 0000000000000..ee2ecc4fb5a56 --- /dev/null +++ b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level.cpp @@ -0,0 +1,50 @@ +//==----------- sycl-opt-level.cpp - DPC++ SYCL on-device test +//---------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// This test verifies the propagation of front-end compiler optimization +// option to the backend. +// API call in device code: +// Following is expected addtion of options: +// Front-end option | OpenCL backend option | L0 backend option +// -O0 | -cl-opt-disable | -ze-opt-disable +// -O1 | /* no option */ | -ze-opt-level=1 +// -O2 | /* no option */ | -ze-opt-level=1 +// -O3 | /* no option */ | -ze-opt-level=2 + +// RUN: %clangxx -O0 -fsycl %s -o %t0.out +// RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 %t0.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK0 +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t0.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL0 +// RUN: %clangxx -O1 -fsycl %s -o %t1.out +// RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 %t1.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK1 +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t1.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL1 +// RUN: %clangxx -O2 -fsycl %s -o %t2.out +// RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 %t2.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK2 +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t2.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL2 +// RUN: %clangxx -O3 -fsycl %s -o %t3.out +// RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 %t3.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK3 +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t3.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL3 + +#include + +int main() { + sycl::queue q; + q.submit([&](sycl::handler &h) { h.single_task([=]() {}); }); + std::cout << "sycl-optlevel test passed\n"; + return 0; +} + +// CHECK-LABEL: ---> piProgramBuild( +// CHECK0: -ze-opt-disable +// CHECKOCL0: -cl-opt-disable +// CHECK1: -ze-opt-level=1 +// CHECKOCL1-NOT: -cl-opt-disable +// CHECK2: -ze-opt-level=1 +// CHECKOCL2-NOT: -cl-opt-disable +// CHECK3: -ze-opt-level=2 +// CHECKOCL3-NOT: -cl-opt-disable +// CHECK: ) ---> pi_result : PI_SUCCESS From aa878ccf25ab695e42a316a7809d08f841cae928 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 29 Mar 2023 08:21:17 -0700 Subject: [PATCH 16/41] Address review comments Signed-off-by: Arvind Sudarsanam --- clang/lib/CodeGen/CodeGenFunction.cpp | 2 +- sycl/source/detail/plugin.hpp | 1 - sycl/source/detail/program_manager/program_manager.cpp | 2 +- 3 files changed, 2 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 9d372f3660409..82424462f47e2 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -591,7 +591,7 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD, // TODO Module identifier is not reliable for this purpose since two modules // can have the same ID, needs improvement - if (getLangOpts().SYCLIsDevice) + if (getLangOpts().SYCLIsDevice) Fn->addFnAttr("sycl-module-id", Fn->getParent()->getModuleIdentifier()); llvm::LLVMContext &Context = getLLVMContext(); diff --git a/sycl/source/detail/plugin.hpp b/sycl/source/detail/plugin.hpp index 4fe3366cc95dc..dcea60814fa13 100644 --- a/sycl/source/detail/plugin.hpp +++ b/sycl/source/detail/plugin.hpp @@ -237,7 +237,6 @@ class plugin { [[maybe_unused]] auto pi_result = call_nocheck( opt_level, backend_option); - return; } // return the index of PiPlatforms. diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 08efadffe436b..82d845ae97726 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -364,7 +364,7 @@ static int getUint32PropAsInt(const RTDeviceBinaryImage &Img, pi_device_binary_property Prop = Img.getProperty(PropName); if (!Prop) return -1; - return (int)(DeviceBinaryProperty(Prop).asUint32()); + return DeviceBinaryProperty(Prop).asUint32(); } static void appendCompileOptionsFromImage(std::string &CompileOpts, From 6e522730ae36c0c03e7c5ed4a48b62e6a08dc3c7 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 29 Mar 2023 08:51:10 -0700 Subject: [PATCH 17/41] Remove warning message Signed-off-by: Arvind Sudarsanam --- sycl/doc/design/PropagateCompilerFlagsToRuntime.md | 3 +-- sycl/source/detail/program_manager/program_manager.cpp | 4 ---- 2 files changed, 1 insertion(+), 6 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 2b0842ebacacf..1ae4490ee7511 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -110,8 +110,7 @@ In the SYCL runtime, the device image properties can be accessed to extract the associated optimization level. Once the optimization level is available, a query is made to the plugin to identify the correct backend option. This backend option is added to the existing list of compiler options and is sent to the -backend. If the plugin returns empty string, then the user is informed of this -via a warning message. +backend. ### Changes to the plugin diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 82d845ae97726..1f73f26a4a291 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -419,10 +419,6 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, if (!CompileOpts.empty()) CompileOpts += " "; CompileOpts += std::string(backend_option); - } else { - // emit warning - if (SYCLConfig::get() >= 2) - std::clog << "Optimization level not propagated to backend"; } if ((Plugin.getBackend() == backend::ext_oneapi_level_zero || Plugin.getBackend() == backend::opencl) && From 41984bea3878c6bedbc4e65bdbb1ea501e8da86a Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Mon, 3 Apr 2023 17:21:32 -0700 Subject: [PATCH 18/41] Update design document and plugin implementation Signed-off-by: Arvind Sudarsanam --- .../design/PropagateCompilerFlagsToRuntime.md | 13 ++++++++---- sycl/plugins/cuda/pi_cuda.cpp | 2 ++ .../esimd_emulator/pi_esimd_emulator.cpp | 2 ++ sycl/plugins/hip/pi_hip.cpp | 2 ++ sycl/plugins/level_zero/pi_level_zero.cpp | 2 +- sycl/plugins/opencl/pi_opencl.cpp | 2 ++ sycl/source/detail/plugin.hpp | 3 ++- .../program_manager/program_manager.cpp | 21 ++++++++++--------- 8 files changed, 31 insertions(+), 16 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 1ae4490ee7511..f926100b1072d 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -6,8 +6,6 @@ proposes a compiler flow that will enable propagation of compiler options specified for front-end to the runtimes and eventually to the backend. Currently, only `O0`/`O1`/`O2`/`O3` options are handled. -**NOTE**: This is not a final version. The document is still in progress. - ## Background When building an application with several source and object files, it is @@ -38,7 +36,7 @@ Here is an example that demonstrates this pain point: clang++ -c test_host.cpp -o test_host.o clang++ -c -fsycl test_device_1.cpp -o test_device_1.o clang++ -c -fsycl -g -O0 test_device_2.cpp -o test_device_2.o -clang++ -g -fsycl -o test test_host.o test_device_1.o test_device_2.o +clang++ -fsycl -g -o test_host.o test_device_1.o test_device_2.o -o test ``` In this scenario, the fat binary is 'test' and there are no compilation flags @@ -114,6 +112,13 @@ backend. ### Changes to the plugin +A new plugin API has been added. It takes the optimization level as input in +integer format and returns `pi_result`. The signature is as follows: + +`pi_result` piPluginGetBackendOptimizationOption(int OptLevel); + In the level-zero and OpenCL plugins, the table provided in the 'Requirements' section is used as a guide to identify the appropriate backend option. For other -plugins (HIP, cuda, and ESIMD emulator), empty string is returned. \ No newline at end of file +plugins (HIP, cuda, and ESIMD emulator), empty string is returned. This API +returns `PI_SUCCESS` for valid inputs (0 <= OptLevel <= 3). For invalid inputs, +it returns `PI_ERROR_INVALID_VALUE`. diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 7075d04b092cd..1dc70b9ed12ea 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -88,6 +88,8 @@ char EmptyStr[2] = ""; // TODO: Determine correct string to be passed. pi_result cuda_piPluginGetBackendOptimizationOption(int opt_level, char **backend_option) { + if ((opt_level < 0) || (opt_level > 3)) + return PI_ERROR_INVALID_VALUE; *backend_option = &EmptyStr[0]; return PI_SUCCESS; } diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index e1cf04f14228f..8face5b1c0e8b 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -175,6 +175,8 @@ char EmptyStr[2] = ""; // TODO: Determine correct string to be passed. pi_result piPluginGetBackendOptimizationOption(int opt_level, char **backend_option) { + if ((opt_level < 0) || (opt_level > 3)) + return PI_ERROR_INVALID_VALUE; *backend_option = &EmptyStr[0]; return PI_SUCCESS; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 4a0aaa6570b77..72904a6493f22 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -140,6 +140,8 @@ char EmptyStr[2] = ""; // TODO: Determine correct string to be passed. pi_result hip_piPluginGetBackendOptimizationOption(int opt_level, char **backend_option) { + if ((opt_level < 0) || (opt_level > 3)) + return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; return PI_SUCCESS; } diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 3b2d34542a0bd..6fdbdcfd2d8fd 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2154,7 +2154,7 @@ pi_result piPluginGetBackendOptimizationOption(int opt_level, break; default: *backend_option = &EmptyStr[0]; - break; + return PI_ERROR_INVALID_VALUE; } return PI_SUCCESS; } diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 1a2b157f0f313..df81858614d44 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -104,6 +104,8 @@ char NoOptStr[16] = "-cl-opt-disable"; // Return '-cl-opt-disable' for opt_level = 0 and '' for others. pi_result piPluginGetBackendOptimizationOption(int opt_level, char **backend_option) { + if ((opt_level < 0) || (opt_level > 3)) + return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; if (opt_level == 0) *backend_option = NoOptStr; diff --git a/sycl/source/detail/plugin.hpp b/sycl/source/detail/plugin.hpp index d2fe6f9385c5e..3f530c07f98b0 100644 --- a/sycl/source/detail/plugin.hpp +++ b/sycl/source/detail/plugin.hpp @@ -238,9 +238,10 @@ class plugin { // Get backend optimization option void getBackendOptimizationOption(int opt_level, char **backend_option) const { - [[maybe_unused]] auto pi_result = + RT::PiResult Err = call_nocheck( opt_level, backend_option); + checkPiResult(Err); } // return the index of PiPlatforms. diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 789398fd1a0d8..9f8537ff17ea3 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -409,16 +409,17 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, // metadata. CompileOpts += isEsimdImage ? "-doubleGRF" : "-ze-opt-large-register-file"; } - // Add optimization flags - // Assigning space for 16 characters. - char *backend_option = nullptr; - // Empty string is returned in backend_option when no appropriate backend - // option is available for a given opt level. - Plugin.getBackendOptimizationOption(optLevel, &backend_option); - if (backend_option && backend_option[0] != '\0') { - if (!CompileOpts.empty()) - CompileOpts += " "; - CompileOpts += std::string(backend_option); + // Add optimization flags. + if (optLevel != -1) { + char *backend_option = nullptr; + // Empty string is returned in backend_option when no appropriate backend + // option is available for a given opt level. + Plugin.getBackendOptimizationOption(optLevel, &backend_option); + if (backend_option && backend_option[0] != '\0') { + if (!CompileOpts.empty()) + CompileOpts += " "; + CompileOpts += std::string(backend_option); + } } if ((Plugin.getBackend() == backend::ext_oneapi_level_zero || Plugin.getBackend() == backend::opencl) && From a34e009a295288e4d54c9a5a8c04015a0fbe2e1a Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Tue, 4 Apr 2023 06:46:24 -0700 Subject: [PATCH 19/41] document correction --- sycl/doc/design/PropagateCompilerFlagsToRuntime.md | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index f926100b1072d..607e7e0ba2b63 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -63,8 +63,6 @@ must be preserved and made use of during the backend compilation. Following are the requirements for this feature. - If the user specifies `-Ox` as a front-end compile option for a particular module, this option must be preserved during backend JIT compilation. -- If the user specifies `-O0` option, SYCL runtime needs to pass the appropriate -backend option to JIT compilation stages. The following table specifies the appropriate backend options for level-zero and OpenCL backends. @@ -115,10 +113,12 @@ backend. A new plugin API has been added. It takes the optimization level as input in integer format and returns `pi_result`. The signature is as follows: -`pi_result` piPluginGetBackendOptimizationOption(int OptLevel); +`pi_result` piPluginGetBackendOptimizationOption(int OptLevel, + char **backend_option); In the level-zero and OpenCL plugins, the table provided in the 'Requirements' -section is used as a guide to identify the appropriate backend option. For other -plugins (HIP, cuda, and ESIMD emulator), empty string is returned. This API -returns `PI_SUCCESS` for valid inputs (0 <= OptLevel <= 3). For invalid inputs, -it returns `PI_ERROR_INVALID_VALUE`. +section is used as a guide to identify the appropriate backend option. +The option is returned in `backend_option`. For other plugins (HIP, cuda, and +ESIMD emulator), empty string is returned. This API returns `PI_SUCCESS` for +valid inputs (0 <= OptLevel <= 3). For invalid inputs, it returns +`PI_ERROR_INVALID_VALUE`. From e14db043791c830b84b33abe93423a123ee3fc58 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Tue, 4 Apr 2023 07:01:49 -0700 Subject: [PATCH 20/41] document correction Signed-off-by: Arvind Sudarsanam --- sycl/doc/design/PropagateCompilerFlagsToRuntime.md | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 607e7e0ba2b63..3c5e3c655662f 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -59,10 +59,11 @@ was not able to set a breakpoint inside device code. In order to support module-level debuggability, the user will compile different module files with different levels of optimization. These optimization levels -must be preserved and made use of during the backend compilation. Following are -the requirements for this feature. +must be preserved and made use of during the backend compilation. Following is a +key requirements for this feature. - If the user specifies `-Ox` as a front-end compile option for a particular -module, this option must be preserved during backend JIT compilation. +module, this option must be converted to appropriate backend option and then +propagated fo use during backend JIT compilation. The following table specifies the appropriate backend options for level-zero and OpenCL backends. From 6359987d865f1e8ae4593d7c467835488d3dfe65 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Tue, 4 Apr 2023 07:03:05 -0700 Subject: [PATCH 21/41] document correction Signed-off-by: Arvind Sudarsanam --- sycl/doc/design/PropagateCompilerFlagsToRuntime.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 3c5e3c655662f..8a2c259bbfd39 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -36,7 +36,7 @@ Here is an example that demonstrates this pain point: clang++ -c test_host.cpp -o test_host.o clang++ -c -fsycl test_device_1.cpp -o test_device_1.o clang++ -c -fsycl -g -O0 test_device_2.cpp -o test_device_2.o -clang++ -fsycl -g -o test_host.o test_device_1.o test_device_2.o -o test +clang++ -fsycl -g test_host.o test_device_1.o test_device_2.o -o test ``` In this scenario, the fat binary is 'test' and there are no compilation flags From b021163babb0ebd1ac602717df928399abef5ab0 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Tue, 4 Apr 2023 19:01:35 -0700 Subject: [PATCH 22/41] Addressing code review changes Signed-off-by: Arvind Sudarsanam --- llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp | 8 +++++--- llvm/tools/sycl-post-link/ModuleSplitter.h | 1 + sycl/doc/design/PropagateCompilerFlagsToRuntime.md | 2 ++ 3 files changed, 8 insertions(+), 3 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp b/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp index afe76f2c9876b..8cea7fd6605ec 100644 --- a/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp @@ -18,10 +18,12 @@ using namespace llvm; PreservedAnalyses SYCLAddOptLevelAttributePass::run(Module &M, ModuleAnalysisManager &MAM) { + // Here, we add a function attribute 'sycl-optlevel' to store the + // optimization level. + assert(OptLevel >= 0 && "Invalid optimization level!"); for (Function &F : M.functions()) { - // Here, we add a function attribute 'sycl-optlevel' to store the - // optimization level. - assert(OptLevel >= 0 && "Invalid optimization level!"); + if (F.isDeclaration()) + continue; F.addFnAttr("sycl-optlevel", std::to_string(OptLevel)); } return PreservedAnalyses::all(); diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index 09b983fe9e5b3..ab5d636fcb398 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -70,6 +70,7 @@ struct EntryPointGroup { Res.UsesLargeGRF = UsesLargeGRF || Other.UsesLargeGRF; // Scope remains global // OptLevel is expected to be the same for both merging EPGs + assert(OptLevel == Other.OptLevel && "OptLevels are not same"); return Res; } }; diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 8a2c259bbfd39..ff7461edf6405 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -5,6 +5,8 @@ compile different modules with different levels of optimization. This document proposes a compiler flow that will enable propagation of compiler options specified for front-end to the runtimes and eventually to the backend. Currently, only `O0`/`O1`/`O2`/`O3` options are handled. +Please note that this document only describes support for JIT path. AOT path +support will be added later. ## Background From f2d98a28526f25eefcf2412b7c83675277545a79 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 5 Apr 2023 12:29:40 -0700 Subject: [PATCH 23/41] Address more review comments Signed-off-by: Arvind Sudarsanam --- .../design/PropagateCompilerFlagsToRuntime.md | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index ff7461edf6405..55036bc10a077 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -3,7 +3,7 @@ In order to ease the process of debugging, there is a user requirement to compile different modules with different levels of optimization. This document proposes a compiler flow that will enable propagation of compiler options -specified for front-end to the runtimes and eventually to the backend. +specified from front-end to the runtimes and eventually to the backend. Currently, only `O0`/`O1`/`O2`/`O3` options are handled. Please note that this document only describes support for JIT path. AOT path support will be added later. @@ -61,8 +61,8 @@ was not able to set a breakpoint inside device code. In order to support module-level debuggability, the user will compile different module files with different levels of optimization. These optimization levels -must be preserved and made use of during the backend compilation. Following is a -key requirements for this feature. +must be preserved and made use of during the backend compilation. The following +is a key requirement for this feature. - If the user specifies `-Ox` as a front-end compile option for a particular module, this option must be converted to appropriate backend option and then propagated fo use during backend JIT compilation. @@ -99,9 +99,9 @@ of sycl-optlevel are bundled into the same device image. See also optional kernel features [design document](https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OptionalDeviceFeatures.md#changes-to-the-post-link-tool). The `sycl-post-link` tool also adds a new property into the `SYCL/misc properties` property set for each device code module. This entry will -be used to store the optimization level. Name of this property is 'optLevel' and -the value is stored as a 32-bit integer. If there is a module where user did not -specify an optimization module, there is no new entry in the property set. +be used to store the optimization level. Name of this property is `optLevel` and +the value is stored as a 32-bit integer. If there is a module where the user did +not specify an optimization module, there is no new entry in the property set. ### Changes to the SYCL runtime @@ -116,8 +116,10 @@ backend. A new plugin API has been added. It takes the optimization level as input in integer format and returns `pi_result`. The signature is as follows: -`pi_result` piPluginGetBackendOptimizationOption(int OptLevel, - char **backend_option); +```C++ +pi_result piPluginGetBackendOptimizationOption(int OptLevel, + char **backend_option); +``` In the level-zero and OpenCL plugins, the table provided in the 'Requirements' section is used as a guide to identify the appropriate backend option. From 22f034b51e63f74b24879388bdd64352c09e9e2e Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 5 Apr 2023 16:28:00 -0700 Subject: [PATCH 24/41] Change char strs in plugin to static const char strs Signed-off-by: Arvind Sudarsanam --- sycl/include/sycl/detail/pi.h | 3 ++- sycl/plugins/cuda/pi_cuda.cpp | 7 ++++--- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 4 ++-- sycl/plugins/hip/pi_hip.cpp | 7 ++++--- sycl/plugins/level_zero/pi_level_zero.cpp | 10 +++++----- sycl/plugins/opencl/pi_opencl.cpp | 6 +++--- sycl/source/detail/plugin.hpp | 2 +- sycl/source/detail/program_manager/program_manager.cpp | 2 +- 8 files changed, 22 insertions(+), 19 deletions(-) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 637b396fae9c2..13d75fa083278 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -2035,7 +2035,8 @@ __SYCL_EXPORT pi_result piPluginGetLastError(char **message); /// \return PI_SUCCESS is returned always. If a valid option is not /// available, an empty string is returned. __SYCL_EXPORT pi_result -piPluginGetBackendOptimizationOption(int opt_level, char **backend_option); +piPluginGetBackendOptimizationOption(int opt_level, + const char **backend_option); /// Queries device for it's global timestamp in nanoseconds, and updates /// HostTime with the value of the host timer at the closest possible point in diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 1dc70b9ed12ea..0bb4935355e3c 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -81,13 +81,14 @@ pi_result cuda_piPluginGetLastError(char **message) { } // Optimization strings -char EmptyStr[2] = ""; +static const char *EmptyStr = ""; // Returns plugin specific backend optimization option. // Return empty string for cuda. // TODO: Determine correct string to be passed. -pi_result cuda_piPluginGetBackendOptimizationOption(int opt_level, - char **backend_option) { +pi_result + cuda_piPluginGetBackendOptimizationOption(int opt_level, + const char **backend_option) { if ((opt_level < 0) || (opt_level > 3)) return PI_ERROR_INVALID_VALUE; *backend_option = &EmptyStr[0]; diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 8face5b1c0e8b..69b8886055060 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -168,13 +168,13 @@ pi_result piPluginGetLastError(char **message) { } // Optimization strings -char EmptyStr[2] = ""; +static const char *EmptyStr = ""; // Returns plugin specific backend optimization option. // Return empty string for esimd emulator. // TODO: Determine correct string to be passed. pi_result piPluginGetBackendOptimizationOption(int opt_level, - char **backend_option) { + const char **backend_option) { if ((opt_level < 0) || (opt_level > 3)) return PI_ERROR_INVALID_VALUE; *backend_option = &EmptyStr[0]; diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 72904a6493f22..1e8e3c97ca4f3 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -133,13 +133,14 @@ pi_result hip_piPluginGetLastError(char **message) { } // Optimization strings -char EmptyStr[2] = ""; +static const char *EmptyStr = ""; // Returns plugin specific backend optimization option. // Return empty string for hip. // TODO: Determine correct string to be passed. -pi_result hip_piPluginGetBackendOptimizationOption(int opt_level, - char **backend_option) { +pi_result + hip_piPluginGetBackendOptimizationOption(int opt_level, + const char **backend_option) { if ((opt_level < 0) || (opt_level > 3)) return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 6fdbdcfd2d8fd..379e11e21681a 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2130,17 +2130,17 @@ pi_result piPluginGetLastError(char **message) { } // Optimization strings -char EmptyStr[2] = ""; -char NoOptStr[16] = "-ze-opt-disable"; -char O1OptStr[16] = "-ze-opt-level=1"; -char O2OptStr[16] = "-ze-opt-level=2"; +static const char *EmptyStr = ""; +static const char *NoOptStr = "-ze-opt-disable"; +static const char *O1OptStr = "-ze-opt-level=1"; +static const char *O2OptStr = "-ze-opt-level=2"; // Returns plugin specific backend optimization option. // Return '-ze-opt-disable' for opt_level = 0. // Return '-ze-opt-level=1' for opt_level = 1/2. // Return '-ze-opt-level=2' for opt_level = 3. pi_result piPluginGetBackendOptimizationOption(int opt_level, - char **backend_option) { + const char **backend_option) { switch (opt_level) { case 0: *backend_option = &NoOptStr[0]; diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index df81858614d44..bafb9a180baf8 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -97,13 +97,13 @@ pi_result piPluginGetLastError(char **message) { } // Optimization strings -char EmptyStr[2] = ""; -char NoOptStr[16] = "-cl-opt-disable"; +static const char *EmptyStr = ""; +static const char *NoOptStr = "-cl-opt-disable"; // Returns plugin specific backend optimization option. // Return '-cl-opt-disable' for opt_level = 0 and '' for others. pi_result piPluginGetBackendOptimizationOption(int opt_level, - char **backend_option) { + const char **backend_option) { if ((opt_level < 0) || (opt_level > 3)) return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; diff --git a/sycl/source/detail/plugin.hpp b/sycl/source/detail/plugin.hpp index 3f530c07f98b0..f0455fc34a8f8 100644 --- a/sycl/source/detail/plugin.hpp +++ b/sycl/source/detail/plugin.hpp @@ -237,7 +237,7 @@ class plugin { // Get backend optimization option void getBackendOptimizationOption(int opt_level, - char **backend_option) const { + const char **backend_option) const { RT::PiResult Err = call_nocheck( opt_level, backend_option); diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 9f8537ff17ea3..532bc0ac71edc 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -411,7 +411,7 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, } // Add optimization flags. if (optLevel != -1) { - char *backend_option = nullptr; + const char *backend_option = nullptr; // Empty string is returned in backend_option when no appropriate backend // option is available for a given opt level. Plugin.getBackendOptimizationOption(optLevel, &backend_option); From ce810eeb67f48c35f532692a0c541a3b8955e16c Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 6 Apr 2023 10:49:05 -0700 Subject: [PATCH 25/41] Address SYCL RT review comments Signed-off-by: Arvind Sudarsanam --- .../design/PropagateCompilerFlagsToRuntime.md | 6 +-- sycl/plugins/cuda/pi_cuda.cpp | 4 +- .../esimd_emulator/pi_esimd_emulator.cpp | 4 +- sycl/plugins/hip/pi_hip.cpp | 2 +- sycl/plugins/level_zero/pi_level_zero.cpp | 10 ++--- sycl/plugins/opencl/pi_opencl.cpp | 2 +- .../program_manager/program_manager.cpp | 4 +- ...evel.cpp => sycl-opt-level-level-zero.cpp} | 32 +++++---------- .../sycl-opt-level-opencl.cpp | 40 +++++++++++++++++++ sycl/unittests/helpers/PiMockPlugin.hpp | 6 ++- 10 files changed, 72 insertions(+), 38 deletions(-) rename sycl/test-e2e/PropagateOptionsToBackend/{sycl-opt-level.cpp => sycl-opt-level-level-zero.cpp} (51%) create mode 100644 sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 55036bc10a077..3bc589318223f 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -86,9 +86,9 @@ pipeline. ### Changes to the clang front-end -For each SYCL kernel, we add a new function attribute that is named -`sycl-optlevel`. Value of this attribute is set to the optimization level which -was used to compile the overlying module. +For each function in SYCL device code, we add a new function attribute that is +named `sycl-optlevel`. Value of this attribute is set to the optimization level +which was used to compile the overlying module. ### Changes to the sycl-post-link tool diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 0bb4935355e3c..66ed98d320632 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -80,7 +80,7 @@ pi_result cuda_piPluginGetLastError(char **message) { return ErrorMessageCode; } -// Optimization strings +// Optimization strings. static const char *EmptyStr = ""; // Returns plugin specific backend optimization option. @@ -91,7 +91,7 @@ pi_result const char **backend_option) { if ((opt_level < 0) || (opt_level > 3)) return PI_ERROR_INVALID_VALUE; - *backend_option = &EmptyStr[0]; + *backend_option = EmptyStr; return PI_SUCCESS; } diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 69b8886055060..a4f292d9b9b3b 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -167,7 +167,7 @@ pi_result piPluginGetLastError(char **message) { return ErrorMessageCode; } -// Optimization strings +// Optimization strings. static const char *EmptyStr = ""; // Returns plugin specific backend optimization option. @@ -177,7 +177,7 @@ pi_result piPluginGetBackendOptimizationOption(int opt_level, const char **backend_option) { if ((opt_level < 0) || (opt_level > 3)) return PI_ERROR_INVALID_VALUE; - *backend_option = &EmptyStr[0]; + *backend_option = EmptyStr; return PI_SUCCESS; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 1e8e3c97ca4f3..5e5b494b018d9 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -132,7 +132,7 @@ pi_result hip_piPluginGetLastError(char **message) { return ErrorMessageCode; } -// Optimization strings +// Optimization strings. static const char *EmptyStr = ""; // Returns plugin specific backend optimization option. diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 379e11e21681a..9d90611ec6a1e 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2129,7 +2129,7 @@ pi_result piPluginGetLastError(char **message) { return pi2ur::piPluginGetLastError(message); } -// Optimization strings +// Optimization strings. static const char *EmptyStr = ""; static const char *NoOptStr = "-ze-opt-disable"; static const char *O1OptStr = "-ze-opt-level=1"; @@ -2143,17 +2143,17 @@ pi_result piPluginGetBackendOptimizationOption(int opt_level, const char **backend_option) { switch (opt_level) { case 0: - *backend_option = &NoOptStr[0]; + *backend_option = NoOptStr; break; case 1: case 2: - *backend_option = &O1OptStr[0]; + *backend_option = O1OptStr; break; case 3: - *backend_option = &O2OptStr[0]; + *backend_option = O2OptStr; break; default: - *backend_option = &EmptyStr[0]; + *backend_option = EmptyStr; return PI_ERROR_INVALID_VALUE; } return PI_SUCCESS; diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index bafb9a180baf8..fba92233333db 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -96,7 +96,7 @@ pi_result piPluginGetLastError(char **message) { return ErrorMessageCode; } -// Optimization strings +// Optimization strings. static const char *EmptyStr = ""; static const char *NoOptStr = "-cl-opt-disable"; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 532bc0ac71edc..c50905ef32ba4 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -410,7 +410,9 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, CompileOpts += isEsimdImage ? "-doubleGRF" : "-ze-opt-large-register-file"; } // Add optimization flags. - if (optLevel != -1) { + // Add only if compile options are not overwritten by environment + // variable + if (!CompileOptsEnv && optLevel != -1) { const char *backend_option = nullptr; // Empty string is returned in backend_option when no appropriate backend // option is available for a given opt level. diff --git a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level.cpp b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-level-zero.cpp similarity index 51% rename from sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level.cpp rename to sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-level-zero.cpp index ee2ecc4fb5a56..2c29a0ba36d38 100644 --- a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level.cpp +++ b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-level-zero.cpp @@ -1,33 +1,23 @@ -//==----------- sycl-opt-level.cpp - DPC++ SYCL on-device test -//---------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// -// This test verifies the propagation of front-end compiler optimization -// option to the backend. -// API call in device code: -// Following is expected addtion of options: -// Front-end option | OpenCL backend option | L0 backend option -// -O0 | -cl-opt-disable | -ze-opt-disable -// -O1 | /* no option */ | -ze-opt-level=1 -// -O2 | /* no option */ | -ze-opt-level=1 -// -O3 | /* no option */ | -ze-opt-level=2 +// REQUIRES: level_zero // RUN: %clangxx -O0 -fsycl %s -o %t0.out // RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 %t0.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK0 -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t0.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL0 // RUN: %clangxx -O1 -fsycl %s -o %t1.out // RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 %t1.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK1 -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t1.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL1 // RUN: %clangxx -O2 -fsycl %s -o %t2.out // RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 %t2.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK2 -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t2.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL2 // RUN: %clangxx -O3 -fsycl %s -o %t3.out // RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 %t3.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK3 -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t3.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL3 + +// This test verifies the propagation of front-end compiler optimization +// option to the backend. +// API call in device code: +// Following is expected addition of options for level_zero backend: +// Front-end option | L0 backend option +// -O0 | -ze-opt-disable +// -O1 | -ze-opt-level=1 +// -O2 | -ze-opt-level=1 +// -O3 | -ze-opt-level=2 #include diff --git a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp new file mode 100644 index 0000000000000..4442242d81630 --- /dev/null +++ b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp @@ -0,0 +1,40 @@ +// REQUIRES: opencl + +// RUN: %clangxx -O0 -fsycl %s -o %t0.out +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t0.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL0 +// RUN: %clangxx -O1 -fsycl %s -o %t1.out +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t1.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL1 +// RUN: %clangxx -O2 -fsycl %s -o %t2.out +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t2.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL2 +// RUN: %clangxx -O3 -fsycl %s -o %t3.out +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t3.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL3 + +// This test verifies the propagation of front-end compiler optimization +// option to the backend. +// API call in device code: +// Following is expected addition of options for opencl backend: +// Front-end option | OpenCL backend option +// -O0 | -cl-opt-disable +// -O1 | /* no option */ +// -O2 | /* no option */ +// -O3 | /* no option */ + +#include + +int main() { + sycl::queue q; + q.submit([&](sycl::handler &h) { h.single_task([=]() {}); }); + std::cout << "sycl-optlevel test passed\n"; + return 0; +} + +// CHECK-LABEL: ---> piProgramBuild( +// CHECK0: -ze-opt-disable +// CHECKOCL0: -cl-opt-disable +// CHECK1: -ze-opt-level=1 +// CHECKOCL1-NOT: -cl-opt-disable +// CHECK2: -ze-opt-level=1 +// CHECKOCL2-NOT: -cl-opt-disable +// CHECK3: -ze-opt-level=2 +// CHECKOCL3-NOT: -cl-opt-disable +// CHECK: ) ---> pi_result : PI_SUCCESS diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index 4a85e1346f6bc..5df05256fb53f 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -1121,8 +1121,10 @@ inline pi_result mock_piPluginGetLastError(char **message) { return PI_SUCCESS; } -inline pi_result mock_piPluginGetBackendOptimizationOption(int opt_level, - char **option) { +inline pi_result + mock_piPluginGetBackendOptimizationOption(int opt_level, + const char **option) { + *option = ""; return PI_SUCCESS; } From a9c01b741fd8b50cd61c0b321eed88c47a37b77d Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 6 Apr 2023 11:34:36 -0700 Subject: [PATCH 26/41] remove unneeded checks in tests Signed-off-by: Arvind Sudarsanam --- .../PropagateOptionsToBackend/sycl-opt-level-level-zero.cpp | 4 ---- .../PropagateOptionsToBackend/sycl-opt-level-opencl.cpp | 4 ---- 2 files changed, 8 deletions(-) diff --git a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-level-zero.cpp b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-level-zero.cpp index 2c29a0ba36d38..c1b1013039b09 100644 --- a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-level-zero.cpp +++ b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-level-zero.cpp @@ -30,11 +30,7 @@ int main() { // CHECK-LABEL: ---> piProgramBuild( // CHECK0: -ze-opt-disable -// CHECKOCL0: -cl-opt-disable // CHECK1: -ze-opt-level=1 -// CHECKOCL1-NOT: -cl-opt-disable // CHECK2: -ze-opt-level=1 -// CHECKOCL2-NOT: -cl-opt-disable // CHECK3: -ze-opt-level=2 -// CHECKOCL3-NOT: -cl-opt-disable // CHECK: ) ---> pi_result : PI_SUCCESS diff --git a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp index 4442242d81630..1a421e1584d2a 100644 --- a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp +++ b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp @@ -29,12 +29,8 @@ int main() { } // CHECK-LABEL: ---> piProgramBuild( -// CHECK0: -ze-opt-disable // CHECKOCL0: -cl-opt-disable -// CHECK1: -ze-opt-level=1 // CHECKOCL1-NOT: -cl-opt-disable -// CHECK2: -ze-opt-level=1 // CHECKOCL2-NOT: -cl-opt-disable -// CHECK3: -ze-opt-level=2 // CHECKOCL3-NOT: -cl-opt-disable // CHECK: ) ---> pi_result : PI_SUCCESS From 15f4abb2c68e9feb3a1c7cd42e8df4e0434eef1a Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 6 Apr 2023 11:38:01 -0700 Subject: [PATCH 27/41] Merge issues Signed-off-by: Arvind Sudarsanam --- sycl/include/sycl/detail/pi.def | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/sycl/include/sycl/detail/pi.def b/sycl/include/sycl/detail/pi.def index c800ad6036d0f..4c4c39cf4e628 100644 --- a/sycl/include/sycl/detail/pi.def +++ b/sycl/include/sycl/detail/pi.def @@ -155,5 +155,10 @@ _PI_API(piGetDeviceAndHostTimer) _PI_API(piextEnqueueDeviceGlobalVariableWrite) _PI_API(piextEnqueueDeviceGlobalVariableRead) +// Queue create and get APIs for immediate commandlists +_PI_API(piextQueueCreate2) +_PI_API(piextQueueGetNativeHandle2) +_PI_API(piextQueueCreateWithNativeHandle2) + _PI_API(piPluginGetBackendOptimizationOption) #undef _PI_API From 14e882a6e1c63efcd99c170f0e6cc13843039143 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 6 Apr 2023 13:31:19 -0700 Subject: [PATCH 28/41] format issue fixed Signed-off-by: Arvind Sudarsanam --- sycl/include/sycl/detail/pi.h | 5 ++--- sycl/plugins/cuda/pi_cuda.cpp | 4 ++-- sycl/plugins/hip/pi_hip.cpp | 4 ++-- sycl/unittests/helpers/PiMockPlugin.hpp | 3 +-- 4 files changed, 7 insertions(+), 9 deletions(-) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 13d75fa083278..294e1f3a29028 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -2034,9 +2034,8 @@ __SYCL_EXPORT pi_result piPluginGetLastError(char **message); /// /// \return PI_SUCCESS is returned always. If a valid option is not /// available, an empty string is returned. -__SYCL_EXPORT pi_result -piPluginGetBackendOptimizationOption(int opt_level, - const char **backend_option); +__SYCL_EXPORT pi_result piPluginGetBackendOptimizationOption( + int opt_level, const char **backend_option); /// Queries device for it's global timestamp in nanoseconds, and updates /// HostTime with the value of the host timer at the closest possible point in diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 66ed98d320632..18e7cc5f5d1c4 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -87,8 +87,8 @@ static const char *EmptyStr = ""; // Return empty string for cuda. // TODO: Determine correct string to be passed. pi_result - cuda_piPluginGetBackendOptimizationOption(int opt_level, - const char **backend_option) { +cuda_piPluginGetBackendOptimizationOption(int opt_level, + const char **backend_option) { if ((opt_level < 0) || (opt_level > 3)) return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 5e5b494b018d9..dfa3a2d703b31 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -139,8 +139,8 @@ static const char *EmptyStr = ""; // Return empty string for hip. // TODO: Determine correct string to be passed. pi_result - hip_piPluginGetBackendOptimizationOption(int opt_level, - const char **backend_option) { +hip_piPluginGetBackendOptimizationOption(int opt_level, + const char **backend_option) { if ((opt_level < 0) || (opt_level > 3)) return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index 5df05256fb53f..3418cf2a06501 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -1122,8 +1122,7 @@ inline pi_result mock_piPluginGetLastError(char **message) { } inline pi_result - mock_piPluginGetBackendOptimizationOption(int opt_level, - const char **option) { +mock_piPluginGetBackendOptimizationOption(int opt_level, const char **option) { *option = ""; return PI_SUCCESS; } From 476a2f0e5dd5761c28deb005555ddb9517567323 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 6 Apr 2023 18:45:15 -0700 Subject: [PATCH 29/41] Modify plugin changes based on user reviews Signed-off-by: Arvind Sudarsanam --- .../design/PropagateCompilerFlagsToRuntime.md | 11 ++++--- sycl/include/sycl/detail/pi.def | 7 +---- sycl/include/sycl/detail/pi.h | 9 +++--- sycl/plugins/cuda/pi_cuda.cpp | 14 ++++----- .../esimd_emulator/pi_esimd_emulator.cpp | 10 +++--- sycl/plugins/hip/pi_hip.cpp | 14 ++++----- sycl/plugins/level_zero/pi_level_zero.cpp | 31 ++++++++----------- sycl/plugins/opencl/pi_opencl.cpp | 17 +++++----- sycl/source/detail/plugin.hpp | 11 +++---- .../program_manager/program_manager.cpp | 30 +++++++++++++----- sycl/test/abi/pi_level_zero_symbol_check.dump | 2 +- sycl/test/abi/pi_opencl_symbol_check.dump | 2 +- sycl/unittests/helpers/PiMockPlugin.hpp | 7 +++-- 13 files changed, 87 insertions(+), 78 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 3bc589318223f..5964c2e7ae1ed 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -113,17 +113,18 @@ backend. ### Changes to the plugin -A new plugin API has been added. It takes the optimization level as input in -integer format and returns `pi_result`. The signature is as follows: +A new plugin API has been added. It takes the frontend option string as input in +string format and returns `pi_result`. The signature is as follows: ```C++ -pi_result piPluginGetBackendOptimizationOption(int OptLevel, - char **backend_option); +pi_result piPluginGetBackendOption(pi_platform platform, + const char *frontend_option, + const char **backend_option); ``` In the level-zero and OpenCL plugins, the table provided in the 'Requirements' section is used as a guide to identify the appropriate backend option. The option is returned in `backend_option`. For other plugins (HIP, cuda, and ESIMD emulator), empty string is returned. This API returns `PI_SUCCESS` for -valid inputs (0 <= OptLevel <= 3). For invalid inputs, it returns +valid inputs (frontend_option != ""). For invalid inputs, it returns `PI_ERROR_INVALID_VALUE`. diff --git a/sycl/include/sycl/detail/pi.def b/sycl/include/sycl/detail/pi.def index 4c4c39cf4e628..0ae3246ac3260 100644 --- a/sycl/include/sycl/detail/pi.def +++ b/sycl/include/sycl/detail/pi.def @@ -155,10 +155,5 @@ _PI_API(piGetDeviceAndHostTimer) _PI_API(piextEnqueueDeviceGlobalVariableWrite) _PI_API(piextEnqueueDeviceGlobalVariableRead) -// Queue create and get APIs for immediate commandlists -_PI_API(piextQueueCreate2) -_PI_API(piextQueueGetNativeHandle2) -_PI_API(piextQueueCreateWithNativeHandle2) - -_PI_API(piPluginGetBackendOptimizationOption) +_PI_API(piPluginGetBackendOption) #undef _PI_API diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 294e1f3a29028..c3ba14303fb25 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -2032,10 +2032,11 @@ __SYCL_EXPORT pi_result piPluginGetLastError(char **message); /// \param backend_option is used to return the backend optimization option /// corresponding to frontend optimization level. /// -/// \return PI_SUCCESS is returned always. If a valid option is not -/// available, an empty string is returned. -__SYCL_EXPORT pi_result piPluginGetBackendOptimizationOption( - int opt_level, const char **backend_option); +/// \return PI_SUCCESS is returned for valid frontend_option. If a valid backend +/// option is not available, an empty string is returned. +__SYCL_EXPORT pi_result piPluginGetBackendOption(pi_platform platform, + const char *frontend_option, + const char **backend_option); /// Queries device for it's global timestamp in nanoseconds, and updates /// HostTime with the value of the host timer at the closest possible point in diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 18e7cc5f5d1c4..9c0386db2f71d 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -83,13 +83,14 @@ pi_result cuda_piPluginGetLastError(char **message) { // Optimization strings. static const char *EmptyStr = ""; -// Returns plugin specific backend optimization option. +// Returns plugin specific backend option. +// Current support is only for optimization options. // Return empty string for cuda. // TODO: Determine correct string to be passed. -pi_result -cuda_piPluginGetBackendOptimizationOption(int opt_level, - const char **backend_option) { - if ((opt_level < 0) || (opt_level > 3)) +pi_result cuda_piPluginGetBackendOption(pi_platform platform, + const char *frontend_option, + const char **backend_option) { + if (frontend_option == nullptr || frontend_option[0] == '\0') return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; return PI_SUCCESS; @@ -5803,8 +5804,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piPluginGetLastError, cuda_piPluginGetLastError) _PI_CL(piTearDown, cuda_piTearDown) _PI_CL(piGetDeviceAndHostTimer, cuda_piGetDeviceAndHostTimer) - _PI_CL(piPluginGetBackendOptimizationOption, - cuda_piPluginGetBackendOptimizationOption) + _PI_CL(piPluginGetBackendOption, cuda_piPluginGetBackendOption) #undef _PI_CL diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index a4f292d9b9b3b..389cf79e69b68 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -170,12 +170,14 @@ pi_result piPluginGetLastError(char **message) { // Optimization strings. static const char *EmptyStr = ""; -// Returns plugin specific backend optimization option. +// Returns plugin specific backend option. +// Current support is only for optimization options. // Return empty string for esimd emulator. // TODO: Determine correct string to be passed. -pi_result piPluginGetBackendOptimizationOption(int opt_level, - const char **backend_option) { - if ((opt_level < 0) || (opt_level > 3)) +pi_result piPluginGetBackendOption(pi_platform platform, + const char *frontend_option, + const char **backend_option) { + if (frontend_option == nullptr || frontend_option[0] == '\0') return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; return PI_SUCCESS; diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index dfa3a2d703b31..f273ef48f2b5c 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -135,13 +135,14 @@ pi_result hip_piPluginGetLastError(char **message) { // Optimization strings. static const char *EmptyStr = ""; -// Returns plugin specific backend optimization option. +// Returns plugin specific backend option. +// Current support is only for optimization options. // Return empty string for hip. // TODO: Determine correct string to be passed. -pi_result -hip_piPluginGetBackendOptimizationOption(int opt_level, - const char **backend_option) { - if ((opt_level < 0) || (opt_level > 3)) +pi_result hip_piPluginGetBackendOption(pi_platform platform, + const char *frontend_option, + const char **backend_option) { + if (frontend_option == nullptr || frontend_option[0] == '\0') return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; return PI_SUCCESS; @@ -5637,8 +5638,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piPluginGetLastError, hip_piPluginGetLastError) _PI_CL(piTearDown, hip_piTearDown) _PI_CL(piGetDeviceAndHostTimer, hip_piGetDeviceAndHostTimer) - _PI_CL(piPluginGetBackendOptimizationOption, - hip_piPluginGetBackendOptimizationOption) + _PI_CL(piPluginGetBackendOption, hip_piPluginGetBackendOption) #undef _PI_CL diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 9d90611ec6a1e..bc760eedb63b5 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2135,27 +2135,22 @@ static const char *NoOptStr = "-ze-opt-disable"; static const char *O1OptStr = "-ze-opt-level=1"; static const char *O2OptStr = "-ze-opt-level=2"; -// Returns plugin specific backend optimization option. -// Return '-ze-opt-disable' for opt_level = 0. -// Return '-ze-opt-level=1' for opt_level = 1/2. -// Return '-ze-opt-level=2' for opt_level = 3. -pi_result piPluginGetBackendOptimizationOption(int opt_level, - const char **backend_option) { - switch (opt_level) { - case 0: +// Returns plugin specific backend option. +// Current support is only for optimization options. +// Return '-ze-opt-disable' for frontend_option = -O0. +// Return '-ze-opt-level=1' for frontend_option = -O1 or -O2. +// Return '-ze-opt-level=2' for frontend_option = -O3. +pi_result piPluginGetBackendOption(pi_platform platform, + const char *frontend_option, + const char **backend_option) { + if (frontend_option == nullptr || frontend_option[0] == '\0') + return PI_ERROR_INVALID_VALUE; + if (!strcmp(frontend_option, "-O0")) *backend_option = NoOptStr; - break; - case 1: - case 2: + if (!strcmp(frontend_option, "-O1") || !strcmp(frontend_option, "-O2")) *backend_option = O1OptStr; - break; - case 3: + if (!strcmp(frontend_option, "-O3")) *backend_option = O2OptStr; - break; - default: - *backend_option = EmptyStr; - return PI_ERROR_INVALID_VALUE; - } return PI_SUCCESS; } diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index fba92233333db..2ce24951379f8 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -100,14 +100,16 @@ pi_result piPluginGetLastError(char **message) { static const char *EmptyStr = ""; static const char *NoOptStr = "-cl-opt-disable"; -// Returns plugin specific backend optimization option. -// Return '-cl-opt-disable' for opt_level = 0 and '' for others. -pi_result piPluginGetBackendOptimizationOption(int opt_level, - const char **backend_option) { - if ((opt_level < 0) || (opt_level > 3)) +// Returns plugin specific backend option. +// Current support is only for optimization options. +// Return '-cl-opt-disable' for frontend_option = -O0 and '' for others. +pi_result piPluginGetBackendOption(pi_platform platform, + const char *frontend_option, + const char **backend_option) { + if (frontend_option == nullptr || frontend_option[0] == '\0') return PI_ERROR_INVALID_VALUE; *backend_option = EmptyStr; - if (opt_level == 0) + if (!strcmp(frontend_option, "-O0")) *backend_option = NoOptStr; return PI_SUCCESS; } @@ -2299,8 +2301,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piPluginGetLastError, piPluginGetLastError) _PI_CL(piTearDown, piTearDown) _PI_CL(piGetDeviceAndHostTimer, piGetDeviceAndHostTimer) - _PI_CL(piPluginGetBackendOptimizationOption, - piPluginGetBackendOptimizationOption) + _PI_CL(piPluginGetBackendOption, piPluginGetBackendOption) #undef _PI_CL diff --git a/sycl/source/detail/plugin.hpp b/sycl/source/detail/plugin.hpp index f0455fc34a8f8..72879bc0f0482 100644 --- a/sycl/source/detail/plugin.hpp +++ b/sycl/source/detail/plugin.hpp @@ -235,12 +235,11 @@ class plugin { void *getLibraryHandle() { return MLibraryHandle; } int unload() { return RT::unloadPlugin(MLibraryHandle); } - // Get backend optimization option - void getBackendOptimizationOption(int opt_level, - const char **backend_option) const { - RT::PiResult Err = - call_nocheck( - opt_level, backend_option); + // Get backend option. + void getBackendOption(pi_platform platform, const char *frontend_option, + const char **backend_option) const { + RT::PiResult Err = call_nocheck( + platform, frontend_option, backend_option); checkPiResult(Err); } diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index c50905ef32ba4..7f07c2ee4266d 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -13,6 +13,7 @@ #include #include #include +#include #include #include #include @@ -359,12 +360,18 @@ static bool getUint32PropAsBool(const RTDeviceBinaryImage &Img, return Prop && (DeviceBinaryProperty(Prop).asUint32() != 0); } -static int getUint32PropAsInt(const RTDeviceBinaryImage &Img, - const char *PropName) { +static const char *getUint32PropAsOptStr(const RTDeviceBinaryImage &Img, + const char *PropName) { pi_device_binary_property Prop = Img.getProperty(PropName); + std::stringstream ss; if (!Prop) - return -1; - return DeviceBinaryProperty(Prop).asUint32(); + return (const char *)(ss.str().c_str()); + int optLevel = DeviceBinaryProperty(Prop).asUint32(); + if (optLevel < 0 || optLevel > 3) + return (const char *)(ss.str().c_str()); + ss << "-O" << optLevel; + std::string temp = ss.str(); + return (const char *)(temp.c_str()); } static void appendCompileOptionsFromImage(std::string &CompileOpts, @@ -389,7 +396,6 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, // TODO: Remove isDoubleGRF check in next ABI break bool isLargeGRF = getUint32PropAsBool(Img, "isLargeGRF") || getUint32PropAsBool(Img, "isDoubleGRF"); - int optLevel = getUint32PropAsInt(Img, "optLevel"); // The -vc-codegen option is always preserved for ESIMD kernels, regardless // of the contents SYCL_PROGRAM_COMPILE_OPTIONS environment variable. if (isEsimdImage) { @@ -412,11 +418,19 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, // Add optimization flags. // Add only if compile options are not overwritten by environment // variable - if (!CompileOptsEnv && optLevel != -1) { + const char *optLevelStr = getUint32PropAsOptStr(Img, "optLevel"); + if (!CompileOptsEnv && optLevelStr != nullptr && optLevelStr[0] == '\0') { + // Making sure all devices have the same platform. + assert(!Devs.empty() && + std::all_of(Devs.begin(), Devs.end(), [&](const device &Dev) { + return Dev.get_platform() == Devs[0].get_platform(); + })); const char *backend_option = nullptr; // Empty string is returned in backend_option when no appropriate backend - // option is available for a given opt level. - Plugin.getBackendOptimizationOption(optLevel, &backend_option); + // option is available for a given frontend option. + Plugin.getBackendOption( + detail::getSyclObjImpl(Devs[0].get_platform())->getHandleRef(), + optLevelStr, &backend_option); if (backend_option && backend_option[0] != '\0') { if (!CompileOpts.empty()) CompileOpts += " "; diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index fe8602450010a..ef0e8fe2b7120 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -59,7 +59,7 @@ piMemRelease piMemRetain piPlatformGetInfo piPlatformsGet -piPluginGetBackendOptimizationOption +piPluginGetBackendOption piPluginGetLastError piPluginInit piProgramBuild diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index c1f8fc1b1d378..d7b9a121c1bad 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -22,7 +22,7 @@ piMemBufferCreate piMemBufferPartition piMemImageCreate piPlatformsGet -piPluginGetBackendOptimizationOption +piPluginGetBackendOption piPluginGetLastError piPluginInit piProgramCreate diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index 3418cf2a06501..11434069f9637 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -1121,9 +1121,10 @@ inline pi_result mock_piPluginGetLastError(char **message) { return PI_SUCCESS; } -inline pi_result -mock_piPluginGetBackendOptimizationOption(int opt_level, const char **option) { - *option = ""; +inline pi_result mock_piPluginGetBackendOption(pi_platform platform, + const char *frontend_option, + const char **backend_option) { + *backend_option = ""; return PI_SUCCESS; } From ad50e96259a8832bbdda5cc67be2df16816df8eb Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 6 Apr 2023 19:14:15 -0700 Subject: [PATCH 30/41] fixing unrelated format change to make tests proceed Signed-off-by: Arvind Sudarsanam --- clang/lib/CodeGen/BackendUtil.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index ded383222848a..d8cef424a9c5b 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -61,8 +61,8 @@ #include "llvm/Support/raw_ostream.h" #include "llvm/Target/TargetMachine.h" #include "llvm/Target/TargetOptions.h" -#include "llvm/Transforms/IPO/DeadArgumentElimination.h" #include "llvm/TargetParser/Triple.h" +#include "llvm/Transforms/IPO/DeadArgumentElimination.h" #include "llvm/Transforms/IPO/LowerTypeTests.h" #include "llvm/Transforms/IPO/ThinLTOBitcodeWriter.h" #include "llvm/Transforms/InstCombine/InstCombine.h" From 36dda040f7b5fe52ad66e6b8de60c44266879ece Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 6 Apr 2023 19:34:37 -0700 Subject: [PATCH 31/41] Fix unused variable warning Signed-off-by: Arvind Sudarsanam --- sycl/plugins/level_zero/pi_level_zero.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index c44cc5503fa83..6b442bff99ba4 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2189,8 +2189,10 @@ static const char *O2OptStr = "-ze-opt-level=2"; pi_result piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { - if (frontend_option == nullptr || frontend_option[0] == '\0') + if (frontend_option == nullptr || frontend_option[0] == '\0') { + *backend_option = EmptyStr; return PI_ERROR_INVALID_VALUE; + } if (!strcmp(frontend_option, "-O0")) *backend_option = NoOptStr; if (!strcmp(frontend_option, "-O1") || !strcmp(frontend_option, "-O2")) From 7f912275af84d73d26e3e5043059afb3086ddeb9 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 6 Apr 2023 23:43:06 -0700 Subject: [PATCH 32/41] Fix test fails Signed-off-by: Arvind Sudarsanam --- sycl/plugins/cuda/pi_cuda.cpp | 11 +++++++--- .../esimd_emulator/pi_esimd_emulator.cpp | 11 +++++++--- sycl/plugins/hip/pi_hip.cpp | 11 +++++++--- sycl/plugins/level_zero/pi_level_zero.cpp | 21 +++++++++++++------ sycl/plugins/opencl/pi_opencl.cpp | 18 ++++++++++++---- .../program_manager/program_manager.cpp | 4 +--- 6 files changed, 54 insertions(+), 22 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 1ba1c09aadf10..2eeddc52c09d2 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -90,10 +90,15 @@ static const char *EmptyStr = ""; pi_result cuda_piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { - if (frontend_option == nullptr || frontend_option[0] == '\0') + if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; - *backend_option = EmptyStr; - return PI_SUCCESS; + if (!strcmp(frontend_option, "-O0") || !strcmp(frontend_option, "-O1") || + !strcmp(frontend_option, "-O2") || !strcmp(frontend_option, "-O3") || + !strcmp(frontend_option, "")) { + *backend_option = EmptyStr; + return PI_SUCCESS; + } + return PI_ERROR_INVALID_VALUE; } // Iterates over the event wait list, returns correct pi_result error codes. diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index f3c73e90b2e61..3c9e58ef33f2b 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -177,10 +177,15 @@ static const char *EmptyStr = ""; pi_result piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { - if (frontend_option == nullptr || frontend_option[0] == '\0') + if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; - *backend_option = EmptyStr; - return PI_SUCCESS; + if (!strcmp(frontend_option, "-O0") || !strcmp(frontend_option, "-O1") || + !strcmp(frontend_option, "-O2") || !strcmp(frontend_option, "-O3") || + !strcmp(frontend_option, "")) { + *backend_option = EmptyStr; + return PI_SUCCESS; + } + return PI_ERROR_INVALID_VALUE; } using IDBuilder = sycl::detail::Builder; diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index f7ad33625c46c..d95e8ee059552 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -142,10 +142,15 @@ static const char *EmptyStr = ""; pi_result hip_piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { - if (frontend_option == nullptr || frontend_option[0] == '\0') + if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; - *backend_option = EmptyStr; - return PI_SUCCESS; + if (!strcmp(frontend_option, "-O0") || !strcmp(frontend_option, "-O1") || + !strcmp(frontend_option, "-O2") || !strcmp(frontend_option, "-O3") || + !strcmp(frontend_option, "")) { + *backend_option = EmptyStr; + return PI_SUCCESS; + } + return PI_ERROR_INVALID_VALUE; } // Iterates over the event wait list, returns correct pi_result error codes. diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 6b442bff99ba4..6306fe20da7a1 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2189,17 +2189,26 @@ static const char *O2OptStr = "-ze-opt-level=2"; pi_result piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { - if (frontend_option == nullptr || frontend_option[0] == '\0') { - *backend_option = EmptyStr; + if (frontend_option == nullptr) { return PI_ERROR_INVALID_VALUE; } - if (!strcmp(frontend_option, "-O0")) + if (!strcmp(frontend_option, "")) { + *backend_option = EmptyStr; + return PI_SUCCESS; + } + if (!strcmp(frontend_option, "-O0")) { *backend_option = NoOptStr; - if (!strcmp(frontend_option, "-O1") || !strcmp(frontend_option, "-O2")) + return PI_SUCCESS; + } + if (!strcmp(frontend_option, "-O1") || !strcmp(frontend_option, "-O2")) { *backend_option = O1OptStr; - if (!strcmp(frontend_option, "-O3")) + return PI_SUCCESS; + } + if (!strcmp(frontend_option, "-O3")) { *backend_option = O2OptStr; - return PI_SUCCESS; + return PI_SUCCESS; + } + return PI_ERROR_INVALID_VALUE; } pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index e84e524df5035..2b42ba9f2c906 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -106,12 +106,22 @@ static const char *NoOptStr = "-cl-opt-disable"; pi_result piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { - if (frontend_option == nullptr || frontend_option[0] == '\0') + if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; - *backend_option = EmptyStr; - if (!strcmp(frontend_option, "-O0")) + if (!strcmp(frontend_option, "")) { + *backend_option = EmptyStr; + return PI_SUCCESS; + } + if (!strcmp(frontend_option, "-O0")) { *backend_option = NoOptStr; - return PI_SUCCESS; + return PI_SUCCESS; + } + if (!strcmp(frontend_option, "-O1") || !strcmp(frontend_option, "-O2") || + !strcmp(frontend_option, "-O3")) { + *backend_option = EmptyStr; + return PI_SUCCESS; + } + return PI_ERROR_INVALID_VALUE; } static cl_int getPlatformVersion(cl_platform_id plat, diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 7f07c2ee4266d..1d98642473ee9 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -416,10 +416,8 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, CompileOpts += isEsimdImage ? "-doubleGRF" : "-ze-opt-large-register-file"; } // Add optimization flags. - // Add only if compile options are not overwritten by environment - // variable const char *optLevelStr = getUint32PropAsOptStr(Img, "optLevel"); - if (!CompileOptsEnv && optLevelStr != nullptr && optLevelStr[0] == '\0') { + if (optLevelStr != nullptr && optLevelStr[0] != '\0') { // Making sure all devices have the same platform. assert(!Devs.empty() && std::all_of(Devs.begin(), Devs.end(), [&](const device &Dev) { From 2f50ba0ed6c8e490a4b05ebe72a3736d64e2e160 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 7 Apr 2023 07:17:26 -0700 Subject: [PATCH 33/41] Turning off option passing for ESIMD images Signed-off-by: Arvind Sudarsanam --- .../program_manager/program_manager.cpp | 4 +++- .../sycl-opt-level-opencl.cpp | 20 +++++++++++++++---- 2 files changed, 19 insertions(+), 5 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 1d98642473ee9..2c435ad8505fc 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -417,7 +417,9 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, } // Add optimization flags. const char *optLevelStr = getUint32PropAsOptStr(Img, "optLevel"); - if (optLevelStr != nullptr && optLevelStr[0] != '\0') { + // TODO: Passing these options to vector compiler causes build failure in + // backend. Will pass the flags once backend compilation issue is resolved. + if (!isEsimdImage && optLevelStr != nullptr && optLevelStr[0] != '\0') { // Making sure all devices have the same platform. assert(!Devs.empty() && std::all_of(Devs.begin(), Devs.end(), [&](const device &Dev) { diff --git a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp index 1a421e1584d2a..604668344a039 100644 --- a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp +++ b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp @@ -1,13 +1,25 @@ // REQUIRES: opencl // RUN: %clangxx -O0 -fsycl %s -o %t0.out -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t0.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL0 +// RUN: %CPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t0.out 2>&1 %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL0 +// RUN: %GPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t0.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL0 // RUN: %clangxx -O1 -fsycl %s -o %t1.out -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t1.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL1 +// RUN: %CPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t1.out 2>&1 %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL1 +// RUN: %GPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t1.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL1 // RUN: %clangxx -O2 -fsycl %s -o %t2.out -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t2.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL2 +// RUN: %CPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t2.out 2>&1 %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL2 +// RUN: %GPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t2.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL2 // RUN: %clangxx -O3 -fsycl %s -o %t3.out -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu SYCL_PI_TRACE=-1 %t3.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL3 +// RUN: %CPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t3.out 2>&1 %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL3 +// RUN: %GPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t3.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL3 + +// requires: cpu, gpu, accelerator +// UNSUPPORTED: hip +// FIXME: enable the test back, see intel/llvm#8146 +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -O0 %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out // This test verifies the propagation of front-end compiler optimization // option to the backend. From e25105c68cd9298fe1b451895fcd053ab2144a59 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 7 Apr 2023 09:09:34 -0700 Subject: [PATCH 34/41] More improvements based on review comments Signed-off-by: Arvind Sudarsanam --- sycl/plugins/cuda/pi_cuda.cpp | 8 +++++--- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 8 +++++--- sycl/plugins/hip/pi_hip.cpp | 8 +++++--- sycl/plugins/level_zero/pi_level_zero.cpp | 10 ++++++---- sycl/plugins/opencl/pi_opencl.cpp | 8 +++++--- .../detail/program_manager/program_manager.cpp | 15 +++++++++------ .../sycl-opt-level-opencl.cpp | 3 --- 7 files changed, 35 insertions(+), 25 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 2eeddc52c09d2..527e82755b14e 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -25,6 +25,7 @@ #include #include #include +#include // Forward declarations void enableCUDATracing(); @@ -90,11 +91,12 @@ static const char *EmptyStr = ""; pi_result cuda_piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { + using namespace std::literals; if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; - if (!strcmp(frontend_option, "-O0") || !strcmp(frontend_option, "-O1") || - !strcmp(frontend_option, "-O2") || !strcmp(frontend_option, "-O3") || - !strcmp(frontend_option, "")) { + if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || + frontend_option == "-O2"sv || frontend_option == "-O3"sv || + frontend_option == ""sv)) { *backend_option = EmptyStr; return PI_SUCCESS; } diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 3c9e58ef33f2b..4265115e50793 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -39,6 +39,7 @@ #include #include #include +#include #include #include @@ -177,11 +178,12 @@ static const char *EmptyStr = ""; pi_result piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { + using namespace std::literals; if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; - if (!strcmp(frontend_option, "-O0") || !strcmp(frontend_option, "-O1") || - !strcmp(frontend_option, "-O2") || !strcmp(frontend_option, "-O3") || - !strcmp(frontend_option, "")) { + if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || + frontend_option == "-O2"sv || frontend_option == "-O3"sv || + frontend_option == ""sv)) { *backend_option = EmptyStr; return PI_SUCCESS; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index d95e8ee059552..f9d93c822caaa 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -24,6 +24,7 @@ #include #include #include +#include #include namespace { @@ -142,11 +143,12 @@ static const char *EmptyStr = ""; pi_result hip_piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { + using namespace std::literals; if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; - if (!strcmp(frontend_option, "-O0") || !strcmp(frontend_option, "-O1") || - !strcmp(frontend_option, "-O2") || !strcmp(frontend_option, "-O3") || - !strcmp(frontend_option, "")) { + if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || + frontend_option == "-O2"sv || frontend_option == "-O3"sv || + frontend_option == ""sv)) { *backend_option = EmptyStr; return PI_SUCCESS; } diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 6306fe20da7a1..228794b7fdf68 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -20,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -2189,22 +2190,23 @@ static const char *O2OptStr = "-ze-opt-level=2"; pi_result piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { + using namespace std::literals; if (frontend_option == nullptr) { return PI_ERROR_INVALID_VALUE; } - if (!strcmp(frontend_option, "")) { + if (frontend_option == ""sv) { *backend_option = EmptyStr; return PI_SUCCESS; } - if (!strcmp(frontend_option, "-O0")) { + if (frontend_option == "-O0"sv) { *backend_option = NoOptStr; return PI_SUCCESS; } - if (!strcmp(frontend_option, "-O1") || !strcmp(frontend_option, "-O2")) { + if (frontend_option == "-O1"sv || frontend_option == "-O2"sv) { *backend_option = O1OptStr; return PI_SUCCESS; } - if (!strcmp(frontend_option, "-O3")) { + if (frontend_option == "-O3"sv) { *backend_option = O2OptStr; return PI_SUCCESS; } diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 2b42ba9f2c906..c8e6e9ddf8319 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -29,6 +29,7 @@ #include #include #include +#include #include #define CHECK_ERR_SET_NULL_RET(err, ptr, reterr) \ @@ -106,9 +107,10 @@ static const char *NoOptStr = "-cl-opt-disable"; pi_result piPluginGetBackendOption(pi_platform platform, const char *frontend_option, const char **backend_option) { + using namespace std::literals; if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; - if (!strcmp(frontend_option, "")) { + if (frontend_option == ""sv) { *backend_option = EmptyStr; return PI_SUCCESS; } @@ -116,8 +118,8 @@ pi_result piPluginGetBackendOption(pi_platform platform, *backend_option = NoOptStr; return PI_SUCCESS; } - if (!strcmp(frontend_option, "-O1") || !strcmp(frontend_option, "-O2") || - !strcmp(frontend_option, "-O3")) { + if (frontend_option == "-O1"sv || frontend_option == "-O2"sv || + frontend_option == "-O3"sv) { *backend_option = EmptyStr; return PI_SUCCESS; } diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 2c435ad8505fc..821f50b0352ac 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -360,18 +360,18 @@ static bool getUint32PropAsBool(const RTDeviceBinaryImage &Img, return Prop && (DeviceBinaryProperty(Prop).asUint32() != 0); } -static const char *getUint32PropAsOptStr(const RTDeviceBinaryImage &Img, +static std::string getUint32PropAsOptStr(const RTDeviceBinaryImage &Img, const char *PropName) { pi_device_binary_property Prop = Img.getProperty(PropName); std::stringstream ss; if (!Prop) - return (const char *)(ss.str().c_str()); + return ""; int optLevel = DeviceBinaryProperty(Prop).asUint32(); if (optLevel < 0 || optLevel > 3) - return (const char *)(ss.str().c_str()); + return ""; ss << "-O" << optLevel; std::string temp = ss.str(); - return (const char *)(temp.c_str()); + return temp; } static void appendCompileOptionsFromImage(std::string &CompileOpts, @@ -416,10 +416,13 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, CompileOpts += isEsimdImage ? "-doubleGRF" : "-ze-opt-large-register-file"; } // Add optimization flags. - const char *optLevelStr = getUint32PropAsOptStr(Img, "optLevel"); + const char *optLevelStr = getUint32PropAsOptStr(Img, "optLevel").c_str(); // TODO: Passing these options to vector compiler causes build failure in // backend. Will pass the flags once backend compilation issue is resolved. - if (!isEsimdImage && optLevelStr != nullptr && optLevelStr[0] != '\0') { + // Update only if compile options are not overwritten by environment + // variable. + if (!isEsimdImage && !CompileOptsEnv && optLevelStr != nullptr && + optLevelStr[0] != '\0') { // Making sure all devices have the same platform. assert(!Devs.empty() && std::all_of(Devs.begin(), Devs.end(), [&](const device &Dev) { diff --git a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp index 604668344a039..52a17f7a2cb13 100644 --- a/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp +++ b/sycl/test-e2e/PropagateOptionsToBackend/sycl-opt-level-opencl.cpp @@ -13,9 +13,6 @@ // RUN: %CPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t3.out 2>&1 %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL3 // RUN: %GPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 %t3.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECKOCL3 -// requires: cpu, gpu, accelerator -// UNSUPPORTED: hip -// FIXME: enable the test back, see intel/llvm#8146 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -O0 %s -o %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out From 8ffca4b9a207e54f9ed35ba03bf25524ed66a815 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 7 Apr 2023 09:23:12 -0700 Subject: [PATCH 35/41] Fix format issues Signed-off-by: Arvind Sudarsanam --- sycl/plugins/cuda/pi_cuda.cpp | 6 +++--- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 6 +++--- sycl/plugins/hip/pi_hip.cpp | 8 ++++---- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 527e82755b14e..553c36e0edb51 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -97,9 +97,9 @@ pi_result cuda_piPluginGetBackendOption(pi_platform platform, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv)) { - *backend_option = EmptyStr; - return PI_SUCCESS; - } + *backend_option = EmptyStr; + return PI_SUCCESS; + } return PI_ERROR_INVALID_VALUE; } diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 4265115e50793..afdd62fd6939b 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -184,9 +184,9 @@ pi_result piPluginGetBackendOption(pi_platform platform, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv)) { - *backend_option = EmptyStr; - return PI_SUCCESS; - } + *backend_option = EmptyStr; + return PI_SUCCESS; + } return PI_ERROR_INVALID_VALUE; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index f9d93c822caaa..cdfff33df7101 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -24,8 +24,8 @@ #include #include #include -#include #include +#include namespace { // Hipify doesn't support cuArrayGetDescriptor, on AMD the hipArray can just be @@ -149,9 +149,9 @@ pi_result hip_piPluginGetBackendOption(pi_platform platform, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv)) { - *backend_option = EmptyStr; - return PI_SUCCESS; - } + *backend_option = EmptyStr; + return PI_SUCCESS; + } return PI_ERROR_INVALID_VALUE; } From f1fadc46144712c5ca2f2cb7ca439075ab7c090f Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 7 Apr 2023 09:46:42 -0700 Subject: [PATCH 36/41] More review comments addressed Signed-off-by: Arvind Sudarsanam --- sycl/include/sycl/detail/pi.h | 8 ++++---- sycl/source/detail/program_manager/program_manager.cpp | 3 ++- 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index e101e6e27c37e..bcc8ec8fe9b73 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -2063,10 +2063,10 @@ __SYCL_EXPORT pi_result piTearDown(void *PluginParameter); /// timestamp __SYCL_EXPORT pi_result piPluginGetLastError(char **message); -/// API to get backend specific optimization option. -/// \param opt_level is an integer that contains frontend optimization level. -/// \param backend_option is used to return the backend optimization option -/// corresponding to frontend optimization level. +/// API to get backend specific option. +/// \param frontend_option is a string that contains frontend option. +/// \param backend_option is used to return the backend option corresponding to +/// frontend option. /// /// \return PI_SUCCESS is returned for valid frontend_option. If a valid backend /// option is not available, an empty string is returned. diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 821f50b0352ac..a84bc290430ea 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -416,7 +416,8 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, CompileOpts += isEsimdImage ? "-doubleGRF" : "-ze-opt-large-register-file"; } // Add optimization flags. - const char *optLevelStr = getUint32PropAsOptStr(Img, "optLevel").c_str(); + auto str = getUint32PropAsOptStr(Img, "optLevel"); + const char *optLevelStr = str.c_str(); // TODO: Passing these options to vector compiler causes build failure in // backend. Will pass the flags once backend compilation issue is resolved. // Update only if compile options are not overwritten by environment From ef009ff8861fed50231beea7ff0ac4e645d8747f Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 7 Apr 2023 10:05:31 -0700 Subject: [PATCH 37/41] Minor changes in plugin code Signed-off-by: Arvind Sudarsanam --- sycl/plugins/cuda/pi_cuda.cpp | 5 +---- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 5 +---- sycl/plugins/hip/pi_hip.cpp | 5 +---- sycl/plugins/level_zero/pi_level_zero.cpp | 14 ++++---------- sycl/plugins/opencl/pi_opencl.cpp | 10 +++------- 5 files changed, 10 insertions(+), 29 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 553c36e0edb51..51e96b12b8035 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -81,9 +81,6 @@ pi_result cuda_piPluginGetLastError(char **message) { return ErrorMessageCode; } -// Optimization strings. -static const char *EmptyStr = ""; - // Returns plugin specific backend option. // Current support is only for optimization options. // Return empty string for cuda. @@ -97,7 +94,7 @@ pi_result cuda_piPluginGetBackendOption(pi_platform platform, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv)) { - *backend_option = EmptyStr; + *backend_option = ""; return PI_SUCCESS; } return PI_ERROR_INVALID_VALUE; diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index afdd62fd6939b..71228c3e5bcc1 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -168,9 +168,6 @@ pi_result piPluginGetLastError(char **message) { return ErrorMessageCode; } -// Optimization strings. -static const char *EmptyStr = ""; - // Returns plugin specific backend option. // Current support is only for optimization options. // Return empty string for esimd emulator. @@ -184,7 +181,7 @@ pi_result piPluginGetBackendOption(pi_platform platform, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv)) { - *backend_option = EmptyStr; + *backend_option = ""; return PI_SUCCESS; } return PI_ERROR_INVALID_VALUE; diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index cdfff33df7101..d650716dc8e42 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -133,9 +133,6 @@ pi_result hip_piPluginGetLastError(char **message) { return ErrorMessageCode; } -// Optimization strings. -static const char *EmptyStr = ""; - // Returns plugin specific backend option. // Current support is only for optimization options. // Return empty string for hip. @@ -149,7 +146,7 @@ pi_result hip_piPluginGetBackendOption(pi_platform platform, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv)) { - *backend_option = EmptyStr; + *backend_option = ""; return PI_SUCCESS; } return PI_ERROR_INVALID_VALUE; diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 228794b7fdf68..bf06eaa70e8e7 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2176,12 +2176,6 @@ pi_result piPluginGetLastError(char **message) { return pi2ur::piPluginGetLastError(message); } -// Optimization strings. -static const char *EmptyStr = ""; -static const char *NoOptStr = "-ze-opt-disable"; -static const char *O1OptStr = "-ze-opt-level=1"; -static const char *O2OptStr = "-ze-opt-level=2"; - // Returns plugin specific backend option. // Current support is only for optimization options. // Return '-ze-opt-disable' for frontend_option = -O0. @@ -2195,19 +2189,19 @@ pi_result piPluginGetBackendOption(pi_platform platform, return PI_ERROR_INVALID_VALUE; } if (frontend_option == ""sv) { - *backend_option = EmptyStr; + *backend_option = ""; return PI_SUCCESS; } if (frontend_option == "-O0"sv) { - *backend_option = NoOptStr; + *backend_option = "-ze-opt-disable"; return PI_SUCCESS; } if (frontend_option == "-O1"sv || frontend_option == "-O2"sv) { - *backend_option = O1OptStr; + *backend_option = "-ze-opt-level=1"; return PI_SUCCESS; } if (frontend_option == "-O3"sv) { - *backend_option = O2OptStr; + *backend_option = "-ze-opt-level=2"; return PI_SUCCESS; } return PI_ERROR_INVALID_VALUE; diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index c8e6e9ddf8319..76d4a962c55bb 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -97,10 +97,6 @@ pi_result piPluginGetLastError(char **message) { return ErrorMessageCode; } -// Optimization strings. -static const char *EmptyStr = ""; -static const char *NoOptStr = "-cl-opt-disable"; - // Returns plugin specific backend option. // Current support is only for optimization options. // Return '-cl-opt-disable' for frontend_option = -O0 and '' for others. @@ -111,16 +107,16 @@ pi_result piPluginGetBackendOption(pi_platform platform, if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; if (frontend_option == ""sv) { - *backend_option = EmptyStr; + *backend_option = ""; return PI_SUCCESS; } if (!strcmp(frontend_option, "-O0")) { - *backend_option = NoOptStr; + *backend_option = "-cl-opt-disable"; return PI_SUCCESS; } if (frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv) { - *backend_option = EmptyStr; + *backend_option = ""; return PI_SUCCESS; } return PI_ERROR_INVALID_VALUE; From b5b1b7548e65bb3cc83a8fe2b0e3c20991bbc19c Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 7 Apr 2023 10:53:40 -0700 Subject: [PATCH 38/41] Removed unused variable and extra parentheses Signed-off-by: Arvind Sudarsanam --- sycl/plugins/cuda/pi_cuda.cpp | 4 ++-- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 5 ++--- sycl/plugins/hip/pi_hip.cpp | 5 ++--- sycl/plugins/level_zero/pi_level_zero.cpp | 3 +-- sycl/plugins/opencl/pi_opencl.cpp | 3 +-- 5 files changed, 8 insertions(+), 12 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 51e96b12b8035..eb80d16e1b30e 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -85,7 +85,7 @@ pi_result cuda_piPluginGetLastError(char **message) { // Current support is only for optimization options. // Return empty string for cuda. // TODO: Determine correct string to be passed. -pi_result cuda_piPluginGetBackendOption(pi_platform platform, +pi_result cuda_piPluginGetBackendOption(pi_platform, const char *frontend_option, const char **backend_option) { using namespace std::literals; @@ -93,7 +93,7 @@ pi_result cuda_piPluginGetBackendOption(pi_platform platform, return PI_ERROR_INVALID_VALUE; if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || - frontend_option == ""sv)) { + frontend_option == ""sv) { *backend_option = ""; return PI_SUCCESS; } diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 71228c3e5bcc1..3c06887de0dd1 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -172,15 +172,14 @@ pi_result piPluginGetLastError(char **message) { // Current support is only for optimization options. // Return empty string for esimd emulator. // TODO: Determine correct string to be passed. -pi_result piPluginGetBackendOption(pi_platform platform, - const char *frontend_option, +pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option, const char **backend_option) { using namespace std::literals; if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || - frontend_option == ""sv)) { + frontend_option == ""sv) { *backend_option = ""; return PI_SUCCESS; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index d650716dc8e42..5b40d51ba43c5 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -137,15 +137,14 @@ pi_result hip_piPluginGetLastError(char **message) { // Current support is only for optimization options. // Return empty string for hip. // TODO: Determine correct string to be passed. -pi_result hip_piPluginGetBackendOption(pi_platform platform, - const char *frontend_option, +pi_result hip_piPluginGetBackendOption(pi_platform, const char *frontend_option, const char **backend_option) { using namespace std::literals; if (frontend_option == nullptr) return PI_ERROR_INVALID_VALUE; if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || - frontend_option == ""sv)) { + frontend_option == ""sv) { *backend_option = ""; return PI_SUCCESS; } diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index bf06eaa70e8e7..8a30716032125 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -2181,8 +2181,7 @@ pi_result piPluginGetLastError(char **message) { // Return '-ze-opt-disable' for frontend_option = -O0. // Return '-ze-opt-level=1' for frontend_option = -O1 or -O2. // Return '-ze-opt-level=2' for frontend_option = -O3. -pi_result piPluginGetBackendOption(pi_platform platform, - const char *frontend_option, +pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option, const char **backend_option) { using namespace std::literals; if (frontend_option == nullptr) { diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 76d4a962c55bb..681c716dfe01b 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -100,8 +100,7 @@ pi_result piPluginGetLastError(char **message) { // Returns plugin specific backend option. // Current support is only for optimization options. // Return '-cl-opt-disable' for frontend_option = -O0 and '' for others. -pi_result piPluginGetBackendOption(pi_platform platform, - const char *frontend_option, +pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option, const char **backend_option) { using namespace std::literals; if (frontend_option == nullptr) From 3937a2a7773cd69586e6ab6b658927189703b2ac Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 7 Apr 2023 11:22:23 -0700 Subject: [PATCH 39/41] Minor format issue Signed-off-by: Arvind Sudarsanam --- sycl/plugins/cuda/pi_cuda.cpp | 6 +++--- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 6 +++--- sycl/plugins/hip/pi_hip.cpp | 6 +++--- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index eb80d16e1b30e..c55bfcf373e52 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -94,9 +94,9 @@ pi_result cuda_piPluginGetBackendOption(pi_platform, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv) { - *backend_option = ""; - return PI_SUCCESS; - } + *backend_option = ""; + return PI_SUCCESS; + } return PI_ERROR_INVALID_VALUE; } diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 3c06887de0dd1..89d8ff14ab4ef 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -180,9 +180,9 @@ pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv) { - *backend_option = ""; - return PI_SUCCESS; - } + *backend_option = ""; + return PI_SUCCESS; + } return PI_ERROR_INVALID_VALUE; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 5b40d51ba43c5..0abe0424db884 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -145,9 +145,9 @@ pi_result hip_piPluginGetBackendOption(pi_platform, const char *frontend_option, if (frontend_option == "-O0"sv || frontend_option == "-O1"sv || frontend_option == "-O2"sv || frontend_option == "-O3"sv || frontend_option == ""sv) { - *backend_option = ""; - return PI_SUCCESS; - } + *backend_option = ""; + return PI_SUCCESS; + } return PI_ERROR_INVALID_VALUE; } From d038d4ee7d648a3aad58c68587e6d6e9099f7e7d Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 7 Apr 2023 14:24:13 -0700 Subject: [PATCH 40/41] Minor typos fixed Signed-off-by: Arvind Sudarsanam --- llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h | 4 ++-- llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h b/llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h index d7d450715b973..4f77568611ec3 100644 --- a/llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h +++ b/llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h @@ -1,4 +1,4 @@ -//===---- SYCLAddOptLevelAttribute.cpp - SYCLAddOptLevelAttribute Pass --===// +//===----- SYCLAddOptLevelAttribute.h - SYCLAddOptLevelAttribute Pass -----===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -30,4 +30,4 @@ class SYCLAddOptLevelAttributePass } // namespace llvm -#endif // LLVM_SYCL_ADD_OPT_LEVEL_ATTRIBUTE_H \ No newline at end of file +#endif // LLVM_SYCL_ADD_OPT_LEVEL_ATTRIBUTE_H diff --git a/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp b/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp index 8cea7fd6605ec..dc0620ccf87e9 100644 --- a/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp @@ -27,4 +27,4 @@ SYCLAddOptLevelAttributePass::run(Module &M, ModuleAnalysisManager &MAM) { F.addFnAttr("sycl-optlevel", std::to_string(OptLevel)); } return PreservedAnalyses::all(); -} \ No newline at end of file +} From a1ce60f82be89ed35e0f4ae74d9510f1003b9353 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 13 Apr 2023 06:53:12 -0700 Subject: [PATCH 41/41] Minor doc changes Signed-off-by: Arvind Sudarsanam --- sycl/doc/design/PropagateCompilerFlagsToRuntime.md | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md index 5964c2e7ae1ed..3981151782140 100644 --- a/sycl/doc/design/PropagateCompilerFlagsToRuntime.md +++ b/sycl/doc/design/PropagateCompilerFlagsToRuntime.md @@ -106,15 +106,19 @@ not specify an optimization module, there is no new entry in the property set. ### Changes to the SYCL runtime In the SYCL runtime, the device image properties can be accessed to extract the -associated optimization level. Once the optimization level is available, a query -is made to the plugin to identify the correct backend option. This backend -option is added to the existing list of compiler options and is sent to the -backend. +associated optimization level. Once the optimization level is available, it is +converted to its equivalent frontend option string +(`-O0`, `-O1`, `-O2`, or `-O3`). This frontend option string is passed into a +query that is made to the plugin to identify the correct backend option. This +backend option is added to the existing list of compiler options and is sent to +the backend. ### Changes to the plugin A new plugin API has been added. It takes the frontend option string as input in -string format and returns `pi_result`. The signature is as follows: +string format and returns `pi_result`. A string format is used for sending the +frontend option so that this API can be used for querying other frontend +options as well. The signature of this API is as follows: ```C++ pi_result piPluginGetBackendOption(pi_platform platform,