[llvm-branch-commits] [llvm] [Offload][Lang] Unify Tests (PR #217722)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Aug 20 11:29:10 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: Sophia Herrmann (jellytabby)
<details>
<summary>Changes</summary>
We need tests to ensure the parsing and computing capability of both cuda and hip flavored inputs. Previously these were separate but same tests, with this PR we follow the same structure as the LanguageRuntime.h files and use macro renaming to generate cuda/hip tests from one common base file. The new tests are in `offload/test/offloading/language/`.
Assisted by GPT-5.5, checked and reviewed manually
---
Patch is 102.04 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/217722.diff
42 Files Affected:
- (removed) offload/test/offloading/CUDA/basic_launch.cu (-30)
- (removed) offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu (-32)
- (removed) offload/test/offloading/CUDA/basic_launch_multi_arg.cu (-39)
- (removed) offload/test/offloading/CUDA/blocking_stream_semantics.cu (-131)
- (removed) offload/test/offloading/CUDA/device_api.cu (-45)
- (removed) offload/test/offloading/CUDA/device_properties.cu (-40)
- (removed) offload/test/offloading/CUDA/devicesync_streams.cu (-98)
- (removed) offload/test/offloading/CUDA/error_kinds.cu (-81)
- (removed) offload/test/offloading/CUDA/get_errs.cu (-82)
- (removed) offload/test/offloading/CUDA/host_alloc.cu (-48)
- (removed) offload/test/offloading/CUDA/launch_tu.cu (-30)
- (removed) offload/test/offloading/CUDA/memcpy_kinds.cu (-51)
- (removed) offload/test/offloading/CUDA/stream_api.cu (-100)
- (removed) offload/test/offloading/CUDA/syncthreads.cu (-41)
- (removed) offload/test/offloading/HIP/basic_launch.hip (-30)
- (removed) offload/test/offloading/HIP/basic_launch_multi_arg.hip (-39)
- (removed) offload/test/offloading/HIP/error_kinds.hip (-81)
- (removed) offload/test/offloading/HIP/get_errs.hip (-81)
- (removed) offload/test/offloading/HIP/host_alloc.hip (-48)
- (removed) offload/test/offloading/HIP/kernel_tu.hip.inc (-1)
- (removed) offload/test/offloading/HIP/memcpy_kinds.hip (-51)
- (removed) offload/test/offloading/HIP/memset.hip (-85)
- (removed) offload/test/offloading/HIP/stream_api.hip (-99)
- (removed) offload/test/offloading/HIP/thread_and_block_id.hip (-46)
- (added) offload/test/offloading/language/Inputs/DefineTestLanguageNames.inc (+18)
- (renamed) offload/test/offloading/language/Inputs/launch_tu_kernel.inc ()
- (renamed) offload/test/offloading/language/basic_launch.cpp (+14-9)
- (renamed) offload/test/offloading/language/basic_launch_blocks_and_threads.cpp (+13-8)
- (added) offload/test/offloading/language/basic_launch_multi_arg.cpp (+44)
- (renamed) offload/test/offloading/language/blocking_stream_semantics.cpp (+36-34)
- (renamed) offload/test/offloading/language/device_api.cpp (+14-9)
- (renamed) offload/test/offloading/language/device_properties.cpp (+13-8)
- (renamed) offload/test/offloading/language/devicesync_streams.cpp (+29-28)
- (added) offload/test/offloading/language/error_kinds.cpp (+95)
- (added) offload/test/offloading/language/get_errs.cpp (+100)
- (added) offload/test/offloading/language/host_alloc.cpp (+52)
- (added) offload/test/offloading/language/launch_tu.cpp (+35)
- (added) offload/test/offloading/language/memcpy_kinds.cpp (+52)
- (renamed) offload/test/offloading/language/memset.cpp (+23-20)
- (added) offload/test/offloading/language/stream_api.cpp (+108)
- (renamed) offload/test/offloading/language/syncthreads.cpp (+12-7)
- (renamed) offload/test/offloading/language/thread_and_block_id.cpp (+13-8)
``````````diff
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();
-...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/217722
More information about the llvm-branch-commits
mailing list