[SYCL][NewOffloadModel] Refactor SYCL pipeline in clang-linker-wrapper by maksimsab · Pull Request #22153 · intel/llvm

maksimsab · 2026-05-28T12:49:52Z

Extract SYCL offloading steps into named functions within the sycl namespace: runSYCLOffloadingPipeline, postLinkProcessModule, compileAndBundle and compileSYCLDevice.

SYCL modifications of linkDevice function has been extracted to the dedicated function - compileSYCLDevice.

Revert linkDevice function to its upstream state. Side effect of this revert: hipspv tests get resolved. Related tracker: CMPLRLLVM-73247

Functions containerizeRawImage and writeOffloadFile are moved so that their location corresponds to the upstream version.

Extract SYCL offloading steps into named functions within the sycl namespace: runSYCLOffloadingPipeline, postLinkProcessModule, compileAndBundle and compileSYCLDevice. SYCL modifications of linkDevice function has been extracted to the dedicated function - compileSYCLDevice. Revert linkDevice function to its upstream state. Side effect of this revert: hipspv tests get resolved. Related tracker: CMPLRLLVM-73247 Functions containerizeRawImage and writeOffloadFile are moved so that their location corresponds to the upstream version.

maksimsab · 2026-05-28T12:56:29Z

Note: this refactoring is done in scope of the adding the support for -fno-sycl-rdc for NewOffloadModel. The added function postLinkProcessModule is going to be invoked for every input module in the following patch as follows (processSYCLModule is the equivalent of postLinkProcessModule):
https://github.com/maksimsab/llvm/blob/19d314d7db175e6e96a05d048126c7c9a9b3cb51/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp#L2018-L2029

mdtoguchi · 2026-05-29T15:12:51Z

+/// For spir[v]{32,64} performs SPIRV translation (JIT case) + possible AOT
+/// compilation (Intel CPU/GPU).
+Expected<StringRef>
+compileSYCLDevice(ArrayRef<StringRef> InputFiles, const ArgList &Args,


Given this is a departure for linkDevice, should this be named to linkSYCLDevice instead? From a overall compilation standpoint, I see 'compile device' as the device compilation step that is performed earlier and not done during link.

There is no actual linking happening for SYCL case. The input is always 1 file:
https://github.com/intel/llvm/pull/22153/changes#diff-f7dc51acede945c7c8481147e773e9c09230a93b937bee9eab08ee52826cfa8eR1869

So, the main functionality here is either invoke clang's backend for nvptx, amdgcn or invoke SPIRV translator + possible aot compilation.

There's lots of different things going on, so maybe a more generic term like processSYCLDevice?

YuriPlyakhin

quite a lot of CI failed. Is it caused by this patch?

YuriPlyakhin · 2026-05-29T21:06:50Z

  case Triple::native_cpu:
-    if (IsSYCLKind)
-      return generic::clang(InputFiles, Args, IsSYCLKind);
+    return generic::clang(InputFiles, Args, /*IsSYCLKind*/ true);


now handling looks same as case Triple::amdgcn, can we just add one more case there, instead of duplicating the same line here?

YuriPlyakhin · 2026-05-29T21:07:52Z

+  default:
    return createStringError(Triple.getArchName() +
-                             " linking is not supported other than for SYCL");
+                             " linking is not supported");


The method is called compileSYCLDevice, yet the error message says linking. Should the error message be updated?

YuriPlyakhin · 2026-05-29T21:10:12Z

+}
+
+/// Function invokes device compilation and bundling for NVPTX and AMDGCN cases.
+Expected<StringRef> compileDeviceAndBundle(StringRef ModuleFilePath,


Comment says, this one is specific for NVPTX and AMDGCN, while function name is generic. What is correct: comment or function name? something needs to be updated.

YuriPlyakhin · 2026-05-29T21:11:04Z

+/// given
+/// \p WrappedOutputCallback).


Suggested change

/// given

/// \p WrappedOutputCallback).

/// given \p WrappedOutputCallback).

YuriPlyakhin · 2026-05-29T21:12:51Z

+  SmallVector<StringRef> InputFilesSYCL = {ModuleFilePath};
+  Expected<std::vector<module_split::SplitModule>> SplitModulesOrErr =
+      UseSYCLPostLinkTool
+          ? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs,


just curious, what's still missing in PostLinkLibrary, that we have to call PostLinkTool?

YuriPlyakhin · 2026-05-29T21:13:54Z

+  }
+
+  // TODO: Take into account Arch values considered as JIT: "native",
+  // "spir64", "spir", "spirv32" and "spirv64" for SPIR targets.


Suggested change

// "spir64", "spir", "spirv32" and "spirv64" for SPIR targets.

// "spir64", "spir", "spirv32" and "spirv64" for SPIR and SPIR-V targets.

YuriPlyakhin · 2026-05-29T21:17:45Z

Revert linkDevice function to its upstream state. Side effect of this revert: hipspv tests get resolved. Related tracker: CMPLRLLVM-73247

Functions containerizeRawImage and writeOffloadFile are moved so that their location corresponds to the upstream version.

thank you for aligning the code to upstream!

srividya-sundaram · 2026-05-29T22:26:51Z

+  bool IsJIT =
+      Triple.isSPIROrSPIRV() && Triple.getSubArch() == llvm::Triple::NoSubArch;
+  if (IsJIT)
+    std::for_each(SplitModules.begin(), SplitModules.end(),


Not a comment, but a question to help with my understanding:
In the old code, CompileOptions/LinkOptions are set inside the loop after the compile call. New code sets them before the compile (compileDeviceAndBundle). Is this an intentional ordering shift?

srividya-sundaram · 2026-05-29T22:34:30Z

+}
+
+/// Function invokes device compilation and bundling for NVPTX and AMDGCN cases.
+Expected<StringRef> compileDeviceAndBundle(StringRef ModuleFilePath,


This function is only called from postLinkProcessModule. Should this be marked as 'static'?

srividya-sundaram · 2026-05-29T22:43:28Z

+  case Triple::aarch64_be:
+  case Triple::ppc64:
+  case Triple::ppc64le:
+  case Triple::spirv64:


L1830 handles SPIRV64 for SYCL offload, what is this SPIRV64 for? Non SYCL offload kind?

maksimsab requested review from a team as code owners May 28, 2026 12:49

maksimsab added the new-offload-model Enables testing with NewOffloadModel. label May 28, 2026

mdtoguchi reviewed May 29, 2026

View reviewed changes

Comment thread clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp Outdated

mdtoguchi reviewed May 29, 2026

View reviewed changes

fix typo and retrigger CI

45b6d83

YuriPlyakhin reviewed May 29, 2026

View reviewed changes

srividya-sundaram reviewed May 29, 2026

View reviewed changes

	/// given
	/// \p WrappedOutputCallback).
	/// given \p WrappedOutputCallback).

	// "spir64", "spir", "spirv32" and "spirv64" for SPIR targets.
	// "spir64", "spir", "spirv32" and "spirv64" for SPIR and SPIR-V targets.

Conversation

maksimsab commented May 28, 2026

Uh oh!

maksimsab commented May 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

YuriPlyakhin left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

YuriPlyakhin commented May 29, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

maksimsab commented May 28, 2026 •

edited

Loading