================
@@ -0,0 +1,82 @@
+// 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 thread_id) {
+ const int N = 1 << 20;
+ size_t bytes = N * sizeof(float);
+
+ float *d_data;
+ cudaMalloc(&d_data, bytes);
+ printError(thread_id, "cudaMalloc", cudaGetLastError());
+
+ thread_id == 1 ? errorKernel<<<4096, 256>>>(d_data)
+ : errorKernel<<<4096, 0>>>(d_data);
+ printError(thread_id, "kernel launch", cudaPeekAtLastError());
+
+ printError(thread_id, "kernel launch get", cudaGetLastError());
+
+ printError(thread_id, "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
----------------
jdoerfert wrote:
If the cuda and hip response is basically the same, we should, at least in the
tests, just check for the important parts. So Success, not cudaSuccess. Now we
duplicate the tests for the prefix to be different only.
https://github.com/llvm/llvm-project/pull/213390
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
