[SYCL] PoC implementation of kernel compiler extension with libtooling and sycl-jit #15701

jopperm · 2024-10-15T12:10:15Z

This PR sets up in-memory compilation for runtime-defined SYCL kernels, via clang's libtooling interface and reusing LLVM-to-SPRIV-translation infrastructure in sycl-jit. I introduced a new, undocumented source language sycljit, which shall be removed again when the proposed approach is ready to replace the current process/file-based implementation for the sycl source language.

Missing features:

Compiler warnings/errors -> build log / exception message
Instantiation of template kernels
Name mangling beyond removing the __sycl_kernel prefix when requesting a kernel from the bundle
Linking of device libs, and sycl-post-link phase, so only very simple kernels are supported
Property handling and integration into program_manager

…oling and sycl-jit Signed-off-by: Julian Oppermann <[email protected]>

Signed-off-by: Julian Oppermann <[email protected]>

sommerlukas · 2024-10-15T12:21:12Z

sycl-jit/jit-compiler/CMakeLists.txt

  SYSTEM PRIVATE
  ${LLVM_MAIN_INCLUDE_DIR}
  ${LLVM_SPIRV_INCLUDE_DIRS}
+  ${CMAKE_SOURCE_DIR}/../clang/include


Does LLVM CMake not define a variable for that?

Changed this to use ${LLVM_EXTERNAL_CLANG_SOURCE_DIR}.

Unfortunately there doesn't seem to be an equivalent for the build directory, from which an .inc file that defines CLANG_VERSION_MAJOR is included.

sycl-jit/jit-compiler/lib/KernelFusion.cpp

sommerlukas · 2024-10-15T12:26:49Z

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

+        NewArgs.push_back((Twine("-resource-dir=") + DPCPPRoot + "/lib/clang/" +
+                           Twine(CLANG_VERSION_MAJOR))
+                              .str());


Do these paths also apply in a packaged release?

Yes, a release icpx follows the same path structure (checked with -print-resource-dir).

sommerlukas · 2024-10-15T12:31:50Z

sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp

+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+namespace detail {
+
+bool SYCLJIT_Compilation_Available() { return false; }
+
+spirv_vec_t
+SYCLJIT_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs,
+                 const std::vector<std::string> &UserArgs, std::string *LogPtr,
+                 const std::vector<std::string> &RegisteredKernelNames) {
+  (void)SYCLSource;
+  (void)IncludePairs;
+  (void)UserArgs;
+  (void)LogPtr;
+  (void)RegisteredKernelNames;
+
+  throw sycl::exception(sycl::errc::build,
+                        "kernel_compiler via sycl-jit is not available");
+}
+


With [[maybe_unused]], we should be able to avoid the double declaration and can just ifdef the body of the function. In any case, we don't need to duplicate the namespace declarations.

Yes, that's neat.

Signed-off-by: Julian Oppermann <[email protected]>

steffenlarsen · 2024-10-15T14:55:29Z

sycl/include/sycl/kernel_bundle_enums.hpp

+  opencl = 0,
+  spirv = 1,
+  sycl = 2 /* cuda */,
+  sycljit = 99


This is part of a public interface. Can we document it somewhere? Also, I don't think you need to make a jump, as long as we don't change it after it's merged.

Side note, I personally prefer sycl_jit.

This is only temporary until we complete functionality in follow-up PRs, so not intended to ever be exposed users.

I renamed the enum value and added a comment that this is temporary.

steffenlarsen · 2024-10-15T15:14:39Z

sycl/source/detail/jit_compiler.cpp

+
+  // TODO: Handle situation.
+  assert(RegisteredKernelNames.empty() &&
+         "Instantiation of kernel templates NYI");


Should we throw instead? In cases where asserts are disabled what would happen if execution continues from here?

Signed-off-by: Julian Oppermann <[email protected]>

cperkinsintel · 2024-10-15T15:48:40Z

sycl-jit/jit-compiler/lib/KernelFusion.cpp

+extern "C" JITResult compileSYCL(const char *SYCLSource,
+                                 View<IncludePair> IncludePairs,
+                                 View<const char *> UserArgs,
+                                 const char *DPCPPRoot) {


DPCPPRoot is not like the other arguments. Doesn't it seem like the routine should be able to figure that out itself, rather than being provided?

Makes sense; it's detected from inside the JIT library now.

cperkinsintel · 2024-10-15T15:51:31Z

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

+  CommandLine.append(UserArgs.begin(), UserArgs.end());
+  clang::tooling::FixedCompilationDatabase DB{"./", CommandLine};
+
+  constexpr auto SourcePath = "rtc.cpp";


rather than hard-code "rtc.cpp" in, can this 'fantasy name' for the file be an argument to the API? It might show up in debug information, so it might be useful to users to be able to disambiguate, rather than having every dynamic device compiled kernel originate with the same fictional file name.

Good idea, will do 👍

Done. I'm passing in a semi-random ID, same as the file-based implementation. There's no property yet in the extension to specify a file name or prefix, correct?

What would be the reason that a user might want to change this name? Why is it better to have a semi-random ID as a default rather than a fixed string?

I can imagine that the name might show up in error / log messages, for example, when there is a syntax error in the source string. If that's the only case the name is visible, it seems like having a fixed string like rtc.cpp would be fine, and probably preferable to a name with a random number.

I'm not opposed to adding a property which allows the user to set this name, but I think it should be an optional property because I think many people will not care what the name is.

I can imagine that the name might show up in error / log messages, for example, when there is a syntax error in the source string. If that's the only case the name is visible, [...]

Yes, that's the only case. I agree that the ID doesn't add much value here because we don't materialise anything on the actual filesystem. I'll keep the plumbing to pass the filename down to the JIT, but will set it rtc.cpp until there's a need and a means to modify it from the extension.

AlexeySachkov · 2024-10-16T09:34:48Z

sycl/source/detail/jit_compiler.cpp

+        sycl::detail::ur::getOsLibraryFuncAddress(LibraryPtr, "compileSYCL"));
+    if (!this->CompileSYCLHandle) {
+      printPerformanceWarning(
+          "Cannot resolve JIT library function entry point");


This sounds more serious than a mere performance warning :)

printPerformanceWarning is the generic error message helper in sycl-jit, but yes, I agree it's a bit of a misnomer when used here (and while attempting to set-up the other entrypoints before).

The reasoning behind that name is that so far, when something went wrong in sycl-jit, we would simply fall back to non-fused execution of kernels (in case of kernel fusion), which still leads to correct results, but potentially lower performance than expected.

But I'm happy to rename this function to something more suitable.

Signed-off-by: Julian Oppermann <[email protected]>

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp

Signed-off-by: Julian Oppermann <[email protected]>

steffenlarsen · 2024-10-18T06:44:00Z

sycl/include/sycl/kernel_bundle_enums.hpp

+  spirv = 1,
+  sycl = 2,
+  /* cuda */
+  sycl_jit /* temporary, alternative implementation for SYCL */


Since it is temporary, I think the previous value given (= 99) is a good idea. Problem is, if someone compiles with this in use and it crosses the library boundary, if we then remove it and another different enum value is added, compatibility could be broken. It's unlikely to happen, but no impossible.

steffenlarsen · 2024-10-18T06:46:26Z

sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp

 // CHECK-DAG: lit.cfg.py
 //
-// CHECK-NUM-MATCHES: 5
+// CHECK-NUM-MATCHES: 6


What has caused this change? I don't see it used in the added test and in general we shouldn't increase this number.

It's in the SYCL source string in the added E2E test:

llvm/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp

Lines 31 to 32 in 407124b

auto constexpr SYCLSource = R"===(

#include <sycl/sycl.hpp>

Signed-off-by: Julian Oppermann <[email protected]>

sommerlukas

Just one nit question

sommerlukas · 2024-10-21T06:57:40Z

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h

+#include <memory>
+
+namespace llvm {
+class Module;


Why is this forward declared and not simply included?

I think it was just a leftover from an early design point. Changed, thanks!

sommerlukas · 2024-10-21T06:59:21Z

sycl/source/detail/jit_compiler.cpp

+        sycl::detail::ur::getOsLibraryFuncAddress(LibraryPtr, "compileSYCL"));
+    if (!this->CompileSYCLHandle) {
+      printPerformanceWarning(
+          "Cannot resolve JIT library function entry point");


The reasoning behind that name is that so far, when something went wrong in sycl-jit, we would simply fall back to non-fused execution of kernels (in case of kernel fusion), which still leads to correct results, but potentially lower performance than expected.

But I'm happy to rename this function to something more suitable.

Signed-off-by: Julian Oppermann <[email protected]>

sommerlukas · 2024-10-24T13:28:53Z

@steffenlarsen @gmlueck @AlexeySachkov Do you want to review this PR again? Or are you fine with us merging this with approval from @cperkinsintel and me?

AlexeySachkov · 2024-10-24T14:41:44Z

@steffenlarsen @gmlueck @AlexeySachkov Do you want to review this PR again? Or are you fine with us merging this with approval from @cperkinsintel and me?

I'm fine with you merging the PR, my comments were very minor, feel free to ignore me here completely :)

steffenlarsen · 2024-10-24T14:50:55Z

@cperkinsintel has worked closer with this feature than I have, so I trust his judgement!

gmlueck · 2024-10-25T13:03:07Z

@steffenlarsen @gmlueck @AlexeySachkov Do you want to review this PR again? Or are you fine with us merging this with approval from @cperkinsintel and me?

Also OK with me. I mostly review interfaces, and I think this PR does not really change the interface to the kernel compiler.

Signed-off-by: Julian Oppermann <[email protected]>

Silences unused variable warning introduced with #15701. Signed-off-by: Julian Oppermann <[email protected]>

jopperm added 2 commits October 15, 2024 11:07

[SYCL] Add plumbing to implement kernel compiler extension with libto…

6e892f9

…oling and sycl-jit Signed-off-by: Julian Oppermann <[email protected]>

Obey SYCL_EXT_JIT_ENABLE.

246bc6d

Signed-off-by: Julian Oppermann <[email protected]>

jopperm requested review from a team as code owners October 15, 2024 12:10

jopperm requested a review from steffenlarsen October 15, 2024 12:10

jopperm had a problem deploying to WindowsCILock October 15, 2024 12:10 — with GitHub Actions Failure

sommerlukas requested review from AlexeySachkov, cperkinsintel and gmlueck October 15, 2024 12:16

sommerlukas reviewed Oct 15, 2024

View reviewed changes

Add missing file.

46e7127

Signed-off-by: Julian Oppermann <[email protected]>

jopperm had a problem deploying to WindowsCILock October 15, 2024 13:36 — with GitHub Actions Error

Tidy up conditional compilation.

7e06ae1

Signed-off-by: Julian Oppermann <[email protected]>

jopperm temporarily deployed to WindowsCILock October 15, 2024 14:02 — with GitHub Actions Inactive

jopperm had a problem deploying to WindowsCILock October 15, 2024 14:37 — with GitHub Actions Error

steffenlarsen reviewed Oct 15, 2024

View reviewed changes

Use CMake variable for clang source dir.

ef90b4b

Signed-off-by: Julian Oppermann <[email protected]>

jopperm temporarily deployed to WindowsCILock October 15, 2024 15:19 — with GitHub Actions Inactive

cperkinsintel reviewed Oct 15, 2024

View reviewed changes

jopperm temporarily deployed to WindowsCILock October 15, 2024 17:17 — with GitHub Actions Inactive

AlexeySachkov reviewed Oct 16, 2024

View reviewed changes

jopperm added 3 commits October 16, 2024 13:09

Detect (and cache) DPCPP root directory inside the JIT library.

bfe9b43

Signed-off-by: Julian Oppermann <[email protected]>

s/sycljit/sycl_jit/

2c2ac32

Signed-off-by: Julian Oppermann <[email protected]>

Account for added RTC test.

f5871cc

Signed-off-by: Julian Oppermann <[email protected]>

jopperm temporarily deployed to WindowsCILock October 16, 2024 12:29 — with GitHub Actions Inactive

jopperm had a problem deploying to WindowsCILock October 16, 2024 12:59 — with GitHub Actions Error

Don't hardcode source file name.

c60a528

Signed-off-by: Julian Oppermann <[email protected]>

jopperm had a problem deploying to WindowsCILock October 16, 2024 13:34 — with GitHub Actions Error

Throw exception instead of assert()-ing.

ce333dc

Signed-off-by: Julian Oppermann <[email protected]>

gmlueck reviewed Oct 17, 2024

View reviewed changes

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp Show resolved Hide resolved

Roll-back use of semi-random ID.

407124b

Signed-off-by: Julian Oppermann <[email protected]>

jopperm temporarily deployed to WindowsCILock October 17, 2024 14:27 — with GitHub Actions Inactive

jopperm temporarily deployed to WindowsCILock October 17, 2024 15:03 — with GitHub Actions Inactive

steffenlarsen reviewed Oct 18, 2024

View reviewed changes

Assign non-consecutive number to sycl_jit language.

8c4597e

Signed-off-by: Julian Oppermann <[email protected]>

jopperm temporarily deployed to WindowsCILock October 18, 2024 08:14 — with GitHub Actions Inactive

jopperm temporarily deployed to WindowsCILock October 18, 2024 08:44 — with GitHub Actions Inactive

cperkinsintel approved these changes Oct 18, 2024

View reviewed changes

sommerlukas approved these changes Oct 21, 2024

View reviewed changes

Don't forward-declare llvm::Module.

4a5c809

Signed-off-by: Julian Oppermann <[email protected]>

jopperm temporarily deployed to WindowsCILock October 21, 2024 08:05 — with GitHub Actions Inactive

jopperm temporarily deployed to WindowsCILock October 21, 2024 08:29 — with GitHub Actions Inactive

jopperm mentioned this pull request Oct 22, 2024

[SYCL] Implement device library linking for runtime compilation #15810

Merged

jopperm self-assigned this Oct 22, 2024

Merge remote-tracking branch 'upstream/sycl' into rtc-via-sycl-jit

863d4b5

Signed-off-by: Julian Oppermann <[email protected]>

jopperm had a problem deploying to WindowsCILock October 28, 2024 08:29 — with GitHub Actions Error

Drop #endif comment.

b727bb1

Signed-off-by: Julian Oppermann <[email protected]>

jopperm temporarily deployed to WindowsCILock October 28, 2024 08:41 — with GitHub Actions Inactive

jopperm temporarily deployed to WindowsCILock October 28, 2024 09:11 — with GitHub Actions Inactive

sommerlukas merged commit 51f8a05 into intel:sycl Oct 28, 2024
13 checks passed

jopperm mentioned this pull request Oct 28, 2024

[SYCL] Fix postcommit failure on macOS #15886

Merged

sommerlukas pushed a commit that referenced this pull request Oct 28, 2024

[SYCL] Fix postcommit failure on macOS (#15886)

8961312

Silences unused variable warning introduced with #15701. Signed-off-by: Julian Oppermann <[email protected]>

jopperm deleted the rtc-via-sycl-jit branch October 28, 2024 14:13

[SYCL] PoC implementation of kernel compiler extension with libtooling and sycl-jit #15701

[SYCL] PoC implementation of kernel compiler extension with libtooling and sycl-jit #15701

Uh oh!

Conversation

jopperm commented Oct 15, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

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!

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!

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!

Choose a reason for hiding this comment

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!

Choose a reason for hiding this comment

Uh oh!

sommerlukas 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!

sommerlukas commented Oct 24, 2024

Uh oh!

AlexeySachkov commented Oct 24, 2024

Uh oh!

steffenlarsen commented Oct 24, 2024

Uh oh!

gmlueck commented Oct 25, 2024

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

jopperm commented Oct 15, 2024 •

edited

Loading