[llvm-branch-commits] [llvm] [Offload][Lang] Add proper DeviceSync (PR #216433)
Sophia Herrmann via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Aug 18 10:38:23 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/216433
>From 7193df90f3c921cd5495181406afb71b2fe5c7d2 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Thu, 13 Aug 2026 16:48:02 -0700
Subject: [PATCH] add proper deviceSync
---
.../languages/kernel/src/LanguageRuntime.cpp | 19 ++--
.../CUDA/basic_launch_blocks_and_threads.cu | 2 +-
.../offloading/CUDA/basic_launch_multi_arg.cu | 2 +-
.../offloading/CUDA/devicesync_streams.cu | 98 +++++++++++++++++++
offload/test/offloading/CUDA/launch_tu.cu | 2 +-
offload/test/offloading/CUDA/syncthreads.cu | 1 +
.../offloading/CUDA/thread_and_block_id.cu | 1 +
.../HIP/basic_launch_blocks_and_threads.hip | 2 +-
.../offloading/HIP/basic_launch_multi_arg.hip | 2 +-
.../offloading/HIP/devicesync_streams.hip | 97 ++++++++++++++++++
offload/test/offloading/HIP/launch_tu.hip | 2 +-
offload/test/offloading/HIP/syncthreads.hip | 1 +
.../offloading/HIP/thread_and_block_id.hip | 1 +
13 files changed, 218 insertions(+), 12 deletions(-)
create mode 100644 offload/test/offloading/CUDA/devicesync_streams.cu
create mode 100644 offload/test/offloading/HIP/devicesync_streams.hip
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 537dcea3bd142..2f6f0573d69e1 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
+#include "llvm/ADT/SmallPtrSet.h"
#ifndef LANGUAGE
#error This file should be included, or used, with a LANGUAGE macro set.
#endif
@@ -23,7 +24,6 @@
#include "Types.h"
#include "OffloadAPI.h"
-#include "llvm/ADT/SmallVector.h"
#include <cassert>
#include <cstdio>
@@ -101,11 +101,18 @@ Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
}
Error_t DeviceSynchronize() {
- // TODO: This is not correct. We likely want to pipe this through to the
- // plugins.
- ol_queue_handle_t Queue = ThreadState::getDefaultQueue();
- ol_result_t Result = olSyncQueue(Queue);
- return convertAndSetLastError(Result);
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ if (!Device)
+ return setLastError(ErrorInvalidDevice);
+
+ llvm::SmallPtrSet<StreamTy *, 8> DeviceStreams =
+ RuntimeState::getDeviceStreams(Device);
+ for (StreamTy *Stream : DeviceStreams) {
+ ol_result_t Result = olSyncQueue(Stream->Queue);
+ if (Result != OL_SUCCESS)
+ return convertAndSetLastError(Result);
+ }
+ return setLastError(Success);
}
Error_t GetDevice(int *DeviceNo) {
diff --git a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
index f82cec6d4692e..6e40fb695c7e1 100644
--- a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
+++ b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
@@ -19,7 +19,6 @@ __global__ void incrementCounter(int *A) {
}
int main(int argc, char **argv) {
- int DevNo = 0;
int *Ptr, I;
cudaMalloc(&Ptr, sizeof(int));
printf("Ptr %p\n", Ptr);
@@ -27,6 +26,7 @@ int main(int argc, char **argv) {
int Zero = 0;
cudaMemcpy(Ptr, &Zero, sizeof(int), cudaMemcpyHostToDevice);
incrementCounter<<<7, 6>>>(Ptr);
+ cudaDeviceSynchronize();
cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
printf("I: %i\n", I);
// CHECK: I: 42
diff --git a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
index 505a9f9379c08..25207536496e7 100644
--- a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
+++ b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
@@ -20,7 +20,6 @@ __global__ void square(int *Dst, short Q, int *Src, short P) {
}
int main(int argc, char **argv) {
- int DevNo = 0;
int *Src, *Ptr;
cudaMalloc(&Ptr, 4);
cudaMalloc(&Src, 8);
@@ -30,6 +29,7 @@ int main(int argc, char **argv) {
cudaMemcpy(Ptr, &I, sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(Src, &HostSrc[0], 2 * sizeof(int), cudaMemcpyHostToDevice);
square<<<1, 1>>>(Ptr, 3, Src, 4);
+ cudaDeviceSynchronize();
cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
cudaMemcpy(&HostSrc[0], Src, 2 * sizeof(int), cudaMemcpyDeviceToHost);
printf("I: %i\n", I);
diff --git a/offload/test/offloading/CUDA/devicesync_streams.cu b/offload/test/offloading/CUDA/devicesync_streams.cu
new file mode 100644
index 0000000000000..a8a547101fbaf
--- /dev/null
+++ b/offload/test/offloading/CUDA/devicesync_streams.cu
@@ -0,0 +1,98 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fgpu-default-stream=legacy -pthread -std=c++17
+// RUN: %t | %fcheck-generic --check-prefix=CHECK
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fgpu-default-stream=per-thread -pthread -std=c++17
+// RUN: %t | %fcheck-generic --check-prefix=CHECK
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <chrono>
+#include <cstdio>
+#include <thread>
+
+__global__ void waitThenSet(volatile int *Gate, volatile int *Out, int Value) {
+ for (unsigned long long I = 0; I < 1000000000ULL && *Gate == 0; ++I)
+ ;
+ *Out = *Gate ? Value : -Value;
+}
+
+int main(int argc, char **argv) {
+ cudaStream_t BlockingStream = nullptr;
+ if (cudaStreamCreateWithFlags(&BlockingStream, cudaStreamDefault) !=
+ cudaSuccess)
+ return 1;
+ cudaStream_t NonBlockingStream = nullptr;
+ if (cudaStreamCreateWithFlags(&NonBlockingStream, cudaStreamNonBlocking) !=
+ cudaSuccess)
+ return 1;
+
+ int *BlockingGate = nullptr;
+ int *NonBlockingGate = nullptr;
+ int *BlockingOutStorage = nullptr;
+ int *NonBlockingOutStorage = nullptr;
+ if (cudaHostAlloc(&BlockingGate, sizeof(int), cudaHostAllocDefault) !=
+ cudaSuccess)
+ return 1;
+ if (cudaHostAlloc(&NonBlockingGate, sizeof(int), cudaHostAllocDefault) !=
+ cudaSuccess)
+ return 1;
+ if (cudaHostAlloc(&BlockingOutStorage, sizeof(int), cudaHostAllocDefault) !=
+ cudaSuccess)
+ return 1;
+ if (cudaHostAlloc(&NonBlockingOutStorage, sizeof(int),
+ cudaHostAllocDefault) != cudaSuccess)
+ return 1;
+
+ volatile int *BlockingOut = BlockingOutStorage;
+ volatile int *NonBlockingOut = NonBlockingOutStorage;
+ *BlockingGate = 0;
+ *NonBlockingGate = 0;
+ *BlockingOut = 0;
+ *NonBlockingOut = 0;
+
+ waitThenSet<<<1, 1, 0, BlockingStream>>>(BlockingGate, BlockingOut, 17);
+ waitThenSet<<<1, 1, 0, NonBlockingStream>>>(NonBlockingGate, NonBlockingOut,
+ 23);
+
+ std::thread Releaser([&]() {
+ std::this_thread::sleep_for(std::chrono::milliseconds(250));
+ *BlockingGate = 1;
+ *NonBlockingGate = 1;
+ });
+
+ cudaError_t SyncResult = cudaDeviceSynchronize();
+
+ if (SyncResult == cudaSuccess) {
+ printf("device sync waited on blocking stream: %d\n", *BlockingOut);
+ // CHECK: device sync waited on blocking stream: 17
+ printf("device sync waited on nonblocking stream: %d\n", *NonBlockingOut);
+ // CHECK: device sync waited on nonblocking stream: 23
+ }
+
+ Releaser.join();
+ if (cudaStreamSynchronize(BlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaStreamSynchronize(NonBlockingStream) != cudaSuccess)
+ return 1;
+ if (SyncResult != cudaSuccess)
+ return 1;
+
+ if (cudaStreamDestroy(BlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaStreamDestroy(NonBlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaFreeHost(BlockingGate) != cudaSuccess)
+ return 1;
+ if (cudaFreeHost(NonBlockingGate) != cudaSuccess)
+ return 1;
+ if (cudaFreeHost(BlockingOutStorage) != cudaSuccess)
+ return 1;
+ if (cudaFreeHost(NonBlockingOutStorage) != cudaSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/CUDA/launch_tu.cu b/offload/test/offloading/CUDA/launch_tu.cu
index 8b92194ba435e..fc24ec1af03b9 100644
--- a/offload/test/offloading/CUDA/launch_tu.cu
+++ b/offload/test/offloading/CUDA/launch_tu.cu
@@ -17,13 +17,13 @@
extern __global__ void square(int *A);
int main(int argc, char **argv) {
- int DevNo = 0;
int *Ptr;
cudaMalloc(&Ptr, 4);
printf("Ptr %p\n", Ptr);
// CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
int I;
+ cudaDeviceSynchronize();
cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
printf("I: %i\n", I);
// CHECK: I: 42
diff --git a/offload/test/offloading/CUDA/syncthreads.cu b/offload/test/offloading/CUDA/syncthreads.cu
index 0c6048c32f824..4c839b85ff768 100644
--- a/offload/test/offloading/CUDA/syncthreads.cu
+++ b/offload/test/offloading/CUDA/syncthreads.cu
@@ -33,6 +33,7 @@ int main(int argc, char **argv) {
int Result = 0;
cudaMalloc(&DevPtr, sizeof(int));
reduceBlock<<<1, 64>>>(DevPtr);
+ cudaDeviceSynchronize();
cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost);
printf("sum: %i\n", Result);
diff --git a/offload/test/offloading/CUDA/thread_and_block_id.cu b/offload/test/offloading/CUDA/thread_and_block_id.cu
index 76c45a7992a46..a56c9ff33e2ab 100644
--- a/offload/test/offloading/CUDA/thread_and_block_id.cu
+++ b/offload/test/offloading/CUDA/thread_and_block_id.cu
@@ -33,6 +33,7 @@ int main(int argc, char **argv) {
printf("DevPtr %p\n", DevPtr);
// CHECK: DevPtr [[DevPtr:0x.*]]
fill<<<NBlocks, NThreads>>>(DevPtr);
+ cudaDeviceSynchronize();
cudaMemcpy(Ptr, DevPtr, Size, cudaMemcpyDeviceToHost);
for (int I = 0; I < NBlocks * NThreads; ++I) {
diff --git a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
index d0b54f6a2c8d3..4f5ce89130052 100644
--- a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
+++ b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
@@ -19,7 +19,6 @@ __global__ void incrementCounter(int *A) {
}
int main(int argc, char **argv) {
- int DevNo = 0;
int *Ptr, I;
hipMalloc(&Ptr, sizeof(int));
printf("Ptr %p\n", Ptr);
@@ -27,6 +26,7 @@ int main(int argc, char **argv) {
int Zero = 0;
hipMemcpy(Ptr, &Zero, sizeof(int), hipMemcpyHostToDevice);
incrementCounter<<<7, 6>>>(Ptr);
+ hipDeviceSynchronize();
hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
printf("I: %i\n", I);
// CHECK: I: 42
diff --git a/offload/test/offloading/HIP/basic_launch_multi_arg.hip b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
index 6e599d6704598..3bca0a7fae484 100644
--- a/offload/test/offloading/HIP/basic_launch_multi_arg.hip
+++ b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
@@ -20,7 +20,6 @@ __global__ void square(int *Dst, short Q, int *Src, short P) {
}
int main(int argc, char **argv) {
- int DevNo = 0;
int *Src, *Ptr;
hipMalloc(&Ptr, 4);
hipMalloc(&Src, 8);
@@ -30,6 +29,7 @@ int main(int argc, char **argv) {
hipMemcpy(Ptr, &I, sizeof(int), hipMemcpyHostToDevice);
hipMemcpy(Src, &HostSrc[0], 2*sizeof(int), hipMemcpyHostToDevice);
square<<<1, 1>>>(Ptr, 3, Src, 4);
+ hipDeviceSynchronize();
hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
hipMemcpy(&HostSrc[0], Src, 2 * sizeof(int), hipMemcpyDeviceToHost);
printf("I: %i\n", I);
diff --git a/offload/test/offloading/HIP/devicesync_streams.hip b/offload/test/offloading/HIP/devicesync_streams.hip
new file mode 100644
index 0000000000000..14ac798f56ea8
--- /dev/null
+++ b/offload/test/offloading/HIP/devicesync_streams.hip
@@ -0,0 +1,97 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fgpu-default-stream=legacy -pthread -std=c++17
+// RUN: %t | %fcheck-generic --check-prefix=CHECK
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fgpu-default-stream=per-thread -pthread -std=c++17
+// RUN: %t | %fcheck-generic --check-prefix=CHECK
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <chrono>
+#include <cstdio>
+#include <thread>
+
+__global__ void waitThenSet(volatile int *Gate, volatile int *Out, int Value) {
+ for (unsigned long long I = 0; I < 1000000000ULL && *Gate == 0; ++I)
+ ;
+ *Out = *Gate ? Value : -Value;
+}
+
+int main(int argc, char **argv) {
+ hipStream_t BlockingStream = nullptr;
+ if (hipStreamCreateWithFlags(&BlockingStream, hipStreamDefault) != hipSuccess)
+ return 1;
+ hipStream_t NonBlockingStream = nullptr;
+ if (hipStreamCreateWithFlags(&NonBlockingStream, hipStreamNonBlocking) !=
+ hipSuccess)
+ return 1;
+
+ int *BlockingGate = nullptr;
+ int *NonBlockingGate = nullptr;
+ int *BlockingOutStorage = nullptr;
+ int *NonBlockingOutStorage = nullptr;
+ if (hipHostAlloc(&BlockingGate, sizeof(int), hipHostAllocDefault) !=
+ hipSuccess)
+ return 1;
+ if (hipHostAlloc(&NonBlockingGate, sizeof(int), hipHostAllocDefault) !=
+ hipSuccess)
+ return 1;
+ if (hipHostAlloc(&BlockingOutStorage, sizeof(int), hipHostAllocDefault) !=
+ hipSuccess)
+ return 1;
+ if (hipHostAlloc(&NonBlockingOutStorage, sizeof(int), hipHostAllocDefault) !=
+ hipSuccess)
+ return 1;
+
+ volatile int *BlockingOut = BlockingOutStorage;
+ volatile int *NonBlockingOut = NonBlockingOutStorage;
+ *BlockingGate = 0;
+ *NonBlockingGate = 0;
+ *BlockingOut = 0;
+ *NonBlockingOut = 0;
+
+ waitThenSet<<<1, 1, 0, BlockingStream>>>(BlockingGate, BlockingOut, 17);
+ waitThenSet<<<1, 1, 0, NonBlockingStream>>>(NonBlockingGate, NonBlockingOut,
+ 23);
+
+ std::thread Releaser([&]() {
+ std::this_thread::sleep_for(std::chrono::milliseconds(250));
+ *BlockingGate = 1;
+ *NonBlockingGate = 1;
+ });
+
+ hipError_t SyncResult = hipDeviceSynchronize();
+
+ if (SyncResult == hipSuccess) {
+ printf("device sync waited on blocking stream: %d\n", *BlockingOut);
+ // CHECK: device sync waited on blocking stream: 17
+ printf("device sync waited on nonblocking stream: %d\n", *NonBlockingOut);
+ // CHECK: device sync waited on nonblocking stream: 23
+ }
+
+ Releaser.join();
+ if (hipStreamSynchronize(BlockingStream) != hipSuccess)
+ return 1;
+ if (hipStreamSynchronize(NonBlockingStream) != hipSuccess)
+ return 1;
+ if (SyncResult != hipSuccess)
+ return 1;
+
+ if (hipStreamDestroy(BlockingStream) != hipSuccess)
+ return 1;
+ if (hipStreamDestroy(NonBlockingStream) != hipSuccess)
+ return 1;
+ if (hipFreeHost(BlockingGate) != hipSuccess)
+ return 1;
+ if (hipFreeHost(NonBlockingGate) != hipSuccess)
+ return 1;
+ if (hipFreeHost(BlockingOutStorage) != hipSuccess)
+ return 1;
+ if (hipFreeHost(NonBlockingOutStorage) != hipSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/HIP/launch_tu.hip b/offload/test/offloading/HIP/launch_tu.hip
index 03073029ca211..20a5d6b0ff6da 100644
--- a/offload/test/offloading/HIP/launch_tu.hip
+++ b/offload/test/offloading/HIP/launch_tu.hip
@@ -17,13 +17,13 @@
extern __global__ void square(int *A);
int main(int argc, char **argv) {
- int DevNo = 0;
int *Ptr;
hipMalloc(&Ptr, 4);
printf("Ptr %p\n", Ptr);
// CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
int I;
+ hipDeviceSynchronize();
hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
printf("I: %i\n", I);
// CHECK: I: 42
diff --git a/offload/test/offloading/HIP/syncthreads.hip b/offload/test/offloading/HIP/syncthreads.hip
index 5962ab5468b86..81e9ed0451ca1 100644
--- a/offload/test/offloading/HIP/syncthreads.hip
+++ b/offload/test/offloading/HIP/syncthreads.hip
@@ -33,6 +33,7 @@ int main(int argc, char **argv) {
int Result = 0;
hipMalloc(&DevPtr, sizeof(int));
reduceBlock<<<1, 64>>>(DevPtr);
+ hipDeviceSynchronize();
hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost);
printf("sum: %i\n", Result);
diff --git a/offload/test/offloading/HIP/thread_and_block_id.hip b/offload/test/offloading/HIP/thread_and_block_id.hip
index 5f9ff157d2a9c..af4daf689e678 100644
--- a/offload/test/offloading/HIP/thread_and_block_id.hip
+++ b/offload/test/offloading/HIP/thread_and_block_id.hip
@@ -33,6 +33,7 @@ int main(int argc, char **argv) {
printf("DevPtr %p\n", DevPtr);
// CHECK: DevPtr [[DevPtr:0x.*]]
fill<<<NBlocks, NThreads>>>(DevPtr);
+ hipDeviceSynchronize();
hipMemcpy(Ptr, DevPtr, Size, hipMemcpyDeviceToHost);
for (int I = 0; I < NBlocks * NThreads; ++I) {
More information about the llvm-branch-commits
mailing list