Skip to content

[NewOffloadModel] Fix transition of compile options from image to SYCL Offload Wrapper #19579

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: sycl
Choose a base branch
from
Open

[NewOffloadModel] Fix transition of compile options from image to SYCL Offload Wrapper #19579

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
47 changes: 34 additions & 13 deletions clang/test/Driver/sycl-linker-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,7 @@
// CHK-CMDS-AOT-GEN-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc
// CHK-CMDS-AOT-GEN-NEXT: "{{.*}}llvm-spirv"{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}}
// CHK-CMDS-AOT-GEN-NEXT: "{{.*}}ocloc"{{.*}} -output_no_suffix -spirv_input -device pvc {{.*}} -output {{.*}} -file {{.*}}
// CHK-CMDS-AOT-GEN-NEXT: offload-wrapper: input: {{.*}}, output: [[WRAPPEROUT:.*]].bc
// CHK-CMDS-AOT-GEN-NEXT: offload-wrapper: input: {{.*}}, output: [[WRAPPEROUT:.*]].bc, compile-opts: , link-opts:
// CHK-CMDS-AOT-GEN-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT:.*]].o [[WRAPPEROUT]].bc
// CHK-CMDS-AOT-GEN-NEXT: "{{.*}}/ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT]].o HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o

Expand Down Expand Up @@ -117,8 +117,8 @@
// CHK-CMDS-AOT-NV-NEXT: "{{.*}}clang"{{.*}} -o [[CLANGOUT:.*]] --target=nvptx64-nvidia-cuda -march={{.*}}
// CHK-CMDS-AOT-NV-NEXT: "{{.*}}ptxas"{{.*}} --output-file [[PTXASOUT:.*]] [[CLANGOUT]]
// CHK-CMDS-AOT-NV-NEXT: "{{.*}}fatbinary"{{.*}} --create [[FATBINOUT:.*]] --image=profile={{.*}},file=[[CLANGOUT]] --image=profile={{.*}},file=[[PTXASOUT]]
// CHK-CMDS-AOT-NV-NEXT: offload-wrapper: input: [[FATBINOUT]], output: [[WRAPPEROUT:.*]]
// CHK-CMDS-AOT-NV-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT:.*]] [[WRAPPEROUT]]
// CHK-CMDS-AOT-NV-NEXT: offload-wrapper: input: [[FATBINOUT]], output: [[WRAPPEROUT:.*]].bc
// CHK-CMDS-AOT-NV-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT:.*]] [[WRAPPEROUT]].bc
// CHK-CMDS-AOT-NV-NEXT: "{{.*}}ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT]] HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o

/// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for AMD)
Expand All @@ -135,8 +135,8 @@
// CHK-CMDS-AOT-AMD-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[FIRSTLLVMLINKOUT]].bc
// CHK-CMDS-AOT-AMD-NEXT: "{{.*}}clang"{{.*}} -o [[CLANGOUT:.*]] --target=amdgcn-amd-amdhsa -mcpu={{.*}}
// CHK-CMDS-AOT-AMD-NEXT: "{{.*}}clang-offload-bundler"{{.*}} -targets=host-x86_64-unknown-linux-gnu,hip-amdgcn-amd-amdhsa--gfx803 -input=/dev/null -input=[[CLANGOUT]] -output=[[BUNDLEROUT:.*]]
// CHK-CMDS-AOT-AMD-NEXT: offload-wrapper: input: [[BUNDLEROUT]], output: [[WRAPPEROUT:.*]]
// CHK-CMDS-AOT-AMD-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT:.*]] [[WRAPPEROUT]]
// CHK-CMDS-AOT-AMD-NEXT: offload-wrapper: input: [[BUNDLEROUT]], output: [[WRAPPEROUT:.*]].bc
// CHK-CMDS-AOT-AMD-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT:.*]] [[WRAPPEROUT]].bc
// CHK-CMDS-AOT-AMD-NEXT: "{{.*}}ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT]] HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o

/// Check for -sycl-embed-ir for standalone clang-linker-wrapper run for sycl (NVPTX)
Expand All @@ -157,13 +157,13 @@
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: offload-wrapper: input: {{.*}}.bc, output: [[WRAPPEROUT1:.*]]
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT1:.*]] [[WRAPPEROUT1]]
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: offload-wrapper: input: {{.*}}.bc, output: [[WRAPPEROUT1:.*]].bc
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT1:.*]] [[WRAPPEROUT1]].bc
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -o [[CLANGOUT:.*]] --target=nvptx64-nvidia-cuda -march={{.*}}
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}ptxas"{{.*}} --output-file [[PTXASOUT:.*]] [[CLANGOUT]]
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}fatbinary"{{.*}} --create [[FATBINOUT:.*]] --image=profile={{.*}},file=[[CLANGOUT]] --image=profile={{.*}},file=[[PTXASOUT]]
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: offload-wrapper: input: [[FATBINOUT]], output: [[WRAPPEROUT:.*]]
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT2:.*]] [[WRAPPEROUT]]
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: offload-wrapper: input: [[FATBINOUT]], output: [[WRAPPEROUT:.*]].bc
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT2:.*]] [[WRAPPEROUT]].bc
// CHK-CMDS-AOT-NV-EMBED-IR-NEXT: "{{.*}}ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT1]] [[LLCOUT2]] HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o

/// Check for -sycl-embed-ir for standalone clang-linker-wrapper run for sycl (AMD)
Expand All @@ -178,12 +178,12 @@
// CHK-CMDS-AOT-AMD-EMBED-IR: "{{.*}}spirv-to-ir-wrapper" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[FIRSTLLVMLINKOUT]].bc
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: offload-wrapper: input: {{.*}}.bc, output: [[WRAPPEROUT1:.*]]
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT1:.*]] [[WRAPPEROUT1]]
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: offload-wrapper: input: {{.*}}.bc, output: [[WRAPPEROUT1:.*]].bc
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT1:.*]] [[WRAPPEROUT1]].bc
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -o [[CLANGOUT:.*]] --target=amdgcn-amd-amdhsa -mcpu={{.*}}
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}clang-offload-bundler"{{.*}} -input=[[CLANGOUT]] -output=[[BUNDLEROUT:.*]]
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: offload-wrapper: input: [[BUNDLEROUT]], output: [[WRAPPEROUT2:.*]]
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT2:.*]] [[WRAPPEROUT2]]
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: offload-wrapper: input: [[BUNDLEROUT]], output: [[WRAPPEROUT2:.*]].bc
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}clang"{{.*}} -c -o [[LLCOUT2:.*]] [[WRAPPEROUT2]].bc
// CHK-CMDS-AOT-AMD-EMBED-IR-NEXT: "{{.*}}ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT1]] [[LLCOUT2]] HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o

// Error handling when --linker-path is not provided for clang-linker-wrapper
Expand Down Expand Up @@ -228,3 +228,24 @@
// CHK-DEVLINK-CMDS-NEXT: "{{.*}}clang"{{.*}} -c -o [[CLANGOUT:.*]] [[WRAPPEROUT]].bc
// CHK-DEVLINK-CMDS-NEXT: "{{.*}}cp"{{.*}} [[CLANGOUT]] a.out
// CHK-DEVLINK-CMDS-NOT: "{{.*}}/ld"

// Check that clang-linker-wrapper recognizes compile/link options and passes them
// over into offload wrapper and ocloc.
// RUN: clang-offload-packager -o %t.packaged_jit.fat "--image=file=%t.o,triple=spir64-unknown-unknown,arch=generic,kind=sycl,compile-opts=aaa aaa,link-opts=ccc ccc"
// RUN: clang-offload-packager -o %t.packaged_aot.fat "--image=file=%t.o,triple=spir64_gen-unknown-unknown,arch=pvc,kind=sycl,compile-opts=aaa aaa,link-opts=ccc ccc"
//
// RUN: %clang -cc1 %s -triple=x86_64-unknown-linux-gnu -emit-obj -o %t.jit.o -fembed-offload-object=%t.packaged_jit.fat
// RUN: %clang -cc1 %s -triple=x86_64-unknown-linux-gnu -emit-obj -o %t.aot.o -fembed-offload-object=%t.packaged_aot.fat
//
// RUN: clang-linker-wrapper --verbose --dry-run "--sycl-backend-compile-options=bbb bbb" "--sycl-target-link-options=ddd ddd" -host-triple=x86_64-unknown-linux-gnu \
// RUN: -sycl-device-libraries=%t.devicelib.o \
// RUN: %t.jit.o -o %t.out 2>&1 --linker-path="/usr/bin/ld" | FileCheck %s --check-prefix=CHECK-COMPILE-LINK-OPTS-JIT-WITH-ARGS
//
// CHECK-COMPILE-LINK-OPTS-JIT-WITH-ARGS: offload-wrapper: {{.*}} compile-opts: aaa aaa bbb bbb, link-opts: ccc ccc ddd ddd

// RUN: clang-linker-wrapper --verbose --dry-run "--sycl-backend-compile-options=bbb bbb" "--sycl-target-link-options=ddd ddd" -host-triple=x86_64-unknown-linux-gnu \
// RUN: -sycl-device-libraries=%t1.devicelib.o \
// RUN: %t.aot.o -o %t.out 2>&1 --linker-path="/usr/bin/ld" | FileCheck %s --check-prefix=CHECK-COMPILE-LINK-OPTS-AOT
//
// CHECK-COMPILE-LINK-OPTS-AOT: "{{.*}}ocloc"{{.*}} -device pvc aaa aaa bbb bbb ccc ccc ddd ddd
// CHECK-COMPILE-LINK-OPTS-AOT: offload-wrapper: {{.*}} compile-opts: , link-opts:
109 changes: 62 additions & 47 deletions clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -158,6 +158,10 @@ static SYCLBIN::BundleState SYCLBINState = SYCLBIN::BundleState::Input;

static SmallString<128> OffloadImageDumpDir;

static std::string SYCLCompileOptionsFromImage;

static std::string SYCLLinkOptionsFromImage;

using OffloadingImage = OffloadBinary::OffloadingImage;

namespace llvm {
Expand Down Expand Up @@ -940,11 +944,14 @@ static Expected<StringRef> runLLVMToSPIRVTranslation(StringRef File,
/// options are generated for AOT compilation targeting Intel GPUs.
static void addBackendOptions(const ArgList &Args,
SmallVector<StringRef, 8> &CmdArgs, bool IsCPU) {
StringRef OptC =
Args.getLastArgValue(OPT_sycl_backend_compile_options_from_image_EQ);
StringRef OptC = SYCLCompileOptionsFromImage;
OptC.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false);
OptC = Args.getLastArgValue(
OPT_sycl_backend_compile_options_EQ); // TODO: test this.
OptC.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false);
StringRef OptL =
Args.getLastArgValue(OPT_sycl_backend_link_options_from_image_EQ);
StringRef OptL = SYCLLinkOptionsFromImage;
OptL.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false);
OptL = Args.getLastArgValue(OPT_sycl_target_link_options_EQ);
OptL.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false);
StringRef OptTool = (IsCPU) ? Args.getLastArgValue(OPT_cpu_tool_arg_EQ)
: Args.getLastArgValue(OPT_gpu_tool_arg_EQ);
Expand Down Expand Up @@ -1052,6 +1059,33 @@ wrapSYCLBinariesFromFile(std::vector<module_split::SplitModule> &SplitModules,
if (!OutputFileOrErr)
return OutputFileOrErr.takeError();

StringRef Target = Args.getLastArgValue(OPT_triple_EQ);
if (Target.empty())
return createStringError(
inconvertibleErrorCode(),
"can't wrap SYCL image. -triple argument is missed.");

// SYCL runtime currently works for spir64 target triple and not for
// spir64-unknown-unknown/spirv64-unknown-unknown/spirv64.
// TODO: Fix SYCL runtime to accept other triples
llvm::Triple T(Target);
bool isJIT = !(T.getSubArch() == llvm::Triple::SPIRSubArch_gen ||
T.getSubArch() == llvm::Triple::SPIRSubArch_x86_64);
SmallString<0> CompileOptions;
SmallString<0> LinkOptions;
if (isJIT) {
CompileOptions =
StringRef(formatv("{0} {1}", SYCLCompileOptionsFromImage,
Args.getLastArgValue(
OPT_sycl_backend_compile_options_EQ, "")))
.trim();
LinkOptions =
StringRef(
formatv("{0} {1}", SYCLLinkOptionsFromImage,
Args.getLastArgValue(OPT_sycl_target_link_options_EQ, "")))
.trim();
}

StringRef OutputFilePath = *OutputFileOrErr;
if (Verbose || DryRun) {
std::string InputFiles;
Expand All @@ -1061,23 +1095,14 @@ wrapSYCLBinariesFromFile(std::vector<module_split::SplitModule> &SplitModules,
InputFiles += ',';
}

errs() << formatv(" offload-wrapper: input: {0}, output: {1}\n", InputFiles,
OutputFilePath);
errs() << formatv(" offload-wrapper: input: {0}, output: {1}, "
"compile-opts: {2}, link-opts: {3}\n",
InputFiles, OutputFilePath, CompileOptions, LinkOptions);
if (DryRun)
return OutputFilePath;
}

StringRef Target = Args.getLastArgValue(OPT_triple_EQ);
if (Target.empty())
return createStringError(
inconvertibleErrorCode(),
"can't wrap SYCL image. -triple argument is missed.");

SmallVector<llvm::offloading::SYCLImage> Images;
// SYCL runtime currently works for spir64 target triple and not for
// spir64-unknown-unknown/spirv64-unknown-unknown/spirv64.
// TODO: Fix SYCL runtime to accept other triples
llvm::Triple T(Target);
std::string EmbeddedIRTarget("llvm_");
EmbeddedIRTarget.append(T.getArchName());
StringRef RegularTarget(T.getArchName());
Expand Down Expand Up @@ -1112,22 +1137,9 @@ wrapSYCLBinariesFromFile(std::vector<module_split::SplitModule> &SplitModules,
M.setTargetTriple(Triple(
Args.getLastArgValue(OPT_host_triple_EQ, sys::getDefaultTargetTriple())));

auto CompileOptionsFromSYCLBackendCompileOptions =
Args.getLastArgValue(OPT_sycl_backend_compile_options_EQ);
auto LinkOptionsFromSYCLTargetLinkOptions =
Args.getLastArgValue(OPT_sycl_target_link_options_EQ);

StringRef CompileOptions(
Args.MakeArgString(CompileOptionsFromSYCLBackendCompileOptions.str()));
StringRef LinkOptions(
Args.MakeArgString(LinkOptionsFromSYCLTargetLinkOptions.str()));
offloading::SYCLWrappingOptions WrappingOptions;
WrappingOptions.CompileOptions = CompileOptions;
WrappingOptions.LinkOptions = LinkOptions;
if (Verbose) {
errs() << formatv(" offload-wrapper: compile-opts: {0}, link-opts: {1}\n",
CompileOptions, LinkOptions);
}
WrappingOptions.CompileOptions = std::string(CompileOptions);
WrappingOptions.LinkOptions = std::string(LinkOptions);
if (Error E = offloading::wrapSYCLBinaries(M, Images, WrappingOptions))
return E;

Expand Down Expand Up @@ -1979,14 +1991,6 @@ DerivedArgList getLinkerArgs(ArrayRef<OffloadFile> Input,
DAL.AddJoinedArg(nullptr, Tbl.getOption(OPT_triple_EQ),
Args.MakeArgString(Input.front().getBinary()->getTriple()));

const auto *Bin = Input.front().getBinary();
DAL.AddJoinedArg(
nullptr, Tbl.getOption(OPT_sycl_backend_compile_options_from_image_EQ),
Args.MakeArgString(Bin->getString(COMPILE_OPTS)));
DAL.AddJoinedArg(nullptr,
Tbl.getOption(OPT_sycl_backend_link_options_from_image_EQ),
Args.MakeArgString(Bin->getString(LINK_OPTS)));

// If every input file is bitcode we have whole program visibility as we
// do only support static linking with bitcode.
auto ContainsBitcode = [](const OffloadFile &F) {
Expand Down Expand Up @@ -2063,6 +2067,23 @@ Error handleOverrideImages(
return Error::success();
}

/// Function looks for compile/link options encoded in SYCL images.
/// If they are found then they are stored in global variables.
void findAndSaveSYCLCompileLinkOptions(ArrayRef<OffloadFile> OffloadFiles) {
for (const OffloadFile &File : OffloadFiles) {
const OffloadBinary &OB = *File.getBinary();
if (OB.getOffloadKind() != OFK_SYCL)
continue;

StringRef CompileOptions = OB.getString("compile-opts");
StringRef LinkOptions = OB.getString("link-opts");
if (!CompileOptions.empty())
SYCLCompileOptionsFromImage = CompileOptions.str();
if (!LinkOptions.empty())
SYCLLinkOptionsFromImage = LinkOptions.str();
}
}

/// Transforms all the extracted offloading input files into an image that can
/// be registered by the runtime.
Expected<SmallVector<StringRef>> linkAndWrapDeviceFiles(
Expand Down Expand Up @@ -2110,14 +2131,8 @@ Expected<SmallVector<StringRef>> linkAndWrapDeviceFiles(
HasNonSYCLOffloadKinds = true;
}

// Write any remaining device inputs to an output file.
SmallVector<StringRef> InputFiles;
for (const OffloadFile &File : Input) {
auto FileNameOrErr = writeOffloadFile(File);
if (!FileNameOrErr)
return FileNameOrErr.takeError();
InputFiles.emplace_back(*FileNameOrErr);
}
if (HasSYCLOffloadKind)
findAndSaveSYCLCompileLinkOptions(Input);

if (HasSYCLOffloadKind) {
SmallVector<StringRef> InputFiles;
Expand Down
13 changes: 2 additions & 11 deletions clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -154,11 +154,11 @@ def sycl_nvptx_device_lib_EQ : CommaJoined<["--", "-"], "sycl-nvptx-device-libra

// Options for SYCL backends and linker options for shared libraries.
def sycl_backend_compile_options_EQ :
Joined<["--", "-"], "sycl-backend-compile-options">,
Joined<["--", "-"], "sycl-backend-compile-options=">,
Flags<[WrapperOnlyOption]>,
HelpText<"Options that are passed to the backend of target device compiler">;
def sycl_target_link_options_EQ :
Joined<["--", "-"], "sycl-target-link-options">,
Joined<["--", "-"], "sycl-target-link-options=">,
Flags<[WrapperOnlyOption]>,
HelpText<"Options that are passed to target linker during a linking of shared device code.">;

Expand Down Expand Up @@ -213,15 +213,6 @@ def cpu_tool_arg_EQ :
Flags<[WrapperOnlyOption]>,
HelpText<"Options that are passed to the backend of target device compiler for Intel CPU during AOT compilation">;

// Hidden options to store backend compile/link options that are stored in
// device images for SYCL offloading
def sycl_backend_compile_options_from_image_EQ : Joined<["--", "-"], "sycl-backend-compile-options-from-image=">,
Flags<[WrapperOnlyOption]>,
HelpText<"Compile options that will be transmitted to the SYCL backend compiler">;
def sycl_backend_link_options_from_image_EQ : Joined<["--", "-"], "sycl-backend-link-options-from-image=">,
Flags<[WrapperOnlyOption]>,
HelpText<"Link options that will be transmitted to the SYCL backend compiler">;

def sycl_thin_lto : Flag<["--", "-"], "sycl-thin-lto">,
Flags<[WrapperOnlyOption]>, HelpText<"Link SYCL device code using thinLTO">;

Expand Down
49 changes: 49 additions & 0 deletions sycl/test-e2e/NewOffloadDriver/compile_options.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,49 @@
// Test checks that compile options are transfared to backend in the New
// Offloading Model. We use -ftarget-register-alloc-mode and
// -ftarget-compile-fast clang options for the testing here.

// REQUIRES: arch-intel_gpu_pvc

// RUN: %{build} --offload-new-driver -ftarget-register-alloc-mode=pvc:auto -o %t_with.out
// RUN: env SYCL_UR_TRACE=2 %{run} %t_with.out 2>&1 | FileCheck --check-prefix=CHECK-OPT %s

// CHECK-OPT: <--- urProgramBuildExp(
// CHECK-SAME-OPT: -ze-intel-enable-auto-large-GRF-mode

// RUN: %{build} --offload-new-driver -Wno-error=unused-command-line-argument -ftarget-compile-fast -o %t_with.out

// RUN: env SYCL_UR_TRACE=2 %{run} %t_with.out 2>&1 | FileCheck --check-prefix=CHECK-INTEL-GPU-WITH %s

// CHECK-INTEL-GPU-WITH: <--- urProgramBuild
// CHECK-INTEL-GPU-WITH-SAME: -igc_opts 'PartitionUnit=1,SubroutineThreshold=50000'

#include <sycl/detail/core.hpp>

int main() {
sycl::buffer<size_t, 1> Buffer(4);

sycl::queue Queue;

sycl::range<1> NumOfWorkItems{Buffer.size()};

Queue.submit([&](sycl::handler &cgh) {
sycl::accessor Accessor{Buffer, cgh, sycl::write_only};
cgh.parallel_for<class FillBuffer>(NumOfWorkItems, [=](sycl::id<1> WIid) {
Accessor[WIid] = WIid.get(0);
});
});

sycl::host_accessor HostAccessor{Buffer, sycl::read_only};

bool MismatchFound = false;
for (size_t I = 0; I < Buffer.size(); ++I) {
if (HostAccessor[I] != I) {
std::cout << "The result is incorrect for element: " << I
<< " , expected: " << I << " , got: " << HostAccessor[I]
<< std::endl;
MismatchFound = true;
}
}

return MismatchFound;
}
Loading