https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/218460
>From ec0766f590d75b18c8fd62a3e3581d1cb303011b Mon Sep 17 00:00:00 2001 From: Sophia Herrmann <[email protected]> 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 | 15 ++++ .../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, 193 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 796ddc9d7cfbb..2008a5039a06c 100644 --- a/offload/languages/kernel/src/LanguageRuntime.cpp +++ b/offload/languages/kernel/src/LanguageRuntime.cpp @@ -26,6 +26,7 @@ #include "OffloadAPI.h" #include <cassert> +#include <cstddef> #include <cstdio> #include <cstdlib> #include <cstring> @@ -86,6 +87,20 @@ 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) { + ThreadStateTy ThreadState = ThreadStateTy::get(); + ol_result_t Result = waitOnBlockingStreams(StateTy::get(), ThreadState); + if (Result != OL_SUCCESS) + return convertAndSetLastError(Result); + + ol_queue_handle_t Queue = ThreadState.getDefaultQueue(); + 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 = ThreadStateTy::get().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; +} _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
