From 61793b3c061ea178be8805675cb80ed5b62b816c Mon Sep 17 00:00:00 2001 From: Michael D Toguchi Date: Fri, 18 Jul 2025 11:35:46 -0700 Subject: [PATCH 1/5] [Driver][SYCL] Update preprocessed file generation When generating preprocessed files for SYCL offloading, create a fully packaged file that contains both the HOST and DEVICE binaries. This will allow for consumption of these binary preprocessed files to be more useful, as opposed to only being able to preprocess and keep the host side of the offloading compilation When the driver encounters preprocessed (file.ii) files on the command line, these are processed in the following way: - Determines if the file is a packaged file (offload binary) - Extracts device side - Compiles device side, packages generated device into offload binary - Extracts host side - Compiles host side, embedding device binary Offload binary determination is performed by checking the magic number associated with the input file. The extraction is done via the clang-offload-packager using a new JobAction. When no output file is given, we will not package the preprocessed files but will just perform the host preprocessing. When output to a file (with an output file option), we will perform the host and device compilation, package and output to that file. --- clang/include/clang/Driver/Action.h | 12 +++ clang/include/clang/Driver/Driver.h | 3 + clang/include/clang/Driver/ToolChain.h | 2 + clang/lib/Driver/Action.cpp | 15 ++++ clang/lib/Driver/Driver.cpp | 93 +++++++++++++++++++- clang/lib/Driver/ToolChain.cpp | 8 ++ clang/lib/Driver/ToolChains/Clang.cpp | 46 ++++++++++ clang/lib/Driver/ToolChains/Clang.h | 13 +++ clang/test/Driver/sycl-int-header-footer.cpp | 2 +- clang/test/Driver/sycl-preprocess.cpp | 32 +++++-- 10 files changed, 217 insertions(+), 9 deletions(-) diff --git a/clang/include/clang/Driver/Action.h b/clang/include/clang/Driver/Action.h index f50aeafdd7697..17db708798c35 100644 --- a/clang/include/clang/Driver/Action.h +++ b/clang/include/clang/Driver/Action.h @@ -76,6 +76,7 @@ class Action { OffloadUnbundlingJobClass, OffloadWrapperJobClass, OffloadPackagerJobClass, + OffloadPackagerExtractJobClass, OffloadDepsJobClass, SPIRVTranslatorJobClass, SYCLPostLinkJobClass, @@ -719,6 +720,17 @@ class OffloadPackagerJobAction : public JobAction { } }; +class OffloadPackagerExtractJobAction : public JobAction { + void anchor() override; + +public: + OffloadPackagerExtractJobAction(ActionList &Inputs, types::ID Type); + + static bool classof(const Action *A) { + return A->getKind() == OffloadPackagerExtractJobClass; + } +}; + class OffloadDepsJobAction final : public JobAction { void anchor() override; diff --git a/clang/include/clang/Driver/Driver.h b/clang/include/clang/Driver/Driver.h index ab9f08f5cfc6f..e916ab8cd5972 100644 --- a/clang/include/clang/Driver/Driver.h +++ b/clang/include/clang/Driver/Driver.h @@ -991,6 +991,9 @@ bool isObjectFile(std::string FileName); /// \return True if the filename has a static archive/lib extension. bool isStaticArchiveFile(const StringRef &FileName); +/// \return True if the filename is an Offload Binary file. +bool isOffloadBinaryFile(const StringRef &FileName); + /// \return True if the argument combination will end up generating remarks. bool willEmitRemarks(const llvm::opt::ArgList &Args); diff --git a/clang/include/clang/Driver/ToolChain.h b/clang/include/clang/Driver/ToolChain.h index a004f057846fa..420cd54b63144 100644 --- a/clang/include/clang/Driver/ToolChain.h +++ b/clang/include/clang/Driver/ToolChain.h @@ -166,6 +166,7 @@ class ToolChain { mutable std::unique_ptr OffloadBundler; mutable std::unique_ptr OffloadWrapper; mutable std::unique_ptr OffloadPackager; + mutable std::unique_ptr OffloadPackagerExtract; mutable std::unique_ptr OffloadDeps; mutable std::unique_ptr SPIRVTranslator; mutable std::unique_ptr SYCLPostLink; @@ -185,6 +186,7 @@ class ToolChain { Tool *getOffloadBundler() const; Tool *getOffloadWrapper() const; Tool *getOffloadPackager() const; + Tool *getOffloadPackagerExtract() const; Tool *getOffloadDeps() const; Tool *getSPIRVTranslator() const; Tool *getSYCLPostLink() const; diff --git a/clang/lib/Driver/Action.cpp b/clang/lib/Driver/Action.cpp index 4fa8c44b2d7db..484b2a15c242c 100644 --- a/clang/lib/Driver/Action.cpp +++ b/clang/lib/Driver/Action.cpp @@ -46,6 +46,8 @@ const char *Action::getClassName(ActionClass AC) { return "clang-offload-wrapper"; case OffloadPackagerJobClass: return "clang-offload-packager"; + case OffloadPackagerExtractJobClass: + return "clang-offload-packager-extract"; case OffloadDepsJobClass: return "clang-offload-deps"; case SPIRVTranslatorJobClass: @@ -86,6 +88,13 @@ void Action::propagateDeviceOffloadInfo(OffloadKind OKind, const char *OArch, // Deps job uses the host kinds. if (Kind == OffloadDepsJobClass) return; + // Packaging actions can use host kinds for preprocessing. + bool hasPreprocessJob = + std::any_of(Inputs.begin(), Inputs.end(), [](const Action *A) { + return A->getKind() == PreprocessJobClass; + }); + if (Kind == OffloadPackagerJobClass && hasPreprocessJob) + return; assert((OffloadingDeviceKind == OKind || OffloadingDeviceKind == OFK_None) && "Setting device kind to a different device??"); @@ -485,6 +494,12 @@ OffloadPackagerJobAction::OffloadPackagerJobAction(ActionList &Inputs, types::ID Type) : JobAction(OffloadPackagerJobClass, Inputs, Type) {} +void OffloadPackagerExtractJobAction::anchor() {} + +OffloadPackagerExtractJobAction::OffloadPackagerExtractJobAction( + ActionList &Inputs, types::ID Type) + : JobAction(OffloadPackagerExtractJobClass, Inputs, Type) {} + void OffloadDepsJobAction::anchor() {} OffloadDepsJobAction::OffloadDepsJobAction( diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 2ecb3729ba798..897dd5b73c1ed 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -7730,7 +7730,8 @@ Action *Driver::BuildOffloadingActions(Compilation &C, Action *HostAction) const { // Don't build offloading actions if explicitly disabled or we do not have a // valid source input. - if (offloadHostOnly() || !types::isSrcFile(Input.first)) + if (offloadHostOnly() || + !(types::isSrcFile(Input.first) || Input.first == types::TY_PP_CXX)) return HostAction; bool HIPNoRDC = @@ -7795,6 +7796,25 @@ Action *Driver::BuildOffloadingActions(Compilation &C, llvm::sort(Sorted); for (StringRef Arch : Sorted) { TCAndArchs.push_back(std::make_pair(TC, Arch)); + // Check if the InputArg is a preprocessed file that is created by the + // clang-offload-packager. + if (InputType == types::TY_PP_CXX && + isOffloadBinaryFile(InputArg->getAsString(Args))) { + // Extract the specific preprocessed file given the current arch + // and triple. Add to DeviceActions if one was extracted. + ActionList PPActions; + OffloadAction::DeviceDependences DDep; + Action *IA = C.MakeAction(*InputArg, InputType, CUID); + PPActions.push_back(IA); + Action *PackagerAction = + C.MakeAction(PPActions, + types::TY_PP_CXX); + DDep.add(*PackagerAction, + *C.getSingleOffloadToolChain(), nullptr, + C.getActiveOffloadKinds()); + DeviceActions.push_back(PackagerAction); + continue; + } DeviceActions.push_back( C.MakeAction(*InputArg, InputType, CUID)); } @@ -7952,6 +7972,41 @@ Action *Driver::BuildOffloadingActions(Compilation &C, DDep.add(*LinkAction, *C.getSingleOffloadToolChain(), nullptr, C.getActiveOffloadKinds()); return C.MakeAction(DDep, types::TY_Nothing); + } else if (C.isOffloadingHostKind(Action::OFK_SYCL) && + isa(HostAction) && + getFinalPhase(Args) == phases::Preprocess && + Args.hasArg(options::OPT_o, options::OPT__SLASH_P, + options::OPT__SLASH_o)) { + // Performing preprocessing only. Take the host and device preprocessed + // files and package them together. + ActionList PackagerActions; + // Only add the preprocess actions from the device side. When one is + // found, add an additional compilation to generate the integration + // header/footer that is used for the host compile. + for (auto OA : OffloadActions) { + if (const OffloadAction *CurOA = dyn_cast(OA)) { + CurOA->doOnEachDependence( + [&](Action *A, const ToolChain *TC, const char *BoundArch) { + assert(TC && "Unknown toolchain"); + if (isa(A)) { + PackagerActions.push_back(OA); + // DDep.add(*A, *TC, BoundArch, Action::OFK_SYCL); + // Action *AA = C.MakeAction(DDep, + // types::TY_PP_CXX); + A->setCannotBeCollapsedWithNextDependentAction(); + Action *CompileAction = + C.MakeAction(A, types::TY_Nothing); + DDeps.add(*CompileAction, *TC, BoundArch, Action::OFK_SYCL); + // PackagerActions.push_back(AA); + } + }); + } + } + PackagerActions.push_back(HostAction); + Action *PackagerAction = C.MakeAction( + PackagerActions, types::TY_PP_CXX); + DDeps.add(*PackagerAction, *C.getSingleOffloadToolChain(), + nullptr, C.getActiveOffloadKinds()); } else if (C.isOffloadingHostKind(Action::OFK_SYCL) && Args.hasArg(options::OPT_fsycl_host_compiler_EQ)) { // -fsycl-host-compiler will create a bundled object instead of an @@ -7974,7 +8029,7 @@ Action *Driver::BuildOffloadingActions(Compilation &C, // add each device output as a host dependency to ensure they are still built. bool SingleDeviceOutput = !llvm::any_of(OffloadActions, [](Action *A) { return A->getType() == types::TY_Nothing; - }) && isa(HostAction); + }) && (isa(HostAction)); OffloadAction::HostDependence HDep( *HostAction, *C.getSingleOffloadToolChain(), /*BoundArch=*/nullptr, SingleDeviceOutput ? DDep : DDeps); @@ -8100,6 +8155,23 @@ Action *Driver::ConstructPhaseAction( return C.MakeAction(Input, types::TY_Nothing); if (Args.hasArg(options::OPT_extract_api)) return C.MakeAction(Input, types::TY_API_INFO); + // New offload driver enabled with a Preprocessed input file - check to make + // sure that the input file is an offload binary - if so, we need to + // extract the actual preprocessed file from the package, and that is what + // we will compile. + if (getUseNewOffloadingDriver() && + TargetDeviceOffloadKind == Action::OFK_None && + Input->getType() == types::TY_PP_CXX) { + const InputAction *IA = dyn_cast(Input); + if (IA && isOffloadBinaryFile(IA->getInputArg().getAsString(Args))) { + ActionList PPActions; + PPActions.push_back(Input); + Action *PackagerAction = C.MakeAction( + PPActions, types::TY_PP_CXX); + return C.MakeAction(PackagerAction, + types::TY_LLVM_BC); + } + } return C.MakeAction(Input, types::TY_LLVM_BC); } case phases::Backend: { @@ -9379,7 +9451,8 @@ const char *Driver::GetNamedOutputPath(Compilation &C, const JobAction &JA, // For /P, preprocess to file named after BaseInput. if (C.getArgs().hasArg(options::OPT__SLASH_P) && ((AtTopLevel && isa(JA)) || - isa(JA))) { + isa(JA) || + isa(JA))) { StringRef BaseName = llvm::sys::path::filename(BaseInput); StringRef NameArg; if (Arg *A = C.getArgs().getLastArg(options::OPT__SLASH_Fi)) @@ -9415,6 +9488,14 @@ const char *Driver::GetNamedOutputPath(Compilation &C, const JobAction &JA, } } + // When generating preprocessed files (-E) with offloading enabled, redirect + // the output to a properly named output file. + if (JA.getType() == types::TY_PP_CXX && isa(JA)) { + if (Arg *FinalOutput = + C.getArgs().getLastArg(options::OPT_o, options::OPT__SLASH_o)) + return C.addResultFile(FinalOutput->getValue(), &JA); + } + // Default to writing to stdout? if (AtTopLevel && !CCGenDiagnostics && HasPreprocessOutput(JA)) { return "-"; @@ -10383,6 +10464,12 @@ bool clang::driver::isStaticArchiveFile(const StringRef &FileName) { return (Magic == llvm::file_magic::archive); } +bool clang::driver::isOffloadBinaryFile(const StringRef &FileName) { + llvm::file_magic Magic; + llvm::identify_magic(FileName, Magic); + return (Magic == llvm::file_magic::offload_binary); +} + bool clang::driver::willEmitRemarks(const ArgList &Args) { // -fsave-optimization-record enables it. if (Args.hasFlag(options::OPT_fsave_optimization_record, diff --git a/clang/lib/Driver/ToolChain.cpp b/clang/lib/Driver/ToolChain.cpp index 3d21c92694458..dbd22936e9006 100644 --- a/clang/lib/Driver/ToolChain.cpp +++ b/clang/lib/Driver/ToolChain.cpp @@ -617,6 +617,12 @@ Tool *ToolChain::getOffloadPackager() const { return OffloadPackager.get(); } +Tool *ToolChain::getOffloadPackagerExtract() const { + if (!OffloadPackagerExtract) + OffloadPackagerExtract.reset(new tools::OffloadPackagerExtract(*this)); + return OffloadPackagerExtract.get(); +} + Tool *ToolChain::getOffloadDeps() const { if (!OffloadDeps) OffloadDeps.reset(new tools::OffloadDeps(*this)); @@ -707,6 +713,8 @@ Tool *ToolChain::getTool(Action::ActionClass AC) const { return getOffloadWrapper(); case Action::OffloadPackagerJobClass: return getOffloadPackager(); + case Action::OffloadPackagerExtractJobClass: + return getOffloadPackagerExtract(); case Action::OffloadDepsJobClass: return getOffloadDeps(); diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index bde06fac4bda3..3d759dfd8fe5a 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10400,6 +10400,8 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, for (const InputInfo &Input : Inputs) { const Action *OffloadAction = Input.getAction(); const ToolChain *TC = OffloadAction->getOffloadingToolChain(); + if (!TC) + TC = &C.getDefaultToolChain(); const ArgList &TCArgs = C.getArgsForToolChain(TC, OffloadAction->getOffloadingArch(), OffloadAction->getOffloadingDeviceKind()); @@ -10481,6 +10483,50 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs, Inputs, Output)); } +// Use the clang-offload-packager to extract binaries from an packaged +// binary. This currently only supports single input/single output. +void OffloadPackagerExtract::ConstructJob(Compilation &C, const JobAction &JA, + const InputInfo &Output, + const InputInfoList &Inputs, + const llvm::opt::ArgList &Args, + const char *LinkingOutput) const { + ArgStringList CmdArgs; + const Action *OffloadAction = Inputs[0].getAction(); + const ToolChain *TC = OffloadAction->getOffloadingToolChain(); + if (!TC) + TC = &C.getDefaultToolChain(); + const ArgList &TCArgs = + C.getArgsForToolChain(TC, OffloadAction->getOffloadingArch(), + OffloadAction->getOffloadingDeviceKind()); + + // Input file name. + StringRef InFile = C.getArgs().MakeArgString(TC->getInputFilename(Inputs[0])); + CmdArgs.push_back(Args.MakeArgString(InFile)); + + // Generated --image option containing the output file name, triple, arch + // and associated offload kind. + assert(Output.isFilename() && "Invalid output."); + StringRef File = Output.getFilename(); + StringRef Arch = OffloadAction->getOffloadingArch() + ? OffloadAction->getOffloadingArch() + : TCArgs.getLastArgValue(options::OPT_march_EQ); + StringRef Kind = + Action::GetOffloadKindName(OffloadAction->getOffloadingDeviceKind()); + + SmallVector Parts{ + "file=" + File.str(), + "triple=" + TC->getTripleString(), + "arch=" + (Arch.empty() ? "generic" : Arch.str()), + "kind=" + Kind.str(), + }; + CmdArgs.push_back(Args.MakeArgString("--image=" + llvm::join(Parts, ","))); + + C.addCommand(std::make_unique( + JA, *this, ResponseFileSupport::None(), + Args.MakeArgString(getToolChain().GetProgramPath(getShortName())), + CmdArgs, Inputs, Output)); +} + // Begin OffloadDeps void OffloadDeps::constructJob(Compilation &C, const JobAction &JA, diff --git a/clang/lib/Driver/ToolChains/Clang.h b/clang/lib/Driver/ToolChains/Clang.h index 538c3715d1c1f..75383cc416077 100644 --- a/clang/lib/Driver/ToolChains/Clang.h +++ b/clang/lib/Driver/ToolChains/Clang.h @@ -190,6 +190,19 @@ class LLVM_LIBRARY_VISIBILITY OffloadPackager final : public Tool { const char *LinkingOutput) const override; }; +/// Offload binary extract tool. +class LLVM_LIBRARY_VISIBILITY OffloadPackagerExtract final : public Tool { +public: + OffloadPackagerExtract(const ToolChain &TC) + : Tool("Offload::PackagerExtract", "clang-offload-packager", TC) {} + + bool hasIntegratedCPP() const override { return false; } + void ConstructJob(Compilation &C, const JobAction &JA, + const InputInfo &Output, const InputInfoList &Inputs, + const llvm::opt::ArgList &TCArgs, + const char *LinkingOutput) const override; +}; + /// Offload deps tool. class LLVM_LIBRARY_VISIBILITY OffloadDeps final : public Tool { void constructJob(Compilation &C, const JobAction &JA, diff --git a/clang/test/Driver/sycl-int-header-footer.cpp b/clang/test/Driver/sycl-int-header-footer.cpp index 8692f29c765e8..d84c645236de0 100644 --- a/clang/test/Driver/sycl-int-header-footer.cpp +++ b/clang/test/Driver/sycl-int-header-footer.cpp @@ -17,7 +17,7 @@ // FOOTER_PREPROC_GEN-SAME: "-dependency-filter" "[[INTHEADER]]" // FOOTER_PREPROC_GEN-SAME: "-include-internal-footer" "[[INTFOOTER]]" // FOOTER_PREPROC_GEN-SAME: "-dependency-filter" "[[INTFOOTER]]" -// FOOTER_PREPROC_GEN-SAME: "-E"{{.*}} "-o" "-" +// FOOTER_PREPROC_GEN-SAME: "-E" /// Preprocessed file use with integration footer // RUN: touch %t.ii diff --git a/clang/test/Driver/sycl-preprocess.cpp b/clang/test/Driver/sycl-preprocess.cpp index 57fe2fbc35270..7977aeb5a056d 100644 --- a/clang/test/Driver/sycl-preprocess.cpp +++ b/clang/test/Driver/sycl-preprocess.cpp @@ -1,9 +1,10 @@ -/// Test preprocessing capabilities when using -fsycl +// Test the behaviors when enabling SYCL offloading with preprocessed files. + /// Creating a preprocessed file is expected to do an integration header /// creation step. // RUN: %clangxx -fsycl --offload-new-driver -E -o %t_output.ii %s -### 2>&1 \ // RUN: | FileCheck -check-prefix PREPROC_ONLY %s -// RUN: %clang_cl -fsycl --offload-new-driver -P -Fi%t_output.ii %s -### 2>&1 \ +// RUN: %clang_cl -fsycl --offload-new-driver -P %s -### 2>&1 \ // RUN: | FileCheck -check-prefix PREPROC_ONLY %s // PREPROC_ONLY: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]" "-fsycl-int-footer=[[INTFOOTER:.+\.h]]"{{.*}} "-E" // PREPROC_ONLY: clang{{.*}} "-fsycl-is-host"{{.*}} "-include-internal-header" "[[INTHEADER]]"{{.*}} "-include-internal-footer" "[[INTFOOTER]]"{{.*}} "-o" "[[HOST_OUT:.+\.ii]]" @@ -14,10 +15,31 @@ // PREPROC_IN-NOT: "-fsycl-int-header={{.*}}" // PREPROC_IN: clang{{.*}} "-fsycl-is-host" -// RUN: %clangxx -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -E %s -ccc-print-phases 2>&1 \ -// RUN: | FileCheck -check-prefix PREPROC_PHASES %s +/// When generating preprocessed files, verify the compilation phases. +// RUN: %clangxx --target=x86_64-unknown-linux-gnu --offload-new-driver -fsycl -E %s -o %t.ii -ccc-print-phases 2>&1 \ +// RUN: | FileCheck %s -check-prefix PREPROC_PHASES +// RUN: %clang_cl --target=x86_64-unknown-linux-gnu --offload-new-driver -fsycl -P %s -Fi%t.ii -ccc-print-phases 2>&1 \ +// RUN: | FileCheck %s -check-prefix PREPROC_PHASES // PREPROC_PHASES: 0: input, "[[INPUT:.+\.cpp]]", c++, (host-sycl) // PREPROC_PHASES: 1: preprocessor, {0}, c++-cpp-output, (host-sycl) // PREPROC_PHASES: 2: input, "[[INPUT]]", c++, (device-sycl) // PREPROC_PHASES: 3: preprocessor, {2}, c++-cpp-output, (device-sycl) -// PREPROC_PHASES: 4: offload, "host-sycl (x86_64-unknown-linux-gnu)" {1}, "device-sycl (spir64-unknown-unknown)" {3}, c++-cpp-output +// PREPROC_PHASES: 4: compiler, {3}, none, (device-sycl) +// PREPROC_PHASES: 5: offload, "device-sycl (spir64-unknown-unknown)" {3}, c++-cpp-output +// PREPROC_PHASES: 6: clang-offload-packager, {5, 1}, c++-cpp-output +// PREPROC_PHASES: 7: offload, "host-sycl (x86_64-unknown-linux-gnu)" {1}, "device-sycl (spir64-unknown-unknown)" {3}, "device-sycl (spir64-unknown-unknown)" {4}, " (x86_64-unknown-linux-gnu)" {6}, c++-cpp-output + +/// When generating preprocessed files, verify the tools called and the expected +/// output file name. +// RUN: %clangxx --offload-new-driver -fsycl -E %s -o sycl-preprocess.ii -### 2>&1 \ +// RUN: | FileCheck %s -check-prefix PREPROC_TOOLS +// RUN: %clang_cl --offload-new-driver -fsycl -P %s -Fisycl-preprocess.ii -### 2>&1 \ +// RUN: | FileCheck %s -check-prefix PREPROC_TOOLS +// RUN: %clang_cl --offload-new-driver -fsycl -E %s -o sycl-preprocess.ii -### 2>&1 \ +// RUN: | FileCheck %s -check-prefix PREPROC_TOOLS +// PREPROC_TOOLS: clang{{.*}} "-fsycl-is-device" +// PREPROC_TOOLS-SAME: "-o" "[[DEVICE_PP_FILE:.+\.ii]] +// PREPROC_TOOLS: clang{{.*}} "-fsycl-is-host" +// PREPROC_TOOLS-SAME: "-o" "[[HOST_PP_FILE:.+\.ii]] +// PREPROC_TOOLS: clang-offload-packager{{.*}} "-o" "sycl-preprocess.ii" +// PREPROC_TOOLS-SAME: "--image=file=[[DEVICE_PP_FILE]],triple=spir64-unknown-unknown,arch=generic,kind=sycl{{.*}}" "--image=file=[[HOST_PP_FILE]],triple={{.*}},arch=generic,kind=host" From 7fdacc4952ed754fa45c4c850ae3712c220b544c Mon Sep 17 00:00:00 2001 From: Michael D Toguchi Date: Thu, 21 Aug 2025 15:55:03 -0700 Subject: [PATCH 2/5] Add E2E test --- .../NewOffloadDriver/preprocess_file.cpp | 42 +++++++++++++++++++ 1 file changed, 42 insertions(+) create mode 100644 sycl/test-e2e/NewOffloadDriver/preprocess_file.cpp diff --git a/sycl/test-e2e/NewOffloadDriver/preprocess_file.cpp b/sycl/test-e2e/NewOffloadDriver/preprocess_file.cpp new file mode 100644 index 0000000000000..e8294e9e93ab4 --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/preprocess_file.cpp @@ -0,0 +1,42 @@ +// REQUIRES: target-spir +// Test with `--offload-new-driver` that exercises the ability to create +// and consume preprocessed files that will perform full offloading compiles. + +// Create the preprocessed file. +// RUN: %{build} --offload-new-driver -E -o %t.ii + +// Compile preprocessed file. +// RUN: %clangxx -Wno-error=unused-command-line-argument -fsycl %{sycl_target_opts} --offload-new-driver %t.ii -o %t.out + +// RUN: %{run} %t.out + +#include + +int main() { + sycl::buffer 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(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; +} From 23b7e9f4297d6bffb554fa2b5fca8f117a73ee8f Mon Sep 17 00:00:00 2001 From: Michael D Toguchi Date: Thu, 21 Aug 2025 15:58:36 -0700 Subject: [PATCH 3/5] Unintended change --- clang/lib/Driver/Driver.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index eaca9c07fc424..a7b6cdbd66387 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -8080,7 +8080,7 @@ Action *Driver::BuildOffloadingActions(Compilation &C, // add each device output as a host dependency to ensure they are still built. bool SingleDeviceOutput = !llvm::any_of(OffloadActions, [](Action *A) { return A->getType() == types::TY_Nothing; - }) && (isa(HostAction)); + }) && isa(HostAction); OffloadAction::HostDependence HDep( *HostAction, *C.getSingleOffloadToolChain(), /*BoundArch=*/nullptr, SingleDeviceOutput ? DDep : DDeps); From fc4489b416b29c42d0c8ef6df758c5bfb182bc82 Mon Sep 17 00:00:00 2001 From: Michael D Toguchi Date: Mon, 25 Aug 2025 17:12:49 -0700 Subject: [PATCH 4/5] Clean out some inadvertent comments --- clang/lib/Driver/Driver.cpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index a7b6cdbd66387..74c0d171fb676 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -8041,14 +8041,10 @@ Action *Driver::BuildOffloadingActions(Compilation &C, assert(TC && "Unknown toolchain"); if (isa(A)) { PackagerActions.push_back(OA); - // DDep.add(*A, *TC, BoundArch, Action::OFK_SYCL); - // Action *AA = C.MakeAction(DDep, - // types::TY_PP_CXX); A->setCannotBeCollapsedWithNextDependentAction(); Action *CompileAction = C.MakeAction(A, types::TY_Nothing); DDeps.add(*CompileAction, *TC, BoundArch, Action::OFK_SYCL); - // PackagerActions.push_back(AA); } }); } From bbc7b1bc7c84525683ee5552eb2f125d6eb6a496 Mon Sep 17 00:00:00 2001 From: Michael D Toguchi Date: Tue, 26 Aug 2025 16:51:51 -0700 Subject: [PATCH 5/5] Address review comments - Update clang-cl based test to use Windows triple - Update some comments --- clang/lib/Driver/Action.cpp | 4 +++- clang/lib/Driver/ToolChains/Clang.cpp | 2 +- clang/test/Driver/sycl-preprocess.cpp | 8 ++++---- 3 files changed, 8 insertions(+), 6 deletions(-) diff --git a/clang/lib/Driver/Action.cpp b/clang/lib/Driver/Action.cpp index 484b2a15c242c..9c1b33e6eb620 100644 --- a/clang/lib/Driver/Action.cpp +++ b/clang/lib/Driver/Action.cpp @@ -88,7 +88,9 @@ void Action::propagateDeviceOffloadInfo(OffloadKind OKind, const char *OArch, // Deps job uses the host kinds. if (Kind == OffloadDepsJobClass) return; - // Packaging actions can use host kinds for preprocessing. + // Packaging actions can use host kinds for preprocessing. When packaging + // preprocessed files, these packaged files will contain both host and device + // files, where the host side does not have any device info to propagate. bool hasPreprocessJob = std::any_of(Inputs.begin(), Inputs.end(), [](const Action *A) { return A->getKind() == PreprocessJobClass; diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 2269fa8a52ebf..ef75ff04320f6 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10391,7 +10391,7 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs, Inputs, Output)); } -// Use the clang-offload-packager to extract binaries from an packaged +// Use the clang-offload-packager to extract binaries from a packaged // binary. This currently only supports single input/single output. void OffloadPackagerExtract::ConstructJob(Compilation &C, const JobAction &JA, const InputInfo &Output, diff --git a/clang/test/Driver/sycl-preprocess.cpp b/clang/test/Driver/sycl-preprocess.cpp index 7977aeb5a056d..65e62eec97e22 100644 --- a/clang/test/Driver/sycl-preprocess.cpp +++ b/clang/test/Driver/sycl-preprocess.cpp @@ -17,9 +17,9 @@ /// When generating preprocessed files, verify the compilation phases. // RUN: %clangxx --target=x86_64-unknown-linux-gnu --offload-new-driver -fsycl -E %s -o %t.ii -ccc-print-phases 2>&1 \ -// RUN: | FileCheck %s -check-prefix PREPROC_PHASES -// RUN: %clang_cl --target=x86_64-unknown-linux-gnu --offload-new-driver -fsycl -P %s -Fi%t.ii -ccc-print-phases 2>&1 \ -// RUN: | FileCheck %s -check-prefix PREPROC_PHASES +// RUN: | FileCheck %s -check-prefix PREPROC_PHASES -DTARGET=x86_64-unknown-linux-gnu +// RUN: %clang_cl --target=x86_64-pc-windows-msvc --offload-new-driver -fsycl -P %s -Fi%t.ii -ccc-print-phases 2>&1 \ +// RUN: | FileCheck %s -check-prefix PREPROC_PHASES -DTARGET=x86_64-pc-windows-msvc // PREPROC_PHASES: 0: input, "[[INPUT:.+\.cpp]]", c++, (host-sycl) // PREPROC_PHASES: 1: preprocessor, {0}, c++-cpp-output, (host-sycl) // PREPROC_PHASES: 2: input, "[[INPUT]]", c++, (device-sycl) @@ -27,7 +27,7 @@ // PREPROC_PHASES: 4: compiler, {3}, none, (device-sycl) // PREPROC_PHASES: 5: offload, "device-sycl (spir64-unknown-unknown)" {3}, c++-cpp-output // PREPROC_PHASES: 6: clang-offload-packager, {5, 1}, c++-cpp-output -// PREPROC_PHASES: 7: offload, "host-sycl (x86_64-unknown-linux-gnu)" {1}, "device-sycl (spir64-unknown-unknown)" {3}, "device-sycl (spir64-unknown-unknown)" {4}, " (x86_64-unknown-linux-gnu)" {6}, c++-cpp-output +// PREPROC_PHASES: 7: offload, "host-sycl ([[TARGET]])" {1}, "device-sycl (spir64-unknown-unknown)" {3}, "device-sycl (spir64-unknown-unknown)" {4}, " ([[TARGET]])" {6}, c++-cpp-output /// When generating preprocessed files, verify the tools called and the expected /// output file name.