[llvm-branch-commits] [llvm] [Offload][Lang] Add Memset to LLVMOffloading library (PR #216825)
Sophia Herrmann via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Aug 18 10:39:34 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/216825
>From e6bd01fbcabac719e65734b9a40f8fdca1e2cf65 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Mon, 17 Aug 2026 13:21:27 -0700
Subject: [PATCH] add Memset
---
.../include/kernel/DefineLanguageNames.inc | 1 +
.../include/kernel/LanguageRuntime.h | 2 +
.../include/kernel/UndefineLanguageNames.inc | 1 +
.../languages/kernel/src/LanguageRuntime.cpp | 19 ++++
.../CUDA/basic_launch_blocks_and_threads.cu | 3 +-
offload/test/offloading/CUDA/memset.cu | 87 +++++++++++++++++++
.../HIP/basic_launch_blocks_and_threads.hip | 3 +-
offload/test/offloading/HIP/memset.hip | 85 ++++++++++++++++++
8 files changed, 197 insertions(+), 4 deletions(-)
create mode 100644 offload/test/offloading/CUDA/memset.cu
create mode 100644 offload/test/offloading/HIP/memset.hip
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index f4732002c1b52..8990fed0241ba 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -14,6 +14,7 @@
#define Malloc COMBINE(LANGUAGE, Malloc)
#define Free COMBINE(LANGUAGE, Free)
#define Memcpy COMBINE(LANGUAGE, Memcpy)
+#define Memset COMBINE(LANGUAGE, Memset)
#define DeviceSynchronize COMBINE(LANGUAGE, DeviceSynchronize)
#define Success COMBINE(LANGUAGE, Success)
#define ErrorInvalidValue COMBINE(LANGUAGE, ErrorInvalidValue)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index b25e735004ee9..796363a11abf5 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -92,6 +92,8 @@ static inline Error_t Memcpy(T *Dst, const T *Src, size_t Size,
}
///}
+Error_t Memset(void *DevPtr, int Value, size_t Count);
+
Error_t DeviceSynchronize();
Error_t GetDevice(int *DeviceNo);
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index 0ec241ead0434..34690d588c3df 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -11,6 +11,7 @@
#undef Malloc
#undef Free
#undef Memcpy
+#undef Memset
#undef DeviceSynchronize
#undef Success
#undef ErrorInvalidValue
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 2f6f0573d69e1..34753699021ff 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -7,6 +7,7 @@
//===----------------------------------------------------------------------===//
#include "llvm/ADT/SmallPtrSet.h"
+#include <cstddef>
#ifndef LANGUAGE
#error This file should be included, or used, with a LANGUAGE macro set.
#endif
@@ -100,6 +101,24 @@ Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
return convertAndSetLastError(Result);
}
+Error_t Memset(void *DevPtr, int Value, size_t Count) {
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ StreamTy *DefaultStream = ThreadState::getDefaultStream();
+ if (!Device || !DefaultStream)
+ return setLastError(ErrorInvalidDevice);
+
+ ol_result_t Result = waitOnBlockingStreams(DefaultStream, Device);
+ if (Result != OL_SUCCESS)
+ return convertAndSetLastError(Result);
+
+ ol_queue_handle_t Queue = DefaultStream->Queue;
+ unsigned char Byte = static_cast<unsigned char>(Value);
+ Result = olMemFill(Queue, DevPtr, 1, &Byte, Count);
+ if (Result == OL_SUCCESS)
+ Result = olSyncQueue(Queue);
+ return convertAndSetLastError(Result);
+}
+
Error_t DeviceSynchronize() {
ol_device_handle_t Device = ThreadState::getDefaultDevice();
if (!Device)
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 6e40fb695c7e1..55f341d6bb90d 100644
--- a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
+++ b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
@@ -23,8 +23,7 @@ int main(int argc, char **argv) {
cudaMalloc(&Ptr, sizeof(int));
printf("Ptr %p\n", Ptr);
// CHECK: Ptr [[Ptr:0x.*]]
- int Zero = 0;
- cudaMemcpy(Ptr, &Zero, sizeof(int), cudaMemcpyHostToDevice);
+ cudaMemset(Ptr, 0, sizeof(int));
incrementCounter<<<7, 6>>>(Ptr);
cudaDeviceSynchronize();
cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
diff --git a/offload/test/offloading/CUDA/memset.cu b/offload/test/offloading/CUDA/memset.cu
new file mode 100644
index 0000000000000..44a5aeba8dcdf
--- /dev/null
+++ b/offload/test/offloading/CUDA/memset.cu
@@ -0,0 +1,87 @@
+// 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=LEGACY
+// 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=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 <chrono>
+#include <cstdio>
+#include <thread>
+
+__global__ void waitThenSet(volatile int *Gate, unsigned char *Out,
+ unsigned char Value) {
+ for (unsigned long long I = 0; I < 1000000000ULL && *Gate == 0; ++I)
+ ;
+ if (*Gate)
+ *Out = Value;
+}
+
+int main(int argc, char **argv) {
+ unsigned char *Dev = nullptr;
+ if (cudaMalloc(&Dev, 4) != cudaSuccess)
+ return 1;
+
+ if (cudaMemset(Dev, 0x2a, 4) != cudaSuccess)
+ return 1;
+
+ unsigned char Host[4] = {};
+ if (cudaMemcpy(Host, Dev, sizeof(Host), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+ printf("memset bytes: %u %u %u %u\n", static_cast<unsigned>(Host[0]),
+ static_cast<unsigned>(Host[1]), static_cast<unsigned>(Host[2]),
+ static_cast<unsigned>(Host[3]));
+ // LEGACY: memset bytes: 42 42 42 42
+ // PERTHREAD: memset bytes: 42 42 42 42
+
+ cudaStream_t BlockingStream = nullptr;
+ if (cudaStreamCreateWithFlags(&BlockingStream, cudaStreamDefault) !=
+ cudaSuccess)
+ return 1;
+
+ int *Gate = nullptr;
+ if (cudaHostAlloc(&Gate, sizeof(int), cudaHostAllocDefault) != cudaSuccess)
+ return 1;
+ *Gate = 0;
+
+ if (cudaMemset(Dev, 0, 1) != cudaSuccess)
+ return 1;
+
+ waitThenSet<<<1, 1, 0, BlockingStream>>>(Gate, Dev, 17);
+
+ std::thread Releaser([&]() {
+ std::this_thread::sleep_for(std::chrono::milliseconds(250));
+ *Gate = 1;
+ });
+
+ cudaError_t MemsetResult = cudaMemset(Dev, 23, 1);
+
+ Releaser.join();
+ if (MemsetResult != cudaSuccess)
+ return 1;
+
+ if (cudaStreamSynchronize(BlockingStream) != cudaSuccess)
+ return 1;
+
+ unsigned char Result = 0;
+ if (cudaMemcpy(&Result, Dev, 1, cudaMemcpyDeviceToHost) != cudaSuccess)
+ return 1;
+ printf("default stream memset result: %u\n", static_cast<unsigned>(Result));
+ // LEGACY: default stream memset result: 23
+ // PERTHREAD: default stream memset result: 17
+
+ if (cudaStreamDestroy(BlockingStream) != cudaSuccess)
+ return 1;
+ if (cudaFreeHost(Gate) != cudaSuccess)
+ return 1;
+ if (cudaFree(Dev) != cudaSuccess)
+ return 1;
+}
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 4f5ce89130052..aa92c0e1db3f0 100644
--- a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
+++ b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
@@ -23,8 +23,7 @@ int main(int argc, char **argv) {
hipMalloc(&Ptr, sizeof(int));
printf("Ptr %p\n", Ptr);
// CHECK: Ptr [[Ptr:0x.*]]
- int Zero = 0;
- hipMemcpy(Ptr, &Zero, sizeof(int), hipMemcpyHostToDevice);
+ hipMemset(Ptr, 0, sizeof(int));
incrementCounter<<<7, 6>>>(Ptr);
hipDeviceSynchronize();
hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
diff --git a/offload/test/offloading/HIP/memset.hip b/offload/test/offloading/HIP/memset.hip
new file mode 100644
index 0000000000000..13bee34670dad
--- /dev/null
+++ b/offload/test/offloading/HIP/memset.hip
@@ -0,0 +1,85 @@
+// 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=LEGACY
+// 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=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 <chrono>
+#include <cstdio>
+#include <thread>
+
+__global__ void waitThenSet(volatile int *Gate, unsigned char *Out,
+ unsigned char Value) {
+ for (unsigned long long I = 0; I < 1000000000ULL && *Gate == 0; ++I)
+ ;
+ if (*Gate)
+ *Out = Value;
+}
+
+int main(int argc, char **argv) {
+ unsigned char *Dev = nullptr;
+ if (hipMalloc(&Dev, 4) != hipSuccess)
+ return 1;
+
+ if (hipMemset(Dev, 0x2a, 4) != hipSuccess)
+ return 1;
+
+ unsigned char Host[4] = {};
+ if (hipMemcpy(Host, Dev, sizeof(Host), hipMemcpyDeviceToHost) != hipSuccess)
+ return 1;
+ printf("memset bytes: %u %u %u %u\n", static_cast<unsigned>(Host[0]),
+ static_cast<unsigned>(Host[1]), static_cast<unsigned>(Host[2]),
+ static_cast<unsigned>(Host[3]));
+ // LEGACY: memset bytes: 42 42 42 42
+ // PERTHREAD: memset bytes: 42 42 42 42
+
+ hipStream_t BlockingStream = nullptr;
+ if (hipStreamCreateWithFlags(&BlockingStream, hipStreamDefault) != hipSuccess)
+ return 1;
+
+ int *Gate = nullptr;
+ if (hipHostAlloc(&Gate, sizeof(int), hipHostAllocDefault) != hipSuccess)
+ return 1;
+ *Gate = 0;
+
+ if (hipMemset(Dev, 0, 1) != hipSuccess)
+ return 1;
+
+ waitThenSet<<<1, 1, 0, BlockingStream>>>(Gate, Dev, 17);
+
+ std::thread Releaser([&]() {
+ std::this_thread::sleep_for(std::chrono::milliseconds(250));
+ *Gate = 1;
+ });
+
+ hipError_t MemsetResult = hipMemset(Dev, 23, 1);
+
+ Releaser.join();
+ if (MemsetResult != hipSuccess)
+ return 1;
+
+ if (hipStreamSynchronize(BlockingStream) != hipSuccess)
+ return 1;
+
+ unsigned char Result = 0;
+ if (hipMemcpy(&Result, Dev, 1, hipMemcpyDeviceToHost) != hipSuccess)
+ return 1;
+ printf("default stream memset result: %u\n", static_cast<unsigned>(Result));
+ // LEGACY: default stream memset result: 23
+ // PERTHREAD: default stream memset result: 17
+
+ if (hipStreamDestroy(BlockingStream) != hipSuccess)
+ return 1;
+ if (hipFreeHost(Gate) != hipSuccess)
+ return 1;
+ if (hipFree(Dev) != hipSuccess)
+ return 1;
+}
More information about the llvm-branch-commits
mailing list