[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 <julian.oppermann@codeplay.com>

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

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

sycl-jit/jit-compiler/CMakeLists.txt

@@ -40,6 +49,8 @@ target_include_directories(sycl-jit
  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.

sommerlukas · 2024-10-15T12:24:06Z

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

+
+  auto *LLVMCtx = &Module->getContext();
+  Module.reset();
+  delete LLVMCtx;


How expensive is it to set up and destroy the LLVMContext on every call to RTC? Would it be an alternative to use the context from JITContext?

If I'm not mistaken, certain things like constants and metadata are stored within LLVMContext and won't be deallocated even if a module which references them is deallocated. Therefore, keeping LLVMContext between RTC call invocations could result in some memory build-up.

At least that is the behavior we discovered a few years ago when we were debugging exceptionally huge memory footprint of sycl-post-link where we had a huge codebase compiled with per-kernel device code split. I don't know if anything has changed since that in the upstream LLVM, but we hadn't proposed any patches back then.

It's certainly possible to pass an existing context into the ToolAction, but that also raises questions of thread safety.

For the performance implications, a) yes, looks like setting up a context does involve a non-trivial amount of work, and b) still seems true that types, constants and metadata are allocated in the context and not freed when the module is destroyed. I'd propose to keep the simple implementation for now, and will look out for the context setup overhead once we start benchmarking the RTC.

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 <julian.oppermann@codeplay.com>

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 <julian.oppermann@codeplay.com>

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).

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

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 <julian.oppermann@codeplay.com>

Obey SYCL_EXT_JIT_ENABLE.

246bc6d

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

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 cperkinsintel, gmlueck and AlexeySachkov October 15, 2024 12:16

sommerlukas reviewed Oct 15, 2024

View reviewed changes

Add missing file.

46e7127

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

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 <julian.oppermann@codeplay.com>

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 <julian.oppermann@codeplay.com>

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 <julian.oppermann@codeplay.com>

s/sycljit/sycl_jit/

2c2ac32

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

Account for added RTC test.

f5871cc

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

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 <julian.oppermann@codeplay.com>

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 <julian.oppermann@codeplay.com>

jopperm temporarily deployed to WindowsCILock October 16, 2024 13:51 — with GitHub Actions Inactive

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

Don't use std::filesystem.

0ad693d

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

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

Clean imports.

7a928b3

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

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

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

Drop __sycl_kernel_ prefix for sycl_jit language, too.

c2eb448

Signed-off-by: Julian Oppermann <julian.oppermann@codeplay.com>

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

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

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 <julian.oppermann@codeplay.com>

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

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[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

jopperm commented Oct 15, 2024 •

edited

Loading

sommerlukas Oct 15, 2024

jopperm Oct 15, 2024

sommerlukas Oct 15, 2024

AlexeySachkov Oct 16, 2024

jopperm Oct 16, 2024

sommerlukas Oct 15, 2024

jopperm Oct 15, 2024

sommerlukas Oct 15, 2024

jopperm Oct 15, 2024

steffenlarsen Oct 15, 2024

sommerlukas Oct 15, 2024

jopperm Oct 16, 2024

steffenlarsen Oct 15, 2024

cperkinsintel Oct 15, 2024

jopperm Oct 16, 2024

cperkinsintel Oct 15, 2024

jopperm Oct 16, 2024

jopperm Oct 16, 2024

gmlueck Oct 16, 2024

jopperm Oct 17, 2024

AlexeySachkov Oct 16, 2024

jopperm Oct 16, 2024

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

Are you sure you want to change the base?

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

Conversation

jopperm commented Oct 15, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

jopperm commented Oct 15, 2024 •

edited

Loading