From ae90952c565438da05cb27402d29e8d725c38f78 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 24 Aug 2026 07:08:03 -0700 Subject: [PATCH 1/5] improve device performance timing kernels only If we do not care about device performance timing for non-kernel commands, skip creating an event for these commands if one is not requested by the application. --- intercept/src/dispatch.cpp | 4 ++-- intercept/src/intercept.h | 22 ++++++++++++++++++++++ 2 files changed, 24 insertions(+), 2 deletions(-) diff --git a/intercept/src/dispatch.cpp b/intercept/src/dispatch.cpp index 4e3ea1a7..a721a024 100644 --- a/intercept/src/dispatch.cpp +++ b/intercept/src/dispatch.cpp @@ -4928,7 +4928,7 @@ CL_API_ENTRY cl_int CL_API_CALL CLIRN(clEnqueueNDRangeKernel)( global_work_offset, global_work_size, local_work_size ); - DEVICE_PERFORMANCE_TIMING_START( event ); + DEVICE_PERFORMANCE_TIMING_START_KERNEL( event ); HOST_PERFORMANCE_TIMING_START(); // ITT_ADD_PARAM_AS_METADATA(command_queue); @@ -5029,7 +5029,7 @@ CL_API_ENTRY cl_int CL_API_CALL CLIRN(clEnqueueTask)( eventWaitListString.c_str()); CHECK_EVENT_LIST( num_events_in_wait_list, event_wait_list, event ); GET_TIMING_TAGS_KERNEL( command_queue, kernel, 0, NULL, NULL, NULL ); - DEVICE_PERFORMANCE_TIMING_START( event ); + DEVICE_PERFORMANCE_TIMING_START_KERNEL( event ); HOST_PERFORMANCE_TIMING_START(); retVal = pIntercept->dispatch().clEnqueueTask( diff --git a/intercept/src/intercept.h b/intercept/src/intercept.h index f153c049..396be90e 100644 --- a/intercept/src/intercept.h +++ b/intercept/src/intercept.h @@ -3383,6 +3383,28 @@ inline bool CLIntercept::checkDevicePerformanceTimingEnqueueLimits( } #define DEVICE_PERFORMANCE_TIMING_START( pEvent ) \ + CLIntercept::clock::time_point queuedTime; \ + cl_event local_event = NULL; \ + bool isLocalEvent = false; \ + bool doDevicePerformanceTiming = \ + ( pIntercept->config().DevicePerformanceTiming || \ + pIntercept->config().ITTPerformanceTiming || \ + pIntercept->config().ChromePerformanceTiming || \ + pIntercept->config().DevicePerfCounterEventBasedSampling ) && \ + !pIntercept->config().DevicePerformanceTimingKernelsOnly && \ + pIntercept->checkDevicePerformanceTimingEnqueueLimits( enqueueCounter ) &&\ + pIntercept->checkConditionalTiming(); \ + if( doDevicePerformanceTiming ) \ + { \ + queuedTime = CLIntercept::clock::now(); \ + if( pEvent == NULL ) \ + { \ + pEvent = &local_event; \ + isLocalEvent = true; \ + } \ + } + +#define DEVICE_PERFORMANCE_TIMING_START_KERNEL( pEvent ) \ CLIntercept::clock::time_point queuedTime; \ cl_event local_event = NULL; \ bool isLocalEvent = false; \ From f72d2a47fb3d0a9b384649721cdd008bd05107b2 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 24 Aug 2026 07:14:08 -0700 Subject: [PATCH 2/5] add a cliloader option to profile kernels only --- cliloader/cliloader.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/cliloader/cliloader.cpp b/cliloader/cliloader.cpp index c85c8a87..0afac23f 100644 --- a/cliloader/cliloader.cpp +++ b/cliloader/cliloader.cpp @@ -517,6 +517,10 @@ static bool parseArguments(int argc, char *argv[]) { checkSetEnv("CLI_HostPerformanceTiming", "1"); } + else if( !strcmp(argv[i], "-ko") || !strcmp(argv[i], "--kernels-only") ) + { + checkSetEnv("CLI_DevicePerformanceTimingKernelsOnly", "1"); + } else if( !strcmp(argv[i], "-l") || !strcmp(argv[i], "--leak-checking") ) { checkSetEnv("CLI_LeakChecking", "1"); @@ -639,6 +643,7 @@ static bool parseArguments(int argc, char *argv[]) " --mdapi-group Choose MDAPI Metrics to Collect (Intel GPU Only)\n" " --mdapi-device Choose MDAPI Device for Metrics (Intel GPU Only)\n" " --host-timing [-h] Report Host API Execution Time\n" + " --kernels-only [-ko] Only Profile Kernels for Device Timing\n" " --capture-enqueue Capture the Specified Kernel Enqueue\n" " --capture-kernel Capture the Specified Kernel Name\n" " --leak-checking [-l] Track and Report OpenCL Leaks\n" From 9cd5cb3621ff798eeb2cf4e5f8ff43231d1ec679 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 26 Aug 2026 13:45:05 -0700 Subject: [PATCH 3/5] a bit more cleanup Since we've already checked for DevicePerformanceTimingKernelsOnly when determining whether to do device performance timing, we do not need to check for it again to determine whether to add the timing event. --- intercept/src/intercept.h | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/intercept/src/intercept.h b/intercept/src/intercept.h index 396be90e..1aed1b4a 100644 --- a/intercept/src/intercept.h +++ b/intercept/src/intercept.h @@ -3429,8 +3429,7 @@ inline bool CLIntercept::checkDevicePerformanceTimingEnqueueLimits( if( doDevicePerformanceTiming && \ ( errorCode == CL_SUCCESS ) && ( pEvent != NULL ) ) \ { \ - if( !pIntercept->config().DevicePerformanceTimingKernelsOnly && \ - ( !pIntercept->config().DevicePerformanceTimingSkipUnmap || \ + if( ( !pIntercept->config().DevicePerformanceTimingSkipUnmap || \ std::string(__FUNCTION__) != "clEnqueueUnmapMemObject" ) ) \ { \ /*TOOL_OVERHEAD_TIMING_START();*/ \ @@ -3454,8 +3453,7 @@ inline bool CLIntercept::checkDevicePerformanceTimingEnqueueLimits( if( doDevicePerformanceTiming && \ ( errorCode == CL_SUCCESS ) && ( pEvent != NULL ) ) \ { \ - if( !pIntercept->config().DevicePerformanceTimingKernelsOnly && \ - ( !pIntercept->config().DevicePerformanceTimingSkipUnmap || \ + if( ( !pIntercept->config().DevicePerformanceTimingSkipUnmap || \ std::string(__FUNCTION__) != "clEnqueueUnmapMemObject" ) ) \ { \ /*TOOL_OVERHEAD_TIMING_START();*/ \ From 6b514619a631953452e34576589265a35414e6c0 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 26 Aug 2026 14:07:59 -0700 Subject: [PATCH 4/5] fix device performance timing for mapped buffers and images --- intercept/src/dispatch.cpp | 19 ++++++++++++++++--- 1 file changed, 16 insertions(+), 3 deletions(-) diff --git a/intercept/src/dispatch.cpp b/intercept/src/dispatch.cpp index a721a024..18338870 100644 --- a/intercept/src/dispatch.cpp +++ b/intercept/src/dispatch.cpp @@ -4550,8 +4550,18 @@ CL_API_ENTRY void* CL_API_CALL CLIRN(clEnqueueMapBuffer)( errcode_ret ); HOST_PERFORMANCE_TIMING_END_WITH_TAG(); - DEVICE_PERFORMANCE_TIMING_END_WITH_TAG( command_queue, retVal, event ); - DUMP_BUFFER_AFTER_MAP( command_queue, buffer, blocking_map, map_flags, retVal, offset, cb ); + DEVICE_PERFORMANCE_TIMING_END_WITH_TAG( + command_queue, + (retVal != NULL ? CL_SUCCESS : CL_INVALID_VALUE), + event ); + DUMP_BUFFER_AFTER_MAP( + command_queue, + buffer, + blocking_map, + map_flags, + retVal, + offset, + cb ); CHECK_ERROR( errcode_ret[0] ); ADD_MAP_POINTER( retVal, map_flags, cb ); ADD_OBJECT_ALLOCATION_EVENT( retVal, event ); @@ -4675,7 +4685,10 @@ CL_API_ENTRY void* CL_API_CALL CLIRN(clEnqueueMapImage)( errcode_ret ); HOST_PERFORMANCE_TIMING_END_WITH_TAG(); - DEVICE_PERFORMANCE_TIMING_END_WITH_TAG( command_queue, retVal, event ); + DEVICE_PERFORMANCE_TIMING_END_WITH_TAG( + command_queue, + (retVal != NULL ? CL_SUCCESS : CL_INVALID_VALUE), + event ); CHECK_ERROR( errcode_ret[0] ); ADD_OBJECT_ALLOCATION_EVENT( retVal, event ); if( pIntercept->config().CallLogging ) From 50ee7f616b56f75ea54f455bd3d0bf19c65234c4 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 26 Aug 2026 14:28:01 -0700 Subject: [PATCH 5/5] a bit more cleanup It's a lot easier to always have the correct error code available for map commands. --- intercept/src/dispatch.cpp | 27 +++++++-------------------- intercept/src/intercept.h | 13 +++++++++---- 2 files changed, 16 insertions(+), 24 deletions(-) diff --git a/intercept/src/dispatch.cpp b/intercept/src/dispatch.cpp index 18338870..bfc28b59 100644 --- a/intercept/src/dispatch.cpp +++ b/intercept/src/dispatch.cpp @@ -4530,7 +4530,7 @@ CL_API_ENTRY void* CL_API_CALL CLIRN(clEnqueueMapBuffer)( cb, eventWaitListString.c_str() ); CHECK_EVENT_LIST( num_events_in_wait_list, event_wait_list, event ); - CHECK_ERROR_INIT( errcode_ret ); + CHECK_ERROR_INIT_MAP( errcode_ret ); GET_TIMING_TAGS_MAP( blocking_map, map_flags, cb ); DEVICE_PERFORMANCE_TIMING_START( event ); HOST_PERFORMANCE_TIMING_START(); @@ -4550,21 +4550,11 @@ CL_API_ENTRY void* CL_API_CALL CLIRN(clEnqueueMapBuffer)( errcode_ret ); HOST_PERFORMANCE_TIMING_END_WITH_TAG(); - DEVICE_PERFORMANCE_TIMING_END_WITH_TAG( - command_queue, - (retVal != NULL ? CL_SUCCESS : CL_INVALID_VALUE), - event ); - DUMP_BUFFER_AFTER_MAP( - command_queue, - buffer, - blocking_map, - map_flags, - retVal, - offset, - cb ); + DEVICE_PERFORMANCE_TIMING_END_WITH_TAG( command_queue, errcode_ret[0], event ); + DUMP_BUFFER_AFTER_MAP( command_queue, buffer, blocking_map, map_flags, retVal, offset, cb ); CHECK_ERROR( errcode_ret[0] ); ADD_MAP_POINTER( retVal, map_flags, cb ); - ADD_OBJECT_ALLOCATION_EVENT( retVal, event ); + ADD_OBJECT_ALLOCATION_EVENT( errcode_ret[0], event ); if( pIntercept->config().CallLogging ) { map_count = 0; @@ -4663,7 +4653,7 @@ CL_API_ENTRY void* CL_API_CALL CLIRN(clEnqueueMapImage)( eventWaitListString.c_str() ); } CHECK_EVENT_LIST( num_events_in_wait_list, event_wait_list, event ); - CHECK_ERROR_INIT( errcode_ret ); + CHECK_ERROR_INIT_MAP( errcode_ret ); GET_TIMING_TAGS_MAP( blocking_map, map_flags, 0 ); DEVICE_PERFORMANCE_TIMING_START( event ); HOST_PERFORMANCE_TIMING_START(); @@ -4685,12 +4675,9 @@ CL_API_ENTRY void* CL_API_CALL CLIRN(clEnqueueMapImage)( errcode_ret ); HOST_PERFORMANCE_TIMING_END_WITH_TAG(); - DEVICE_PERFORMANCE_TIMING_END_WITH_TAG( - command_queue, - (retVal != NULL ? CL_SUCCESS : CL_INVALID_VALUE), - event ); + DEVICE_PERFORMANCE_TIMING_END_WITH_TAG( command_queue, errcode_ret[0], event ); CHECK_ERROR( errcode_ret[0] ); - ADD_OBJECT_ALLOCATION_EVENT( retVal, event ); + ADD_OBJECT_ALLOCATION_EVENT( errcode_ret[0], event ); if( pIntercept->config().CallLogging ) { map_count = 0; diff --git a/intercept/src/intercept.h b/intercept/src/intercept.h index 1aed1b4a..1a9a594f 100644 --- a/intercept/src/intercept.h +++ b/intercept/src/intercept.h @@ -2244,11 +2244,16 @@ inline CObjectTracker& CLIntercept::objectTracker() pErrorCode = &localErrorCode; \ } +// For map APIs, setup the error code pointer unconditionally: +#define CHECK_ERROR_INIT_MAP( pErrorCode ) \ + cl_int localErrorCode = CL_SUCCESS; \ + if( pErrorCode == NULL ) \ + { \ + pErrorCode = &localErrorCode; \ + } + #define CHECK_ERROR( errorCode ) \ - if( ( pIntercept->config().ErrorLogging || \ - pIntercept->config().ErrorAssert || \ - pIntercept->config().NoErrors ) && \ - ( errorCode != CL_SUCCESS ) ) \ + if( errorCode != CL_SUCCESS ) \ { \ if( pIntercept->config().ErrorLogging ) \ { \