[flang-commits] [flang] [llvm] [flang][cuda][NFC] Add filename and line number in error reporting (PR #185516)

[Valentin Clement バレンタイン クレメン via flang-commits]

https://github.com/clementval created https://github.com/llvm/llvm-project/pull/185516

Some entry points carry over filename and line number for error reporting. Use this information when reporting cuda error. 

>From 6d7e29c85c521ff7444d5534618c6609e565c8a2 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Mon, 9 Mar 2026 13:50:56 -0700
Subject: [PATCH] [flang][cuda] Add filename and line number in error reporting

---
 flang-rt/lib/cuda/descriptor.cpp          |  8 +++++---
 flang-rt/lib/cuda/memory.cpp              | 23 ++++++++++++++---------
 flang/include/flang/Runtime/CUDA/common.h | 11 +++++++++++
 3 files changed, 30 insertions(+), 12 deletions(-)

diff --git a/flang-rt/lib/cuda/descriptor.cpp b/flang-rt/lib/cuda/descriptor.cpp
index aa75d4eff0511..2cf795181ea7b 100644
--- a/flang-rt/lib/cuda/descriptor.cpp
+++ b/flang-rt/lib/cuda/descriptor.cpp
@@ -33,7 +33,8 @@ void *RTDEF(CUFGetDeviceAddress)(
     void *hostPtr, const char *sourceFile, int sourceLine) {
   Terminator terminator{sourceFile, sourceLine};
   void *p;
-  CUDA_REPORT_IF_ERROR(cudaGetSymbolAddress((void **)&p, hostPtr));
+  CUDA_REPORT_IF_ERROR_LOC(
+      cudaGetSymbolAddress((void **)&p, hostPtr), sourceFile, sourceLine);
   if (!p) {
     terminator.Crash("Could not retrieve symbol's address");
   }
@@ -43,8 +44,9 @@ void *RTDEF(CUFGetDeviceAddress)(
 void RTDEF(CUFDescriptorSync)(Descriptor *dst, const Descriptor *src,
     const char *sourceFile, int sourceLine) {
   std::size_t count{src->SizeInBytes()};
-  CUDA_REPORT_IF_ERROR(cudaMemcpy(
-      (void *)dst, (const void *)src, count, cudaMemcpyHostToDevice));
+  CUDA_REPORT_IF_ERROR_LOC(
+      cudaMemcpy((void *)dst, (const void *)src, count, cudaMemcpyHostToDevice),
+      sourceFile, sourceLine);
 }
 
 void RTDEF(CUFSyncGlobalDescriptor)(
diff --git a/flang-rt/lib/cuda/memory.cpp b/flang-rt/lib/cuda/memory.cpp
index 78270fef07c36..494703bc09e74 100644
--- a/flang-rt/lib/cuda/memory.cpp
+++ b/flang-rt/lib/cuda/memory.cpp
@@ -28,16 +28,20 @@ void *RTDEF(CUFMemAlloc)(
   bytes = bytes ? bytes : 1;
   if (type == kMemTypeDevice) {
     if (Fortran::runtime::executionEnvironment.cudaDeviceIsManaged) {
-      CUDA_REPORT_IF_ERROR(
-          cudaMallocManaged((void **)&ptr, bytes, cudaMemAttachGlobal));
+      CUDA_REPORT_IF_ERROR_LOC(
+          cudaMallocManaged((void **)&ptr, bytes, cudaMemAttachGlobal),
+          sourceFile, sourceLine);
     } else {
-      CUDA_REPORT_IF_ERROR(cudaMalloc((void **)&ptr, bytes));
+      CUDA_REPORT_IF_ERROR_LOC(
+          cudaMalloc((void **)&ptr, bytes), sourceFile, sourceLine);
     }
   } else if (type == kMemTypeManaged || type == kMemTypeUnified) {
-    CUDA_REPORT_IF_ERROR(
-        cudaMallocManaged((void **)&ptr, bytes, cudaMemAttachGlobal));
+    CUDA_REPORT_IF_ERROR_LOC(
+        cudaMallocManaged((void **)&ptr, bytes, cudaMemAttachGlobal),
+        sourceFile, sourceLine);
   } else if (type == kMemTypePinned) {
-    CUDA_REPORT_IF_ERROR(cudaMallocHost((void **)&ptr, bytes));
+    CUDA_REPORT_IF_ERROR_LOC(
+        cudaMallocHost((void **)&ptr, bytes), sourceFile, sourceLine);
   } else {
     Terminator terminator{sourceFile, sourceLine};
     terminator.Crash("unsupported memory type");
@@ -51,9 +55,9 @@ void RTDEF(CUFMemFree)(
     return;
   if (type == kMemTypeDevice || type == kMemTypeManaged ||
       type == kMemTypeUnified) {
-    CUDA_REPORT_IF_ERROR(cudaFree(ptr));
+    CUDA_REPORT_IF_ERROR_LOC(cudaFree(ptr), sourceFile, sourceLine);
   } else if (type == kMemTypePinned) {
-    CUDA_REPORT_IF_ERROR(cudaFreeHost(ptr));
+    CUDA_REPORT_IF_ERROR_LOC(cudaFreeHost(ptr), sourceFile, sourceLine);
   } else {
     Terminator terminator{sourceFile, sourceLine};
     terminator.Crash("unsupported memory type");
@@ -81,7 +85,8 @@ void RTDEF(CUFDataTransferPtrPtr)(void *dst, void *src, std::size_t bytes,
     terminator.Crash("host to host copy not supported");
   }
   // TODO: Use cudaMemcpyAsync when we have support for stream.
-  CUDA_REPORT_IF_ERROR(cudaMemcpy(dst, src, bytes, kind));
+  CUDA_REPORT_IF_ERROR_LOC(
+      cudaMemcpy(dst, src, bytes, kind), sourceFile, sourceLine);
 }
 
 void RTDEF(CUFDataTransferPtrDesc)(void *addr, Descriptor *desc,
diff --git a/flang/include/flang/Runtime/CUDA/common.h b/flang/include/flang/Runtime/CUDA/common.h
index 474f8e6578b89..36e47b0c13655 100644
--- a/flang/include/flang/Runtime/CUDA/common.h
+++ b/flang/include/flang/Runtime/CUDA/common.h
@@ -34,4 +34,15 @@ static constexpr unsigned kDeviceToDevice = 2;
     terminator.Crash("'%s' failed with '%s'", #expr, name); \
   }(expr)
 
+#define CUDA_REPORT_IF_ERROR_LOC(expr, file, line) \
+  [](cudaError_t err, const char *sourceFile, int sourceLine) { \
+    if (err == cudaSuccess) \
+      return; \
+    const char *name = cudaGetErrorName(err); \
+    if (!name) \
+      name = "<unknown>"; \
+    Fortran::runtime::Terminator terminator{sourceFile, sourceLine}; \
+    terminator.Crash("'%s' failed with '%s'", #expr, name); \
+  }(expr, sourceFile, sourceLine)
+
 #endif // FORTRAN_RUNTIME_CUDA_COMMON_H_