[llvm] [OpenMP][offload] Add LIBOMPTARGET_KERNEL_EXE_TIME (PR #222105)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 8 12:11:23 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Robert Imschweiler (ro-i)
<details>
<summary>Changes</summary>
If set to 1, prints kernel-specific execution time.
This is a port of the original downstream ROCm patch, extended with the inclusion of the kernel name in the printed line.
Claude assisted with this patch.
---
Full diff: https://github.com/llvm/llvm-project/pull/222105.diff
3 Files Affected:
- (modified) offload/plugins-nextgen/amdgpu/src/rtl.cpp (+81-2)
- (modified) offload/plugins-nextgen/common/include/PluginInterface.h (+19)
- (added) offload/test/env/kernel_exe_time.c (+24)
``````````diff
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index 281b9e3795a54..6358541e0abbd 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -12,6 +12,7 @@
#include <atomic>
#include <cassert>
+#include <cinttypes>
#include <cstddef>
#include <cstdint>
#include <deque>
@@ -1040,6 +1041,19 @@ struct AMDGPUStreamTy {
AMDGPUSignalManagerTy *SignalManager;
};
+ /// Utility struct holding arguments for tracing a kernel's duration.
+ struct KernelDurationTracingArgsTy {
+ hsa_agent_t Agent;
+ AMDGPUSignalTy *Signal;
+ double TicksToTime;
+ int32_t DeviceId;
+ uint32_t LaunchId;
+ uint32_t NumTeams;
+ uint32_t NumThreads;
+ /// Owned by the kernel, which outlives the launch this traces.
+ const char *Name;
+ };
+
using AMDGPUStreamCallbackTy = Error(void *Data);
/// The stream is composed of N stream's slots. The struct below represents
@@ -1067,6 +1081,7 @@ struct AMDGPUStreamTy {
MemcpyArgsTy MemcpyArgs;
ReleaseBufferArgsTy ReleaseBufferArgs;
ReleaseSignalArgsTy ReleaseSignalArgs;
+ KernelDurationTracingArgsTy KernelDurationTracingArgs;
void *CallbackArgs;
};
@@ -1104,6 +1119,18 @@ struct AMDGPUStreamTy {
return Plugin::success();
}
+ /// Schedule a kernel duration tracing action on the slot.
+ Error schedKernelDurationTracing(hsa_agent_t Agent, AMDGPUSignalTy *Signal,
+ double TicksToTime, int32_t DeviceId,
+ uint32_t LaunchId, uint32_t NumTeams,
+ uint32_t NumThreads, const char *Name) {
+ Callbacks.emplace_back(kernelDurationTracingAction);
+ ActionArgs.emplace_back().KernelDurationTracingArgs =
+ KernelDurationTracingArgsTy{Agent, Signal, TicksToTime, DeviceId,
+ LaunchId, NumTeams, NumThreads, Name};
+ return Plugin::success();
+ }
+
/// Register a callback to be called on compleition
Error schedCallback(AMDGPUStreamCallbackTy *Func, void *Data) {
Callbacks.emplace_back(Func);
@@ -1129,6 +1156,9 @@ struct AMDGPUStreamTy {
} else if (Callback == releaseSignalAction) {
if (auto Err = releaseSignalAction(&ActionArg))
return Err;
+ } else if (Callback == kernelDurationTracingAction) {
+ if (auto Err = kernelDurationTracingAction(&ActionArg))
+ return Err;
} else if (Callback) {
if (auto Err = Callback(ActionArg.CallbackArgs))
return Err;
@@ -1146,6 +1176,10 @@ struct AMDGPUStreamTy {
/// The device agent where the stream was created.
hsa_agent_t Agent;
+ /// Factor converting HSA clock ticks into nanoseconds. Only used when kernel
+ /// duration tracing is enabled.
+ double TicksToTime;
+
/// The queue that the stream uses to launch kernels.
AMDGPUQueueTy *Queue;
@@ -1383,6 +1417,32 @@ struct AMDGPUStreamTy {
return Plugin::success();
}
+ /// Report the duration of a completed kernel. This is a post completion
+ /// action, taken while the kernel's signal is still alive.
+ static Error kernelDurationTracingAction(void *Data) {
+ KernelDurationTracingArgsTy *Args =
+ reinterpret_cast<KernelDurationTracingArgsTy *>(Data);
+ assert(Args && "Invalid arguments");
+ assert(Args->Signal && "Invalid signal");
+
+ hsa_amd_profiling_dispatch_time_t TimeRec = {};
+ hsa_status_t Status = hsa_amd_profiling_get_dispatch_time(
+ Args->Agent, Args->Signal->get(), &TimeRec);
+ if (auto Err = Plugin::check(
+ Status, "error in hsa_amd_profiling_get_dispatch_time: %s"))
+ return Err;
+
+ uint64_t Duration = static_cast<uint64_t>(
+ static_cast<double>(TimeRec.end - TimeRec.start) * Args->TicksToTime);
+
+ INFO_MESSAGE(
+ Args->DeviceId,
+ "LaunchID: %2u TeamsXthrds:(%4uX%4u) Duration(ns): %" PRIu64 " n:%s\n",
+ Args->LaunchId, Args->NumTeams, Args->NumThreads, Duration, Args->Name);
+
+ return Plugin::success();
+ }
+
public:
/// Create an empty stream associated with a specific device.
AMDGPUStreamTy(AMDGPUDeviceTy &Device);
@@ -1421,6 +1481,15 @@ struct AMDGPUStreamTy {
if (auto Err = Slots[Curr].schedReleaseBuffer(KernelArgs, MemoryManager))
return Err;
+ // When LIBOMPTARGET_KERNEL_EXE_TIME is set, register a post action that
+ // reports the kernel's duration once it has completed.
+ if (Device.enableKernelDurationTracing())
+ if (auto Err = Slots[Curr].schedKernelDurationTracing(
+ Agent, OutputSignal, TicksToTime, Device.getDeviceId(),
+ Device.getAndIncrementLaunchId(), NumBlocks[0], NumThreads[0],
+ Kernel.getName()))
+ return Err;
+
// If we are running an RPC server we want to wake up the server thread
// whenever there is a kernel running and let it sleep otherwise.
if (Device.getRPCServer())
@@ -3813,9 +3882,19 @@ Error AMDGPUResourceRef<ResourceTy>::create(GenericDeviceTy &Device) {
return Resource->init();
}
+/// Compute the factor converting the device's clock ticks into nanoseconds.
+/// Returns zero if the system timestamp frequency is unavailable, in which case
+/// all reported durations are zero.
+static double getTicksToTime(AMDGPUDeviceTy &Device) {
+ uint64_t TicksPerSecond = Device.getSystemTimestampFrequency();
+ if (TicksPerSecond == 0)
+ return 0.0;
+ return 1e9 / static_cast<double>(TicksPerSecond);
+}
+
AMDGPUStreamTy::AMDGPUStreamTy(AMDGPUDeviceTy &Device)
- : Agent(Device.getAgent()), Queue(nullptr),
- SignalManager(Device.getSignalManager()), Device(Device),
+ : Agent(Device.getAgent()), TicksToTime(getTicksToTime(Device)),
+ Queue(nullptr), SignalManager(Device.getSignalManager()), Device(Device),
// Initialize the std::deque with some empty positions.
Slots(32), NextSlot(0), SyncCycle(0),
StreamBusyWaitMicroseconds(Device.getStreamBusyWaitMicroseconds()),
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 29513661867b1..463b1a04549ed 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -11,6 +11,7 @@
#ifndef OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_PLUGININTERFACE_H
#define OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_PLUGININTERFACE_H
+#include <atomic>
#include <cstddef>
#include <cstdint>
#include <deque>
@@ -1258,6 +1259,16 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
return OMPX_ReuseBlocksForHighTripCount;
}
+ /// Whether the user asked for a trace of every kernel launch's duration.
+ /// @see OMPX_KernelDurationTracing
+ bool enableKernelDurationTracing() const {
+ return OMPX_KernelDurationTracing;
+ }
+
+ /// Get a unique, monotonically increasing identifier for the next kernel
+ /// launch on this device.
+ uint32_t getAndIncrementLaunchId() { return LaunchId.fetch_add(1); }
+
/// Get the total amount of hardware parallelism supported by the target
/// device. This is the total amount of warps or wavefronts that can be
/// resident on the device simultaneously.
@@ -1446,6 +1457,14 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
BoolEnvar OMPX_ReuseBlocksForHighTripCount =
BoolEnvar("LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT", true);
+ /// Environment variable to trace the duration of every kernel launch.
+ BoolEnvar OMPX_KernelDurationTracing =
+ BoolEnvar("LIBOMPTARGET_KERNEL_EXE_TIME", false);
+
+ /// Counter handing out a unique identifier to each kernel launch on this
+ /// device.
+ std::atomic<uint32_t> LaunchId{0};
+
/// Indicate whether mapped host buffers should be locked automatically.
bool LockMappedBuffers;
diff --git a/offload/test/env/kernel_exe_time.c b/offload/test/env/kernel_exe_time.c
new file mode 100644
index 0000000000000..04e14e7acc6e4
--- /dev/null
+++ b/offload/test/env/kernel_exe_time.c
@@ -0,0 +1,24 @@
+// Check that LIBOMPTARGET_KERNEL_EXE_TIME reports one trace line per kernel
+// launch, carrying a monotonically increasing launch id and the name of the
+// kernel that was launched.
+//
+// RUN: %libomptarget-compile-generic && \
+// RUN: env LIBOMPTARGET_KERNEL_EXE_TIME=1 %libomptarget-run-generic 2>&1 | \
+// RUN: %fcheck-generic
+//
+// REQUIRES: amdgpu
+
+int main(void) {
+ int X = 0;
+
+#pragma omp target map(tofrom : X)
+ X = 1;
+
+#pragma omp target map(tofrom : X)
+ X += 1;
+
+ return X == 2 ? 0 : 1;
+}
+
+// CHECK: device {{[0-9]+}} info: LaunchID: [[#ID:]] TeamsXthrds:({{.*}}) Duration(ns): {{[0-9]+}} n:__omp_offloading_{{.*}}_main_l{{[0-9]+}}
+// CHECK: device {{[0-9]+}} info: LaunchID: [[#ID+1]] TeamsXthrds:({{.*}}) Duration(ns): {{[0-9]+}} n:__omp_offloading_{{.*}}_main_l{{[0-9]+}}
``````````
</details>
https://github.com/llvm/llvm-project/pull/222105
More information about the llvm-commits
mailing list