[SYCL][Graph][OpenCL] Map copy/fill to SVM (#18177)

EwanC · web-flow · commit 4ed8534f20ea · 2025-04-28T09:39:48.000+02:00
There is currently no cl_khr_command_buffer or
cl_intel_unified_shared_memory entry-point for appending USM copy or
fill commands to a command-buffer. This prevents these commands from
being added to a graph for the OpenCL backend.

The long term solution to this is the OpenCL USVM extension which will
align USM and SVM, allowing the existing SVM entry-points to be used. To
prepare for this, map the UR entry-points to the cl_khr_command_buffer
SVM copy/fill commands.

This will work on OpenCL implementations that share an implementation
for USM and SVM.

Signed-off-by: Ewan Crawford &lt;ewan@codeplay.com&gt;
diff --git a/sycl/doc/design/CommandGraph.md b/sycl/doc/design/CommandGraph.md
@@ -626,8 +626,8 @@ adapter where there is matching support for each function in the list.
 | urCommandBufferReleaseExp | clReleaseCommandBufferKHR | Yes |
 | urCommandBufferFinalizeExp | clFinalizeCommandBufferKHR | Yes |
 | urCommandBufferAppendKernelLaunchExp | clCommandNDRangeKernelKHR | Yes |
-| urCommandBufferAppendUSMMemcpyExp |  | No |
-| urCommandBufferAppendUSMFillExp |  | No |
+| urCommandBufferAppendUSMMemcpyExp | clCommandSVMMemcpyKHR | Partial, [see below](#unsupported-command-types) |
+| urCommandBufferAppendUSMFillExp | clCommandSVMMemFillKHR | Partial, [see below](#unsupported-command-types) |
 | urCommandBufferAppendMembufferCopyExp | clCommandCopyBufferKHR | Yes |
 | urCommandBufferAppendMemBufferWriteExp |  | No |
 | urCommandBufferAppendMemBufferReadExp |  | No |
@@ -643,8 +643,6 @@ adapter where there is matching support for each function in the list.
 |  | clCommandCopyImageToBufferKHR | No |
 |  | clCommandFillImageKHR | No |
 |  | clGetCommandBufferInfoKHR | No |
-|  | clCommandSVMMemcpyKHR | No |
-|  | clCommandSVMMemFillKHR | No |
 | urCommandBufferUpdateKernelLaunchExp | clUpdateMutableCommandsKHR | Partial [See Update Section](#update-support) |
 
 We are looking to address these gaps in the future so that SYCL-Graph can be
@@ -664,17 +662,21 @@ terminate called after throwing an instance of 'sycl::_V1::exception'
 what():  USM copy command not supported by graph backend
 ```
 
-The types of commands which are unsupported, and lead to this exception are:
+The types of commands which are unsupported, and may lead to this exception are:
 * `handler::copy(src, dest)` - Where `src` is an accessor and `dest` is a pointer.
    This corresponds to a memory buffer read command.
 * `handler::copy(src, dest)` - Where `src` is an pointer and `dest` is an accessor.
   This corresponds to a memory buffer write command.
 * `handler::copy(src, dest)` or `handler::memcpy(dest, src)` - Where both `src` and
-   `dest` are USM pointers. This corresponds to a USM copy command.
+   `dest` are USM pointers. This corresponds to a USM copy command that is
+   mapped to `clCommandSVMMemcpyKHR`, which will only work on OpenCL devices
+   which don't differentiate between USM and SVM.
 * `handler::fill(ptr, pattern, count)` - This corresponds to a USM memory
-  fill command.
+  fill command that is mapped to `clCommandSVMMemFillKHR`, which will only work
+  on OpenCL devices which don't differentiate between USM and SVM.
 * `handler::memset(ptr, value, numBytes)` - This corresponds to a USM memory
-  fill command.
+  fill command that is mapped to `clCommandSVMMemFillKHR`, which will only work
+  on OpenCL devices which don't differentiate between USM and SVM.
 * `handler::prefetch()`.
 * `handler::mem_advise()`.
 
diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp
@@ -1567,9 +1567,17 @@ void MemoryManager::ext_oneapi_fill_usm_cmd_buffer(
                     "NULL pointer argument in memory fill operation.");
 
   const AdapterPtr &Adapter = Context->getAdapter();
-  Adapter->call<UrApiKind::urCommandBufferAppendUSMFillExp>(
-      CommandBuffer, DstMem, Pattern.data(), Pattern.size(), Len, Deps.size(),
-      Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr);
+  ur_result_t Result =
+      Adapter->call_nocheck<UrApiKind::urCommandBufferAppendUSMFillExp>(
+          CommandBuffer, DstMem, Pattern.data(), Pattern.size(), Len,
+          Deps.size(), Deps.data(), 0, nullptr, OutSyncPoint, nullptr, nullptr);
+  if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
+    throw sycl::exception(
+        sycl::make_error_code(sycl::errc::feature_not_supported),
+        "USM fill command not supported by graph backend");
+  } else {
+    Adapter->checkUrResult(Result);
+  }
 }
 
 void MemoryManager::ext_oneapi_fill_cmd_buffer(
diff --git a/sycl/test-e2e/Graph/Explicit/usm_copy.cpp b/sycl/test-e2e/Graph/Explicit/usm_copy.cpp
@@ -4,10 +4,6 @@
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
-//
-//
-// Intended - USM copy command not supported for OpenCL
-// UNSUPPORTED: opencl
 
 #define GRAPH_E2E_EXPLICIT
 
diff --git a/sycl/test-e2e/Graph/Explicit/usm_fill.cpp b/sycl/test-e2e/Graph/Explicit/usm_fill.cpp
@@ -4,9 +4,6 @@
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
-//
-// Intended - USM fill command not supported for OpenCL
-// UNSUPPORTED: opencl
 
 #define GRAPH_E2E_EXPLICIT
 
diff --git a/sycl/test-e2e/Graph/Explicit/usm_fill_host.cpp b/sycl/test-e2e/Graph/Explicit/usm_fill_host.cpp
@@ -7,9 +7,6 @@
 
 // REQUIRES: aspect-usm_host_allocations
 
-// Intended - USM fill command not supported for OpenCL
-// UNSUPPORTED: opencl
-
 #define GRAPH_E2E_EXPLICIT
 
 #include "../Inputs/usm_fill_host.cpp"
diff --git a/sycl/test-e2e/Graph/Explicit/usm_fill_shared.cpp b/sycl/test-e2e/Graph/Explicit/usm_fill_shared.cpp
@@ -7,9 +7,6 @@
 
 // REQUIRES: aspect-usm_shared_allocations
 
-// Intended - USM fill command not supported for OpenCL
-// UNSUPPORTED: opencl
-
 #define GRAPH_E2E_EXPLICIT
 
 #include "../Inputs/usm_fill_shared.cpp"
diff --git a/sycl/test-e2e/Graph/Explicit/usm_memset.cpp b/sycl/test-e2e/Graph/Explicit/usm_memset.cpp
@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// Intended - USM memset command not supported for OpenCL
-// UNSUPPORTED: opencl
-
 #define GRAPH_E2E_EXPLICIT
 
 #include "../Inputs/usm_memset.cpp"
diff --git a/sycl/test-e2e/Graph/RecordReplay/barrier_multi_graph.cpp b/sycl/test-e2e/Graph/RecordReplay/barrier_multi_graph.cpp
@@ -4,9 +4,6 @@
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
-//
-// UNSUPPORTED: opencl
-// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
 
 #include "../graph_common.hpp"
 
diff --git a/sycl/test-e2e/Graph/RecordReplay/barrier_multi_queue.cpp b/sycl/test-e2e/Graph/RecordReplay/barrier_multi_queue.cpp
@@ -4,9 +4,6 @@
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
-//
-// UNSUPPORTED: opencl
-// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
 
 #include "../graph_common.hpp"
 
diff --git a/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions.cpp b/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions.cpp
@@ -7,9 +7,6 @@
 
 // Tests the enqueue free function kernel shortcuts.
 
-// UNSUPPORTED: opencl
-// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
-
 #include "../graph_common.hpp"
 #include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
 #include <sycl/properties/all_properties.hpp>
diff --git a/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions_submit_with_event.cpp b/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions_submit_with_event.cpp
@@ -8,9 +8,6 @@
 // Tests the enqueue free function using USM and submit_with_event for
 // dependencies
 
-// UNSUPPORTED: opencl
-// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
-
 #include "../graph_common.hpp"
 #include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
 
diff --git a/sycl/test-e2e/Graph/RecordReplay/transitive_queue_barrier.cpp b/sycl/test-e2e/Graph/RecordReplay/transitive_queue_barrier.cpp
@@ -9,9 +9,6 @@
 // event is passed as a dependency to a barrier operation in a different
 // in-order queue.
 
-// UNSUPPORTED: opencl
-// UNSUPPORTED-INTENDED: USM fill command not supported for OpenCL
-
 #include "../graph_common.hpp"
 #include <sycl/properties/all_properties.hpp>
 
diff --git a/sycl/test-e2e/Graph/RecordReplay/usm_copy.cpp b/sycl/test-e2e/Graph/RecordReplay/usm_copy.cpp
@@ -4,9 +4,6 @@
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
-//
-// Intended - USM copy command not supported for OpenCL
-// UNSUPPORTED: opencl
 
 #define GRAPH_E2E_RECORD_REPLAY
 
diff --git a/sycl/test-e2e/Graph/RecordReplay/usm_copy_in_order.cpp b/sycl/test-e2e/Graph/RecordReplay/usm_copy_in_order.cpp
@@ -4,10 +4,6 @@
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
-//
-//
-// Intended - USM copy command not supported for OpenCL
-// UNSUPPORTED: opencl
 
 // Tests memcpy operation using device USM and an in-order queue.
 
diff --git a/sycl/test-e2e/Graph/RecordReplay/usm_fill.cpp b/sycl/test-e2e/Graph/RecordReplay/usm_fill.cpp
@@ -4,9 +4,6 @@
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
-//
-// Intended - USM fill command not supported for OpenCL
-// UNSUPPORTED: opencl
 
 #define GRAPH_E2E_RECORD_REPLAY
 
diff --git a/sycl/test-e2e/Graph/RecordReplay/usm_fill_host.cpp b/sycl/test-e2e/Graph/RecordReplay/usm_fill_host.cpp
@@ -7,9 +7,6 @@
 
 // REQUIRES: aspect-usm_host_allocations
 
-// Intended - USM fill command not supported for OpenCL
-// UNSUPPORTED: opencl
-
 #define GRAPH_E2E_RECORD_REPLAY
 
 #include "../Inputs/usm_fill_host.cpp"
diff --git a/sycl/test-e2e/Graph/RecordReplay/usm_fill_shared.cpp b/sycl/test-e2e/Graph/RecordReplay/usm_fill_shared.cpp
@@ -7,9 +7,6 @@
 
 // REQUIRES: aspect-usm_shared_allocations
 
-// Intended - USM fill command not supported for OpenCL
-// UNSUPPORTED: opencl
-
 #define GRAPH_E2E_RECORD_REPLAY
 
 #include "../Inputs/usm_fill_shared.cpp"
diff --git a/sycl/test-e2e/Graph/RecordReplay/usm_memset.cpp b/sycl/test-e2e/Graph/RecordReplay/usm_memset.cpp
@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// Intended - USM memset command not supported for OpenCL
-// UNSUPPORTED: opencl
-
 #define GRAPH_E2E_RECORD_REPLAY
 
 #include "../Inputs/usm_memset.cpp"
diff --git a/sycl/test-e2e/Graph/RecordReplay/usm_memset_shortcut.cpp b/sycl/test-e2e/Graph/RecordReplay/usm_memset_shortcut.cpp
@@ -3,9 +3,6 @@
 // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
 // RUN: %if level_zero %{ %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 //
-// Intended - USM fill command not supported for OpenCL
-// UNSUPPORTED: opencl
-//
 // Tests adding a USM memset queue shortcut operation as a graph node.
 
 #include "../graph_common.hpp"
diff --git a/sycl/test-e2e/Graph/ValidUsage/linear_graph_copy.cpp b/sycl/test-e2e/Graph/ValidUsage/linear_graph_copy.cpp
@@ -12,9 +12,6 @@
 // Tests that the optimization to use the L0 Copy Engine for memory commands
 // does not interfere with the linear graph optimization
 
-// UNSUPPORTED: opencl
-// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
-
 #include "../graph_common.hpp"
 
 #include <sycl/properties/queue_properties.hpp>
diff --git a/sycl/test-e2e/Graph/ValidUsage/linear_graph_copy_D2H.cpp b/sycl/test-e2e/Graph/ValidUsage/linear_graph_copy_D2H.cpp
@@ -13,8 +13,6 @@
 // does not interfere with the linear graph optimization
 //
 // REQUIRES: aspect-usm_host_allocations
-// UNSUPPORTED: opencl
-// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
 
 #include "../graph_common.hpp"
 
diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp
@@ -54,7 +54,7 @@
 // tests to match the required format and in that case you should just update
 // (i.e. reduce) the number and the list below.
 //
-// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 273
+// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 261
 //
 // List of improperly UNSUPPORTED tests.
 // Remove the CHECK once the test has been properly UNSUPPORTED.
@@ -165,11 +165,6 @@
 // CHECK-NEXT: Graph/Explicit/prefetch.cpp
 // CHECK-NEXT: Graph/Explicit/spec_constants_handler_api.cpp
 // CHECK-NEXT: Graph/Explicit/spec_constants_kernel_bundle_api.cpp
-// CHECK-NEXT: Graph/Explicit/usm_copy.cpp
-// CHECK-NEXT: Graph/Explicit/usm_fill.cpp
-// CHECK-NEXT: Graph/Explicit/usm_fill_host.cpp
-// CHECK-NEXT: Graph/Explicit/usm_fill_shared.cpp
-// CHECK-NEXT: Graph/Explicit/usm_memset.cpp
 // CHECK-NEXT: Graph/Explicit/work_group_size_prop.cpp
 // CHECK-NEXT: Graph/RecordReplay/buffer_copy_host2target.cpp
 // CHECK-NEXT: Graph/RecordReplay/buffer_copy_host2target_2d.cpp
@@ -185,13 +180,6 @@
 // CHECK-NEXT: Graph/RecordReplay/prefetch.cpp
 // CHECK-NEXT: Graph/RecordReplay/spec_constants_handler_api.cpp
 // CHECK-NEXT: Graph/RecordReplay/spec_constants_kernel_bundle_api.cpp
-// CHECK-NEXT: Graph/RecordReplay/usm_copy.cpp
-// CHECK-NEXT: Graph/RecordReplay/usm_copy_in_order.cpp
-// CHECK-NEXT: Graph/RecordReplay/usm_fill.cpp
-// CHECK-NEXT: Graph/RecordReplay/usm_fill_host.cpp
-// CHECK-NEXT: Graph/RecordReplay/usm_fill_shared.cpp
-// CHECK-NEXT: Graph/RecordReplay/usm_memset.cpp
-// CHECK-NEXT: Graph/RecordReplay/usm_memset_shortcut.cpp
 // CHECK-NEXT: Graph/RecordReplay/work_group_size_prop.cpp
 // CHECK-NEXT: Graph/UnsupportedDevice/device_query.cpp
 // CHECK-NEXT: GroupAlgorithm/SYCL2020/reduce_over_group_size.cpp
diff --git a/unified-runtime/source/adapters/opencl/command_buffer.cpp b/unified-runtime/source/adapters/opencl/command_buffer.cpp
@@ -204,33 +204,59 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendKernelLaunchExp(
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendUSMMemcpyExp(
-    [[maybe_unused]] ur_exp_command_buffer_handle_t hCommandBuffer,
-    [[maybe_unused]] void *pDst, [[maybe_unused]] const void *pSrc,
-    [[maybe_unused]] size_t size,
-    [[maybe_unused]] uint32_t numSyncPointsInWaitList,
-    [[maybe_unused]] const ur_exp_command_buffer_sync_point_t
-        *pSyncPointWaitList,
-    [[maybe_unused]] uint32_t numEventsInWaitList,
-    [[maybe_unused]] const ur_event_handle_t *phEventWaitList,
-    [[maybe_unused]] ur_exp_command_buffer_sync_point_t *pSyncPoint,
-    [[maybe_unused]] ur_event_handle_t *phEvent,
-    [[maybe_unused]] ur_exp_command_buffer_command_handle_t *phCommand) {
-  return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+    ur_exp_command_buffer_handle_t hCommandBuffer, void *pDst, const void *pSrc,
+    size_t size, uint32_t numSyncPointsInWaitList,
+    const ur_exp_command_buffer_sync_point_t *pSyncPointWaitList, uint32_t,
+    const ur_event_handle_t *, ur_exp_command_buffer_sync_point_t *pSyncPoint,
+    ur_event_handle_t *, ur_exp_command_buffer_command_handle_t *) {
+  // No extension entry-point exists for USM memcpy, use SVM memcpy in
+  // preparation for USVM.
+  cl_context CLContext = hCommandBuffer->hContext->CLContext;
+  cl_ext::clCommandSVMMemcpyKHR_fn clCommandSVMMemcpyKHR = nullptr;
+  UR_RETURN_ON_FAILURE(
+      cl_ext::getExtFuncFromContext<decltype(clCommandSVMMemcpyKHR)>(
+          CLContext, ur::cl::getAdapter()->fnCache.clCommandSVMMemcpyKHRCache,
+          cl_ext::CommandSVMMemcpyName, &clCommandSVMMemcpyKHR));
+
+  const bool IsInOrder = hCommandBuffer->IsInOrder;
+  cl_sync_point_khr *RetSyncPoint = IsInOrder ? nullptr : pSyncPoint;
+  const cl_sync_point_khr *SyncPointWaitList =
+      IsInOrder ? nullptr : pSyncPointWaitList;
+  uint32_t WaitListSize = IsInOrder ? 0 : numSyncPointsInWaitList;
+  CL_RETURN_ON_FAILURE(clCommandSVMMemcpyKHR(
+      hCommandBuffer->CLCommandBuffer, nullptr, nullptr, pDst, pSrc, size,
+      WaitListSize, SyncPointWaitList, RetSyncPoint, nullptr));
+
+  return UR_RESULT_SUCCESS;
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendUSMFillExp(
-    [[maybe_unused]] ur_exp_command_buffer_handle_t hCommandBuffer,
-    [[maybe_unused]] void *pMemory, [[maybe_unused]] const void *pPattern,
-    [[maybe_unused]] size_t patternSize, [[maybe_unused]] size_t size,
-    [[maybe_unused]] uint32_t numSyncPointsInWaitList,
-    [[maybe_unused]] const ur_exp_command_buffer_sync_point_t
-        *pSyncPointWaitList,
-    [[maybe_unused]] uint32_t numEventsInWaitList,
-    [[maybe_unused]] const ur_event_handle_t *phEventWaitList,
-    [[maybe_unused]] ur_exp_command_buffer_sync_point_t *pSyncPoint,
-    [[maybe_unused]] ur_event_handle_t *phEvent,
-    [[maybe_unused]] ur_exp_command_buffer_command_handle_t *phCommand) {
-  return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+    ur_exp_command_buffer_handle_t hCommandBuffer, void *pMemory,
+    const void *pPattern, size_t patternSize, size_t size,
+    uint32_t numSyncPointsInWaitList,
+    const ur_exp_command_buffer_sync_point_t *pSyncPointWaitList, uint32_t,
+    const ur_event_handle_t *, ur_exp_command_buffer_sync_point_t *pSyncPoint,
+    ur_event_handle_t *, ur_exp_command_buffer_command_handle_t *) {
+  // No extension entry-point exists for USM fill, use SVM fill in preparation
+  // for USVM.
+  cl_context CLContext = hCommandBuffer->hContext->CLContext;
+  cl_ext::clCommandSVMMemFillKHR_fn clCommandSVMMemFillKHR = nullptr;
+  UR_RETURN_ON_FAILURE(
+      cl_ext::getExtFuncFromContext<decltype(clCommandSVMMemFillKHR)>(
+          CLContext, ur::cl::getAdapter()->fnCache.clCommandSVMMemFillKHRCache,
+          cl_ext::CommandSVMMemFillName, &clCommandSVMMemFillKHR));
+
+  const bool IsInOrder = hCommandBuffer->IsInOrder;
+  cl_sync_point_khr *RetSyncPoint = IsInOrder ? nullptr : pSyncPoint;
+  const cl_sync_point_khr *SyncPointWaitList =
+      IsInOrder ? nullptr : pSyncPointWaitList;
+  uint32_t WaitListSize = IsInOrder ? 0 : numSyncPointsInWaitList;
+  CL_RETURN_ON_FAILURE(
+      clCommandSVMMemFillKHR(hCommandBuffer->CLCommandBuffer, nullptr, nullptr,
+                             pMemory, pPattern, patternSize, size, WaitListSize,
+                             SyncPointWaitList, RetSyncPoint, nullptr));
+
+  return UR_RESULT_SUCCESS;
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendMemBufferCopyExp(
diff --git a/unified-runtime/source/adapters/opencl/common.hpp b/unified-runtime/source/adapters/opencl/common.hpp
diff --git a/unified-runtime/source/adapters/opencl/extension_functions.def b/unified-runtime/source/adapters/opencl/extension_functions.def
diff --git a/unified-runtime/test/conformance/exp_command_buffer/commands.cpp b/unified-runtime/test/conformance/exp_command_buffer/commands.cpp
diff --git a/unified-runtime/test/conformance/exp_command_buffer/fill.cpp b/unified-runtime/test/conformance/exp_command_buffer/fill.cpp
