Skip to content

improve device performance timing for kernels only #478

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
bashbaug merged 5 commits into from
Aug 27, 2026
Merged

improve device performance timing for kernels only #478

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
ae90952
improve device performance timing kernels only
bashbaug Aug 24, 2026
f72d2a4
add a cliloader option to profile kernels only
bashbaug Aug 24, 2026
9cd5cb3
a bit more cleanup
bashbaug Aug 26, 2026
6b51461
fix device performance timing for mapped buffers and images
bashbaug Aug 26, 2026
50ee7f6
a bit more cleanup
bashbaug Aug 26, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions cliloader/cliloader.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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");
Expand Down Expand Up @@ -639,6 +643,7 @@ static bool parseArguments(int argc, char *argv[])
" --mdapi-group <NAME> Choose MDAPI Metrics to Collect (Intel GPU Only)\n"
" --mdapi-device <INDEX> 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 <NUMBER> Capture the Specified Kernel Enqueue\n"
" --capture-kernel <NAME> Capture the Specified Kernel Name\n"
" --leak-checking [-l] Track and Report OpenCL Leaks\n"
Expand Down
16 changes: 8 additions & 8 deletions intercept/src/dispatch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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();
Expand All @@ -4550,11 +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, event );
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;
Expand Down Expand Up @@ -4653,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();
Expand All @@ -4675,9 +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, 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;
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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(
Expand Down
41 changes: 33 additions & 8 deletions intercept/src/intercept.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 ) \
{ \
Expand Down Expand Up @@ -3383,6 +3388,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; \
Expand All @@ -3407,8 +3434,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();*/ \
Expand All @@ -3432,8 +3458,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();*/ \
Expand Down
Loading