From 591291679e8c03a515be48b2d6a575c8ae8d92eb Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Tue, 17 Jun 2025 12:25:58 +0200
Subject: [PATCH] Add device-side exception reporting.

---
 src/device/runtime.jl | 21 ++++++++++++++++++---
 1 file changed, 18 insertions(+), 3 deletions(-)

diff --git a/src/device/runtime.jl b/src/device/runtime.jl
index b2a78a29..76309f31 100644
--- a/src/device/runtime.jl
+++ b/src/device/runtime.jl
@@ -7,8 +7,23 @@ malloc(sz) = C_NULL
 
 report_oom(sz) = return
 
-report_exception(ex) = return
+function report_exception(ex)
+    @printf(
+        "ERROR: a %s was thrown during kernel execution on thread (%d, %d, %d).\n",
+        ex, get_global_id(UInt32(0)), get_global_id(UInt32(1)), get_global_id(UInt32(2))
+    )
+    return
+end
 
-report_exception_name(ex) = return
+function report_exception_name(ex)
+    @printf(
+        "ERROR: a %s was thrown during kernel execution on thread (%d, %d, %d).\nStacktrace:\n",
+        ex, get_global_id(UInt32(0)), get_global_id(UInt32(1)), get_global_id(UInt32(2))
+    )
+    return
+end
 
-report_exception_frame(idx, func, file, line) = return
+function report_exception_frame(idx, func, file, line)
+    @printf(" [%d] %s at %s:%d\n", idx, func, file, line)
+    return
+end