[clang] [llvm] [LLVMOffload] Get LastError for bad Kernels (PR #213390)

[Johannes Doerfert via llvm-commits]

================
@@ -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