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 <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"
diff --git a/intercept/src/dispatch.cpp b/intercept/src/dispatch.cpp
index 4e3ea1a7..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,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;
@@ -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();
@@ -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;
@@ -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..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 )                             \
         {                                                                   \
@@ -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;                                       \
@@ -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();*/                               \
@@ -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();*/                               \