https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/218050
>From 4f93033b4abceddd2db4927b0ac5a34df84fb539 Mon Sep 17 00:00:00 2001 From: Sophia Herrmann <[email protected]> Date: Thu, 13 Aug 2026 16:48:02 -0700 Subject: [PATCH] add proper deviceSync --- .../languages/kernel/src/LanguageRuntime.cpp | 20 ++-- .../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(+), 13 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 75a00ae88a77f..796ddc9d7cfbb 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> @@ -87,12 +87,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. - StreamTy *DefaultStream = ThreadStateTy::get().getDefaultStream(); - ol_result_t Result = - DefaultStream ? DefaultStream->sync() : olSyncQueue(nullptr); - return convertAndSetLastError(Result); + ol_device_handle_t Device = ThreadStateTy::get().getDefaultDevice(); + if (!Device) + return setLastError(ErrorInvalidDevice); + + llvm::SmallPtrSet<StreamTy *, 8> DeviceStreams = + StateTy::get().getDeviceStreams(Device); + for (StreamTy *Stream : DeviceStreams) { + ol_result_t Result = Stream->sync(); + 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) { _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
