[llvm-branch-commits] [clang] [llvm] [Offload][Lang] Add blocking to LaunchKernel and Memcpy (PR #216430)
Sophia Herrmann via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Aug 18 10:37:01 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/216430
>From 5257b5626251646ff2f98175732116db2969d499 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Thu, 13 Aug 2026 16:47:23 -0700
Subject: [PATCH] add blocking semantics to LaunchKernel and Memcpy
---
clang/lib/CodeGen/CGCUDANV.cpp | 4 +-
clang/lib/Driver/ToolChains/Clang.cpp | 1 +
.../languages/kernel/include/LanguageUtils.h | 54 +++++++-
.../languages/kernel/src/LanguageLaunch.cpp | 20 ++-
.../languages/kernel/src/LanguageRuntime.cpp | 25 +++-
.../CUDA/blocking_stream_semantics.cu | 131 ++++++++++++++++++
offload/test/offloading/CUDA/stream_api.cu | 6 +-
.../HIP/blocking_stream_semantics.hip | 125 +++++++++++++++++
offload/test/offloading/HIP/stream_api.hip | 6 +-
9 files changed, 355 insertions(+), 17 deletions(-)
create mode 100644 offload/test/offloading/CUDA/blocking_stream_semantics.cu
create mode 100644 offload/test/offloading/HIP/blocking_stream_semantics.hip
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index e03b7e754ab3f..14904c52b98ff 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -442,7 +442,9 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
std::string KernelLaunchAPI = "LaunchKernel";
if (CGF.getLangOpts().GPUDefaultStream ==
LangOptions::GPUDefaultStreamKind::PerThread) {
- if (CGF.getLangOpts().HIP)
+ if (CGF.getLangOpts().OffloadViaLLVM)
+ KernelLaunchAPI = KernelLaunchAPI + "";
+ else if (CGF.getLangOpts().HIP)
KernelLaunchAPI = KernelLaunchAPI + "_spt";
else if (CGF.getLangOpts().CUDA)
KernelLaunchAPI = KernelLaunchAPI + "_ptsz";
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 54583fe3abbd8..8b21400ab959f 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -8396,6 +8396,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
if (IsHIP) {
CmdArgs.push_back("-fcuda-allow-variadic-functions");
+ /// TODO: Why is this not forwarded when IsCUDA?
Args.AddLastArg(CmdArgs, options::OPT_fgpu_default_stream_EQ);
}
diff --git a/offload/languages/kernel/include/LanguageUtils.h b/offload/languages/kernel/include/LanguageUtils.h
index 730857e00461f..e7f95e0d0b485 100644
--- a/offload/languages/kernel/include/LanguageUtils.h
+++ b/offload/languages/kernel/include/LanguageUtils.h
@@ -13,6 +13,11 @@
#include "OffloadAPI.h"
#include "State.h"
#include "Stream.h"
+#include "llvm/ADT/SmallVector.h"
+
+using RuntimeState = llvm::offload::StateTy;
+using ThreadState = llvm::offload::ThreadStateTy;
+using StreamTy = llvm::offload::StreamTy;
/// Convert an ol_result_t to the active language's Error_t.
static inline Error_t convertResult(ol_result_t Result) {
@@ -40,8 +45,7 @@ static inline Error_t convertResult(ol_result_t Result) {
/// Set the last error for the current thread and return it.
static inline Error_t setLastError(Error_t Error) {
// TODO: find a more efficient way to set last error
- return static_cast<Error_t>(
- llvm::offload::ThreadStateTy::setLastError(Error));
+ return static_cast<Error_t>(ThreadState::setLastError(Error));
}
/// Convert an ol_result_t to the active language's Error_t and set it as the
@@ -51,12 +55,54 @@ static inline Error_t convertAndSetLastError(ol_result_t Result) {
}
/// Convert between the language-facing opaque stream and the internal stream.
-static inline Stream_t makeLanguageStream(llvm::offload::StreamTy *Stream) {
+static inline Stream_t makeLanguageStream(StreamTy *Stream) {
return reinterpret_cast<Stream_t>(Stream);
}
static inline llvm::offload::StreamTy *getInternalStream(Stream_t Stream) {
- return reinterpret_cast<llvm::offload::StreamTy *>(Stream);
+ return reinterpret_cast<StreamTy *>(Stream);
+}
+
+/// Wait for blocking streams before executing on the legacy default stream.
+static inline ol_result_t waitOnBlockingStreams(StreamTy *LegacyDefaultStream,
+ ol_device_handle_t Device) {
+ if (!LegacyDefaultStream ||
+ LegacyDefaultStream->Kind != llvm::offload::QueueKind::LegacyDefault)
+ return OL_SUCCESS;
+
+ llvm::SmallVector<ol_event_handle_t, 8> Events;
+ for (StreamTy *BlockingStream : RuntimeState::getBlockingStreams(Device)) {
+ ol_event_handle_t Event = nullptr;
+ ol_result_t Result =
+ olCreateEvent(BlockingStream->Queue, OL_EVENT_FLAGS_NONE, &Event);
+ if (Result != OL_SUCCESS)
+ return Result;
+ Events.push_back(Event);
+ }
+
+ if (Events.empty())
+ return OL_SUCCESS;
+ return olWaitEvents(LegacyDefaultStream->Queue, Events.data(), Events.size());
+}
+
+/// Wait for the legacy default stream to complete before launching a kernel on
+/// a blocking stream.
+static inline ol_result_t waitOnLegacyDefaultStream(StreamTy *SourceStream,
+ ol_device_handle_t Device) {
+ if (!RuntimeState::hasLegacyDefaultStream(Device))
+ return OL_SUCCESS;
+
+ StreamTy *DefaultStream = ThreadState::getDefaultStream();
+ if (!DefaultStream ||
+ DefaultStream->Kind != llvm::offload::QueueKind::LegacyDefault)
+ return OL_SUCCESS;
+
+ ol_event_handle_t Event = nullptr;
+ ol_result_t Result =
+ olCreateEvent(DefaultStream->Queue, OL_EVENT_FLAGS_NONE, &Event);
+ if (Result != OL_SUCCESS)
+ return Result;
+ return olWaitEvents(SourceStream->Queue, &Event, 1);
}
/// Convert a Stream_t to an ol_queue_handle_t.
diff --git a/offload/languages/kernel/src/LanguageLaunch.cpp b/offload/languages/kernel/src/LanguageLaunch.cpp
index cbbe9e41291b8..4ca8b8dee6430 100644
--- a/offload/languages/kernel/src/LanguageLaunch.cpp
+++ b/offload/languages/kernel/src/LanguageLaunch.cpp
@@ -8,6 +8,7 @@
#include "LanguageLaunch.h"
#include "LanguageUtils.h"
+#include "OffloadAPI.h"
#include "State.h"
#include "Stream.h"
#include <cstdio>
@@ -56,8 +57,21 @@ ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
LaunchSizeArgs.GroupSize.z = BlockDim.z;
LaunchSizeArgs.DynSharedMemory = DynamicSharedMem;
- ol_queue_handle_t Queue = Stream ? reinterpret_cast<StreamTy *>(Stream)->Queue
- : ThreadState::getDefaultQueue();
+ StreamTy *LaunchStream = Stream ? reinterpret_cast<StreamTy *>(Stream)
+ : ThreadState::getDefaultStream();
+ if (!LaunchStream || !RuntimeState::isStreamRegistered(LaunchStream) ||
+ LaunchStream->Device != Device)
+ return &InvalidConfigurationError;
+
+ if (LaunchStream->Kind == llvm::offload::QueueKind::LegacyDefault) {
+ ol_result_t Result = waitOnBlockingStreams(LaunchStream, Device);
+ if (Result != OL_SUCCESS)
+ return Result;
+ } else if (LaunchStream->Kind == llvm::offload::QueueKind::ExplicitBlocking) {
+ ol_result_t Result = waitOnLegacyDefaultStream(LaunchStream, Device);
+ if (Result != OL_SUCCESS)
+ return Result;
+ }
struct OffloadKernelArgs {
void **Args;
@@ -73,7 +87,7 @@ ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
if (!OKA->Args[I] || OKA->ArgSizes[I] == 0)
return &InvalidArgumentError;
- return olLaunchKernel(Queue, Device, Kernel, &LaunchSizeArgs,
+ return olLaunchKernel(LaunchStream->Queue, Device, Kernel, &LaunchSizeArgs,
/*Properties=*/nullptr, OKA->NumArgs, OKA->Args,
OKA->ArgSizes);
}
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index f0298302b2889..537dcea3bd142 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -23,6 +23,7 @@
#include "Types.h"
#include "OffloadAPI.h"
+#include "llvm/ADT/SmallVector.h"
#include <cassert>
#include <cstdio>
@@ -45,7 +46,22 @@ Error_t Free(void *DevPtr) {
}
Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
- ol_queue_handle_t Queue = ThreadState::getDefaultQueue();
+ ol_device_handle_t Device = nullptr;
+ StreamTy *DefaultStream = nullptr;
+ ol_queue_handle_t Queue = nullptr;
+
+ if (Kind != MemcpyHostToHost) {
+ Device = ThreadState::getDefaultDevice();
+ DefaultStream = ThreadState::getDefaultStream();
+ if (!Device || !DefaultStream)
+ return setLastError(ErrorInvalidDevice);
+
+ ol_result_t Result = waitOnBlockingStreams(DefaultStream, Device);
+ if (Result != OL_SUCCESS)
+ return convertAndSetLastError(Result);
+
+ Queue = DefaultStream->Queue;
+ }
ol_result_t Result;
switch (Kind) {
@@ -55,21 +71,17 @@ Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
break;
}
case MemcpyHostToDevice: {
- ol_device_handle_t Device = ThreadState::getDefaultDevice();
ol_device_handle_t Host = RuntimeState::getHostDevice();
Result = olMemcpy(Queue, Dst, Device, const_cast<void *>(Src), Host, Size);
break;
}
case MemcpyDeviceToHost: {
- ol_device_handle_t Device = ThreadState::getDefaultDevice();
ol_device_handle_t Host = RuntimeState::getHostDevice();
Result = olMemcpy(Queue, Dst, Host, const_cast<void *>(Src), Device, Size);
break;
}
case MemcpyDeviceToDevice: {
- ol_device_handle_t Device = ThreadState::getDefaultDevice();
-
Result =
olMemcpy(Queue, Dst, Device, const_cast<void *>(Src), Device, Size);
break;
@@ -81,6 +93,9 @@ Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
if (Result != OL_SUCCESS)
return convertAndSetLastError(Result);
+ if (!Queue)
+ return convertAndSetLastError(Result);
+
Result = olSyncQueue(Queue);
return convertAndSetLastError(Result);
}
diff --git a/offload/test/offloading/CUDA/blocking_stream_semantics.cu b/offload/test/offloading/CUDA/blocking_stream_semantics.cu
new file mode 100644
index 0000000000000..83aa3383571f6
--- /dev/null
+++ b/offload/test/offloading/CUDA/blocking_stream_semantics.cu
@@ -0,0 +1,131 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fgpu-default-stream=per-thread
+// RUN: %t | %fcheck-generic --check-prefix=PERTHREAD
+// 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 <stdio.h>
+
+__global__ void delayedSetValue(int *Out, int Value) {
+ volatile unsigned long long Delay = 0;
+ for (unsigned I = 0; I < 1000000; ++I)
+ Delay += I;
+ if (Delay)
+ *Out = Value;
+}
+
+__global__ void copyValue(int *In, int *Out) { *Out = *In; }
+
+__global__ void waitThenSetValue(int *Gate, int *Out, int Value) {
+ volatile int *VolatileGate = Gate;
+ for (unsigned I = 0; I < 100000000 && *VolatileGate == 0; ++I)
+ ;
+ *Out = Value;
+}
+
+__global__ void copyValueAndRelease(int *In, int *Out, int *Gate) {
+ *Out = *In;
+ volatile int *VolatileGate = Gate;
+ *VolatileGate = 1;
+}
+
+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 *In = nullptr;
+ int *Out = nullptr;
+ int *Gate = nullptr;
+ if (cudaMalloc(&In, sizeof(int)) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&Out, sizeof(int)) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&Gate, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ int Initial = 0;
+ int Result = 0;
+ if (cudaMemcpy(In, &Initial, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(Out, &Initial, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+
+ delayedSetValue<<<1, 1, 0, BlockingStream>>>(In, 99);
+ copyValue<<<1, 1>>>(In, Out);
+ if (cudaMemcpy(&Result, Out, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("legacy default waited on blocking stream: %d\n", Result);
+ // LEGACY: legacy default waited on blocking stream: 99
+ // PERTHREAD: legacy default waited on blocking stream: 0
+
+ Result = 0;
+ if (cudaMemcpy(Out, &Initial, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+
+ delayedSetValue<<<1, 1>>>(In, 123);
+ copyValue<<<1, 1, 0, BlockingStream>>>(In, Out);
+ if (cudaStreamSynchronize(BlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaMemcpy(&Result, Out, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("blocking stream waited on legacy default: %d\n", Result);
+ // LEGACY: blocking stream waited on legacy default: 123
+ // PERTHREAD: blocking stream waited on legacy default: 99
+
+ Result = 0;
+ if (cudaMemcpy(In, &Initial, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(Out, &Initial, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(Gate, &Initial, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+
+ waitThenSetValue<<<1, 1>>>(Gate, In, 321);
+ copyValueAndRelease<<<1, 1, 0, NonBlockingStream>>>(In, Out, Gate);
+ if (cudaStreamSynchronize(NonBlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaMemcpy(&Result, Out, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("nonblocking stream did not wait on legacy default: %d\n", Result);
+ // LEGACY: nonblocking stream did not wait on legacy default: 0
+ // PERTHREAD: nonblocking stream did not wait on legacy default: 0
+
+ if (cudaStreamDestroy(BlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaStreamDestroy(NonBlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaFree(In) != cudaSuccess)
+ return 1;
+ if (cudaFree(Out) != cudaSuccess)
+ return 1;
+ if (cudaFree(Gate) != cudaSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/CUDA/stream_api.cu b/offload/test/offloading/CUDA/stream_api.cu
index c5b1caaa5315c..0e1c328cc5f1b 100644
--- a/offload/test/offloading/CUDA/stream_api.cu
+++ b/offload/test/offloading/CUDA/stream_api.cu
@@ -70,8 +70,6 @@ int main(int argc, char **argv) {
if (cudaStreamSynchronize(Stream) != cudaSuccess)
return 1;
- if (cudaDeviceSynchronize() != cudaSuccess)
- return 1;
if (cudaMemcpy(&StreamResult, StreamPtr, sizeof(int),
cudaMemcpyDeviceToHost) != cudaSuccess)
return 1;
@@ -86,6 +84,10 @@ int main(int argc, char **argv) {
if (cudaStreamDestroy(Stream) != cudaSuccess)
return 1;
+ if (cudaStreamDestroy(BlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaStreamDestroy(NonBlockingStream) != cudaSuccess)
+ return 1;
print_error("destroyed stream destroy", cudaStreamDestroy(Stream));
// CHECK: destroyed stream destroy value: 4
// CHECK: destroyed stream destroy name: cudaErrorInvalidResourceHandle
diff --git a/offload/test/offloading/HIP/blocking_stream_semantics.hip b/offload/test/offloading/HIP/blocking_stream_semantics.hip
new file mode 100644
index 0000000000000..8c28c7b1b2ee8
--- /dev/null
+++ b/offload/test/offloading/HIP/blocking_stream_semantics.hip
@@ -0,0 +1,125 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fgpu-default-stream=per-thread
+// RUN: %t | %fcheck-generic --check-prefix=PERTHREAD
+// 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 <stdio.h>
+
+__global__ void delayedSetValue(int *Out, int Value) {
+ volatile unsigned long long Delay = 0;
+ for (unsigned I = 0; I < 1000000; ++I)
+ Delay += I;
+ if (Delay)
+ *Out = Value;
+}
+
+__global__ void copyValue(int *In, int *Out) { *Out = *In; }
+
+__global__ void waitThenSetValue(int *Gate, int *Out, int Value) {
+ volatile int *VolatileGate = Gate;
+ for (unsigned I = 0; I < 100000000 && *VolatileGate == 0; ++I)
+ ;
+ *Out = Value;
+}
+
+__global__ void copyValueAndRelease(int *In, int *Out, int *Gate) {
+ *Out = *In;
+ volatile int *VolatileGate = Gate;
+ *VolatileGate = 1;
+}
+
+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 *In = nullptr;
+ int *Out = nullptr;
+ int *Gate = nullptr;
+ if (hipMalloc(&In, sizeof(int)) != hipSuccess)
+ return 1;
+ if (hipMalloc(&Out, sizeof(int)) != hipSuccess)
+ return 1;
+ if (hipMalloc(&Gate, sizeof(int)) != hipSuccess)
+ return 1;
+
+ int Initial = 0;
+ int Result = 0;
+ if (hipMemcpy(In, &Initial, sizeof(int), hipMemcpyHostToDevice) != hipSuccess)
+ return 1;
+ if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
+ hipSuccess)
+ return 1;
+
+ delayedSetValue<<<1, 1, 0, BlockingStream>>>(In, 99);
+ copyValue<<<1, 1>>>(In, Out);
+ if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
+ return 1;
+
+ printf("legacy default waited on blocking stream: %d\n", Result);
+ // LEGACY: legacy default waited on blocking stream: 99
+ // PERTHREAD: legacy default waited on blocking stream: 0
+
+ Result = 0;
+ if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
+ hipSuccess)
+ return 1;
+
+ delayedSetValue<<<1, 1>>>(In, 123);
+ copyValue<<<1, 1, 0, BlockingStream>>>(In, Out);
+ if (hipStreamSynchronize(BlockingStream) != hipSuccess)
+ return 1;
+ if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
+ return 1;
+
+ printf("blocking stream waited on legacy default: %d\n", Result);
+ // LEGACY: blocking stream waited on legacy default: 123
+ // PERTHREAD: blocking stream waited on legacy default: 99
+
+ Result = 0;
+ if (hipMemcpy(In, &Initial, sizeof(int), hipMemcpyHostToDevice) != hipSuccess)
+ return 1;
+ if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
+ hipSuccess)
+ return 1;
+ if (hipMemcpy(Gate, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
+ hipSuccess)
+ return 1;
+
+ waitThenSetValue<<<1, 1>>>(Gate, In, 321);
+ copyValueAndRelease<<<1, 1, 0, NonBlockingStream>>>(In, Out, Gate);
+ if (hipStreamSynchronize(NonBlockingStream) != hipSuccess)
+ return 1;
+ if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
+ return 1;
+
+ printf("nonblocking stream did not wait on legacy default: %d\n", Result);
+ // LEGACY: nonblocking stream did not wait on legacy default: 0
+ // PERTHREAD: nonblocking stream did not wait on legacy default: 0
+
+ if (hipStreamDestroy(BlockingStream) != hipSuccess)
+ return 1;
+ if (hipStreamDestroy(NonBlockingStream) != hipSuccess)
+ return 1;
+ if (hipFree(In) != hipSuccess)
+ return 1;
+ if (hipFree(Out) != hipSuccess)
+ return 1;
+ if (hipFree(Gate) != hipSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/HIP/stream_api.hip b/offload/test/offloading/HIP/stream_api.hip
index fbfca230ee697..3460a6a0c351c 100644
--- a/offload/test/offloading/HIP/stream_api.hip
+++ b/offload/test/offloading/HIP/stream_api.hip
@@ -69,8 +69,6 @@ int main(int argc, char **argv) {
if (hipStreamSynchronize(Stream) != hipSuccess)
return 1;
- if (hipDeviceSynchronize() != hipSuccess)
- return 1;
if (hipMemcpy(&StreamResult, StreamPtr, sizeof(int), hipMemcpyDeviceToHost) !=
hipSuccess)
return 1;
@@ -85,6 +83,10 @@ int main(int argc, char **argv) {
if (hipStreamDestroy(Stream) != hipSuccess)
return 1;
+ if (hipStreamDestroy(BlockingStream) != hipSuccess)
+ return 1;
+ if (hipStreamDestroy(NonBlockingStream) != hipSuccess)
+ return 1;
print_error("destroyed stream destroy", hipStreamDestroy(Stream));
// CHECK: destroyed stream destroy value: 4
// CHECK: destroyed stream destroy name: hipErrorInvalidResourceHandle
More information about the llvm-branch-commits
mailing list