diff --git a/clang/test/Driver/clang-offload-wrapper.c b/clang/test/Driver/clang-offload-wrapper.c
index af04576846b50..84506d1fa2b57 100644
--- a/clang/test/Driver/clang-offload-wrapper.c
+++ b/clang/test/Driver/clang-offload-wrapper.c
@@ -5,52 +5,64 @@
//
// RUN: clang-offload-wrapper --help | FileCheck %s --check-prefix CHECK-HELP
-// CHECK-HELP: {{.*}}OVERVIEW: A tool to create a wrapper bitcode for offload target binaries.
-// CHECK-HELP: {{.*}}Takes offload target binaries and optional manifest files as input
-// CHECK-HELP: {{.*}}and produces bitcode file containing target binaries packaged as data
-// CHECK-HELP: {{.*}}and initialization code which registers target binaries in the offload
-// CHECK-HELP: {{.*}}runtime. Manifest files format and contents are not restricted and are
-// CHECK-HELP: {{.*}}a subject of agreement between the device compiler and the native
-// CHECK-HELP: {{.*}}runtime for that device. When present, manifest file name should
-// CHECK-HELP: {{.*}}immediately follow the corresponding device image filename on the
-// CHECK-HELP: {{.*}}command line. Options annotating a device binary have effect on all
-// CHECK-HELP: {{.*}}subsequent input, until redefined. For example:
-// CHECK-HELP: {{.*}}$clang-offload-wrapper -host x86_64-pc-linux-gnu \
-// CHECK-HELP: {{.*}} -kind=sycl -target=spir64 -format=spirv -build-opts=-g \
-// CHECK-HELP: {{.*}} a.spv a_mf.txt \
-// CHECK-HELP: {{.*}} -target=xxx -format=native -build-opts="" \
-// CHECK-HELP: {{.*}} b.bin b_mf.txt \
-// CHECK-HELP: {{.*}} -kind=openmp \
-// CHECK-HELP: {{.*}} c.bin
-// CHECK-HELP: {{.*}}will generate an x86 wrapper object (.bc) enclosing the following
-// CHECK-HELP: {{.*}}tuples describing a single device binary each ('-' means 'none')
-// CHECK-HELP: {{.*}}offload kind | target | data format | data | manifest | build options:
-// CHECK-HELP: {{.*}}----------------------------------------------------------------------
-// CHECK-HELP: {{.*}} sycl | spir64 | spirv | a.spv| a_mf.txt | -g
-// CHECK-HELP: {{.*}} sycl | xxx | native | b.bin| b_mf.txt | -
-// CHECK-HELP: {{.*}} openmp | xxx | native | c.bin| - | -
-// CHECK-HELP: {{.*}}USAGE: clang-offload-wrapper [options]
-// CHECK-HELP: {{.*}}OPTIONS:
-// CHECK-HELP: {{.*}}clang-offload-wrapper options:
-// CHECK-HELP: {{.*}} --build-opts= - build options passed to the offload runtime
-// CHECK-HELP: {{.*}} --desc-name= - Specifies offload descriptor symbol name: '..', and makes it globally visible
-// CHECK-HELP: {{.*}} --emit-reg-funcs - Emit [un-]registration functions
-// CHECK-HELP: {{.*}} --entries= - File listing all offload function entries, SYCL offload only
-// CHECK-HELP: {{.*}} --format= - device binary image formats:
-// CHECK-HELP: {{.*}} =none - not set
-// CHECK-HELP: {{.*}} =native - unknown or native
-// CHECK-HELP: {{.*}} =spirv - SPIRV binary
-// CHECK-HELP: {{.*}} =llvmbc - LLVMIR bitcode
-// CHECK-HELP: {{.*}} --host= - Target triple for the output module
-// CHECK-HELP: {{.*}} --kind= - offload kind:
-// CHECK-HELP: {{.*}} =unknown - unknown
-// CHECK-HELP: {{.*}} =host - host
-// CHECK-HELP: {{.*}} =openmp - OpenMP
-// CHECK-HELP: {{.*}} =hip - HIP
-// CHECK-HELP: {{.*}} =sycl - SYCL
-// CHECK-HELP: {{.*}} -o= - Output filename
-// CHECK-HELP: {{.*}} --target= - offload target triple
-// CHECK-HELP: {{.*}} -v - verbose output
+// CHECK-HELP: A tool to create a wrapper bitcode for offload target binaries.
+// CHECK-HELP: Takes offload target binaries and optional manifest files as input
+// CHECK-HELP: and produces bitcode file containing target binaries packaged as data
+// CHECK-HELP: and initialization code which registers target binaries in the offload
+// CHECK-HELP: runtime. Manifest files format and contents are not restricted and are
+// CHECK-HELP: a subject of agreement between the device compiler and the native
+// CHECK-HELP: runtime for that device. When present, manifest file name should
+// CHECK-HELP: immediately follow the corresponding device image filename on the
+// CHECK-HELP: command line. Options annotating a device binary have effect on all
+// CHECK-HELP: subsequent input, until redefined.
+// CHECK-HELP: For example:
+// CHECK-HELP: clang-offload-wrapper
+// CHECK-HELP: -host x86_64-pc-linux-gnu
+// CHECK-HELP: -kind=sycl
+// CHECK-HELP: -target=spir64
+// CHECK-HELP: -format=spirv
+// CHECK-HELP: -compile-opts=-g
+// CHECK-HELP: -link-opts=-cl-denorms-are-zero
+// CHECK-HELP: a.spv
+// CHECK-HELP: a_mf.txt
+// CHECK-HELP: -target=xxx
+// CHECK-HELP: -format=native
+// CHECK-HELP: -compile-opts=""
+// CHECK-HELP: -link-opts=""
+// CHECK-HELP: b.bin
+// CHECK-HELP: b_mf.txt
+// CHECK-HELP: -kind=openmp
+// CHECK-HELP: c.bin
+// CHECK-HELP: This command generates an x86 wrapper object (.bc) enclosing the
+// CHECK-HELP: following tuples describing a single device binary each:
+// CHECK-HELP: offload kind | target | data format | data | manifest | build options:
+// CHECK-HELP: ----------------------------------------------------------------------
+// CHECK-HELP: sycl | spir64 | spirv | a.spv| a_mf.txt | -g
+// CHECK-HELP: sycl | xxx | native | b.bin| b_mf.txt | -
+// CHECK-HELP: openmp | xxx | native | c.bin| n/a | -
+// CHECK-HELP: USAGE: clang-offload-wrapper [options]
+// CHECK-HELP: OPTIONS:
+// CHECK-HELP: clang-offload-wrapper options:
+// CHECK-HELP: --compile-opts= - compile options passed to the offload runtime
+// CHECK-HELP: --desc-name= - Specifies offload descriptor symbol name: '..', and makes it globally visible
+// CHECK-HELP: --emit-reg-funcs - Emit [un-]registration functions
+// CHECK-HELP: --entries= - File listing all offload function entries, SYCL offload only
+// CHECK-HELP: --format= - device binary image formats:
+// CHECK-HELP: =none - not set
+// CHECK-HELP: =native - unknown or native
+// CHECK-HELP: =spirv - SPIRV binary
+// CHECK-HELP: =llvmbc - LLVMIR bitcode
+// CHECK-HELP: --host= - Target triple for the output module
+// CHECK-HELP: --kind= - offload kind:
+// CHECK-HELP: =unknown - unknown
+// CHECK-HELP: =host - host
+// CHECK-HELP: =openmp - OpenMP
+// CHECK-HELP: =hip - HIP
+// CHECK-HELP: =sycl - SYCL
+// CHECK-HELP: --link-opts= - link options passed to the offload runtime
+// CHECK-HELP: -o= - Output filename
+// CHECK-HELP: --target= - offload target triple
+// CHECK-HELP: -v - verbose output
// -------
// Generate files to wrap.
@@ -66,8 +78,10 @@
// RUN: clang-offload-wrapper \
// RUN: -host=x86_64-pc-linux-gnu \
// RUN: -kind=openmp -target=tg2 -format=native %t3.tgt %t1_mf.txt \
-// RUN: -kind=sycl -target=tg1 -build-opts=-g -format spirv %t1.tgt \
-// RUN: -target=tg2 -build-opts= -format native %t2.tgt \
+// RUN: -kind=sycl -target=tg1 -compile-opts=-g -link-opts=-cl-denorms-are-zero \
+// RUN: -format spirv %t1.tgt \
+// RUN: -target=tg2 -compile-opts= -link-opts= \
+// RUN: -format native %t2.tgt \
// RUN: -o %t.wrapper.bc
// RUN: llvm-dis %t.wrapper.bc -o - | FileCheck %s --check-prefix CHECK-IR
@@ -77,7 +91,7 @@
// CHECK-IR-DAG: [[IMAGETY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
// CHECK-IR-DAG: [[DESCTY:%.+]] = type { i32, [[IMAGETY]]*, [[ENTTY]]*, [[ENTTY]]* }
-// CHECK-IR-DAG: [[SYCL_IMAGETY:%.+]] = type { i16, i8, i8, i8*, i8*, i8*, i8*, i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
+// CHECK-IR-DAG: [[SYCL_IMAGETY:%.+]] = type { i16, i8, i8, i8*, i8*, i8*, i8*, i8*, i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
// CHECK-IR-DAG: [[SYCL_DESCTY:%.+]] = type { i16, i16, [[SYCL_IMAGETY]]*, [[ENTTY]]*, [[ENTTY]]* }
// CHECK-IR: [[ENTBEGIN:@.+]] = external hidden constant [[ENTTY]]
@@ -92,14 +106,16 @@
// CHECK-IR: [[OMP_DESC:@.+]] = internal constant [[DESCTY]] { i32 1, [[IMAGETY]]* getelementptr inbounds ([1 x [[IMAGETY]]], [1 x [[IMAGETY]]]* [[OMP_IMAGES]], i64 0, i64 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }
// CHECK-IR: [[SYCL_TGT0:@.+]] = internal unnamed_addr constant [4 x i8] c"tg1\00"
-// CHECK-IR: [[SYCL_OPTS0:@.+]] = internal unnamed_addr constant [3 x i8] c"-g\00"
+// CHECK-IR: [[SYCL_COMPILE_OPTS0:@.+]] = internal unnamed_addr constant [3 x i8] c"-g\00"
+// CHECK-IR: [[SYCL_LINK_OPTS0:@.+]] = internal unnamed_addr constant [21 x i8] c"-cl-denorms-are-zero\00"
// CHECK-IR: [[SYCL_BIN0:@.+]] = internal unnamed_addr constant [[SYCL_BIN0TY:\[[0-9]+ x i8\]]] c"Content of device file1{{.+}}"
// CHECK-IR: [[SYCL_TGT1:@.+]] = internal unnamed_addr constant [4 x i8] c"tg2\00"
-// CHECK-IR: [[SYCL_OPTS1:@.+]] = internal unnamed_addr constant [1 x i8] zeroinitializer
+// CHECK-IR: [[SYCL_COMPILE_OPTS1:@.+]] = internal unnamed_addr constant [1 x i8] zeroinitializer
+// CHECK-IR: [[SYCL_LINK_OPTS1:@.+]] = internal unnamed_addr constant [1 x i8] zeroinitializer
// CHECK-IR: [[SYCL_BIN1:@.+]] = internal unnamed_addr constant [[SYCL_BIN1TY:\[[0-9]+ x i8\]]] c"Content of device file2{{.+}}"
-// CHECK-IR: [[SYCL_IMAGES:@.+]] = internal unnamed_addr constant [2 x [[SYCL_IMAGETY]]] [{{.+}} { i16 1, i8 4, i8 2, i8* getelementptr inbounds ([4 x i8], [4 x i8]* [[SYCL_TGT0]], i64 0, i64 0), i8* getelementptr inbounds ([3 x i8], [3 x i8]* [[SYCL_OPTS0]], i64 0, i64 0), i8* null, i8* null, i8* getelementptr inbounds ([[SYCL_BIN0TY]], [[SYCL_BIN0TY]]* [[SYCL_BIN0]], i64 0, i64 0), i8* getelementptr inbounds ([[SYCL_BIN0TY]], [[SYCL_BIN0TY]]* [[SYCL_BIN0]], i64 1, i64 0), [[ENTTY]]* null, [[ENTTY]]* null }, [[SYCL_IMAGETY]] { i16 1, i8 4, i8 1, i8* getelementptr inbounds ([4 x i8], [4 x i8]* [[SYCL_TGT1]], i64 0, i64 0), i8* getelementptr inbounds ([1 x i8], [1 x i8]* [[SYCL_OPTS1]], i64 0, i64 0), i8* null, i8* null, i8* getelementptr inbounds ([[SYCL_BIN1TY]], [[SYCL_BIN1TY]]* [[SYCL_BIN1]], i64 0, i64 0), i8* getelementptr inbounds ([[SYCL_BIN1TY]], [[SYCL_BIN1TY]]* [[SYCL_BIN1]], i64 1, i64 0), [[ENTTY]]* null, [[ENTTY]]* null }]
+// CHECK-IR: [[SYCL_IMAGES:@.+]] = internal unnamed_addr constant [2 x [[SYCL_IMAGETY]]] [{{.+}} { i16 1, i8 4, i8 2, i8* getelementptr inbounds ([4 x i8], [4 x i8]* [[SYCL_TGT0]], i64 0, i64 0), i8* getelementptr inbounds ([3 x i8], [3 x i8]* [[SYCL_COMPILE_OPTS0]], i64 0, i64 0), i8* getelementptr inbounds ([21 x i8], [21 x i8]* [[SYCL_LINK_OPTS0]], i64 0, i64 0), i8* null, i8* null, i8* getelementptr inbounds ([[SYCL_BIN0TY]], [[SYCL_BIN0TY]]* [[SYCL_BIN0]], i64 0, i64 0), i8* getelementptr inbounds ([[SYCL_BIN0TY]], [[SYCL_BIN0TY]]* [[SYCL_BIN0]], i64 1, i64 0), [[ENTTY]]* null, [[ENTTY]]* null }, [[SYCL_IMAGETY]] { i16 1, i8 4, i8 1, i8* getelementptr inbounds ([4 x i8], [4 x i8]* [[SYCL_TGT1]], i64 0, i64 0), i8* getelementptr inbounds ([1 x i8], [1 x i8]* [[SYCL_COMPILE_OPTS1]], i64 0, i64 0), i8* getelementptr inbounds ([1 x i8], [1 x i8]* [[SYCL_LINK_OPTS1]], i64 0, i64 0), i8* null, i8* null, i8* getelementptr inbounds ([[SYCL_BIN1TY]], [[SYCL_BIN1TY]]* [[SYCL_BIN1]], i64 0, i64 0), i8* getelementptr inbounds ([[SYCL_BIN1TY]], [[SYCL_BIN1TY]]* [[SYCL_BIN1]], i64 1, i64 0), [[ENTTY]]* null, [[ENTTY]]* null }]
// CHECK-IR: [[SYCL_DESC:@.+]] = internal constant [[SYCL_DESCTY]] { i16 1, i16 2, [[SYCL_IMAGETY]]* getelementptr inbounds ([2 x [[SYCL_IMAGETY]]], [2 x [[SYCL_IMAGETY]]]* [[SYCL_IMAGES]], i64 0, i64 0), [[ENTTY]]* null, [[ENTTY]]* null }
@@ -138,7 +154,7 @@
//
// RUN: clang-offload-wrapper -kind sycl -host=x86_64-pc-linux-gnu -emit-reg-funcs=0 -desc-name=lalala -o - %t.tgt | llvm-dis | FileCheck %s --check-prefix CHECK-IR1
// CHECK-IR1: source_filename = "offload.wrapper.object"
-// CHECK-IR1: [[IMAGETY:%.+]] = type { i16, i8, i8, i8*, i8*, i8*, i8*, i8*, i8*, %__tgt_offload_entry*, %__tgt_offload_entry* }
+// CHECK-IR1: [[IMAGETY:%.+]] = type { i16, i8, i8, i8*, i8*, i8*, i8*, i8*, i8*, i8*, %__tgt_offload_entry*, %__tgt_offload_entry* }
// CHECK-IR1: [[ENTTY:%.+]] = type { i8*, i8*, i64, i32, i32 }
// CHECK-IR1: [[DESCTY:%.+]] = type { i16, i16, [[IMAGETY]]*, [[ENTTY]]*, [[ENTTY]]* }
// CHECK-IR1-NOT: @llvm.global_ctors
diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp
index 489b51ae37e8b..aad2afa9abbe2 100644
--- a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp
+++ b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp
@@ -131,12 +131,19 @@ static cl::list Targets("target", cl::ZeroOrMore,
cl::cat(ClangOffloadWrapperCategory),
cl::cat(ClangOffloadWrapperCategory));
-/// Sets build options for device binary image.
+/// Sets compile options for device binary image.
static cl::list
- Options("build-opts", cl::ZeroOrMore,
- cl::desc("build options passed to the offload runtime"),
- cl::cat(ClangOffloadWrapperCategory),
- cl::cat(ClangOffloadWrapperCategory));
+ CompileOptions("compile-opts", cl::ZeroOrMore,
+ cl::desc("compile options passed to the offload runtime"),
+ cl::cat(ClangOffloadWrapperCategory),
+ cl::cat(ClangOffloadWrapperCategory));
+
+/// Sets link options for device binary image.
+static cl::list
+ LinkOptions("link-opts", cl::ZeroOrMore,
+ cl::desc("link options passed to the offload runtime"),
+ cl::cat(ClangOffloadWrapperCategory),
+ cl::cat(ClangOffloadWrapperCategory));
/// Sets the name of the file containing offload function entries
static cl::list Entries(
@@ -214,8 +221,10 @@ class BinaryWrapper {
public:
Image(const llvm::StringRef File_, const llvm::StringRef Manif_,
const llvm::StringRef Tgt_, BinaryImageFormat Fmt_,
- const llvm::StringRef Opts_, const llvm::StringRef EntriesFile_)
- : File(File_), Manif(Manif_), Tgt(Tgt_), Fmt(Fmt_), Opts(Opts_),
+ const llvm::StringRef CompileOpts_, const llvm::StringRef LinkOpts_,
+ const llvm::StringRef EntriesFile_)
+ : File(File_), Manif(Manif_), Tgt(Tgt_), Fmt(Fmt_),
+ CompileOpts(CompileOpts_), LinkOpts(LinkOpts_),
EntriesFile(EntriesFile_) {}
/// Name of the file with actual contents
@@ -226,8 +235,10 @@ class BinaryWrapper {
const llvm::StringRef Tgt;
/// Format
const BinaryImageFormat Fmt;
- /// Build options
- const llvm::StringRef Opts;
+ /// Compile options
+ const llvm::StringRef CompileOpts;
+ /// Link options
+ const llvm::StringRef LinkOpts;
/// File listing contained entries
const llvm::StringRef EntriesFile;
@@ -258,13 +269,13 @@ class BinaryWrapper {
public:
void addImage(const OffloadKind Kind, llvm::StringRef File,
llvm::StringRef Manif, llvm::StringRef Tgt,
- const BinaryImageFormat Fmt, llvm::StringRef Opts,
- llvm::StringRef EntriesFile) {
+ const BinaryImageFormat Fmt, llvm::StringRef CompileOpts,
+ llvm::StringRef LinkOpts, llvm::StringRef EntriesFile) {
std::unique_ptr &Pack = Packs[Kind];
if (!Pack)
Pack.reset(new SameKindPack());
- Pack->emplace_back(
- std::make_unique(File, Manif, Tgt, Fmt, Opts, EntriesFile));
+ Pack->emplace_back(std::make_unique(
+ File, Manif, Tgt, Fmt, CompileOpts, LinkOpts, EntriesFile));
}
private:
@@ -347,8 +358,11 @@ class BinaryWrapper {
// /// architecture
// const char *DeviceTargetSpec;
// /// a null-terminated string; target- and compiler-specific options
- // /// which are suggested to use to "build" program at runtime
- // const char *BuildOptions;
+ // /// which are suggested to use to "compile" program at runtime
+ // const char *CompileOptions;
+ // /// a null-terminated string; target- and compiler-specific options
+ // /// which are suggested to use to "link" program at runtime
+ // const char *LinkOptions;
// /// Pointer to the manifest data start
// const unsigned char *ManifestStart;
// /// Pointer to the manifest data end
@@ -370,7 +384,8 @@ class BinaryWrapper {
Type::getInt8Ty(C), // OffloadKind
Type::getInt8Ty(C), // Format
Type::getInt8PtrTy(C), // DeviceTargetSpec
- Type::getInt8PtrTy(C), // BuildOptions
+ Type::getInt8PtrTy(C), // CompileOptions
+ Type::getInt8PtrTy(C), // LinkOptions
Type::getInt8PtrTy(C), // ManifestStart
Type::getInt8PtrTy(C), // ManifestEnd
Type::getInt8PtrTy(C), // ImageStart
@@ -658,8 +673,12 @@ class BinaryWrapper {
auto *Ffmt = ConstantInt::get(Type::getInt8Ty(C), Img.Fmt);
auto *Ftgt = addStringToModule(
Img.Tgt, Twine(OffloadKindTag) + Twine("target.") + Twine(ImgId));
- auto *Fopt = addStringToModule(
- Img.Opts, Twine(OffloadKindTag) + Twine("opts.") + Twine(ImgId));
+ auto *Foptcompile = addStringToModule(
+ Img.CompileOpts,
+ Twine(OffloadKindTag) + Twine("opts.compile.") + Twine(ImgId));
+ auto *Foptlink = addStringToModule(Img.LinkOpts, Twine(OffloadKindTag) +
+ Twine("opts.link.") +
+ Twine(ImgId));
std::pair FMnf;
if (Img.Manif.empty()) {
@@ -694,9 +713,9 @@ class BinaryWrapper {
return EntriesOrErr.takeError();
std::pair ImageEntriesPtrs = *EntriesOrErr;
ImagesInits.push_back(ConstantStruct::get(
- getSyclDeviceImageTy(), Fver, Fknd, Ffmt, Ftgt, Fopt, FMnf.first,
- FMnf.second, Fbin.first, Fbin.second, ImageEntriesPtrs.first,
- ImageEntriesPtrs.second));
+ getSyclDeviceImageTy(), Fver, Fknd, Ffmt, Ftgt, Foptcompile,
+ Foptlink, FMnf.first, FMnf.second, Fbin.first, Fbin.second,
+ ImageEntriesPtrs.first, ImageEntriesPtrs.second));
} else
ImagesInits.push_back(ConstantStruct::get(
getDeviceImageTy(), Fbin.first, Fbin.second, EntriesB, EntriesE));
@@ -830,7 +849,10 @@ llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
Out << " manifest = " << (Img.Manif.empty() ? "-" : Img.Manif) << "\n";
Out << " format = " << formatToString(Img.Fmt) << "\n";
Out << " target = " << (Img.Tgt.empty() ? "-" : Img.Tgt) << "\n";
- Out << " options = " << (Img.Opts.empty() ? "-" : Img.Opts) << "\n";
+ Out << " compile options = "
+ << (Img.CompileOpts.empty() ? "-" : Img.CompileOpts) << "\n";
+ Out << " link options = " << (Img.LinkOpts.empty() ? "-" : Img.LinkOpts)
+ << "\n";
Out << "}\n";
return Out;
}
@@ -990,21 +1012,35 @@ int main(int argc, const char **argv) {
"runtime for that device. When present, manifest file name should\n"
"immediately follow the corresponding device image filename on the\n"
"command line. Options annotating a device binary have effect on all\n"
- "subsequent input, until redefined. For example:\n"
- "$clang-offload-wrapper -host x86_64-pc-linux-gnu \\\n"
- " -kind=sycl -target=spir64 -format=spirv -build-opts=-g \\\n"
- " a.spv a_mf.txt \\\n"
- " -target=xxx -format=native -build-opts=\"\" \\\n"
- " b.bin b_mf.txt \\\n"
- " -kind=openmp \\\n"
- " c.bin\n"
- "will generate an x86 wrapper object (.bc) enclosing the following\n"
- "tuples describing a single device binary each ('-' means 'none')\n\n"
+ "subsequent input, until redefined.\n"
+ "\n"
+ "For example:\n"
+ " clang-offload-wrapper \\\n"
+ " -host x86_64-pc-linux-gnu \\\n"
+ " -kind=sycl \\\n"
+ " -target=spir64 \\\n"
+ " -format=spirv \\\n"
+ " -compile-opts=-g \\\n"
+ " -link-opts=-cl-denorms-are-zero \\\n"
+ " a.spv \\\n"
+ " a_mf.txt \\\n"
+ " -target=xxx \\\n"
+ " -format=native \\\n"
+ " -compile-opts=\"\" \\\n"
+ " -link-opts=\"\" \\\n"
+ " b.bin \\\n"
+ " b_mf.txt \\\n"
+ " -kind=openmp \\\n"
+ " c.bin\\n"
+ "\n"
+ "This command generates an x86 wrapper object (.bc) enclosing the\n"
+ "following tuples describing a single device binary each:\n"
+ "\n"
"offload kind | target | data format | data | manifest | build options:\n"
"----------------------------------------------------------------------\n"
" sycl | spir64 | spirv | a.spv| a_mf.txt | -g\n"
" sycl | xxx | native | b.bin| b_mf.txt | -\n"
- " openmp | xxx | native | c.bin| - | -\n");
+ " openmp | xxx | native | c.bin| n/a | -\n");
if (Help) {
cl::PrintHelpMessage();
@@ -1028,13 +1064,16 @@ int main(int argc, const char **argv) {
OffloadKind Knd = OffloadKind::Unknown;
llvm::StringRef Tgt = "";
BinaryImageFormat Fmt = BinaryImageFormat::none;
- llvm::StringRef Opts = "";
+ llvm::StringRef CompileOpts = "";
+ llvm::StringRef LinkOpts = "";
llvm::StringRef EntriesFile = "";
llvm::SmallVector CurInputPair;
ListArgsSequencer
- ArgSeq((size_t)argc, Inputs, Kinds, Formats, Targets, Options, Entries);
+ decltype(Targets), decltype(CompileOptions),
+ decltype(LinkOptions), decltype(Entries)>
+ ArgSeq((size_t)argc, Inputs, Kinds, Formats, Targets, CompileOptions,
+ LinkOptions, Entries);
int ID = -1;
do {
@@ -1058,7 +1097,8 @@ int main(int argc, const char **argv) {
}
StringRef File = CurInputPair[0];
StringRef Manif = CurInputPair.size() > 1 ? CurInputPair[1] : "";
- Wr.addImage(Knd, File, Manif, Tgt, Fmt, Opts, EntriesFile);
+ Wr.addImage(Knd, File, Manif, Tgt, Fmt, CompileOpts, LinkOpts,
+ EntriesFile);
CurInputPair.clear();
}
}
@@ -1077,11 +1117,14 @@ int main(int argc, const char **argv) {
case 3: // Targets
Tgt = *(ArgSeq.template get<3>());
break;
- case 4: // Options
- Opts = *(ArgSeq.template get<4>());
+ case 4: // CompileOptions
+ CompileOpts = *(ArgSeq.template get<4>());
+ break;
+ case 5: // LinkOptions
+ LinkOpts = *(ArgSeq.template get<5>());
break;
- case 5: // Entries
- EntriesFile = *(ArgSeq.template get<5>());
+ case 6: // Entries
+ EntriesFile = *(ArgSeq.template get<6>());
break;
default:
llvm_unreachable("bad option class ID");
diff --git a/sycl/doc/SYCLEnvironmentVariables.md b/sycl/doc/SYCLEnvironmentVariables.md
index 700b0fac6fc6b..209a1e3160552 100644
--- a/sycl/doc/SYCLEnvironmentVariables.md
+++ b/sycl/doc/SYCLEnvironmentVariables.md
@@ -13,7 +13,8 @@ subject to change. Do not rely on these variables in production code.
| SYCL_PI_TRACE | Any(*) | Force tracing of PI calls to stderr. |
| SYCL_BE | PI_OPENCL, PI_OTHER | When SYCL RT is built with PI this controls which plugin to use. Default value is PI_OPENCL. |
| SYCL_DEVICE_TYPE | One of: CPU, GPU, ACC, HOST | Force SYCL to use the specified device type. If unset, default selection rules are applied. If set to any unlisted value, this control has no effect. If the requested device type is not found, a `cl::sycl::runtime_error` exception is thrown. If a non-default device selector is used, a device must satisfy both the selector and this control to be chosen. This control only has effect on devices created with a selector. |
-| SYCL_PROGRAM_BUILD_OPTIONS | String of valid OpenCL build options | Override build options for all programs. |
+| SYCL_PROGRAM_COMPILE_OPTIONS | String of valid OpenCL compile options | Override compile options for all programs. |
+| SYCL_PROGRAM_LINK_OPTIONS | String of valid OpenCL link options | Override link options for all programs. |
| SYCL_USE_KERNEL_SPV | Path to the SPIR-V binary | Load device image from the specified file. If runtime is unable to read the file, `cl::sycl::runtime_error` exception is thrown.|
| SYCL_DUMP_IMAGES | Any(*) | Dump device image binaries to file. Control has no effect if SYCL_USE_KERNEL_SPV is set. |
| SYCL_PRINT_EXECUTION_GRAPH | Described [below](#sycl_print_execution_graph-options) | Print execution graph to DOT text file. |
diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h
index f7cf7cb56a31d..d5797c5b81ebc 100644
--- a/sycl/include/CL/sycl/detail/pi.h
+++ b/sycl/include/CL/sycl/detail/pi.h
@@ -302,8 +302,11 @@ struct pi_device_binary_struct {
/// PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA - 64-bit OpenCL FPGA device
const char *DeviceTargetSpec;
/// a null-terminated string; target- and compiler-specific options
- /// which are suggested to use to "build" program at runtime
- const char *BuildOptions;
+ /// which are suggested to use to "compile" program at runtime
+ const char *CompileOptions;
+ /// a null-terminated string; target- and compiler-specific options
+ /// which are suggested to use to "link" program at runtime
+ const char *LinkOptions;
/// Pointer to the manifest data start
const char *ManifestStart;
/// Pointer to the manifest data end
diff --git a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp
index 884fae473d4fd..63a50cace5a01 100644
--- a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp
+++ b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp
@@ -77,7 +77,8 @@ class ProgramManager {
using ProgramPtr = unique_ptr_class,
decltype(&::piProgramRelease)>;
ProgramPtr build(ProgramPtr Program, RT::PiContext Context,
- const string_class &Options,
+ const string_class &CompileOptions,
+ const string_class &LinkOptions,
const std::vector &Devices,
std::map &CachedLibPrograms,
bool LinkDeviceLibs = false);
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 3a40c5c114eda..c3f1445244cb3 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -318,8 +318,9 @@ ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context,
[](const device Dev) { return getRawSyclObjImpl(Dev)->getHandleRef(); });
ProgramPtr BuiltProgram =
- build(std::move(ProgramManaged), PiContext, Img.BuildOptions, PiDevices,
- ContextImpl->getCachedLibPrograms(), LinkDeviceLibs);
+ build(std::move(ProgramManaged), PiContext, Img.CompileOptions,
+ Img.LinkOptions, PiDevices, ContextImpl->getCachedLibPrograms(),
+ LinkDeviceLibs);
return BuiltProgram.release();
};
@@ -508,7 +509,8 @@ ProgramManager::ProgramManager() {
ImgPtr->Version = PI_DEVICE_BINARY_VERSION;
ImgPtr->Kind = PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL;
ImgPtr->DeviceTargetSpec = PI_DEVICE_BINARY_TARGET_UNKNOWN;
- ImgPtr->BuildOptions = "";
+ ImgPtr->CompileOptions = "";
+ ImgPtr->LinkOptions = "";
ImgPtr->ManifestStart = nullptr;
ImgPtr->ManifestEnd = nullptr;
ImgPtr->BinaryStart = Data.release();
@@ -614,19 +616,26 @@ static std::vector getDeviceLibPrograms(
ProgramManager::ProgramPtr
ProgramManager::build(ProgramPtr Program, RT::PiContext Context,
- const string_class &Options,
+ const string_class &CompileOptions,
+ const string_class &LinkOptions,
const std::vector &Devices,
std::map &CachedLibPrograms,
bool LinkDeviceLibs) {
if (DbgProgMgr > 0) {
std::cerr << ">>> ProgramManager::build(" << Program.get() << ", "
- << Options << ", ... " << Devices.size() << ")\n";
+ << CompileOptions << ", " << LinkOptions << ", ... "
+ << Devices.size() << ")\n";
}
- const char *Opts = std::getenv("SYCL_PROGRAM_BUILD_OPTIONS");
- if (!Opts)
- Opts = Options.c_str();
+ const char *CompileOpts = std::getenv("SYCL_PROGRAM_COMPILE_OPTIONS");
+ if (!CompileOpts) {
+ CompileOpts = CompileOptions.c_str();
+ }
+ const char *LinkOpts = std::getenv("SYCL_PROGRAM_LINK_OPTIONS");
+ if (!LinkOpts) {
+ LinkOpts = LinkOptions.c_str();
+ }
std::vector LinkPrograms;
if (LinkDeviceLibs) {
@@ -634,21 +643,27 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context,
}
if (LinkPrograms.empty()) {
+ std::string Opts(CompileOpts);
+ Opts += " ";
+ Opts += LinkOpts;
+
RT::PiResult Error = PI_CALL_NOCHECK(piProgramBuild)(
- Program.get(), Devices.size(), Devices.data(), Opts, nullptr, nullptr);
+ Program.get(), Devices.size(), Devices.data(), Opts.c_str(), nullptr,
+ nullptr);
if (Error != PI_SUCCESS)
throw compile_program_error(getProgramBuildLog(Program.get()));
return Program;
}
// Include the main program and compile/link everything together
- PI_CALL(piProgramCompile)(Program.get(), Devices.size(), Devices.data(), Opts,
- 0, nullptr, nullptr, nullptr, nullptr);
+ PI_CALL(piProgramCompile)(Program.get(), Devices.size(), Devices.data(),
+ CompileOpts, 0, nullptr, nullptr, nullptr, nullptr);
LinkPrograms.push_back(Program.get());
RT::PiProgram LinkedProg = nullptr;
+
RT::PiResult Error = PI_CALL_NOCHECK(piProgramLink)(
- Context, Devices.size(), Devices.data(), Opts, LinkPrograms.size(),
+ Context, Devices.size(), Devices.data(), LinkOpts, LinkPrograms.size(),
LinkPrograms.data(), nullptr, nullptr, &LinkedProg);
// Link program call returns a new program object if all parameters are valid,
@@ -719,12 +734,14 @@ void ProgramManager::debugDumpBinaryImage(const DeviceImage *Img) const {
std::cerr << " --- Image " << Img << "\n";
if (!Img)
return;
- std::cerr << " Version : " << (int)Img->Version << "\n";
- std::cerr << " Kind : " << (int)Img->Kind << "\n";
- std::cerr << " Format : " << (int)Img->Format << "\n";
- std::cerr << " Target : " << Img->DeviceTargetSpec << "\n";
- std::cerr << " Options : "
- << (Img->BuildOptions ? Img->BuildOptions : "NULL") << "\n";
+ std::cerr << " Version : " << (int)Img->Version << "\n";
+ std::cerr << " Kind : " << (int)Img->Kind << "\n";
+ std::cerr << " Format : " << (int)Img->Format << "\n";
+ std::cerr << " Target : " << Img->DeviceTargetSpec << "\n";
+ std::cerr << " Compile options : "
+ << (Img->CompileOptions ? Img->CompileOptions : "NULL") << "\n";
+ std::cerr << " Link options : "
+ << (Img->LinkOptions ? Img->LinkOptions : "NULL") << "\n";
std::cerr << " Bin size : "
<< ((intptr_t)Img->BinaryEnd - (intptr_t)Img->BinaryStart) << "\n";
std::cerr << " Entries : ";
diff --git a/sycl/test/program_manager/env_vars.cpp b/sycl/test/program_manager/env_vars.cpp
new file mode 100644
index 0000000000000..e747eab855e12
--- /dev/null
+++ b/sycl/test/program_manager/env_vars.cpp
@@ -0,0 +1,50 @@
+// RUN: %clangxx -O0 -fsycl %s -o %t.out -lsycl
+//
+// Deprecated SYCL_PROGRAM_BUILD_OPTIONS should work as an alias to
+// SYCL_PROGRAM_COMPILE_OPTIONS:
+// RUN: %CPU_RUN_PLACEHOLDER SYCL_PROGRAM_BUILD_OPTIONS="-g" %t.out
+// RUN: %GPU_RUN_PLACEHOLDER SYCL_PROGRAM_BUILD_OPTIONS="-g" %t.out
+// RUN: %CPU_RUN_PLACEHOLDER SYCL_PROGRAM_COMPILE_OPTIONS="-g" %t.out
+// RUN: %GPU_RUN_PLACEHOLDER SYCL_PROGRAM_COMPILE_OPTIONS="-g" %t.out
+// RUN: %CPU_RUN_PLACEHOLDER SYCL_PROGRAM_LINK_OPTIONS="-enable-link-options -cl-denorms-are-zero" %t.out
+// RUN: %GPU_RUN_PLACEHOLDER SYCL_PROGRAM_LINK_OPTIONS="-enable-link-options -cl-denorms-are-zero" %t.out
+//
+// Now test for invalid options to make sure they are really passed to
+// a device compiler. Intel GPU runtime doesn't give an error for
+// invalid options, so we don't test it here.
+//
+// RUN: %CPU_RUN_PLACEHOLDER SYCL_PROGRAM_COMPILE_OPTIONS="-enable-link-options -cl-denorms-are-zero" SHOULD_CRASH=1 %t.out
+// RUN: %CPU_RUN_PLACEHOLDER SYCL_PROGRAM_LINK_OPTIONS="-g" SHOULD_CRASH=1 %t.out
+
+
+#include
+#include
+#include
+
+using namespace cl::sycl;
+
+int main() {
+ int data = 5;
+ buffer buf(&data, range<1>(1));
+ queue myQueue;
+ if (getenv("SHOULD_CRASH")) {
+ try {
+ myQueue.submit([&](handler &cgh) {
+ auto B = buf.get_access(cgh);
+ cgh.single_task([=]() { B[0] = 0; });
+ });
+ } catch (cl::sycl::runtime_error &e) {
+ // Exit immediately, otherwise the buffer destructor may actually try to
+ // enqueue the command once again, and throw another exception.
+ exit(0);
+ } catch (cl::sycl::compile_program_error &e) {
+ exit(0);
+ }
+ assert(0 && "Expected exception was *not* thrown");
+ } else {
+ myQueue.submit([&](handler &cgh) {
+ auto B = buf.get_access(cgh);
+ cgh.single_task([=]() { B[0] = 0; });
+ });
+ }
+}