[OpenMP] Build OpenMP DeviceRTL on SPIR-V #121600

sarnex · 2025-01-03T20:31:56Z

We actually can build it today with the current OMP FE.

We want to minimize device-specific code with SPIR-V, so I used the spirv64 triple instead of spirv64-intel which is currently the only planned SPIR-V based OMP offloading triple. The existing driver code finds the library fine.

I also made a minor change in the RPC code called when building the RTL, __has_builtin(__builtin_ia32_pause) was returning true and then we were getting an assert saying it's not supported on this target, it seems there are no features associated with __builtin_ia32_pause so __has_builtin always returns true on SPIR-V, so just add a arch check so it does nothing as there's no sleep we can use.

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

llvmbot · 2025-01-03T21:48:19Z

@llvm/pr-subscribers-offload

@llvm/pr-subscribers-libc

Author: Nick Sarnie (sarnex)

Changes

We actually can build it today with the current OMP FE.

We want to minimize device-specific code with SPIR-V, so I used the spirv64 triple instead of spirv64-intel which is currently the only planned SPIR-V based OMP offloading triple. The existing driver code finds the library fine.

I also made a minor change in the RPC code called when building the RTL, __has_builtin(__builtin_ia32_pause) was returning true and then we were getting an assert saying it's not supported on this target, it seems there are no features associated with __builtin_ia32_pause so __has_builtin always returns true on SPIR-V, so just add a arch check so it does nothing as there's no sleep we can use.

Full diff: https://github.com/llvm/llvm-project/pull/121600.diff

2 Files Affected:

(modified) libc/shared/rpc_util.h (+1-1)
(modified) offload/DeviceRTL/CMakeLists.txt (+3)

diff --git a/libc/shared/rpc_util.h b/libc/shared/rpc_util.h
index 9406de59f63b71..6b9df43776edac 100644
--- a/libc/shared/rpc_util.h
+++ b/libc/shared/rpc_util.h
@@ -157,7 +157,7 @@ RPC_ATTRS void sleep_briefly() {
     asm("nanosleep.u32 64;" ::: "memory");
 #elif defined(__AMDGPU__) && defined(RPC_TARGET_IS_GPU)
   __builtin_amdgcn_s_sleep(2);
-#elif __has_builtin(__builtin_ia32_pause)
+#elif __has_builtin(__builtin_ia32_pause) && !defined(__SPIRV__)
   __builtin_ia32_pause();
 #elif __has_builtin(__builtin_arm_isb)
   __builtin_arm_isb(0xf);
diff --git a/offload/DeviceRTL/CMakeLists.txt b/offload/DeviceRTL/CMakeLists.txt
index 22940264f9b19a..e707429a5196cf 100644
--- a/offload/DeviceRTL/CMakeLists.txt
+++ b/offload/DeviceRTL/CMakeLists.txt
@@ -264,6 +264,9 @@ compileDeviceRTLLibrary(amdgpu amdgcn-amd-amdhsa -Xclang -mcode-object-version=n
 add_custom_target(omptarget.devicertl.nvptx)
 compileDeviceRTLLibrary(nvptx nvptx64-nvidia-cuda --cuda-feature=+ptx63)
 
+add_custom_target(omptarget.devicertl.spirv64)
+compileDeviceRTLLibrary(spirv64 spirv64)
+
 # Archive all the object files generated above into a static library
 add_library(omptarget.devicertl STATIC)
 set_target_properties(omptarget.devicertl PROPERTIES

sarnex · 2025-01-03T21:51:16Z

Joseph said he's going to rework the DeviceRTL files soon, so let me know if I should hold off on/rework this.

Also let me know if I should put this under a CMake option or something while it's experimental.

jhuber6 · 2025-01-04T16:30:58Z

libc/shared/rpc_util.h

@@ -157,7 +157,7 @@ RPC_ATTRS void sleep_briefly() {
    asm("nanosleep.u32 64;" ::: "memory");
 #elif defined(__AMDGPU__) && defined(RPC_TARGET_IS_GPU)
  __builtin_amdgcn_s_sleep(2);
-#elif __has_builtin(__builtin_ia32_pause)
+#elif __has_builtin(__builtin_ia32_pause) && !defined(__SPIRV__)


Why is this necessary? This seems like a massive bug that needs to be fixed.

Let me debug more into it, will reply here when I have a better answer.

If this is showing up it means that the x64 builtins are being emitted, which is very bad. I'm going to guess something got crossed while defining the target info.

I debugged into it quickly before and found the problem was that it seems there's some map in the preprocessor between builtins and features those builtins require, and that feature map was empty and when it's empty __has_builtin returns true, but I don't actually know what that feature map repesents or why it's empty, will do a real investigation on this which I should have done to begin with :P

This turned out to be pretty annoying, I have a draft PR here that attempts to fix it, will mark it ready for review once CI passes.

Do you think it's okay to have this target-check workaround until we resolve the discussion on the __has_builtin PR?

The way you have it not just implicitly assumes the host is x64. If we're dedicated to working around that then you'd just need to put an empty ifdef at the top.

Thx, I just pushed a commit that hopefully fixes what you meant. If you feel strongly about not doing the workaround I'm happy to wait, but since IMO it's relatively simple and self-contained I was thinking we could not block the PR with that.

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

jhuber6 · 2025-01-07T22:09:50Z

libc/shared/rpc_util.h

@@ -152,12 +152,14 @@ template <typename T> class optional {

 /// Suspend the thread briefly to assist the thread scheduler during busy loops.
 RPC_ATTRS void sleep_briefly() {
-#if defined(__NVPTX__) && defined(RPC_TARGET_IS_GPU)
+#if defined(__SPIRV__)
+  // It's unsupported and __has_builtin doesn't always work as we expect.


I'd prefer a FIXME: __has_builtin does not work on offloading targets. or something along those lines.

done, thanks

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

jhuber6

It's definitely easier to support this now that I merged all the architecture into one. We'll need a tester at some point but as long as we're making forward progress.

sarnex · 2025-01-07T22:20:23Z

Thanks for that, we had an earlier version of this internally before your patches and it was definitely something.

sarnex · 2025-01-08T16:10:23Z

Hi @jdoerfert @shiltian, do you guys have any feedback on this? If not, I'll merge this tomorrow. Thanks!

jhuber6 · 2025-01-08T17:24:05Z

You know I'm surprised this works, considering we don't even have spirv64 hooked up to emit CGOpenMPRuntimeGPU. Do we have anything that confirms the code generated here makes sense? Maybe we should hold off on landing before then.

sarnex · 2025-01-08T17:32:40Z

You know I'm surprised this works, considering we don't even have spirv64 hooked up to emit CGOpenMPRuntimeGPU. Do we have anything that confirms the code generated here makes sense? Maybe we should hold off on landing before then.

Right, that's the kind of stuff I was referencing when I was saying we need work in the FE. The DeviceRTL won't be able to look up/write to variables like kernel_environment because the FE doesn't generate them yet.

I agree this won't have use to anyone trying to seriously use SPIR-V with OpenMP, but it is useful for LLVM developers working on OMP/SPIR-V. Maybe it's better to put this under a default-off CMake option?

jhuber6 · 2025-01-08T19:15:08Z

Unsure, I might just remove -fopenmp from the DeviceRTL while I'm at it, but I don't know how SPIRV will handle that since it works fine for AMD / NVIDIA.

sarnex · 2025-01-08T19:18:24Z

Sure, I can test it if you end up trying that.

jhuber6 · 2025-01-09T15:34:26Z

Also this probably wouldn't do anything because we don't have any of the SPIR-V builtins to do much of anything, so we should probably hold off on landing it since it can't be used.

sarnex · 2025-01-09T15:36:51Z

Sounds good, will mark as draft for now.

[OpenMP] Build OpenMP DeviceRTL on SPIRV64

b2d1534

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

sarnex changed the title ~~[OpenMP] Build OpenMP DeviceRTL on SPIRV64~~ [OpenMP] Build OpenMP DeviceRTL on SPIR-V Jan 3, 2025

sarnex marked this pull request as ready for review January 3, 2025 21:47

llvmbot added libc offload labels Jan 3, 2025

lntue requested a review from jhuber6 January 4, 2025 16:24

jhuber6 reviewed Jan 4, 2025

View reviewed changes

jdoerfert requested review from jdoerfert and shiltian January 6, 2025 18:15

Fix __has_builtin workaround

c40b5b6

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

jhuber6 reviewed Jan 7, 2025

View reviewed changes

Update comment

f0cc084

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

jhuber6 approved these changes Jan 7, 2025

View reviewed changes

Merge remote-tracking branch 'upstream/main' into HEAD

acb36a1

shiltian approved these changes Jan 8, 2025

View reviewed changes

sarnex marked this pull request as draft January 9, 2025 15:36

[OpenMP] Build OpenMP DeviceRTL on SPIR-V #121600

Are you sure you want to change the base?

[OpenMP] Build OpenMP DeviceRTL on SPIR-V #121600

Uh oh!

Conversation

sarnex commented Jan 3, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jan 3, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

sarnex commented Jan 3, 2025 • 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!

Choose a reason for hiding this comment

Uh oh!

sarnex Jan 6, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

sarnex Jan 6, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

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!

jhuber6 left a comment

Choose a reason for hiding this comment

Uh oh!

sarnex commented Jan 7, 2025

Uh oh!

sarnex commented Jan 8, 2025

Uh oh!

jhuber6 commented Jan 8, 2025

Uh oh!

sarnex commented Jan 8, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

jhuber6 commented Jan 8, 2025

Uh oh!

sarnex commented Jan 8, 2025

Uh oh!

jhuber6 commented Jan 9, 2025

Uh oh!

sarnex commented Jan 9, 2025

Uh oh!

Uh oh!

sarnex commented Jan 3, 2025 •

edited

Loading

llvmbot commented Jan 3, 2025 •

edited

Loading

sarnex commented Jan 3, 2025 •

edited

Loading

sarnex Jan 6, 2025 •

edited

Loading

sarnex Jan 6, 2025 •

edited

Loading

sarnex commented Jan 8, 2025 •

edited

Loading