[llvm-branch-commits] [llvm] [Offload][Lang] Unify Tests (PR #218515)
Sophia Herrmann via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Aug 25 16:55:04 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/218515
>From 7e652763f0d38f635c93d7737a475365ff985425 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 ------
.../HIP/blocking_stream_semantics.hip | 125 -----------------
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/launch_tu.hip | 30 ----
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
.../test/offloading/language/basic_launch.cpp | 37 +++++
.../basic_launch_blocks_and_threads.cpp} | 23 +--
.../language/basic_launch_multi_arg.cpp | 46 ++++++
.../language/blocking_stream_semantics.cpp | 129 +++++++++++++++++
.../device_api.cpp} | 25 ++--
.../device_properties.cpp} | 23 +--
.../devicesync_streams.cpp} | 59 ++++----
.../test/offloading/language/error_kinds.cpp | 97 +++++++++++++
offload/test/offloading/language/get_errs.cpp | 102 ++++++++++++++
.../test/offloading/language/host_alloc.cpp | 54 ++++++++
.../test/offloading/language/launch_tu.cpp | 37 +++++
.../test/offloading/language/memcpy_kinds.cpp | 54 ++++++++
.../{CUDA/memset.cu => language/memset.cpp} | 45 +++---
.../test/offloading/language/stream_api.cpp | 110 +++++++++++++++
.../syncthreads.cpp} | 21 ++-
.../thread_and_block_id.cpp} | 23 +--
44 files changed, 815 insertions(+), 1652 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/blocking_stream_semantics.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/launch_tu.hip
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%)
create mode 100644 offload/test/offloading/language/basic_launch.cpp
rename offload/test/offloading/{HIP/basic_launch_blocks_and_threads.hip => language/basic_launch_blocks_and_threads.cpp} (50%)
create mode 100644 offload/test/offloading/language/basic_launch_multi_arg.cpp
create mode 100644 offload/test/offloading/language/blocking_stream_semantics.cpp
rename offload/test/offloading/{HIP/device_api.hip => language/device_api.cpp} (53%)
rename offload/test/offloading/{HIP/device_properties.hip => language/device_properties.cpp} (57%)
rename offload/test/offloading/{HIP/devicesync_streams.hip => language/devicesync_streams.cpp} (50%)
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} (51%)
create mode 100644 offload/test/offloading/language/stream_api.cpp
rename offload/test/offloading/{HIP/syncthreads.hip => language/syncthreads.cpp} (54%)
rename offload/test/offloading/{CUDA/thread_and_block_id.cu => language/thread_and_block_id.cpp} (59%)
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/blocking_stream_semantics.hip b/offload/test/offloading/HIP/blocking_stream_semantics.hip
deleted file mode 100644
index 8c28c7b1b2ee8..0000000000000
--- a/offload/test/offloading/HIP/blocking_stream_semantics.hip
+++ /dev/null
@@ -1,125 +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) {
- hipStream_t BlockingStream = nullptr;
- if (hipStreamCreateWithFlags(&BlockingStream, hipStreamDefault) != hipSuccess)
- return 1;
- hipStream_t NonBlockingStream = nullptr;
- if (hipStreamCreateWithFlags(&NonBlockingStream, hipStreamNonBlocking) !=
- hipSuccess)
- return 1;
-
- int *In = nullptr;
- int *Out = nullptr;
- int *Gate = nullptr;
- if (hipMalloc(&In, sizeof(int)) != hipSuccess)
- return 1;
- if (hipMalloc(&Out, sizeof(int)) != hipSuccess)
- return 1;
- if (hipMalloc(&Gate, sizeof(int)) != hipSuccess)
- return 1;
-
- int Initial = 0;
- int Result = 0;
- if (hipMemcpy(In, &Initial, sizeof(int), hipMemcpyHostToDevice) != hipSuccess)
- return 1;
- if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
- return 1;
-
- delayedSetValue<<<1, 1, 0, BlockingStream>>>(In, 99);
- copyValue<<<1, 1>>>(In, Out);
- if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
- return 1;
-
- printf("legacy default waited on blocking stream: %d\n", Result);
- // LEGACY: legacy default waited on blocking stream: 99
- // PERTHREAD: legacy default waited on blocking stream: 0
-
- Result = 0;
- if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
- return 1;
-
- delayedSetValue<<<1, 1>>>(In, 123);
- copyValue<<<1, 1, 0, BlockingStream>>>(In, Out);
- if (hipStreamSynchronize(BlockingStream) != hipSuccess)
- return 1;
- if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
- return 1;
-
- printf("blocking stream waited on legacy default: %d\n", Result);
- // LEGACY: blocking stream waited on legacy default: 123
- // PERTHREAD: blocking stream waited on legacy default: 99
-
- Result = 0;
- if (hipMemcpy(In, &Initial, sizeof(int), hipMemcpyHostToDevice) != hipSuccess)
- return 1;
- if (hipMemcpy(Out, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
- return 1;
- if (hipMemcpy(Gate, &Initial, sizeof(int), hipMemcpyHostToDevice) !=
- hipSuccess)
- return 1;
-
- waitThenSetValue<<<1, 1>>>(Gate, In, 321);
- copyValueAndRelease<<<1, 1, 0, NonBlockingStream>>>(In, Out, Gate);
- if (hipStreamSynchronize(NonBlockingStream) != hipSuccess)
- return 1;
- if (hipMemcpy(&Result, Out, sizeof(int), hipMemcpyDeviceToHost) != hipSuccess)
- return 1;
-
- printf("nonblocking stream did not wait on legacy default: %d\n", Result);
- // LEGACY: nonblocking stream did not wait on legacy default: 0
- // PERTHREAD: nonblocking stream did not wait on legacy default: 0
-
- if (hipStreamDestroy(BlockingStream) != hipSuccess)
- return 1;
- if (hipStreamDestroy(NonBlockingStream) != hipSuccess)
- return 1;
- if (hipFree(In) != hipSuccess)
- return 1;
- if (hipFree(Out) != hipSuccess)
- return 1;
- if (hipFree(Gate) != hipSuccess)
- return 1;
-}
diff --git a/offload/test/offloading/HIP/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/launch_tu.hip b/offload/test/offloading/HIP/launch_tu.hip
deleted file mode 100644
index 20a5d6b0ff6da..0000000000000
--- a/offload/test/offloading/HIP/launch_tu.hip
+++ /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 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
-// 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;
- hipMalloc(&Ptr, 4);
- printf("Ptr %p\n", Ptr);
- // CHECK: Ptr [[Ptr:0x.*]]
- square<<<1, 1>>>(Ptr);
- int I;
- hipDeviceSynchronize();
- hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
- printf("I: %i\n", I);
- // CHECK: I: 42
-}
diff --git a/offload/test/offloading/HIP/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/language/basic_launch.cpp b/offload/test/offloading/language/basic_launch.cpp
new file mode 100644
index 0000000000000..7067c6cc8a9d0
--- /dev/null
+++ b/offload/test/offloading/language/basic_launch.cpp
@@ -0,0 +1,37 @@
+// 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
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+__global__ void square(int *A) { *A = 42; }
+
+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 = 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 50%
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..17243df360024 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
@@ -12,7 +16,10 @@
// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
+// clang-format off
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
__global__ void incrementCounter(int *A) {
__scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
@@ -20,13 +27,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..e538897f9906d
--- /dev/null
+++ b/offload/test/offloading/language/basic_launch_multi_arg.cpp
@@ -0,0 +1,46 @@
+// 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
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+__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/language/blocking_stream_semantics.cpp b/offload/test/offloading/language/blocking_stream_semantics.cpp
new file mode 100644
index 0000000000000..bf0fec4d858e3
--- /dev/null
+++ b/offload/test/offloading/language/blocking_stream_semantics.cpp
@@ -0,0 +1,129 @@
+// clang-format off
+// 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
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+__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) {
+ Stream_t BlockingStream = nullptr;
+ if (StreamCreateWithFlags(&BlockingStream, StreamDefault) != Success)
+ return 1;
+ Stream_t NonBlockingStream = nullptr;
+ if (StreamCreateWithFlags(&NonBlockingStream, StreamNonBlocking) != Success)
+ return 1;
+
+ int *In = nullptr;
+ int *Out = nullptr;
+ int *Gate = nullptr;
+ if (Malloc(&In, sizeof(int)) != Success)
+ return 1;
+ if (Malloc(&Out, sizeof(int)) != Success)
+ return 1;
+ if (Malloc(&Gate, sizeof(int)) != Success)
+ return 1;
+
+ int Initial = 0;
+ int Result = 0;
+ if (Memcpy(In, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
+ return 1;
+ if (Memcpy(Out, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
+ return 1;
+
+ delayedSetValue<<<1, 1, 0, BlockingStream>>>(In, 99);
+ copyValue<<<1, 1>>>(In, Out);
+ if (Memcpy(&Result, Out, sizeof(int), MemcpyDeviceToHost) != Success)
+ 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 (Memcpy(Out, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
+ return 1;
+
+ delayedSetValue<<<1, 1>>>(In, 123);
+ copyValue<<<1, 1, 0, BlockingStream>>>(In, Out);
+ if (StreamSynchronize(BlockingStream) != Success)
+ return 1;
+ if (Memcpy(&Result, Out, sizeof(int), MemcpyDeviceToHost) != Success)
+ 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 (Memcpy(In, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
+ return 1;
+ if (Memcpy(Out, &Initial, sizeof(int), MemcpyHostToDevice) != Success)
+ return 1;
+ 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 (StreamSynchronize(NonBlockingStream) != Success)
+ return 1;
+ 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 (StreamDestroy(BlockingStream) != Success)
+ return 1;
+ if (StreamDestroy(NonBlockingStream) != Success)
+ return 1;
+ if (Free(In) != Success)
+ return 1;
+ if (Free(Out) != Success)
+ return 1;
+ 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 53%
rename from offload/test/offloading/HIP/device_api.hip
rename to offload/test/offloading/language/device_api.cpp
index 5fb66e6e45eeb..27e09deda0d5e 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
@@ -12,34 +16,37 @@
// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
+// clang-format off
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
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 57%
rename from offload/test/offloading/HIP/device_properties.hip
rename to offload/test/offloading/language/device_properties.cpp
index 1a9b9a70f8ea9..2b3ea3004526c 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
@@ -12,13 +16,16 @@
// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
+// clang-format off
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
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 50%
rename from offload/test/offloading/HIP/devicesync_streams.hip
rename to offload/test/offloading/language/devicesync_streams.cpp
index 14ac798f56ea8..617a077363ae5 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
@@ -12,9 +16,12 @@
// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
+// clang-format off
#include <chrono>
#include <cstdio>
#include <thread>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
__global__ void waitThenSet(volatile int *Gate, volatile int *Out, int Value) {
for (unsigned long long I = 0; I < 1000000000ULL && *Gate == 0; ++I)
@@ -23,29 +30,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 +68,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 +78,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..deb4047b4502f
--- /dev/null
+++ b/offload/test/offloading/language/error_kinds.cpp
@@ -0,0 +1,97 @@
+// 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
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+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..f9411397841aa
--- /dev/null
+++ b/offload/test/offloading/language/get_errs.cpp
@@ -0,0 +1,102 @@
+// 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
+
+// clang-format off
+#include <cstdio>
+#include <mutex>
+#include <thread>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+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..55e66eb43b0b1
--- /dev/null
+++ b/offload/test/offloading/language/host_alloc.cpp
@@ -0,0 +1,54 @@
+// 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
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+__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..2e52d338346af
--- /dev/null
+++ b/offload/test/offloading/language/launch_tu.cpp
@@ -0,0 +1,37 @@
+// 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
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+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..4be36740855ef
--- /dev/null
+++ b/offload/test/offloading/language/memcpy_kinds.cpp
@@ -0,0 +1,54 @@
+// 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
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+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 51%
rename from offload/test/offloading/CUDA/memset.cu
rename to offload/test/offloading/language/memset.cpp
index 44a5aeba8dcdf..17756d57f435f 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
@@ -12,9 +16,12 @@
// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
+// clang-format off
#include <chrono>
#include <cstdio>
#include <thread>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
__global__ void waitThenSet(volatile int *Gate, unsigned char *Out,
unsigned char Value) {
@@ -26,15 +33,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 +48,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 +67,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..98fde30e3fd08
--- /dev/null
+++ b/offload/test/offloading/language/stream_api.cpp
@@ -0,0 +1,110 @@
+// 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
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+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 54%
rename from offload/test/offloading/HIP/syncthreads.hip
rename to offload/test/offloading/language/syncthreads.cpp
index 81e9ed0451ca1..8563280bf9148 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
@@ -12,7 +16,10 @@
// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
+// clang-format off
#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
__global__ void reduceBlock(int *Out) {
__shared__ int Scratch[64];
@@ -31,10 +38,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 59%
rename from offload/test/offloading/CUDA/thread_and_block_id.cu
rename to offload/test/offloading/language/thread_and_block_id.cpp
index a56c9ff33e2ab..8b79a5ad98334 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
@@ -14,8 +18,11 @@
// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
+// clang-format off
#include <stdio.h>
#include <stdlib.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
__global__ void fill(int *A) {
int tid = threadIdx.x + blockDim.x * blockIdx.x;
@@ -28,13 +35,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