add control to force non-profiling command queues (#162)

This can be used for performance analysis, but may cause errors if
an application requires event profiling.  If this control is set
along with another control that requires event profiling, such as
DevicePerformanceTiming, then the other control will override this
control, but this behavior may change in the future and should not
be relied upon!

Author: bashbaug

docs/controls.md

@@ -584,6 +584,10 @@ If set to a nonzero value, the Intercept Layer for OpenCL Applications inserts a
 
 If set to a nonzero value, the Intercept Layer for OpenCL Applications will force all queues to be created in-order.  This can be used for performance analysis, but may lead to deadlocks in some cases.
 
+##### `NoProfilingQueue` (bool)
+
+If set to a nonzero value, the Intercept Layer for OpenCL Applications will force all queues to be created without event profiling support.  This can be used for performance analysis, but may lead to errors if the application requires event profiling.
+
 ##### `NullEnqueue` (bool)
 
 If set to a nonzero value, the Intercept Layer for OpenCL Applications will silently ignore any enqueue.  This can be used for performance analysis, but will likely cause errors if the application relies on any sort of information from OpenCL events and should be used carefully.

intercept/src/controls.h

@@ -156,6 +156,7 @@ CLI_CONTROL( bool,          FinishAfterEnqueue,                     false, "If s
 CLI_CONTROL( bool,          FlushAfterEnqueue,                      false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications inserts a call to clFlush() after every enqueue.  The command queue that the command was just enqueued to is passed to clFlush().  This can also be used to debug possible timing or resource management issues and is slightly less obtrusive than FinishAfterEnqueue but still will likely impact performance.  If both FinishAfterEnqueue and FlushAfterEnqueue are nonzero then the Intercept Layer for OpenCL Applications will only insert a call to clFinish() after every enqueue, because clFinish() implies clFlush()." )
 CLI_CONTROL( bool,          FlushAfterEnqueueBarrier,               false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications inserts a call to clFlush() after every barrier enqueue.  The command queue that the command was just enqueued to is passed to clFlush().  This has been useful to debug out-of-order queue issues." )
 CLI_CONTROL( bool,          InOrderQueue,                           false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will force all queues to be created in-order.  This can be used for performance analysis, but may lead to deadlocks in some cases." )
+CLI_CONTROL( bool,          NoProfilingQueue,                       false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will force all queues to be created without event profiling support.  This can be used for performance analysis, but may lead to errors if the application requires event profiling." )
 CLI_CONTROL( bool,          NullEnqueue,                            false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will silently ignore any enqueue.  This can be used for performance analysis, but will likely cause errors if the application relies on any sort of information from OpenCL events and should be used carefully." )
 CLI_CONTROL( bool,          NullLocalWorkSize,                      false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will force the local work size argument to clEnqueueNDRangeKernel() to be NULL, which causes the OpenCL implementation to pick the local work size. Note that this control takes effect before NullLocalWorkSizeX / NullLocalWorkSizeY / NullLocalWorkSizeZ (see below), so enabling both controls will have the effect of forcing a specific local work size." )
 CLI_CONTROL( size_t,        NullLocalWorkSizeX,                     0,     "If set to a nonzero value, the Intercept Layer for OpenCL Applications will set the local work size that will be used if an application passes NULL as the local work size to clEnqueueNDRangeKernel().  1D dispatches will only look at NullLocalWorkSizeX, 2D dispatches will only look at NullLocalWorkSizeX and NullLocalWorkSizeY, while 3D dispatches will look at NullLocalWorkSizeX, NullLocalWorkSizeY, and NullLocalWorkSizeZ.  If the specified values for NullLocalWorkSize do not evenly divide the global work size then the specified values of NullLocalWorkSize will not take effect." )

intercept/src/intercept.cpp

@@ -4495,6 +4495,14 @@ void CLIntercept::updateHostTimingStats(
 void CLIntercept::modifyCommandQueueProperties(
     cl_command_queue_properties& props ) const
 {
+    if( config().InOrderQueue )
+    {
+        props &= ~(cl_command_queue_properties)CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
+    }
+    if( config().NoProfilingQueue )
+    {
+        props &= ~(cl_command_queue_properties)CL_QUEUE_PROFILING_ENABLE;
+    }
     if( config().DevicePerformanceTiming ||
         config().ITTPerformanceTiming ||
         config().ChromePerformanceTiming ||
@@ -4503,10 +4511,6 @@ void CLIntercept::modifyCommandQueueProperties(
     {
         props |= (cl_command_queue_properties)CL_QUEUE_PROFILING_ENABLE;
     }
-    if( config().InOrderQueue )
-    {
-        props &= ~(cl_command_queue_properties)CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
-    }
 }
 
 ///////////////////////////////////////////////////////////////////////////////

intercept/src/intercept.h

@@ -2172,6 +2172,7 @@ inline bool CLIntercept::checkDevicePerformanceTimingEnqueueLimits() const
         pIntercept->config().SIMDSurvey ||                                  \
         pIntercept->config().DevicePerfCounterEventBasedSampling ||         \
         pIntercept->config().InOrderQueue ||                                \
+        pIntercept->config().NoProfilingQueue ||                            \
         pIntercept->config().DefaultQueuePriorityHint ||                    \
         pIntercept->config().DefaultQueueThrottleHint )                     \
     {                                                                       \
@@ -2188,6 +2189,7 @@ inline bool CLIntercept::checkDevicePerformanceTimingEnqueueLimits() const
         pIntercept->config().SIMDSurvey ||                                  \
         pIntercept->config().DevicePerfCounterEventBasedSampling ||         \
         pIntercept->config().InOrderQueue ||                                \
+        pIntercept->config().NoProfilingQueue ||                            \
         pIntercept->config().DefaultQueuePriorityHint ||                    \
         pIntercept->config().DefaultQueueThrottleHint )                     \
     {                                                                       \