Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
wants to merge 4 commits into
base: main
Choose a base branch
from
Draft

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
b2d1534
[OpenMP] Build OpenMP DeviceRTL on SPIRV64
sarnex Jan 3, 2025
c40b5b6
Fix __has_builtin workaround
sarnex Jan 7, 2025
f0cc084
Update comment
sarnex Jan 7, 2025
acb36a1
Merge remote-tracking branch 'upstream/main' into HEAD
sarnex Jan 7, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion libc/shared/rpc_util.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -152,7 +152,9 @@ template <typename T> class optional {

/// Suspend the thread briefly to assist the thread scheduler during busy loops.
RPC_ATTRS void sleep_briefly() {
#if __has_builtin(__nvvm_reflect)
#if defined(__SPIRV__)
// FIXME: __has_builtin does not work on offloading targets.
#elif __has_builtin(__nvvm_reflect)
if (__nvvm_reflect("__CUDA_ARCH") >= 700)
asm("nanosleep.u32 64;" ::: "memory");
#elif __has_builtin(__builtin_amdgcn_s_sleep)
Expand Down
3 changes: 3 additions & 0 deletions offload/DeviceRTL/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
Loading