[llvm-branch-commits] [llvm] [Offload][Lang] Unify Tests (PR #217722)
Sophia Herrmann via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Aug 20 11:28:27 PDT 2026
https://github.com/jellytabby created https://github.com/llvm/llvm-project/pull/217722
We need tests to ensure the parsing and computing capability of both cuda and hip flavored inputs. Previously these were separate but same tests, with this PR we follow the same structure as the LanguageRuntime.h files and use macro renaming to generate cuda/hip tests from one common base file. The new tests are in `offload/test/offloading/language/`.
Assisted by GPT-5.5, checked and reviewed manually
>From b5f07499ac006d352908e76f59848ff87e926bca Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Thu, 20 Aug 2026 11:21:21 -0700
Subject: [PATCH] unify tests
---
offload/test/offloading/CUDA/basic_launch.cu | 30 ----
.../CUDA/basic_launch_blocks_and_threads.cu | 32 -----
.../offloading/CUDA/basic_launch_multi_arg.cu | 39 ------
.../CUDA/blocking_stream_semantics.cu | 131 ------------------
offload/test/offloading/CUDA/device_api.cu | 45 ------
.../test/offloading/CUDA/device_properties.cu | 40 ------
.../offloading/CUDA/devicesync_streams.cu | 98 -------------
offload/test/offloading/CUDA/error_kinds.cu | 81 -----------
offload/test/offloading/CUDA/get_errs.cu | 82 -----------
offload/test/offloading/CUDA/host_alloc.cu | 48 -------
offload/test/offloading/CUDA/launch_tu.cu | 30 ----
offload/test/offloading/CUDA/memcpy_kinds.cu | 51 -------
offload/test/offloading/CUDA/stream_api.cu | 100 -------------
offload/test/offloading/CUDA/syncthreads.cu | 41 ------
offload/test/offloading/HIP/basic_launch.hip | 30 ----
.../offloading/HIP/basic_launch_multi_arg.hip | 39 ------
offload/test/offloading/HIP/error_kinds.hip | 81 -----------
offload/test/offloading/HIP/get_errs.hip | 81 -----------
offload/test/offloading/HIP/host_alloc.hip | 48 -------
offload/test/offloading/HIP/kernel_tu.hip.inc | 1 -
offload/test/offloading/HIP/memcpy_kinds.hip | 51 -------
offload/test/offloading/HIP/memset.hip | 85 ------------
offload/test/offloading/HIP/stream_api.hip | 99 -------------
.../offloading/HIP/thread_and_block_id.hip | 46 ------
.../Inputs/DefineTestLanguageNames.inc | 18 +++
.../Inputs/launch_tu_kernel.inc} | 0
.../basic_launch.cpp} | 23 +--
.../basic_launch_blocks_and_threads.cpp} | 21 +--
.../language/basic_launch_multi_arg.cpp | 44 ++++++
.../blocking_stream_semantics.cpp} | 70 +++++-----
.../device_api.cpp} | 23 +--
.../device_properties.cpp} | 21 +--
.../devicesync_streams.cpp} | 57 ++++----
.../test/offloading/language/error_kinds.cpp | 95 +++++++++++++
offload/test/offloading/language/get_errs.cpp | 100 +++++++++++++
.../test/offloading/language/host_alloc.cpp | 52 +++++++
.../test/offloading/language/launch_tu.cpp | 35 +++++
.../test/offloading/language/memcpy_kinds.cpp | 52 +++++++
.../{CUDA/memset.cu => language/memset.cpp} | 43 +++---
.../test/offloading/language/stream_api.cpp | 108 +++++++++++++++
.../syncthreads.cpp} | 19 ++-
.../thread_and_block_id.cpp} | 21 +--
42 files changed, 671 insertions(+), 1540 deletions(-)
delete mode 100644 offload/test/offloading/CUDA/basic_launch.cu
delete mode 100644 offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
delete mode 100644 offload/test/offloading/CUDA/basic_launch_multi_arg.cu
delete mode 100644 offload/test/offloading/CUDA/blocking_stream_semantics.cu
delete mode 100644 offload/test/offloading/CUDA/device_api.cu
delete mode 100644 offload/test/offloading/CUDA/device_properties.cu
delete mode 100644 offload/test/offloading/CUDA/devicesync_streams.cu
delete mode 100644 offload/test/offloading/CUDA/error_kinds.cu
delete mode 100644 offload/test/offloading/CUDA/get_errs.cu
delete mode 100644 offload/test/offloading/CUDA/host_alloc.cu
delete mode 100644 offload/test/offloading/CUDA/launch_tu.cu
delete mode 100644 offload/test/offloading/CUDA/memcpy_kinds.cu
delete mode 100644 offload/test/offloading/CUDA/stream_api.cu
delete mode 100644 offload/test/offloading/CUDA/syncthreads.cu
delete mode 100644 offload/test/offloading/HIP/basic_launch.hip
delete mode 100644 offload/test/offloading/HIP/basic_launch_multi_arg.hip
delete mode 100644 offload/test/offloading/HIP/error_kinds.hip
delete mode 100644 offload/test/offloading/HIP/get_errs.hip
delete mode 100644 offload/test/offloading/HIP/host_alloc.hip
delete mode 100644 offload/test/offloading/HIP/kernel_tu.hip.inc
delete mode 100644 offload/test/offloading/HIP/memcpy_kinds.hip
delete mode 100644 offload/test/offloading/HIP/memset.hip
delete mode 100644 offload/test/offloading/HIP/stream_api.hip
delete mode 100644 offload/test/offloading/HIP/thread_and_block_id.hip
create mode 100644 offload/test/offloading/language/Inputs/DefineTestLanguageNames.inc
rename offload/test/offloading/{CUDA/kernel_tu.cu.inc => language/Inputs/launch_tu_kernel.inc} (100%)
rename offload/test/offloading/{HIP/launch_tu.hip => language/basic_launch.cpp} (50%)
rename offload/test/offloading/{HIP/basic_launch_blocks_and_threads.hip => language/basic_launch_blocks_and_threads.cpp} (51%)
create mode 100644 offload/test/offloading/language/basic_launch_multi_arg.cpp
rename offload/test/offloading/{HIP/blocking_stream_semantics.hip => language/blocking_stream_semantics.cpp} (50%)
rename offload/test/offloading/{HIP/device_api.hip => language/device_api.cpp} (54%)
rename offload/test/offloading/{HIP/device_properties.hip => language/device_properties.cpp} (58%)
rename offload/test/offloading/{HIP/devicesync_streams.hip => language/devicesync_streams.cpp} (51%)
create mode 100644 offload/test/offloading/language/error_kinds.cpp
create mode 100644 offload/test/offloading/language/get_errs.cpp
create mode 100644 offload/test/offloading/language/host_alloc.cpp
create mode 100644 offload/test/offloading/language/launch_tu.cpp
create mode 100644 offload/test/offloading/language/memcpy_kinds.cpp
rename offload/test/offloading/{CUDA/memset.cu => language/memset.cpp} (52%)
create mode 100644 offload/test/offloading/language/stream_api.cpp
rename offload/test/offloading/{HIP/syncthreads.hip => language/syncthreads.cpp} (55%)
rename offload/test/offloading/{CUDA/thread_and_block_id.cu => language/thread_and_block_id.cpp} (61%)
diff --git a/offload/test/offloading/CUDA/basic_launch.cu b/offload/test/offloading/CUDA/basic_launch.cu
deleted file mode 100644
index 5ecfc3e9d5601..0000000000000
--- a/offload/test/offloading/CUDA/basic_launch.cu
+++ /dev/null
@@ -1,30 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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 square(int *A) { *A = 42; }
-
-int main(int argc, char **argv) {
- int *Ptr;
- cudaMalloc(&Ptr, 4);
- printf("Ptr %p\n", Ptr);
- // CHECK: Ptr [[Ptr:0x.*]]
- square<<<1, 1>>>(Ptr);
- int I = 0;
- cudaDeviceSynchronize();
- cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
- printf("I: %i\n", I);
- // CHECK: I: 42
-}
diff --git a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
deleted file mode 100644
index 55f341d6bb90d..0000000000000
--- a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
+++ /dev/null
@@ -1,32 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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 incrementCounter(int *A) {
- __scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
-}
-
-int main(int argc, char **argv) {
- int *Ptr, I;
- cudaMalloc(&Ptr, sizeof(int));
- printf("Ptr %p\n", Ptr);
- // CHECK: Ptr [[Ptr:0x.*]]
- cudaMemset(Ptr, 0, sizeof(int));
- 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
deleted file mode 100644
index 25207536496e7..0000000000000
--- a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
+++ /dev/null
@@ -1,39 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// clang-format on
-
-// REQUIRES: gpu
-// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
-// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
-// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
-// UNSUPPORTED: intelgpu
-
-#include <stdio.h>
-
-__global__ void square(int *Dst, short Q, int *Src, short P) {
- *Dst = (Src[0] + Src[1]) * (Q + P);
- Src[0] = Q;
- Src[1] = P;
-}
-
-int main(int argc, char **argv) {
- int *Src, *Ptr;
- cudaMalloc(&Ptr, 4);
- cudaMalloc(&Src, 8);
-
- int I = 7;
- int HostSrc[2] = {-2, 8};
- 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);
- // CHECK: I: 42
- printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
- // CHECK: Src: 3, 4
-}
diff --git a/offload/test/offloading/CUDA/blocking_stream_semantics.cu b/offload/test/offloading/CUDA/blocking_stream_semantics.cu
deleted file mode 100644
index 83aa3383571f6..0000000000000
--- a/offload/test/offloading/CUDA/blocking_stream_semantics.cu
+++ /dev/null
@@ -1,131 +0,0 @@
-// 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/device_api.cu b/offload/test/offloading/CUDA/device_api.cu
deleted file mode 100644
index af2b046eee397..0000000000000
--- a/offload/test/offloading/CUDA/device_api.cu
+++ /dev/null
@@ -1,45 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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>
-
-int main(int argc, char **argv) {
- int Count = 0;
- if (cudaGetDeviceCount(&Count) != cudaSuccess)
- return 1;
-
- printf("device count: %d\n", Count);
- // CHECK: device count: {{[1-9][0-9]*}}
-
- int Device = -1;
- if (cudaGetDevice(&Device) != cudaSuccess)
- return 1;
-
- printf("device: %d\n", Device);
- // CHECK: device: {{[0-9]+}}
-
- if (cudaSetDevice(Device) != cudaSuccess)
- return 1;
-
- int After = -1;
- if (cudaGetDevice(&After) != cudaSuccess)
- return 1;
-
- printf("device after set: %d\n", After);
- // CHECK: device after set: {{[0-9]+}}
-
- cudaError_t Err = cudaSetDevice(-1);
- printf("set invalid device: %u\n", Err);
- // CHECK: set invalid device: 2
-}
diff --git a/offload/test/offloading/CUDA/device_properties.cu b/offload/test/offloading/CUDA/device_properties.cu
deleted file mode 100644
index 8f625f6ccabe1..0000000000000
--- a/offload/test/offloading/CUDA/device_properties.cu
+++ /dev/null
@@ -1,40 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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>
-
-int main(int argc, char **argv) {
- cudaDeviceProp Prop = {};
- cudaError_t Err = cudaGetDeviceProperties(&Prop, 0);
- if (Err != cudaSuccess) {
- printf("cudaGetDeviceProperties failed: %u\n", Err);
- return 1;
- }
-
- printf("Device name: %s\n", Prop.name);
- // CHECK: Device name:
- printf("Total global memory: %zu\n", Prop.totalGlobalMem);
- // CHECK: Total global memory:
- printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
- // CHECK: Multiprocessors:
- printf("Warp size: %i\n", Prop.warpSize);
- // CHECK: Warp size:
-
- if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
- !Prop.warpSize)
- return 1;
-
- printf("Device properties are populated.\n");
- // CHECK: Device properties are populated.
-}
diff --git a/offload/test/offloading/CUDA/devicesync_streams.cu b/offload/test/offloading/CUDA/devicesync_streams.cu
deleted file mode 100644
index a8a547101fbaf..0000000000000
--- a/offload/test/offloading/CUDA/devicesync_streams.cu
+++ /dev/null
@@ -1,98 +0,0 @@
-// 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/error_kinds.cu b/offload/test/offloading/CUDA/error_kinds.cu
deleted file mode 100644
index c5c2a1d83680d..0000000000000
--- a/offload/test/offloading/CUDA/error_kinds.cu
+++ /dev/null
@@ -1,81 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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>
-
-static void print_error(const char *Label, cudaError_t Error) {
- printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
- printf("%s name: %s\n", Label, cudaGetErrorName(Error));
- printf("%s string: %s\n", Label, cudaGetErrorString(Error));
-}
-
-int main() {
- print_error("success", cudaSuccess);
- // CHECK: success value: 0
- // CHECK: success name: cudaSuccess
- // CHECK: success string: No error
-
- print_error("invalid value", cudaErrorInvalidValue);
- // CHECK: invalid value value: 1
- // CHECK: invalid value name: cudaErrorInvalidValue
- // CHECK: invalid value string: Invalid argument value
-
- print_error("invalid device", cudaErrorInvalidDevice);
- // CHECK: invalid device value: 2
- // CHECK: invalid device name: cudaErrorInvalidDevice
- // CHECK: invalid device string: Invalid device number
-
- print_error("unknown", cudaErrorUnknown);
- // CHECK: unknown value: 3
- // CHECK: unknown name: Unrecognized error
- // CHECK: unknown string: Unknown error
-
- print_error("invalid resource handle", cudaErrorInvalidResourceHandle);
- // CHECK: invalid resource handle value: 4
- // CHECK: invalid resource handle name: cudaErrorInvalidResourceHandle
- // CHECK: invalid resource handle string: Invalid resource handle
-
- print_error("invalid configuration", cudaErrorInvalidConfiguration);
- // CHECK: invalid configuration value: 5
- // CHECK: invalid configuration name: cudaErrorInvalidConfiguration
- // CHECK: invalid configuration string: Invalid configuration argument
-
- cudaError_t Unrecognized = static_cast<cudaError_t>(999);
- print_error("unrecognized", Unrecognized);
- // CHECK: unrecognized value: 999
- // CHECK: unrecognized name: Unrecognized error
- // CHECK: unrecognized string: Unrecognized error
-
- print_error("set invalid device", cudaSetDevice(-1));
- // CHECK: set invalid device value: 2
- // CHECK: set invalid device name: cudaErrorInvalidDevice
- // CHECK: set invalid device string: Invalid device number
-
- print_error("get last error", cudaGetLastError());
- // CHECK: get last error value: 2
- // CHECK: get last error name: cudaErrorInvalidDevice
- // CHECK: get last error string: Invalid device number
-
- print_error("cleared last error", cudaGetLastError());
- // CHECK: cleared last error value: 0
- // CHECK: cleared last error name: cudaSuccess
- // CHECK: cleared last error string: No error
-
- print_error("null stream destroy", cudaStreamDestroy(nullptr));
- // CHECK: null stream destroy value: 1
- // CHECK: null stream destroy name: cudaErrorInvalidValue
- // CHECK: null stream destroy string: Invalid argument value
-
- return 0;
-}
diff --git a/offload/test/offloading/CUDA/get_errs.cu b/offload/test/offloading/CUDA/get_errs.cu
deleted file mode 100644
index f29c4665de3f3..0000000000000
--- a/offload/test/offloading/CUDA/get_errs.cu
+++ /dev/null
@@ -1,82 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -pthread -std=c++17
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp -pthread -std=c++17
-// RUN: %t | %fcheck-generic
-// 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 <cstdio>
-#include <cuda_runtime.h>
-#include <mutex>
-#include <thread>
-
-static std::mutex PrintMutex;
-
-static void printError(int ThreadId, const char *Label, cudaError_t Error) {
- std::lock_guard<std::mutex> Lock(PrintMutex);
- printf("thread %d %s: %s\n", ThreadId, Label, cudaGetErrorName(Error));
- std::fflush(stdout);
-}
-
-__global__ void errorKernel(float *d_out) {
- int idx = blockIdx.x * blockDim.x + threadIdx.x;
- d_out[idx] = idx * 0.5f;
-}
-
-void runTask(int ThreadId) {
- const int N = 1 << 20;
- size_t bytes = N * sizeof(float);
-
- float *d_data;
- cudaMalloc(&d_data, bytes);
- printError(ThreadId, "cudaMalloc", cudaGetLastError());
-
- ThreadId == 1 ? errorKernel<<<4096, 256>>>(d_data)
- : errorKernel<<<4096, 0>>>(d_data);
- printError(ThreadId, "kernel launch", cudaPeekAtLastError());
-
- printError(ThreadId, "kernel launch get", cudaGetLastError());
-
- printError(ThreadId, "kernel launch get again", cudaGetLastError());
-
- cudaDeviceSynchronize();
- cudaFree(d_data);
-}
-
-int main() {
- printError(0, "initial", cudaPeekAtLastError());
- // CHECK: thread 0 initial: cudaSuccess
-
- std::thread t1(runTask, 1);
- std::thread t2(runTask, 2);
-
- t1.join();
- t2.join();
- // CHECK-DAG: thread 1 cudaMalloc: cudaSuccess
- // CHECK-DAG: thread 2 cudaMalloc: cudaSuccess
- // CHECK-DAG: thread 1 kernel launch: cudaSuccess
- // CHECK-DAG: thread 2 kernel launch: cudaErrorInvalidConfiguration
- // CHECK-DAG: thread 1 kernel launch get: cudaSuccess
- // CHECK-DAG: thread 2 kernel launch get: cudaErrorInvalidConfiguration
- // CHECK-DAG: thread 1 kernel launch get again: cudaSuccess
- // CHECK-DAG: thread 2 kernel launch get again: cudaSuccess
-
- std::thread t3(runTask, 3);
- t3.join();
- // CHECK: thread 3 cudaMalloc: cudaSuccess
- // CHECK: thread 3 kernel launch: cudaErrorInvalidConfiguration
- // CHECK: thread 3 kernel launch get: cudaErrorInvalidConfiguration
- // CHECK: thread 3 kernel launch get again: cudaSuccess
-
- printError(0, "joined", cudaGetLastError());
- // CHECK: thread 0 joined: cudaSuccess
-
- return 0;
-}
diff --git a/offload/test/offloading/CUDA/host_alloc.cu b/offload/test/offloading/CUDA/host_alloc.cu
deleted file mode 100644
index 1440b276c6602..0000000000000
--- a/offload/test/offloading/CUDA/host_alloc.cu
+++ /dev/null
@@ -1,48 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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 add(int *Ptr, int Value) { *Ptr += Value; }
-
-int main(int argc, char **argv) {
- int *HostAllocPtr = nullptr;
- if (cudaHostAlloc(&HostAllocPtr, sizeof(int), cudaHostAllocDefault) !=
- cudaSuccess)
- return 1;
-
- *HostAllocPtr = 17;
- add<<<1, 1>>>(HostAllocPtr, 5);
- if (cudaDeviceSynchronize() != cudaSuccess)
- return 1;
- printf("cudaHostAlloc value: %d\n", *HostAllocPtr);
- // CHECK: cudaHostAlloc value: 22
-
- if (cudaFreeHost(HostAllocPtr) != cudaSuccess)
- return 1;
-
- int *MallocHostPtr = nullptr;
- if (cudaMallocHost(&MallocHostPtr, sizeof(int)) != cudaSuccess)
- return 1;
-
- *MallocHostPtr = 23;
- add<<<1, 1>>>(MallocHostPtr, 7);
- if (cudaDeviceSynchronize() != cudaSuccess)
- return 1;
- printf("cudaMallocHost value: %d\n", *MallocHostPtr);
- // CHECK: cudaMallocHost value: 30
-
- if (cudaFreeHost(MallocHostPtr) != cudaSuccess)
- return 1;
-}
diff --git a/offload/test/offloading/CUDA/launch_tu.cu b/offload/test/offloading/CUDA/launch_tu.cu
deleted file mode 100644
index fc24ec1af03b9..0000000000000
--- a/offload/test/offloading/CUDA/launch_tu.cu
+++ /dev/null
@@ -1,30 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda %S/kernel_tu.cu.inc -o %t.kernel_tu.o -c
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
-// RUN: %t | %fcheck-generic
-// 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>
-
-extern __global__ void square(int *A);
-
-int main(int argc, char **argv) {
- 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/memcpy_kinds.cu b/offload/test/offloading/CUDA/memcpy_kinds.cu
deleted file mode 100644
index a4288ee51ee3b..0000000000000
--- a/offload/test/offloading/CUDA/memcpy_kinds.cu
+++ /dev/null
@@ -1,51 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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>
-
-int main(int argc, char **argv) {
- int HostSrc = 11;
- int HostDst = 0;
- if (cudaMemcpy(&HostDst, &HostSrc, sizeof(int), cudaMemcpyHostToHost) !=
- cudaSuccess)
- return 1;
-
- printf("host to host: %d\n", HostDst);
- // CHECK: host to host: 11
-
- int *DevSrc = nullptr;
- int *DevDst = nullptr;
- int Result = 0;
- if (cudaMalloc(&DevSrc, sizeof(int)) != cudaSuccess)
- return 1;
- if (cudaMalloc(&DevDst, sizeof(int)) != cudaSuccess)
- return 1;
-
- HostSrc = 42;
- if (cudaMemcpy(DevSrc, &HostSrc, sizeof(int), cudaMemcpyHostToDevice) !=
- cudaSuccess)
- return 1;
- if (cudaMemcpy(DevDst, DevSrc, sizeof(int), cudaMemcpyDeviceToDevice) !=
- cudaSuccess)
- return 1;
- if (cudaMemcpy(&Result, DevDst, sizeof(int), cudaMemcpyDeviceToHost) !=
- cudaSuccess)
- return 1;
-
- printf("device to device: %d\n", Result);
- // CHECK: device to device: 42
-
- cudaFree(DevSrc);
- cudaFree(DevDst);
-}
diff --git a/offload/test/offloading/CUDA/stream_api.cu b/offload/test/offloading/CUDA/stream_api.cu
deleted file mode 100644
index 0e1c328cc5f1b..0000000000000
--- a/offload/test/offloading/CUDA/stream_api.cu
+++ /dev/null
@@ -1,100 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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>
-
-static void print_error(const char *Label, cudaError_t Error) {
- printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
- printf("%s name: %s\n", Label, cudaGetErrorName(Error));
-}
-
-__global__ void setValue(int *Out, int Value) { *Out = Value; }
-
-int main(int argc, char **argv) {
- print_error("null stream create", cudaStreamCreate(nullptr));
- // CHECK: null stream create value: 1
- // CHECK: null stream create name: cudaErrorInvalidValue
- print_error("null flags stream create",
- cudaStreamCreateWithFlags(nullptr, cudaStreamDefault));
- // CHECK: null flags stream create value: 1
- // CHECK: null flags stream create name: cudaErrorInvalidValue
-
- cudaStream_t InvalidFlagsStream = nullptr;
- print_error("invalid stream flags",
- cudaStreamCreateWithFlags(&InvalidFlagsStream, ~0u));
- // CHECK: invalid stream flags value: 1
- // CHECK: invalid stream flags name: cudaErrorInvalidValue
- printf("invalid flags stream: %d\n", InvalidFlagsStream == nullptr);
- // CHECK: invalid flags stream: 1
-
- cudaStream_t Stream = nullptr;
- if (cudaStreamCreate(&Stream) != cudaSuccess)
- return 1;
- cudaStream_t BlockingStream = nullptr;
- if (cudaStreamCreateWithFlags(&BlockingStream, cudaStreamDefault) !=
- cudaSuccess)
- return 1;
- cudaStream_t NonBlockingStream = nullptr;
- if (cudaStreamCreateWithFlags(&NonBlockingStream, cudaStreamNonBlocking) !=
- cudaSuccess)
- return 1;
-
- printf("stream created: %d\n", Stream != nullptr);
- // CHECK: stream created: 1
- printf("stream flags created: %d %d\n", BlockingStream != nullptr,
- NonBlockingStream != nullptr);
- // CHECK: stream flags created: 1 1
-
- int *StreamPtr = nullptr;
- int *DefaultPtr = nullptr;
- int StreamResult = 0;
- int DefaultResult = 0;
- if (cudaMalloc(&StreamPtr, sizeof(int)) != cudaSuccess)
- return 1;
- if (cudaMalloc(&DefaultPtr, sizeof(int)) != cudaSuccess)
- return 1;
-
- setValue<<<1, 1, 0, Stream>>>(StreamPtr, 42);
- setValue<<<1, 1>>>(DefaultPtr, 17);
-
- if (cudaStreamSynchronize(Stream) != cudaSuccess)
- return 1;
- if (cudaMemcpy(&StreamResult, StreamPtr, sizeof(int),
- cudaMemcpyDeviceToHost) != cudaSuccess)
- return 1;
- if (cudaMemcpy(&DefaultResult, DefaultPtr, sizeof(int),
- cudaMemcpyDeviceToHost) != cudaSuccess)
- return 1;
-
- printf("stream result: %d\n", StreamResult);
- // CHECK: stream result: 42
- printf("default result: %d\n", DefaultResult);
- // CHECK: default result: 17
-
- 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
- print_error("destroyed stream synchronize", cudaStreamSynchronize(Stream));
- // CHECK: destroyed stream synchronize value: 4
- // CHECK: destroyed stream synchronize name: cudaErrorInvalidResourceHandle
-
- cudaFree(StreamPtr);
- cudaFree(DefaultPtr);
-}
diff --git a/offload/test/offloading/CUDA/syncthreads.cu b/offload/test/offloading/CUDA/syncthreads.cu
deleted file mode 100644
index 4c839b85ff768..0000000000000
--- a/offload/test/offloading/CUDA/syncthreads.cu
+++ /dev/null
@@ -1,41 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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 reduceBlock(int *Out) {
- __shared__ int Scratch[64];
- int Tid = threadIdx.x;
- Scratch[Tid] = Tid;
- __syncthreads();
-
- if (Tid == 0) {
- int Sum = 0;
- for (int I = 0; I < 64; ++I)
- Sum += Scratch[I];
- Out[0] = Sum;
- }
-}
-
-int main(int argc, char **argv) {
- int *DevPtr;
- int Result = 0;
- cudaMalloc(&DevPtr, sizeof(int));
- reduceBlock<<<1, 64>>>(DevPtr);
- cudaDeviceSynchronize();
- cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost);
-
- printf("sum: %i\n", Result);
- // CHECK: sum: 2016
-}
diff --git a/offload/test/offloading/HIP/basic_launch.hip b/offload/test/offloading/HIP/basic_launch.hip
deleted file mode 100644
index bd2f2a6078671..0000000000000
--- a/offload/test/offloading/HIP/basic_launch.hip
+++ /dev/null
@@ -1,30 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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 square(int *A) { *A = 42; }
-
-int main(int argc, char **argv) {
- int *Ptr;
- hipMalloc(&Ptr, 4);
- printf("Ptr %p\n", Ptr);
- // CHECK: Ptr [[Ptr:0x.*]]
- square<<<1, 1>>>(Ptr);
- int I = 0;
- 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
deleted file mode 100644
index 3bca0a7fae484..0000000000000
--- a/offload/test/offloading/HIP/basic_launch_multi_arg.hip
+++ /dev/null
@@ -1,39 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// clang-format on
-
-// REQUIRES: gpu
-// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
-// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
-// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
-// UNSUPPORTED: intelgpu
-
-#include <stdio.h>
-
-__global__ void square(int *Dst, short Q, int *Src, short P) {
- *Dst = (Src[0] + Src[1]) * (Q + P);
- Src[0] = Q;
- Src[1] = P;
-}
-
-int main(int argc, char **argv) {
- int *Src, *Ptr;
- hipMalloc(&Ptr, 4);
- hipMalloc(&Src, 8);
-
- int I = 7;
- int HostSrc[2] = {-2,8};
- 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);
- // CHECK: I: 42
- printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
- // CHECK: Src: 3, 4
-}
diff --git a/offload/test/offloading/HIP/error_kinds.hip b/offload/test/offloading/HIP/error_kinds.hip
deleted file mode 100644
index 2830d9de9b9dc..0000000000000
--- a/offload/test/offloading/HIP/error_kinds.hip
+++ /dev/null
@@ -1,81 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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>
-
-static void print_error(const char *Label, hipError_t Error) {
- printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
- printf("%s name: %s\n", Label, hipGetErrorName(Error));
- printf("%s string: %s\n", Label, hipGetErrorString(Error));
-}
-
-int main() {
- print_error("success", hipSuccess);
- // CHECK: success value: 0
- // CHECK: success name: hipSuccess
- // CHECK: success string: No error
-
- print_error("invalid value", hipErrorInvalidValue);
- // CHECK: invalid value value: 1
- // CHECK: invalid value name: hipErrorInvalidValue
- // CHECK: invalid value string: Invalid argument value
-
- print_error("invalid device", hipErrorInvalidDevice);
- // CHECK: invalid device value: 2
- // CHECK: invalid device name: hipErrorInvalidDevice
- // CHECK: invalid device string: Invalid device number
-
- print_error("unknown", hipErrorUnknown);
- // CHECK: unknown value: 3
- // CHECK: unknown name: Unrecognized error
- // CHECK: unknown string: Unknown error
-
- print_error("invalid resource handle", hipErrorInvalidResourceHandle);
- // CHECK: invalid resource handle value: 4
- // CHECK: invalid resource handle name: hipErrorInvalidResourceHandle
- // CHECK: invalid resource handle string: Invalid resource handle
-
- print_error("invalid configuration", hipErrorInvalidConfiguration);
- // CHECK: invalid configuration value: 5
- // CHECK: invalid configuration name: hipErrorInvalidConfiguration
- // CHECK: invalid configuration string: Invalid configuration argument
-
- hipError_t Unrecognized = static_cast<hipError_t>(999);
- print_error("unrecognized", Unrecognized);
- // CHECK: unrecognized value: 999
- // CHECK: unrecognized name: Unrecognized error
- // CHECK: unrecognized string: Unrecognized error
-
- print_error("set invalid device", hipSetDevice(-1));
- // CHECK: set invalid device value: 2
- // CHECK: set invalid device name: hipErrorInvalidDevice
- // CHECK: set invalid device string: Invalid device number
-
- print_error("get last error", hipGetLastError());
- // CHECK: get last error value: 2
- // CHECK: get last error name: hipErrorInvalidDevice
- // CHECK: get last error string: Invalid device number
-
- print_error("cleared last error", hipGetLastError());
- // CHECK: cleared last error value: 0
- // CHECK: cleared last error name: hipSuccess
- // CHECK: cleared last error string: No error
-
- print_error("null stream destroy", hipStreamDestroy(nullptr));
- // CHECK: null stream destroy value: 1
- // CHECK: null stream destroy name: hipErrorInvalidValue
- // CHECK: null stream destroy string: Invalid argument value
-
- return 0;
-}
diff --git a/offload/test/offloading/HIP/get_errs.hip b/offload/test/offloading/HIP/get_errs.hip
deleted file mode 100644
index c400c5e371cda..0000000000000
--- a/offload/test/offloading/HIP/get_errs.hip
+++ /dev/null
@@ -1,81 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -pthread -std=c++17
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp -pthread -std=c++17
-// RUN: %t | %fcheck-generic
-// 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 <cstdio>
-#include <mutex>
-#include <thread>
-
-static std::mutex PrintMutex;
-
-static void printError(int ThreadId, const char *Label, hipError_t Error) {
- std::lock_guard<std::mutex> Lock(PrintMutex);
- printf("thread %d %s: %s\n", ThreadId, Label, hipGetErrorName(Error));
- std::fflush(stdout);
-}
-
-__global__ void errorKernel(float *d_out) {
- int idx = blockIdx.x * blockDim.x + threadIdx.x;
- d_out[idx] = idx * 0.5f;
-}
-
-void runTask(int ThreadId) {
- const int N = 1 << 20;
- size_t Bytes = N * sizeof(float);
-
- float *d_data;
- hipMalloc(&d_data, Bytes);
- printError(ThreadId, "hipMalloc", hipGetLastError());
-
- ThreadId == 1 ? errorKernel<<<4096, 256>>>(d_data)
- : errorKernel<<<4096, 0>>>(d_data);
- printError(ThreadId, "kernel launch", hipPeekAtLastError());
-
- printError(ThreadId, "kernel launch get", hipGetLastError());
-
- printError(ThreadId, "kernel launch get again", hipGetLastError());
-
- hipDeviceSynchronize();
- hipFree(d_data);
-}
-
-int main() {
- printError(0, "initial", hipPeekAtLastError());
- // CHECK: thread 0 initial: hipSuccess
-
- std::thread t1(runTask, 1);
- std::thread t2(runTask, 2);
-
- t1.join();
- t2.join();
- // CHECK-DAG: thread 1 hipMalloc: hipSuccess
- // CHECK-DAG: thread 2 hipMalloc: hipSuccess
- // CHECK-DAG: thread 1 kernel launch: hipSuccess
- // CHECK-DAG: thread 2 kernel launch: hipErrorInvalidConfiguration
- // CHECK-DAG: thread 1 kernel launch get: hipSuccess
- // CHECK-DAG: thread 2 kernel launch get: hipErrorInvalidConfiguration
- // CHECK-DAG: thread 1 kernel launch get again: hipSuccess
- // CHECK-DAG: thread 2 kernel launch get again: hipSuccess
-
- std::thread t3(runTask, 3);
- t3.join();
- // CHECK: thread 3 hipMalloc: hipSuccess
- // CHECK: thread 3 kernel launch: hipErrorInvalidConfiguration
- // CHECK: thread 3 kernel launch get: hipErrorInvalidConfiguration
- // CHECK: thread 3 kernel launch get again: hipSuccess
-
- printError(0, "joined", hipGetLastError());
- // CHECK: thread 0 joined: hipSuccess
-
- return 0;
-}
diff --git a/offload/test/offloading/HIP/host_alloc.hip b/offload/test/offloading/HIP/host_alloc.hip
deleted file mode 100644
index 38a83ee64f333..0000000000000
--- a/offload/test/offloading/HIP/host_alloc.hip
+++ /dev/null
@@ -1,48 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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 add(int *Ptr, int Value) { *Ptr += Value; }
-
-int main(int argc, char **argv) {
- int *HostAllocPtr = nullptr;
- if (hipHostAlloc(&HostAllocPtr, sizeof(int), hipHostAllocDefault) !=
- hipSuccess)
- return 1;
-
- *HostAllocPtr = 17;
- add<<<1, 1>>>(HostAllocPtr, 5);
- if (hipDeviceSynchronize() != hipSuccess)
- return 1;
- printf("hipHostAlloc value: %d\n", *HostAllocPtr);
- // CHECK: hipHostAlloc value: 22
-
- if (hipFreeHost(HostAllocPtr) != hipSuccess)
- return 1;
-
- int *MallocHostPtr = nullptr;
- if (hipMallocHost(&MallocHostPtr, sizeof(int)) != hipSuccess)
- return 1;
-
- *MallocHostPtr = 23;
- add<<<1, 1>>>(MallocHostPtr, 7);
- if (hipDeviceSynchronize() != hipSuccess)
- return 1;
- printf("hipMallocHost value: %d\n", *MallocHostPtr);
- // CHECK: hipMallocHost value: 30
-
- if (hipFreeHost(MallocHostPtr) != hipSuccess)
- return 1;
-}
diff --git a/offload/test/offloading/HIP/kernel_tu.hip.inc b/offload/test/offloading/HIP/kernel_tu.hip.inc
deleted file mode 100644
index d7d28a109dfc5..0000000000000
--- a/offload/test/offloading/HIP/kernel_tu.hip.inc
+++ /dev/null
@@ -1 +0,0 @@
-__global__ void square(int *A) { *A = 42; }
diff --git a/offload/test/offloading/HIP/memcpy_kinds.hip b/offload/test/offloading/HIP/memcpy_kinds.hip
deleted file mode 100644
index 6755a55aa0794..0000000000000
--- a/offload/test/offloading/HIP/memcpy_kinds.hip
+++ /dev/null
@@ -1,51 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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>
-
-int main(int argc, char **argv) {
- int HostSrc = 11;
- int HostDst = 0;
- if (hipMemcpy(&HostDst, &HostSrc, sizeof(int), hipMemcpyHostToHost) !=
- hipSuccess)
- return 1;
-
- printf("host to host: %d\n", HostDst);
- // CHECK: host to host: 11
-
- int *DevSrc = nullptr;
- int *DevDst = nullptr;
- int Result = 0;
- if (hipMalloc(&DevSrc, sizeof(int)) != hipSuccess)
- return 1;
- if (hipMalloc(&DevDst, sizeof(int)) != hipSuccess)
- return 1;
-
- HostSrc = 42;
- if (hipMemcpy(DevSrc, &HostSrc, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
- return 1;
- if (hipMemcpy(DevDst, DevSrc, sizeof(int), hipMemcpyDeviceToDevice) !=
- hipSuccess)
- return 1;
- if (hipMemcpy(&Result, DevDst, sizeof(int), hipMemcpyDeviceToHost) !=
- hipSuccess)
- return 1;
-
- printf("device to device: %d\n", Result);
- // CHECK: device to device: 42
-
- hipFree(DevSrc);
- hipFree(DevDst);
-}
diff --git a/offload/test/offloading/HIP/memset.hip b/offload/test/offloading/HIP/memset.hip
deleted file mode 100644
index 13bee34670dad..0000000000000
--- a/offload/test/offloading/HIP/memset.hip
+++ /dev/null
@@ -1,85 +0,0 @@
-// 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;
-}
diff --git a/offload/test/offloading/HIP/stream_api.hip b/offload/test/offloading/HIP/stream_api.hip
deleted file mode 100644
index 3460a6a0c351c..0000000000000
--- a/offload/test/offloading/HIP/stream_api.hip
+++ /dev/null
@@ -1,99 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// 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>
-
-static void print_error(const char *Label, hipError_t Error) {
- printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
- printf("%s name: %s\n", Label, hipGetErrorName(Error));
-}
-
-__global__ void setValue(int *Out, int Value) { *Out = Value; }
-
-int main(int argc, char **argv) {
- print_error("null stream create", hipStreamCreate(nullptr));
- // CHECK: null stream create value: 1
- // CHECK: null stream create name: hipErrorInvalidValue
- print_error("null flags stream create",
- hipStreamCreateWithFlags(nullptr, hipStreamDefault));
- // CHECK: null flags stream create value: 1
- // CHECK: null flags stream create name: hipErrorInvalidValue
-
- hipStream_t InvalidFlagsStream = nullptr;
- print_error("invalid stream flags",
- hipStreamCreateWithFlags(&InvalidFlagsStream, ~0u));
- // CHECK: invalid stream flags value: 1
- // CHECK: invalid stream flags name: hipErrorInvalidValue
- printf("invalid flags stream: %d\n", InvalidFlagsStream == nullptr);
- // CHECK: invalid flags stream: 1
-
- hipStream_t Stream = nullptr;
- if (hipStreamCreate(&Stream) != hipSuccess)
- return 1;
- hipStream_t BlockingStream = nullptr;
- if (hipStreamCreateWithFlags(&BlockingStream, hipStreamDefault) != hipSuccess)
- return 1;
- hipStream_t NonBlockingStream = nullptr;
- if (hipStreamCreateWithFlags(&NonBlockingStream, hipStreamNonBlocking) !=
- hipSuccess)
- return 1;
-
- printf("stream created: %d\n", Stream != nullptr);
- // CHECK: stream created: 1
- printf("stream flags created: %d %d\n", BlockingStream != nullptr,
- NonBlockingStream != nullptr);
- // CHECK: stream flags created: 1 1
-
- int *StreamPtr = nullptr;
- int *DefaultPtr = nullptr;
- int StreamResult = 0;
- int DefaultResult = 0;
- if (hipMalloc(&StreamPtr, sizeof(int)) != hipSuccess)
- return 1;
- if (hipMalloc(&DefaultPtr, sizeof(int)) != hipSuccess)
- return 1;
-
- setValue<<<1, 1, 0, Stream>>>(StreamPtr, 42);
- setValue<<<1, 1>>>(DefaultPtr, 17);
-
- if (hipStreamSynchronize(Stream) != hipSuccess)
- return 1;
- if (hipMemcpy(&StreamResult, StreamPtr, sizeof(int), hipMemcpyDeviceToHost) !=
- hipSuccess)
- return 1;
- if (hipMemcpy(&DefaultResult, DefaultPtr, sizeof(int),
- hipMemcpyDeviceToHost) != hipSuccess)
- return 1;
-
- printf("stream result: %d\n", StreamResult);
- // CHECK: stream result: 42
- printf("default result: %d\n", DefaultResult);
- // CHECK: default result: 17
-
- 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
- print_error("destroyed stream synchronize", hipStreamSynchronize(Stream));
- // CHECK: destroyed stream synchronize value: 4
- // CHECK: destroyed stream synchronize name: hipErrorInvalidResourceHandle
-
- hipFree(StreamPtr);
- hipFree(DefaultPtr);
-}
diff --git a/offload/test/offloading/HIP/thread_and_block_id.hip b/offload/test/offloading/HIP/thread_and_block_id.hip
deleted file mode 100644
index af4daf689e678..0000000000000
--- a/offload/test/offloading/HIP/thread_and_block_id.hip
+++ /dev/null
@@ -1,46 +0,0 @@
-// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
-// clang-format on
-
-// UNSUPPORTED: aarch64-unknown-linux-gnu
-// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
-// UNSUPPORTED: x86_64-unknown-linux-gnu
-// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
-// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
-// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
-// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
-// UNSUPPORTED: intelgpu
-
-#include <stdio.h>
-#include <stdlib.h>
-
-__global__ void fill(int *A) {
- int tid = threadIdx.x + blockDim.x * blockIdx.x;
- A[tid] = 42;
-}
-
-int main(int argc, char **argv) {
- int NThreads = 128;
- int NBlocks = 512;
- int Size = sizeof(int) * NThreads * NBlocks;
- int *Ptr = (int*)calloc(1, Size);
- int *DevPtr;
- hipMalloc(&DevPtr, Size);
- hipMemcpy(DevPtr, Ptr, Size, hipMemcpyHostToDevice);
- 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) {
- if (Ptr[I] == 42)
- continue;
- printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
- return 1;
- }
- return 0;
-}
diff --git a/offload/test/offloading/language/Inputs/DefineTestLanguageNames.inc b/offload/test/offloading/language/Inputs/DefineTestLanguageNames.inc
new file mode 100644
index 0000000000000..f088eae3a1e64
--- /dev/null
+++ b/offload/test/offloading/language/Inputs/DefineTestLanguageNames.inc
@@ -0,0 +1,18 @@
+//===-- DefineTestLanguageNames.inc - Test language API renaming ---------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OFFLOAD_TEST_LANGUAGE
+#error "OFFLOAD_TEST_LANGUAGE must be defined to cuda or hip"
+#endif
+
+#define LANGUAGE OFFLOAD_TEST_LANGUAGE
+#include "kernel/DefineLanguageNames.inc"
+
+#define OFFLOAD_TEST_STRINGIZE_IMPL(X) #X
+#define OFFLOAD_TEST_STRINGIZE(X) OFFLOAD_TEST_STRINGIZE_IMPL(X)
+#define OFFLOAD_TEST_LANGUAGE_NAME OFFLOAD_TEST_STRINGIZE(OFFLOAD_TEST_LANGUAGE)
diff --git a/offload/test/offloading/CUDA/kernel_tu.cu.inc b/offload/test/offloading/language/Inputs/launch_tu_kernel.inc
similarity index 100%
rename from offload/test/offloading/CUDA/kernel_tu.cu.inc
rename to offload/test/offloading/language/Inputs/launch_tu_kernel.inc
diff --git a/offload/test/offloading/HIP/launch_tu.hip b/offload/test/offloading/language/basic_launch.cpp
similarity index 50%
rename from offload/test/offloading/HIP/launch_tu.hip
rename to offload/test/offloading/language/basic_launch.cpp
index 20a5d6b0ff6da..44251d08df934 100644
--- a/offload/test/offloading/HIP/launch_tu.hip
+++ b/offload/test/offloading/language/basic_launch.cpp
@@ -1,8 +1,12 @@
// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip %S/kernel_tu.hip.inc -o %t.kernel_tu.o -c
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
-// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -13,18 +17,19 @@
// UNSUPPORTED: intelgpu
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
-extern __global__ void square(int *A);
+__global__ void square(int *A) { *A = 42; }
int main(int argc, char **argv) {
int *Ptr;
- hipMalloc(&Ptr, 4);
+ Malloc(&Ptr, 4);
printf("Ptr %p\n", Ptr);
// CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
- int I;
- hipDeviceSynchronize();
- hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ int I = 0;
+ DeviceSynchronize();
+ Memcpy(&I, Ptr, sizeof(int), MemcpyDeviceToHost);
printf("I: %i\n", I);
// CHECK: I: 42
}
diff --git a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip b/offload/test/offloading/language/basic_launch_blocks_and_threads.cpp
similarity index 51%
rename from offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
rename to offload/test/offloading/language/basic_launch_blocks_and_threads.cpp
index aa92c0e1db3f0..0c2e19a0f0cc4 100644
--- a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
+++ b/offload/test/offloading/language/basic_launch_blocks_and_threads.cpp
@@ -1,8 +1,12 @@
// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -13,6 +17,7 @@
// UNSUPPORTED: intelgpu
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
__global__ void incrementCounter(int *A) {
__scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
@@ -20,13 +25,13 @@ __global__ void incrementCounter(int *A) {
int main(int argc, char **argv) {
int *Ptr, I;
- hipMalloc(&Ptr, sizeof(int));
+ Malloc(&Ptr, sizeof(int));
printf("Ptr %p\n", Ptr);
// CHECK: Ptr [[Ptr:0x.*]]
- hipMemset(Ptr, 0, sizeof(int));
+ Memset(Ptr, 0, sizeof(int));
incrementCounter<<<7, 6>>>(Ptr);
- hipDeviceSynchronize();
- hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ DeviceSynchronize();
+ Memcpy(&I, Ptr, sizeof(int), MemcpyDeviceToHost);
printf("I: %i\n", I);
// CHECK: I: 42
}
diff --git a/offload/test/offloading/language/basic_launch_multi_arg.cpp b/offload/test/offloading/language/basic_launch_multi_arg.cpp
new file mode 100644
index 0000000000000..71839e938c895
--- /dev/null
+++ b/offload/test/offloading/language/basic_launch_multi_arg.cpp
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
+// clang-format on
+
+// REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+
+__global__ void square(int *Dst, short Q, int *Src, short P) {
+ *Dst = (Src[0] + Src[1]) * (Q + P);
+ Src[0] = Q;
+ Src[1] = P;
+}
+
+int main(int argc, char **argv) {
+ int *Src, *Ptr;
+ Malloc(&Ptr, 4);
+ Malloc(&Src, 8);
+
+ int I = 7;
+ int HostSrc[2] = {-2, 8};
+ Memcpy(Ptr, &I, sizeof(int), MemcpyHostToDevice);
+ Memcpy(Src, &HostSrc[0], 2 * sizeof(int), MemcpyHostToDevice);
+ square<<<1, 1>>>(Ptr, 3, Src, 4);
+ DeviceSynchronize();
+ Memcpy(&I, Ptr, sizeof(int), MemcpyDeviceToHost);
+ Memcpy(&HostSrc[0], Src, 2 * sizeof(int), MemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+ printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+ // CHECK: Src: 3, 4
+}
diff --git a/offload/test/offloading/HIP/blocking_stream_semantics.hip b/offload/test/offloading/language/blocking_stream_semantics.cpp
similarity index 50%
rename from offload/test/offloading/HIP/blocking_stream_semantics.hip
rename to offload/test/offloading/language/blocking_stream_semantics.cpp
index 8c28c7b1b2ee8..003f22a18d869 100644
--- a/offload/test/offloading/HIP/blocking_stream_semantics.hip
+++ b/offload/test/offloading/language/blocking_stream_semantics.cpp
@@ -1,10 +1,16 @@
// 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
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.legacy
+// RUN: %t.cuda.legacy | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.legacy.omp -fopenmp
+// RUN: %t.cuda.legacy.omp | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.perthread -fgpu-default-stream=per-thread
+// RUN: %t.cuda.perthread | %fcheck-generic --check-prefix=PERTHREAD
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.legacy
+// RUN: %t.hip.legacy | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.legacy.omp -fopenmp
+// RUN: %t.hip.legacy.omp | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.perthread -fgpu-default-stream=per-thread
+// RUN: %t.hip.perthread | %fcheck-generic --check-prefix=PERTHREAD
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -15,6 +21,7 @@
// UNSUPPORTED: intelgpu
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
__global__ void delayedSetValue(int *Out, int Value) {
volatile unsigned long long Delay = 0;
@@ -40,35 +47,33 @@ __global__ void copyValueAndRelease(int *In, int *Out, int *Gate) {
}
int main(int argc, char **argv) {
- hipStream_t BlockingStream = nullptr;
- if (hipStreamCreateWithFlags(&BlockingStream, hipStreamDefault) != hipSuccess)
+ Stream_t BlockingStream = nullptr;
+ if (StreamCreateWithFlags(&BlockingStream, StreamDefault) != Success)
return 1;
- hipStream_t NonBlockingStream = nullptr;
- if (hipStreamCreateWithFlags(&NonBlockingStream, hipStreamNonBlocking) !=
- hipSuccess)
+ Stream_t NonBlockingStream = nullptr;
+ if (StreamCreateWithFlags(&NonBlockingStream, StreamNonBlocking) != Success)
return 1;
int *In = nullptr;
int *Out = nullptr;
int *Gate = nullptr;
- if (hipMalloc(&In, sizeof(int)) != hipSuccess)
+ if (Malloc(&In, sizeof(int)) != Success)
return 1;
- if (hipMalloc(&Out, sizeof(int)) != hipSuccess)
+ if (Malloc(&Out, sizeof(int)) != Success)
return 1;
- if (hipMalloc(&Gate, sizeof(int)) != hipSuccess)
+ if (Malloc(&Gate, sizeof(int)) != Success)
return 1;
int Initial = 0;
int Result = 0;
- if (hipMemcpy(In, &Initial, sizeof(int), hipMemcpyHostToDevice) != hipSuccess)
+ if (Memcpy(In, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
return 1;
- if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
+ if (Memcpy(Out, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
return 1;
delayedSetValue<<<1, 1, 0, BlockingStream>>>(In, 99);
copyValue<<<1, 1>>>(In, Out);
- if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
+ if (Memcpy(&Result, Out, sizeof(int), MemcpyDeviceToHost) != Success)
return 1;
printf("legacy default waited on blocking stream: %d\n", Result);
@@ -76,15 +81,14 @@ int main(int argc, char **argv) {
// PERTHREAD: legacy default waited on blocking stream: 0
Result = 0;
- if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
+ if (Memcpy(Out, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
return 1;
delayedSetValue<<<1, 1>>>(In, 123);
copyValue<<<1, 1, 0, BlockingStream>>>(In, Out);
- if (hipStreamSynchronize(BlockingStream) != hipSuccess)
+ if (StreamSynchronize(BlockingStream) != Success)
return 1;
- if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
+ if (Memcpy(&Result, Out, sizeof(int), MemcpyDeviceToHost) != Success)
return 1;
printf("blocking stream waited on legacy default: %d\n", Result);
@@ -92,34 +96,32 @@ int main(int argc, char **argv) {
// PERTHREAD: blocking stream waited on legacy default: 99
Result = 0;
- if (hipMemcpy(In, &Initial, sizeof(int), hipMemcpyHostToDevice) != hipSuccess)
+ if (Memcpy(In, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
return 1;
- if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
+ if (Memcpy(Out, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
return 1;
- if (hipMemcpy(Gate, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
+ if (Memcpy(Gate, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
return 1;
waitThenSetValue<<<1, 1>>>(Gate, In, 321);
copyValueAndRelease<<<1, 1, 0, NonBlockingStream>>>(In, Out, Gate);
- if (hipStreamSynchronize(NonBlockingStream) != hipSuccess)
+ if (StreamSynchronize(NonBlockingStream) != Success)
return 1;
- if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
+ if (Memcpy(&Result, Out, sizeof(int), MemcpyDeviceToHost) != Success)
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)
+ if (StreamDestroy(BlockingStream) != Success)
return 1;
- if (hipStreamDestroy(NonBlockingStream) != hipSuccess)
+ if (StreamDestroy(NonBlockingStream) != Success)
return 1;
- if (hipFree(In) != hipSuccess)
+ if (Free(In) != Success)
return 1;
- if (hipFree(Out) != hipSuccess)
+ if (Free(Out) != Success)
return 1;
- if (hipFree(Gate) != hipSuccess)
+ if (Free(Gate) != Success)
return 1;
}
diff --git a/offload/test/offloading/HIP/device_api.hip b/offload/test/offloading/language/device_api.cpp
similarity index 54%
rename from offload/test/offloading/HIP/device_api.hip
rename to offload/test/offloading/language/device_api.cpp
index 5fb66e6e45eeb..2c19555be2ead 100644
--- a/offload/test/offloading/HIP/device_api.hip
+++ b/offload/test/offloading/language/device_api.cpp
@@ -1,8 +1,12 @@
// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -13,33 +17,34 @@
// UNSUPPORTED: intelgpu
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
int main(int argc, char **argv) {
int Count = 0;
- if (hipGetDeviceCount(&Count) != hipSuccess)
+ if (GetDeviceCount(&Count) != Success)
return 1;
printf("device count: %d\n", Count);
// CHECK: device count: {{[1-9][0-9]*}}
int Device = -1;
- if (hipGetDevice(&Device) != hipSuccess)
+ if (GetDevice(&Device) != Success)
return 1;
printf("device: %d\n", Device);
// CHECK: device: {{[0-9]+}}
- if (hipSetDevice(Device) != hipSuccess)
+ if (SetDevice(Device) != Success)
return 1;
int After = -1;
- if (hipGetDevice(&After) != hipSuccess)
+ if (GetDevice(&After) != Success)
return 1;
printf("device after set: %d\n", After);
// CHECK: device after set: {{[0-9]+}}
- hipError_t Err = hipSetDevice(-1);
+ Error_t Err = SetDevice(-1);
printf("set invalid device: %u\n", Err);
// CHECK: set invalid device: 2
}
diff --git a/offload/test/offloading/HIP/device_properties.hip b/offload/test/offloading/language/device_properties.cpp
similarity index 58%
rename from offload/test/offloading/HIP/device_properties.hip
rename to offload/test/offloading/language/device_properties.cpp
index 1a9b9a70f8ea9..229770b974891 100644
--- a/offload/test/offloading/HIP/device_properties.hip
+++ b/offload/test/offloading/language/device_properties.cpp
@@ -1,8 +1,12 @@
// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -13,12 +17,13 @@
// UNSUPPORTED: intelgpu
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
int main(int argc, char **argv) {
- hipDeviceProp_t Prop = {};
- hipError_t Err = hipGetDeviceProperties(&Prop, 0);
- if (Err != hipSuccess) {
- printf("hipGetDeviceProperties failed: %u\n", Err);
+ DeviceProp_t Prop = {};
+ Error_t Err = GetDeviceProperties(&Prop, 0);
+ if (Err != Success) {
+ printf("GetDeviceProperties failed: %u\n", Err);
return 1;
}
diff --git a/offload/test/offloading/HIP/devicesync_streams.hip b/offload/test/offloading/language/devicesync_streams.cpp
similarity index 51%
rename from offload/test/offloading/HIP/devicesync_streams.hip
rename to offload/test/offloading/language/devicesync_streams.cpp
index 14ac798f56ea8..b72daa7184d49 100644
--- a/offload/test/offloading/HIP/devicesync_streams.hip
+++ b/offload/test/offloading/language/devicesync_streams.cpp
@@ -1,8 +1,12 @@
// 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
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.legacy -fgpu-default-stream=legacy -pthread -std=c++17
+// RUN: %t.cuda.legacy | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.perthread -fgpu-default-stream=per-thread -pthread -std=c++17
+// RUN: %t.cuda.perthread | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.legacy -fgpu-default-stream=legacy -pthread -std=c++17
+// RUN: %t.hip.legacy | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.perthread -fgpu-default-stream=per-thread -pthread -std=c++17
+// RUN: %t.hip.perthread | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -15,6 +19,7 @@
#include <chrono>
#include <cstdio>
#include <thread>
+#include "Inputs/DefineTestLanguageNames.inc"
__global__ void waitThenSet(volatile int *Gate, volatile int *Out, int Value) {
for (unsigned long long I = 0; I < 1000000000ULL && *Gate == 0; ++I)
@@ -23,29 +28,25 @@ __global__ void waitThenSet(volatile int *Gate, volatile int *Out, int Value) {
}
int main(int argc, char **argv) {
- hipStream_t BlockingStream = nullptr;
- if (hipStreamCreateWithFlags(&BlockingStream, hipStreamDefault) != hipSuccess)
+ Stream_t BlockingStream = nullptr;
+ if (StreamCreateWithFlags(&BlockingStream, StreamDefault) != Success)
return 1;
- hipStream_t NonBlockingStream = nullptr;
- if (hipStreamCreateWithFlags(&NonBlockingStream, hipStreamNonBlocking) !=
- hipSuccess)
+ Stream_t NonBlockingStream = nullptr;
+ if (StreamCreateWithFlags(&NonBlockingStream, StreamNonBlocking) != Success)
return 1;
int *BlockingGate = nullptr;
int *NonBlockingGate = nullptr;
int *BlockingOutStorage = nullptr;
int *NonBlockingOutStorage = nullptr;
- if (hipHostAlloc(&BlockingGate, sizeof(int), hipHostAllocDefault) !=
- hipSuccess)
+ if (HostAlloc(&BlockingGate, sizeof(int), HostAllocDefault) != Success)
return 1;
- if (hipHostAlloc(&NonBlockingGate, sizeof(int), hipHostAllocDefault) !=
- hipSuccess)
+ if (HostAlloc(&NonBlockingGate, sizeof(int), HostAllocDefault) != Success)
return 1;
- if (hipHostAlloc(&BlockingOutStorage, sizeof(int), hipHostAllocDefault) !=
- hipSuccess)
+ if (HostAlloc(&BlockingOutStorage, sizeof(int), HostAllocDefault) != Success)
return 1;
- if (hipHostAlloc(&NonBlockingOutStorage, sizeof(int), hipHostAllocDefault) !=
- hipSuccess)
+ if (HostAlloc(&NonBlockingOutStorage, sizeof(int), HostAllocDefault) !=
+ Success)
return 1;
volatile int *BlockingOut = BlockingOutStorage;
@@ -65,9 +66,9 @@ int main(int argc, char **argv) {
*NonBlockingGate = 1;
});
- hipError_t SyncResult = hipDeviceSynchronize();
+ Error_t SyncResult = DeviceSynchronize();
- if (SyncResult == hipSuccess) {
+ if (SyncResult == Success) {
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);
@@ -75,23 +76,23 @@ int main(int argc, char **argv) {
}
Releaser.join();
- if (hipStreamSynchronize(BlockingStream) != hipSuccess)
+ if (StreamSynchronize(BlockingStream) != Success)
return 1;
- if (hipStreamSynchronize(NonBlockingStream) != hipSuccess)
+ if (StreamSynchronize(NonBlockingStream) != Success)
return 1;
- if (SyncResult != hipSuccess)
+ if (SyncResult != Success)
return 1;
- if (hipStreamDestroy(BlockingStream) != hipSuccess)
+ if (StreamDestroy(BlockingStream) != Success)
return 1;
- if (hipStreamDestroy(NonBlockingStream) != hipSuccess)
+ if (StreamDestroy(NonBlockingStream) != Success)
return 1;
- if (hipFreeHost(BlockingGate) != hipSuccess)
+ if (FreeHost(BlockingGate) != Success)
return 1;
- if (hipFreeHost(NonBlockingGate) != hipSuccess)
+ if (FreeHost(NonBlockingGate) != Success)
return 1;
- if (hipFreeHost(BlockingOutStorage) != hipSuccess)
+ if (FreeHost(BlockingOutStorage) != Success)
return 1;
- if (hipFreeHost(NonBlockingOutStorage) != hipSuccess)
+ if (FreeHost(NonBlockingOutStorage) != Success)
return 1;
}
diff --git a/offload/test/offloading/language/error_kinds.cpp b/offload/test/offloading/language/error_kinds.cpp
new file mode 100644
index 0000000000000..96671b3a4ad3b
--- /dev/null
+++ b/offload/test/offloading/language/error_kinds.cpp
@@ -0,0 +1,95 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic --check-prefixes=CHECK,CUDA
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic --check-prefixes=CHECK,CUDA
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic --check-prefixes=CHECK,HIP
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic --check-prefixes=CHECK,HIP
+// 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>
+#include "Inputs/DefineTestLanguageNames.inc"
+
+static void print_error(const char *Label, Error_t Error) {
+ printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+ printf("%s name: %s\n", Label, GetErrorName(Error));
+ printf("%s string: %s\n", Label, GetErrorString(Error));
+}
+
+int main() {
+ print_error("success", Success);
+ // CHECK: success value: 0
+ // CUDA: success name: cudaSuccess
+ // HIP: success name: hipSuccess
+ // CHECK: success string: No error
+
+ print_error("invalid value", ErrorInvalidValue);
+ // CHECK: invalid value value: 1
+ // CUDA: invalid value name: cudaErrorInvalidValue
+ // HIP: invalid value name: hipErrorInvalidValue
+ // CHECK: invalid value string: Invalid argument value
+
+ print_error("invalid device", ErrorInvalidDevice);
+ // CHECK: invalid device value: 2
+ // CUDA: invalid device name: cudaErrorInvalidDevice
+ // HIP: invalid device name: hipErrorInvalidDevice
+ // CHECK: invalid device string: Invalid device number
+
+ print_error("unknown", ErrorUnknown);
+ // CHECK: unknown value: 3
+ // CHECK: unknown name: Unrecognized error
+ // CHECK: unknown string: Unknown error
+
+ print_error("invalid resource handle", ErrorInvalidResourceHandle);
+ // CHECK: invalid resource handle value: 4
+ // CUDA: invalid resource handle name: cudaErrorInvalidResourceHandle
+ // HIP: invalid resource handle name: hipErrorInvalidResourceHandle
+ // CHECK: invalid resource handle string: Invalid resource handle
+
+ print_error("invalid configuration", ErrorInvalidConfiguration);
+ // CHECK: invalid configuration value: 5
+ // CUDA: invalid configuration name: cudaErrorInvalidConfiguration
+ // HIP: invalid configuration name: hipErrorInvalidConfiguration
+ // CHECK: invalid configuration string: Invalid configuration argument
+
+ Error_t Unrecognized = static_cast<Error_t>(999);
+ print_error("unrecognized", Unrecognized);
+ // CHECK: unrecognized value: 999
+ // CHECK: unrecognized name: Unrecognized error
+ // CHECK: unrecognized string: Unrecognized error
+
+ print_error("set invalid device", SetDevice(-1));
+ // CHECK: set invalid device value: 2
+ // CUDA: set invalid device name: cudaErrorInvalidDevice
+ // HIP: set invalid device name: hipErrorInvalidDevice
+ // CHECK: set invalid device string: Invalid device number
+
+ print_error("get last error", GetLastError());
+ // CHECK: get last error value: 2
+ // CUDA: get last error name: cudaErrorInvalidDevice
+ // HIP: get last error name: hipErrorInvalidDevice
+ // CHECK: get last error string: Invalid device number
+
+ print_error("cleared last error", GetLastError());
+ // CHECK: cleared last error value: 0
+ // CUDA: cleared last error name: cudaSuccess
+ // HIP: cleared last error name: hipSuccess
+ // CHECK: cleared last error string: No error
+
+ print_error("null stream destroy", StreamDestroy(nullptr));
+ // CHECK: null stream destroy value: 1
+ // CUDA: null stream destroy name: cudaErrorInvalidValue
+ // HIP: null stream destroy name: hipErrorInvalidValue
+ // CHECK: null stream destroy string: Invalid argument value
+
+ return 0;
+}
diff --git a/offload/test/offloading/language/get_errs.cpp b/offload/test/offloading/language/get_errs.cpp
new file mode 100644
index 0000000000000..5371369d86ef4
--- /dev/null
+++ b/offload/test/offloading/language/get_errs.cpp
@@ -0,0 +1,100 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda -pthread -std=c++17
+// RUN: %t.cuda | %fcheck-generic --check-prefix=CUDA
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp -pthread -std=c++17
+// RUN: %t.cuda.omp | %fcheck-generic --check-prefix=CUDA
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip -pthread -std=c++17
+// RUN: %t.hip | %fcheck-generic --check-prefix=HIP
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp -pthread -std=c++17
+// RUN: %t.hip.omp | %fcheck-generic --check-prefix=HIP
+// 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 <cstdio>
+#include <mutex>
+#include <thread>
+#include "Inputs/DefineTestLanguageNames.inc"
+
+static std::mutex PrintMutex;
+
+static void printError(int ThreadId, const char *Label, Error_t Error) {
+ std::lock_guard<std::mutex> Lock(PrintMutex);
+ printf("thread %d %s: %s\n", ThreadId, Label, GetErrorName(Error));
+ std::fflush(stdout);
+}
+
+__global__ void errorKernel(float *d_out) {
+ int idx = blockIdx.x * blockDim.x + threadIdx.x;
+ d_out[idx] = idx * 0.5f;
+}
+
+void runTask(int ThreadId) {
+ const int N = 1 << 20;
+ size_t bytes = N * sizeof(float);
+
+ float *d_data;
+ Malloc(&d_data, bytes);
+ printError(ThreadId, "Malloc", GetLastError());
+
+ ThreadId == 1 ? errorKernel<<<4096, 256>>>(d_data)
+ : errorKernel<<<4096, 0>>>(d_data);
+ printError(ThreadId, "kernel launch", PeekAtLastError());
+
+ printError(ThreadId, "kernel launch get", GetLastError());
+
+ printError(ThreadId, "kernel launch get again", GetLastError());
+
+ DeviceSynchronize();
+ Free(d_data);
+}
+
+int main() {
+ printError(0, "initial", PeekAtLastError());
+ // CUDA: thread 0 initial: cudaSuccess
+ // HIP: thread 0 initial: hipSuccess
+
+ std::thread t1(runTask, 1);
+ std::thread t2(runTask, 2);
+
+ t1.join();
+ t2.join();
+ // CUDA-DAG: thread 1 Malloc: cudaSuccess
+ // HIP-DAG: thread 1 Malloc: hipSuccess
+ // CUDA-DAG: thread 2 Malloc: cudaSuccess
+ // HIP-DAG: thread 2 Malloc: hipSuccess
+ // CUDA-DAG: thread 1 kernel launch: cudaSuccess
+ // HIP-DAG: thread 1 kernel launch: hipSuccess
+ // CUDA-DAG: thread 2 kernel launch: cudaErrorInvalidConfiguration
+ // HIP-DAG: thread 2 kernel launch: hipErrorInvalidConfiguration
+ // CUDA-DAG: thread 1 kernel launch get: cudaSuccess
+ // HIP-DAG: thread 1 kernel launch get: hipSuccess
+ // CUDA-DAG: thread 2 kernel launch get: cudaErrorInvalidConfiguration
+ // HIP-DAG: thread 2 kernel launch get: hipErrorInvalidConfiguration
+ // CUDA-DAG: thread 1 kernel launch get again: cudaSuccess
+ // HIP-DAG: thread 1 kernel launch get again: hipSuccess
+ // CUDA-DAG: thread 2 kernel launch get again: cudaSuccess
+ // HIP-DAG: thread 2 kernel launch get again: hipSuccess
+
+ std::thread t3(runTask, 3);
+ t3.join();
+ // CUDA: thread 3 Malloc: cudaSuccess
+ // HIP: thread 3 Malloc: hipSuccess
+ // CUDA: thread 3 kernel launch: cudaErrorInvalidConfiguration
+ // HIP: thread 3 kernel launch: hipErrorInvalidConfiguration
+ // CUDA: thread 3 kernel launch get: cudaErrorInvalidConfiguration
+ // HIP: thread 3 kernel launch get: hipErrorInvalidConfiguration
+ // CUDA: thread 3 kernel launch get again: cudaSuccess
+ // HIP: thread 3 kernel launch get again: hipSuccess
+
+ printError(0, "joined", GetLastError());
+ // CUDA: thread 0 joined: cudaSuccess
+ // HIP: thread 0 joined: hipSuccess
+
+ return 0;
+}
diff --git a/offload/test/offloading/language/host_alloc.cpp b/offload/test/offloading/language/host_alloc.cpp
new file mode 100644
index 0000000000000..a61f618fd9a6c
--- /dev/null
+++ b/offload/test/offloading/language/host_alloc.cpp
@@ -0,0 +1,52 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
+// 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>
+#include "Inputs/DefineTestLanguageNames.inc"
+
+__global__ void add(int *Ptr, int Value) { *Ptr += Value; }
+
+int main(int argc, char **argv) {
+ int *HostAllocPtr = nullptr;
+ if (HostAlloc(&HostAllocPtr, sizeof(int), HostAllocDefault) != Success)
+ return 1;
+
+ *HostAllocPtr = 17;
+ add<<<1, 1>>>(HostAllocPtr, 5);
+ if (DeviceSynchronize() != Success)
+ return 1;
+ printf("HostAlloc value: %d\n", *HostAllocPtr);
+ // CHECK: HostAlloc value: 22
+
+ if (FreeHost(HostAllocPtr) != Success)
+ return 1;
+
+ int *MallocHostPtr = nullptr;
+ if (MallocHost(&MallocHostPtr, sizeof(int)) != Success)
+ return 1;
+
+ *MallocHostPtr = 23;
+ add<<<1, 1>>>(MallocHostPtr, 7);
+ if (DeviceSynchronize() != Success)
+ return 1;
+ printf("MallocHost value: %d\n", *MallocHostPtr);
+ // CHECK: MallocHost value: 30
+
+ if (FreeHost(MallocHostPtr) != Success)
+ return 1;
+}
diff --git a/offload/test/offloading/language/launch_tu.cpp b/offload/test/offloading/language/launch_tu.cpp
new file mode 100644
index 0000000000000..2711a072597fb
--- /dev/null
+++ b/offload/test/offloading/language/launch_tu.cpp
@@ -0,0 +1,35 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.launch_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %S/Inputs/launch_tu_kernel.inc -o %t.cuda.kernel_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.cuda.launch_tu.o %t.cuda.kernel_tu.o -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.launch_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %S/Inputs/launch_tu_kernel.inc -o %t.hip.kernel_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.hip.launch_tu.o %t.hip.kernel_tu.o -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// 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>
+#include "Inputs/DefineTestLanguageNames.inc"
+
+extern __global__ void square(int *A);
+
+int main(int argc, char **argv) {
+ int *Ptr;
+ Malloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<1, 1>>>(Ptr);
+ int I;
+ DeviceSynchronize();
+ Memcpy(&I, Ptr, sizeof(int), MemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/language/memcpy_kinds.cpp b/offload/test/offloading/language/memcpy_kinds.cpp
new file mode 100644
index 0000000000000..bc1b4fc6614b5
--- /dev/null
+++ b/offload/test/offloading/language/memcpy_kinds.cpp
@@ -0,0 +1,52 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
+// 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>
+#include "Inputs/DefineTestLanguageNames.inc"
+
+int main(int argc, char **argv) {
+ int HostSrc = 11;
+ int HostDst = 0;
+ if (Memcpy(&HostDst, &HostSrc, sizeof(int), MemcpyHostToHost) != Success)
+ return 1;
+
+ printf("host to host: %d\n", HostDst);
+ // CHECK: host to host: 11
+
+ int *DevSrc = nullptr;
+ int *DevDst = nullptr;
+ int Result = 0;
+ if (Malloc(&DevSrc, sizeof(int)) != Success)
+ return 1;
+ if (Malloc(&DevDst, sizeof(int)) != Success)
+ return 1;
+
+ HostSrc = 42;
+ if (Memcpy(DevSrc, &HostSrc, sizeof(int), MemcpyHostToDevice) != Success)
+ return 1;
+ if (Memcpy(DevDst, DevSrc, sizeof(int), MemcpyDeviceToDevice) != Success)
+ return 1;
+ if (Memcpy(&Result, DevDst, sizeof(int), MemcpyDeviceToHost) != Success)
+ return 1;
+
+ printf("device to device: %d\n", Result);
+ // CHECK: device to device: 42
+
+ Free(DevSrc);
+ Free(DevDst);
+}
diff --git a/offload/test/offloading/CUDA/memset.cu b/offload/test/offloading/language/memset.cpp
similarity index 52%
rename from offload/test/offloading/CUDA/memset.cu
rename to offload/test/offloading/language/memset.cpp
index 44a5aeba8dcdf..156d9eabb0dec 100644
--- a/offload/test/offloading/CUDA/memset.cu
+++ b/offload/test/offloading/language/memset.cpp
@@ -1,8 +1,12 @@
// 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
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.legacy -fgpu-default-stream=legacy -pthread -std=c++17
+// RUN: %t.cuda.legacy | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.perthread -fgpu-default-stream=per-thread -pthread -std=c++17
+// RUN: %t.cuda.perthread | %fcheck-generic --check-prefix=PERTHREAD
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.legacy -fgpu-default-stream=legacy -pthread -std=c++17
+// RUN: %t.hip.legacy | %fcheck-generic --check-prefix=LEGACY
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.perthread -fgpu-default-stream=per-thread -pthread -std=c++17
+// RUN: %t.hip.perthread | %fcheck-generic --check-prefix=PERTHREAD
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -15,6 +19,7 @@
#include <chrono>
#include <cstdio>
#include <thread>
+#include "Inputs/DefineTestLanguageNames.inc"
__global__ void waitThenSet(volatile int *Gate, unsigned char *Out,
unsigned char Value) {
@@ -26,15 +31,14 @@ __global__ void waitThenSet(volatile int *Gate, unsigned char *Out,
int main(int argc, char **argv) {
unsigned char *Dev = nullptr;
- if (cudaMalloc(&Dev, 4) != cudaSuccess)
+ if (Malloc(&Dev, 4) != Success)
return 1;
- if (cudaMemset(Dev, 0x2a, 4) != cudaSuccess)
+ if (Memset(Dev, 0x2a, 4) != Success)
return 1;
unsigned char Host[4] = {};
- if (cudaMemcpy(Host, Dev, sizeof(Host), cudaMemcpyDeviceToHost) !=
- cudaSuccess)
+ if (Memcpy(Host, Dev, sizeof(Host), MemcpyDeviceToHost) != Success)
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]),
@@ -42,17 +46,16 @@ int main(int argc, char **argv) {
// LEGACY: memset bytes: 42 42 42 42
// PERTHREAD: memset bytes: 42 42 42 42
- cudaStream_t BlockingStream = nullptr;
- if (cudaStreamCreateWithFlags(&BlockingStream, cudaStreamDefault) !=
- cudaSuccess)
+ Stream_t BlockingStream = nullptr;
+ if (StreamCreateWithFlags(&BlockingStream, StreamDefault) != Success)
return 1;
int *Gate = nullptr;
- if (cudaHostAlloc(&Gate, sizeof(int), cudaHostAllocDefault) != cudaSuccess)
+ if (HostAlloc(&Gate, sizeof(int), HostAllocDefault) != Success)
return 1;
*Gate = 0;
- if (cudaMemset(Dev, 0, 1) != cudaSuccess)
+ if (Memset(Dev, 0, 1) != Success)
return 1;
waitThenSet<<<1, 1, 0, BlockingStream>>>(Gate, Dev, 17);
@@ -62,26 +65,26 @@ int main(int argc, char **argv) {
*Gate = 1;
});
- cudaError_t MemsetResult = cudaMemset(Dev, 23, 1);
+ Error_t MemsetResult = Memset(Dev, 23, 1);
Releaser.join();
- if (MemsetResult != cudaSuccess)
+ if (MemsetResult != Success)
return 1;
- if (cudaStreamSynchronize(BlockingStream) != cudaSuccess)
+ if (StreamSynchronize(BlockingStream) != Success)
return 1;
unsigned char Result = 0;
- if (cudaMemcpy(&Result, Dev, 1, cudaMemcpyDeviceToHost) != cudaSuccess)
+ if (Memcpy(&Result, Dev, 1, MemcpyDeviceToHost) != Success)
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)
+ if (StreamDestroy(BlockingStream) != Success)
return 1;
- if (cudaFreeHost(Gate) != cudaSuccess)
+ if (FreeHost(Gate) != Success)
return 1;
- if (cudaFree(Dev) != cudaSuccess)
+ if (Free(Dev) != Success)
return 1;
}
diff --git a/offload/test/offloading/language/stream_api.cpp b/offload/test/offloading/language/stream_api.cpp
new file mode 100644
index 0000000000000..04bd5122502f6
--- /dev/null
+++ b/offload/test/offloading/language/stream_api.cpp
@@ -0,0 +1,108 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic --check-prefixes=CHECK,CUDA
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic --check-prefixes=CHECK,CUDA
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic --check-prefixes=CHECK,HIP
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic --check-prefixes=CHECK,HIP
+// 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>
+#include "Inputs/DefineTestLanguageNames.inc"
+
+static void print_error(const char *Label, Error_t Error) {
+ printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+ printf("%s name: %s\n", Label, GetErrorName(Error));
+}
+
+__global__ void setValue(int *Out, int Value) { *Out = Value; }
+
+int main(int argc, char **argv) {
+ print_error("null stream create", StreamCreate(nullptr));
+ // CHECK: null stream create value: 1
+ // CUDA: null stream create name: cudaErrorInvalidValue
+ // HIP: null stream create name: hipErrorInvalidValue
+ print_error("null flags stream create",
+ StreamCreateWithFlags(nullptr, StreamDefault));
+ // CHECK: null flags stream create value: 1
+ // CUDA: null flags stream create name: cudaErrorInvalidValue
+ // HIP: null flags stream create name: hipErrorInvalidValue
+
+ Stream_t InvalidFlagsStream = nullptr;
+ print_error("invalid stream flags",
+ StreamCreateWithFlags(&InvalidFlagsStream, ~0u));
+ // CHECK: invalid stream flags value: 1
+ // CUDA: invalid stream flags name: cudaErrorInvalidValue
+ // HIP: invalid stream flags name: hipErrorInvalidValue
+ printf("invalid flags stream: %d\n", InvalidFlagsStream == nullptr);
+ // CHECK: invalid flags stream: 1
+
+ Stream_t Stream = nullptr;
+ if (StreamCreate(&Stream) != Success)
+ return 1;
+ Stream_t BlockingStream = nullptr;
+ if (StreamCreateWithFlags(&BlockingStream, StreamDefault) != Success)
+ return 1;
+ Stream_t NonBlockingStream = nullptr;
+ if (StreamCreateWithFlags(&NonBlockingStream, StreamNonBlocking) != Success)
+ return 1;
+
+ printf("stream created: %d\n", Stream != nullptr);
+ // CHECK: stream created: 1
+ printf("stream flags created: %d %d\n", BlockingStream != nullptr,
+ NonBlockingStream != nullptr);
+ // CHECK: stream flags created: 1 1
+
+ int *StreamPtr = nullptr;
+ int *DefaultPtr = nullptr;
+ int StreamResult = 0;
+ int DefaultResult = 0;
+ if (Malloc(&StreamPtr, sizeof(int)) != Success)
+ return 1;
+ if (Malloc(&DefaultPtr, sizeof(int)) != Success)
+ return 1;
+
+ setValue<<<1, 1, 0, Stream>>>(StreamPtr, 42);
+ setValue<<<1, 1>>>(DefaultPtr, 17);
+
+ if (StreamSynchronize(Stream) != Success)
+ return 1;
+ if (Memcpy(&StreamResult, StreamPtr, sizeof(int), MemcpyDeviceToHost) !=
+ Success)
+ return 1;
+ if (Memcpy(&DefaultResult, DefaultPtr, sizeof(int), MemcpyDeviceToHost) !=
+ Success)
+ return 1;
+
+ printf("stream result: %d\n", StreamResult);
+ // CHECK: stream result: 42
+ printf("default result: %d\n", DefaultResult);
+ // CHECK: default result: 17
+
+ if (StreamDestroy(Stream) != Success)
+ return 1;
+ if (StreamDestroy(BlockingStream) != Success)
+ return 1;
+ if (StreamDestroy(NonBlockingStream) != Success)
+ return 1;
+ print_error("destroyed stream destroy", StreamDestroy(Stream));
+ // CHECK: destroyed stream destroy value: 4
+ // CUDA: destroyed stream destroy name: cudaErrorInvalidResourceHandle
+ // HIP: destroyed stream destroy name: hipErrorInvalidResourceHandle
+ print_error("destroyed stream synchronize", StreamSynchronize(Stream));
+ // CHECK: destroyed stream synchronize value: 4
+ // CUDA: destroyed stream synchronize name: cudaErrorInvalidResourceHandle
+ // HIP: destroyed stream synchronize name: hipErrorInvalidResourceHandle
+
+ Free(StreamPtr);
+ Free(DefaultPtr);
+}
diff --git a/offload/test/offloading/HIP/syncthreads.hip b/offload/test/offloading/language/syncthreads.cpp
similarity index 55%
rename from offload/test/offloading/HIP/syncthreads.hip
rename to offload/test/offloading/language/syncthreads.cpp
index 81e9ed0451ca1..71781a3a3aea8 100644
--- a/offload/test/offloading/HIP/syncthreads.hip
+++ b/offload/test/offloading/language/syncthreads.cpp
@@ -1,8 +1,12 @@
// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -13,6 +17,7 @@
// UNSUPPORTED: intelgpu
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
__global__ void reduceBlock(int *Out) {
__shared__ int Scratch[64];
@@ -31,10 +36,10 @@ __global__ void reduceBlock(int *Out) {
int main(int argc, char **argv) {
int *DevPtr;
int Result = 0;
- hipMalloc(&DevPtr, sizeof(int));
+ Malloc(&DevPtr, sizeof(int));
reduceBlock<<<1, 64>>>(DevPtr);
- hipDeviceSynchronize();
- hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost);
+ DeviceSynchronize();
+ Memcpy(&Result, DevPtr, sizeof(int), MemcpyDeviceToHost);
printf("sum: %i\n", Result);
// CHECK: sum: 2016
diff --git a/offload/test/offloading/CUDA/thread_and_block_id.cu b/offload/test/offloading/language/thread_and_block_id.cpp
similarity index 61%
rename from offload/test/offloading/CUDA/thread_and_block_id.cu
rename to offload/test/offloading/language/thread_and_block_id.cpp
index a56c9ff33e2ab..014dbda9e36f8 100644
--- a/offload/test/offloading/CUDA/thread_and_block_id.cu
+++ b/offload/test/offloading/language/thread_and_block_id.cpp
@@ -1,8 +1,12 @@
// clang-format off
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
-// RUN: %t | %fcheck-generic
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
-// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
@@ -16,6 +20,7 @@
#include <stdio.h>
#include <stdlib.h>
+#include "Inputs/DefineTestLanguageNames.inc"
__global__ void fill(int *A) {
int tid = threadIdx.x + blockDim.x * blockIdx.x;
@@ -28,13 +33,13 @@ int main(int argc, char **argv) {
int Size = sizeof(int) * NThreads * NBlocks;
int *Ptr = (int *)calloc(1, Size);
int *DevPtr;
- cudaMalloc(&DevPtr, Size);
- cudaMemcpy(DevPtr, Ptr, Size, cudaMemcpyHostToDevice);
+ Malloc(&DevPtr, Size);
+ Memcpy(DevPtr, Ptr, Size, MemcpyHostToDevice);
printf("DevPtr %p\n", DevPtr);
// CHECK: DevPtr [[DevPtr:0x.*]]
fill<<<NBlocks, NThreads>>>(DevPtr);
- cudaDeviceSynchronize();
- cudaMemcpy(Ptr, DevPtr, Size, cudaMemcpyDeviceToHost);
+ DeviceSynchronize();
+ Memcpy(Ptr, DevPtr, Size, MemcpyDeviceToHost);
for (int I = 0; I < NBlocks * NThreads; ++I) {
if (Ptr[I] == 42)
More information about the llvm-branch-commits
mailing list