HIP: record direct atomic PSF timings

src/frame.c

@@ -134,6 +134,7 @@ typedef struct {
 #ifdef PSF_BACKEND_HIP
   HipPsfSink *hip;
   char hip_message[256];
+  HipPsfTiming hip_timing;
 #endif
   int failed;
 } PsfEventSink;
@@ -227,6 +228,10 @@ static int psf_event_sink_resume_gpu(PsfEventSink *sink) {
 static int psf_event_sink_destroy(PsfEventSink *sink) {
   int result = psf_event_sink_finish_for_cpu(sink);
 #ifdef PSF_BACKEND_HIP
+  if (sink->hip != NULL && hip_psf_sink_get_timing(sink->hip, &sink->hip_timing)) {
+    psf_event_sink_mark_failed(sink, "timing collection");
+    result = -1;
+  }
   hip_psf_sink_destroy(sink->hip);
 #endif
   free(sink->events);
@@ -998,6 +1003,8 @@ typedef struct {
   size_t direct_fallbacks;
   size_t cached_wing_clipped;
   size_t discarded_below_min_y;
+  size_t gpu_event_count, gpu_batch_count, gpu_timed_batch_count;
+  double gpu_upload_seconds, gpu_kernel_seconds, gpu_download_seconds;
   int failed;
 #ifdef GR_DEBUG
   double max_raw_magnification;
@@ -1016,6 +1023,12 @@ static void copy_psf_splat_stats(PsfSplatStats *destination,
                                  .cached_wing_clipped = source.cached_wing_clipped,
                                  .direct_fallbacks = source.direct_fallbacks,
                                  .discarded_below_min_y = source.discarded_below_min_y,
+                                 .gpu_event_count = source.gpu_event_count,
+                                 .gpu_batch_count = source.gpu_batch_count,
+                                 .gpu_timed_batch_count = source.gpu_timed_batch_count,
+                                 .gpu_upload_seconds = source.gpu_upload_seconds,
+                                 .gpu_kernel_seconds = source.gpu_kernel_seconds,
+                                 .gpu_download_seconds = source.gpu_download_seconds,
 #ifdef GR_DEBUG
                                  .max_raw_magnification = source.max_raw_magnification,
                                  .magnification_clamped_triangles =
@@ -1162,8 +1175,18 @@ static CatalogSplatStats splat_catalog_triangles(
     stats.cached_wing_clipped += context.cached_wing_clipped;
     stats.discarded_below_min_y += context.discarded_below_min_y;
   }
-  if (owns_sink && psf_event_sink_destroy(&owned_sink))
-    stats.failed = 1;
+  if (owns_sink) {
+    if (psf_event_sink_destroy(&owned_sink))
+      stats.failed = 1;
+#ifdef PSF_BACKEND_HIP
+    stats.gpu_event_count = owned_sink.hip_timing.event_count;
+    stats.gpu_batch_count = owned_sink.hip_timing.batch_count;
+    stats.gpu_timed_batch_count = owned_sink.hip_timing.timed_batch_count;
+    stats.gpu_upload_seconds = owned_sink.hip_timing.upload_seconds;
+    stats.gpu_kernel_seconds = owned_sink.hip_timing.kernel_seconds;
+    stats.gpu_download_seconds = owned_sink.hip_timing.download_seconds;
+#endif
+  }
   return stats;
 }
 

src/hip_psf.h

@@ -13,6 +13,11 @@ extern "C" {
  * for PsfCachedEvent chunks, immutable cache weights, and double RGB HDR. */
 typedef struct HipPsfSink HipPsfSink;
 
+typedef struct {
+  size_t event_count, batch_count, timed_batch_count;
+  double upload_seconds, kernel_seconds, download_seconds;
+} HipPsfTiming;
+
 int hip_psf_available(char *message, size_t message_size);
 
 /* Creates a zeroed GPU HDR framebuffer. event_capacity is the reusable upload
@@ -35,6 +40,9 @@ int hip_psf_sink_submit(HipPsfSink *sink, const PsfCachedEvent *events,
 /* Completes all work and overwrites hdr with double linear RGB device HDR. */
 int hip_psf_sink_finish(HipPsfSink *sink, double *hdr, char *message,
                         size_t message_size);
+/* Valid after hip_psf_sink_finish(); timings are accumulated without a
+ * per-batch host synchronization. */
+int hip_psf_sink_get_timing(const HipPsfSink *sink, HipPsfTiming *timing);
 void hip_psf_sink_destroy(HipPsfSink *sink);
 
 #ifdef __cplusplus

src/hip_psf.hip

@@ -5,6 +5,14 @@
 #include <cmath>
 #include <cstdio>
 #include <limits>
+#include <vector>
+
+struct HipPsfBatchTiming {
+  hipEvent_t upload_start = nullptr, upload_end = nullptr;
+  hipEvent_t kernel_start = nullptr, kernel_end = nullptr;
+};
+
+enum { HIP_PSF_MAX_TIMED_BATCHES = 64 };
 
 struct HipPsfSink {
   PsfCachedEvent *events = nullptr;
@@ -14,6 +22,10 @@ struct HipPsfSink {
   size_t event_capacity = 0, hdr_values = 0;
   int width = 0, height = 0, phase_resolution = 0, radius_pixels = 0;
   double max_radius_pixels = 0.0;
+  hipEvent_t download_start = nullptr, download_end = nullptr;
+  std::vector<HipPsfBatchTiming> batches;
+  size_t measured_batches = 0;
+  HipPsfTiming timing = {};
 };
 
 static int report(hipError_t status, char *message, size_t message_size) {
@@ -25,6 +37,15 @@ static void ok(char *message, size_t message_size) {
   if (message && message_size) std::snprintf(message, message_size, "ok");
 }
 
+static void destroy_batch_timing(HipPsfBatchTiming *timing) {
+  if (!timing) return;
+  (void)hipEventDestroy(timing->upload_start);
+  (void)hipEventDestroy(timing->upload_end);
+  (void)hipEventDestroy(timing->kernel_start);
+  (void)hipEventDestroy(timing->kernel_end);
+  *timing = {};
+}
+
 __device__ static size_t weight_index(int phase_resolution, int radius_pixels,
                                       int phase_x, int phase_y, int offset_x, int offset_y) {
   const size_t nodes = (size_t)phase_resolution + 1;
@@ -118,6 +139,8 @@ extern "C" int hip_psf_sink_create(HipPsfSink **out, int width, int height,
   sink->radius_pixels = cache->radius_pixels; sink->max_radius_pixels = cache->max_radius_pixels;
   const size_t weight_count = nodes * nodes * side * side;
   if (report(hipStreamCreate(&sink->stream), message, message_size) ||
+      report(hipEventCreate(&sink->download_start), message, message_size) ||
+      report(hipEventCreate(&sink->download_end), message, message_size) ||
       report(hipMalloc(&sink->events, event_capacity * sizeof *sink->events), message, message_size) ||
       report(hipMalloc(&sink->weights, weight_count * sizeof *sink->weights), message, message_size) ||
       report(hipMalloc(&sink->hdr, sink->hdr_values * sizeof *sink->hdr), message, message_size) ||
@@ -137,13 +160,35 @@ extern "C" int hip_psf_sink_submit(HipPsfSink *sink, const PsfCachedEvent *event
   if (!event_count) { ok(message, message_size); return 0; }
   const unsigned int threads = 128;
   const size_t blocks = (event_count + threads - 1) / threads;
+  HipPsfBatchTiming timing;
+  const int timed = sink->batches.size() < HIP_PSF_MAX_TIMED_BATCHES;
   if (blocks > std::numeric_limits<unsigned int>::max() ||
-      report(hipMemcpyAsync(sink->events, events, event_count * sizeof *events, hipMemcpyHostToDevice, sink->stream), message, message_size))
+      (timed && (report(hipEventCreate(&timing.upload_start), message, message_size) ||
+                 report(hipEventCreate(&timing.upload_end), message, message_size) ||
+                 report(hipEventCreate(&timing.kernel_start), message, message_size) ||
+                 report(hipEventCreate(&timing.kernel_end), message, message_size) ||
+                 report(hipEventRecord(timing.upload_start, sink->stream), message, message_size))) ||
+      report(hipMemcpyAsync(sink->events, events, event_count * sizeof *events,
+                            hipMemcpyHostToDevice, sink->stream), message, message_size) ||
+      (timed && (report(hipEventRecord(timing.upload_end, sink->stream), message, message_size) ||
+                 report(hipEventRecord(timing.kernel_start, sink->stream), message, message_size)))) {
+    destroy_batch_timing(&timing);
     return -1;
+  }
   hipLaunchKernelGGL(splat_kernel, dim3((unsigned int)blocks), dim3(threads), 0, sink->stream,
                      sink->events, event_count, sink->width, sink->height, sink->weights,
                      sink->phase_resolution, sink->radius_pixels, sink->max_radius_pixels, sink->hdr);
-  if (report(hipGetLastError(), message, message_size)) return -1;
+  if (report(hipGetLastError(), message, message_size) ||
+      (timed && report(hipEventRecord(timing.kernel_end, sink->stream), message, message_size))) {
+    destroy_batch_timing(&timing);
+    return -1;
+  }
+  if (timed) {
+    sink->batches.push_back(timing);
+    ++sink->timing.timed_batch_count;
+  }
+  ++sink->timing.batch_count;
+  sink->timing.event_count += event_count;
   ok(message, message_size); return 0;
 }
 
@@ -161,13 +206,39 @@ extern "C" int hip_psf_sink_load_hdr(HipPsfSink *sink, const double *hdr,
 
 extern "C" int hip_psf_sink_finish(HipPsfSink *sink, double *hdr, char *message, size_t message_size) {
   if (!sink || !hdr) { if (message && message_size) std::snprintf(message, message_size, "invalid HIP PSF HDR download"); return -1; }
-  if (report(hipMemcpyAsync(hdr, sink->hdr, sink->hdr_values * sizeof *hdr, hipMemcpyDeviceToHost, sink->stream), message, message_size) ||
+  if (report(hipEventRecord(sink->download_start, sink->stream), message, message_size) ||
+      report(hipMemcpyAsync(hdr, sink->hdr, sink->hdr_values * sizeof *hdr, hipMemcpyDeviceToHost, sink->stream), message, message_size) ||
+      report(hipEventRecord(sink->download_end, sink->stream), message, message_size) ||
       report(hipStreamSynchronize(sink->stream), message, message_size)) return -1;
+  float milliseconds = 0.0f;
+  for (size_t i = sink->measured_batches; i < sink->batches.size(); ++i) {
+    const HipPsfBatchTiming &batch = sink->batches[i];
+    if (report(hipEventElapsedTime(&milliseconds, batch.upload_start, batch.upload_end),
+               message, message_size)) return -1;
+    sink->timing.upload_seconds += milliseconds * 1e-3;
+    if (report(hipEventElapsedTime(&milliseconds, batch.kernel_start, batch.kernel_end),
+               message, message_size)) return -1;
+    sink->timing.kernel_seconds += milliseconds * 1e-3;
+  }
+  sink->measured_batches = sink->batches.size();
+  if (report(hipEventElapsedTime(&milliseconds, sink->download_start, sink->download_end),
+             message, message_size)) return -1;
+  sink->timing.download_seconds += milliseconds * 1e-3;
   ok(message, message_size); return 0;
 }
 
+extern "C" int hip_psf_sink_get_timing(const HipPsfSink *sink, HipPsfTiming *timing) {
+  if (!sink || !timing) return -1;
+  *timing = sink->timing;
+  return 0;
+}
+
 extern "C" void hip_psf_sink_destroy(HipPsfSink *sink) {
   if (!sink) return;
+  for (HipPsfBatchTiming &timing : sink->batches)
+    destroy_batch_timing(&timing);
+  (void)hipEventDestroy(sink->download_start);
+  (void)hipEventDestroy(sink->download_end);
   (void)hipFree(sink->events); (void)hipFree(sink->weights); (void)hipFree(sink->hdr);
   (void)hipStreamDestroy(sink->stream); delete sink;
 }

src/optics.c

@@ -301,6 +301,13 @@ void psf_kernel_cache_report(const PsfKernelCache *cache,
           stats == NULL ? 0u : stats->cached_wing_clipped,
           stats == NULL ? 0u : stats->direct_fallbacks,
           stats == NULL ? 0u : stats->discarded_below_min_y);
+  if (stats != NULL && stats->gpu_batch_count != 0)
+    fprintf(stream, "HIP PSF: %zu events in %zu batches (%zu timed); upload %.6f s, "
+                    "kernel %.6f s, download %.6f s\n",
+            stats->gpu_event_count, stats->gpu_batch_count,
+            stats->gpu_timed_batch_count,
+            stats->gpu_upload_seconds, stats->gpu_kernel_seconds,
+            stats->gpu_download_seconds);
   if (stats != NULL && stats->discarded_below_min_y != 0)
     fputs("Warning: --psf-min-y discarded one or more PSF events.\n", stream);
 #ifdef GR_DEBUG

src/optics.h

@@ -26,6 +26,8 @@ typedef struct {
   size_t cached_wing_clipped;
   size_t direct_fallbacks;
   size_t discarded_below_min_y;
+  size_t gpu_event_count, gpu_batch_count, gpu_timed_batch_count;
+  double gpu_upload_seconds, gpu_kernel_seconds, gpu_download_seconds;
 #ifdef GR_DEBUG
   double max_raw_magnification;
   size_t magnification_clamped_triangles;