Merge branch 'main' into cl_khr_unified_svm

bashbaug · bashbaug · commit e3e943c07528 · 2025-11-06T14:44:19.000-08:00
diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml
@@ -18,7 +18,7 @@ jobs:
   check:
     runs-on: ubuntu-latest
     steps:
-    - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+    - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
 
     - name: Check Tabs
       run: |
@@ -39,7 +39,7 @@ jobs:
     runs-on: ${{matrix.os}}
 
     steps:
-    - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+    - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
 
     - name: Get Ubuntu OpenGL Dependencies
       if: matrix.os == 'ubuntu-latest'
@@ -48,29 +48,29 @@ jobs:
         sudo apt-get install -y libglfw3-dev
 
     - name: Get OpenCL Headers
-      uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+      uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
       with:
         repository: KhronosGroup/OpenCL-Headers
         ref: cl_khr_unified_svm
         path: external/OpenCL-Headers
 
     - name: Get OpenCL ICD Loader
-      uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+      uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
       with:
         repository: KhronosGroup/OpenCL-ICD-Loader
         path: external/opencl-icd-loader
 
     - name: Get OpenCL Extension Loader
       if: matrix.ext == 'YES'
-      uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+      uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
       with:
         repository: bashbaug/opencl-extension-loader
         ref: cl_khr_unified_svm
         path: external/opencl-extension-loader
 
     - name: Get SPIR-V Headers
       if: matrix.ext == 'YES'
-      uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
+      uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
       with:
         repository: KhronosGroup/SPIRV-Headers
         path: external/SPIRV-Headers
diff --git a/layers/10_cmdbufemu/README.md b/layers/10_cmdbufemu/README.md
@@ -33,6 +33,7 @@ The following environment variables can modify the behavior of the command buffe
 | Environment Variable | Behavior |  Example Format |
 |----------------------|----------|-----------------|
 | `CMDBUFEMU_EnhancedErrorChecking` | Enables additional error checking when commands are added to a command buffer using a command buffer "test queue".  By default, the additional error checking is disabled. | `export CMDBUFEMU_EnhancedErrorChecking=1`<br/><br/>`set CMDBUFEMU_EnhancedErrorChecking=1` |
+| `CMDBUFEMU_KernelForProfiling` | Enables use of an empty kernel for event profiling instead of event profiling on a command-queue barrier.  By default, to minimize overhead, the empty kernel is not used. | `export CMDBUFEMU_KernelForProfiling=1`<br/><br/>`set CMDBUFEMU_KernelForProfiling=1` |
 
 ## Known Limitations
 
diff --git a/layers/10_cmdbufemu/emulate.cpp b/layers/10_cmdbufemu/emulate.cpp
@@ -39,6 +39,26 @@ const cl_mutable_dispatch_fields_khr g_MutableDispatchCaps =
     CL_MUTABLE_DISPATCH_ARGUMENTS_KHR |
     CL_MUTABLE_DISPATCH_EXEC_INFO_KHR;
 
+static cl_int enqueueProfilingKernel(
+    cl_command_queue queue,
+    cl_kernel kernel,
+    cl_uint num_events_in_wait_list,
+    const cl_event* event_wait_list,
+    cl_event* event )
+{
+    const size_t one = 1;
+    return g_pNextDispatch->clEnqueueNDRangeKernel(
+        queue,
+        kernel,
+        1,
+        nullptr,
+        &one,
+        nullptr,
+        num_events_in_wait_list,
+        event_wait_list,
+        event );
+}
+
 typedef struct _cl_mutable_command_khr
 {
     static bool isValid( cl_mutable_command_khr command )
@@ -215,7 +235,7 @@ struct BarrierWithWaitList : Command
         return g_pNextDispatch->clEnqueueBarrierWithWaitList(
             queue,
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -271,7 +291,7 @@ struct CopyBuffer : Command
             dst_offset,
             size,
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -345,7 +365,7 @@ struct CopyBufferRect : Command
             dst_row_pitch,
             dst_slice_pitch,
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -411,7 +431,7 @@ struct CopyBufferToImage : Command
             dst_origin.data(),
             region.data(),
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -473,7 +493,7 @@ struct CopyImage : Command
             dst_origin.data(),
             region.data(),
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -535,7 +555,7 @@ struct CopyImageToBuffer : Command
             region.data(),
             dst_offset,
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -601,7 +621,7 @@ struct FillBuffer : Command
             offset,
             size,
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -678,7 +698,7 @@ struct FillImage : Command
             origin.data(),
             region.data(),
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -725,7 +745,7 @@ struct SVMMemcpy : Command
             src_ptr,
             size,
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -779,7 +799,7 @@ struct SVMMemFill : Command
             pattern.size(),
             size,
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -1073,7 +1093,7 @@ struct NDRangeKernel : Command
             global_work_size.data(),
             local_work_size.size() ? local_work_size.data() : nullptr,
             static_cast<cl_uint>(wait_list.size()),
-            wait_list.data(),
+            wait_list.size() ? wait_list.data() : nullptr,
             signal);
     }
 
@@ -1229,6 +1249,7 @@ typedef struct _cl_command_buffer_khr
                     (props & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE) == 0 );
 
                 cmdbuf->setupTestQueue(queue);
+                cmdbuf->setupProfilingKernel(queue);
             }
         }
 
@@ -1254,6 +1275,11 @@ typedef struct _cl_command_buffer_khr
         {
             g_pNextDispatch->clReleaseCommandQueue(queue);
         }
+
+        for( auto kernel : ProfilingKernels )
+        {
+            g_pNextDispatch->clReleaseKernel(kernel);
+        }
     }
 
     static bool isValid( cl_command_buffer_khr cmdbuf )
@@ -1297,20 +1323,17 @@ typedef struct _cl_command_buffer_khr
 
     cl_command_queue    getQueue() const
     {
-        if( Queues.size() > 0 )
-        {
-            return Queues[0];
-        }
-        return nullptr;
+        return Queues.empty() ? nullptr : Queues[0];
     }
 
     cl_command_queue    getTestQueue() const
     {
-        if( TestQueues.size() > 0 )
-        {
-            return TestQueues[0];
-        }
-        return nullptr;
+        return TestQueues.empty() ? nullptr : TestQueues[0];
+    }
+
+    cl_kernel   getProfilingKernel() const
+    {
+        return ProfilingKernels.empty() ? nullptr : ProfilingKernels[0];
     }
 
     cl_mutable_dispatch_asserts_khr getMutableDispatchAsserts() const
@@ -1510,10 +1533,13 @@ typedef struct _cl_command_buffer_khr
             NextSyncPoint.fetch_add(1, std::memory_order_relaxed) :
             0;
 
-        command->addDependencies(
-            num_sync_points,
-            wait_list,
-            syncPoint);
+        // We only need to add dependencies if there is more than one queue (so
+        // we have possible cross-queue dependencies) or the queue is an
+        // out-of-order queue (so we have possible intra-queue dependencies).
+        if( Queues.size() > 1 || !IsInOrder[0] )
+        {
+            command->addDependencies(num_sync_points, wait_list, syncPoint);
+        }
 
         if( sync_point != nullptr )
         {
@@ -1671,6 +1697,7 @@ typedef struct _cl_command_buffer_khr
     std::vector<bool>   IsInOrder;
     std::vector<cl_command_queue>   TestQueues;
     std::vector<cl_event>   BlockingEvents;
+    std::vector<cl_kernel>  ProfilingKernels;
 
     std::vector<std::unique_ptr<Command>> Commands;
     std::atomic<uint32_t> NextSyncPoint;
@@ -1747,6 +1774,52 @@ typedef struct _cl_command_buffer_khr
         }
     }
 
+    void setupProfilingKernel(cl_command_queue queue)
+    {
+        if( g_KernelForProfiling )
+        {
+            cl_context context = nullptr;
+            g_pNextDispatch->clGetCommandQueueInfo(
+                queue,
+                CL_QUEUE_CONTEXT,
+                sizeof(context),
+                &context,
+                nullptr );
+
+            cl_device_id device = nullptr;
+            g_pNextDispatch->clGetCommandQueueInfo(
+                queue,
+                CL_QUEUE_DEVICE,
+                sizeof(device),
+                &device,
+                nullptr );
+
+            const char* kernelString = "kernel void Empty() {}";
+            cl_program program = g_pNextDispatch->clCreateProgramWithSource(
+                context,
+                1,
+                &kernelString,
+                nullptr,
+                nullptr );
+            g_pNextDispatch->clBuildProgram(
+                program,
+                1,
+                &device,
+                nullptr,
+                nullptr,
+                nullptr );
+
+            cl_kernel kernel = g_pNextDispatch->clCreateKernel(
+                program,
+                "Empty",
+                nullptr );
+            g_pNextDispatch->clReleaseProgram(
+                program );
+
+            ProfilingKernels.push_back(kernel);
+        }
+    }
+
     _cl_command_buffer_khr(
             cl_command_buffer_flags_khr flags,
             cl_mutable_dispatch_asserts_khr mutableDispatchAsserts) :
@@ -1993,7 +2066,16 @@ cl_int CL_API_CALL clEnqueueCommandBufferKHR_EMU(
             queue,
             num_events_in_wait_list,
             event_wait_list,
-            event ? &startEvent : nullptr);
+            event == nullptr || g_KernelForProfiling ? nullptr : &startEvent );
+        if( errorCode == CL_SUCCESS && event && g_KernelForProfiling )
+        {
+            errorCode = enqueueProfilingKernel(
+                queue,
+                cmdbuf->getProfilingKernel(),
+                0,
+                nullptr,
+                &startEvent );
+        }
     }
 
     if( errorCode == CL_SUCCESS )
@@ -2007,7 +2089,16 @@ cl_int CL_API_CALL clEnqueueCommandBufferKHR_EMU(
             queue,
             0,
             nullptr,
-            event );
+            g_KernelForProfiling ? nullptr : event );
+        if( errorCode == CL_SUCCESS && g_KernelForProfiling )
+        {
+            errorCode = enqueueProfilingKernel(
+                queue,
+                cmdbuf->getProfilingKernel(),
+                0,
+                nullptr,
+                event );
+        }
     }
 
     if( event )
diff --git a/layers/10_cmdbufemu/emulate.h b/layers/10_cmdbufemu/emulate.h
@@ -10,6 +10,7 @@
 #include <map>
 
 extern bool g_EnhancedErrorChecking;
+extern bool g_KernelForProfiling;
 
 extern const struct _cl_icd_dispatch* g_pNextDispatch;
 
diff --git a/layers/10_cmdbufemu/main.cpp b/layers/10_cmdbufemu/main.cpp
@@ -34,6 +34,11 @@
 
 bool g_EnhancedErrorChecking = false;
 
+// Using kernels for profiling can fix issues with some implementations
+// that do not properly support event profiling on barrkers.
+
+bool g_KernelForProfiling = false;
+
 const struct _cl_icd_dispatch* g_pNextDispatch = NULL;
 
 static cl_int CL_API_CALL
@@ -283,6 +288,7 @@ CL_API_ENTRY cl_int CL_API_CALL clInitLayer(
     _init_dispatch();
 
     getControl("CMDBUFEMU_EnhancedErrorChecking", g_EnhancedErrorChecking);
+    getControl("CMDBUFEMU_KernelForProfiling", g_KernelForProfiling);
 
     g_pNextDispatch = target_dispatch;
 
diff --git a/samples/python/requirements.txt b/samples/python/requirements.txt
@@ -3,4 +3,4 @@ numpy==1.26.3
 pyopencl==2023.1.4
 pillow==10.3.0
 
-jupyterlab==4.2.5
+jupyterlab==4.4.8
