commit 0a0483e997bea993cdf26e62ef5c68a043011e1b
author Andrew Woloszyn <andrew.woloszyn@gmail.com> Fri Dec 20 12:48:40 2024 -0800
committer GitHub <noreply@github.com> Fri Dec 20 20:48:40 2024 +0000
tree ba94d6ec488f11b2b639f40216df2be797f9664e
parent 47ccd93ca3f1a55f08e983e92f5960a8d26f577c

[hip] Add trace zones to copy/fill/update buffer commands. (#19544)

They were missing in the hip stream case.

Signed-off-by: Andrew Woloszyn <andrew.woloszyn@gmail.com>

runtime/src/iree/hal/drivers/hip/stream_command_buffer.c

diff --git a/runtime/src/iree/hal/drivers/hip/stream_command_buffer.c b/runtime/src/iree/hal/drivers/hip/stream_command_buffer.c
index d0dcea1..6ccc93c 100644
--- a/runtime/src/iree/hal/drivers/hip/stream_command_buffer.c
+++ b/runtime/src/iree/hal/drivers/hip/stream_command_buffer.c
@@ -328,6 +328,9 @@
       iree_hal_buffer_byte_offset(target_ref.buffer) + target_ref.offset;
   hipDeviceptr_t dst = (uint8_t*)target_device_buffer + target_offset;
   size_t num_elements = target_ref.length / pattern_length;
+  IREE_HAL_STREAM_TRACE_ZONE_BEGIN(command_buffer->tracing_context,
+                                   &command_buffer->tracing_event_list,
+                                   IREE_HAL_STREAM_TRACING_VERBOSITY_FINE);
 
   switch (pattern_length) {
     case 4: {
@@ -359,7 +362,9 @@
       return iree_make_status(IREE_STATUS_INTERNAL,
                               "unsupported fill pattern length");
   }
-
+  IREE_HAL_STREAM_TRACE_ZONE_END(command_buffer->tracing_context,
+                                 &command_buffer->tracing_event_list,
+                                 IREE_HAL_STREAM_TRACING_VERBOSITY_FINE);
   IREE_TRACE_ZONE_END(z0);
   return iree_ok_status();
 }
@@ -397,12 +402,17 @@
   hipDeviceptr_t dst = (uint8_t*)target_device_buffer +
                        iree_hal_buffer_byte_offset(target_ref.buffer) +
                        target_ref.offset;
+  IREE_HAL_STREAM_TRACE_ZONE_BEGIN(command_buffer->tracing_context,
+                                   &command_buffer->tracing_event_list,
+                                   IREE_HAL_STREAM_TRACING_VERBOSITY_FINE);
   IREE_HIP_RETURN_AND_END_ZONE_IF_ERROR(
       z0, command_buffer->hip_symbols,
       hipMemcpyHtoDAsync(dst, (void*)src, target_ref.length,
                          command_buffer->hip_stream),
       "hipMemcpyHtoDAsync");
-
+  IREE_HAL_STREAM_TRACE_ZONE_END(command_buffer->tracing_context,
+                                 &command_buffer->tracing_event_list,
+                                 IREE_HAL_STREAM_TRACING_VERBOSITY_FINE);
   IREE_TRACE_ZONE_END(z0);
   return iree_ok_status();
 }
@@ -417,6 +427,9 @@
 
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_hal_hip_stream_command_buffer_flush_collectives(command_buffer));
+  IREE_HAL_STREAM_TRACE_ZONE_BEGIN(command_buffer->tracing_context,
+                                   &command_buffer->tracing_event_list,
+                                   IREE_HAL_STREAM_TRACING_VERBOSITY_FINE);
 
   hipDeviceptr_t target_device_buffer = iree_hal_hip_buffer_device_pointer(
       iree_hal_buffer_allocated_buffer(target_ref.buffer));
@@ -435,6 +448,10 @@
                      command_buffer->hip_stream),
       "hipMemcpyAsync");
 
+  IREE_HAL_STREAM_TRACE_ZONE_END(command_buffer->tracing_context,
+                                 &command_buffer->tracing_event_list,
+                                 IREE_HAL_STREAM_TRACING_VERBOSITY_FINE);
+
   IREE_TRACE_ZONE_END(z0);
   return iree_ok_status();
 }