From 20522910d617702364208122f191aa57e1fa0de4 Mon Sep 17 00:00:00 2001 From: Mark Abraham Date: Thu, 19 Mar 2026 15:04:56 +0100 Subject: [PATCH 1/2] Provide example for shared system USM using malloc() It's important to be clear to users that with suitable hardware and kernel, one does not have to change anything to do with memory allocation. --- source/examples/usm-system-shared.cpp | 34 +++++++++++++++++++++++++++ source/iface/usm_basic_concept.rst | 24 +++++++++++++++---- 2 files changed, 53 insertions(+), 5 deletions(-) create mode 100644 source/examples/usm-system-shared.cpp diff --git a/source/examples/usm-system-shared.cpp b/source/examples/usm-system-shared.cpp new file mode 100644 index 00000000..6e55d826 --- /dev/null +++ b/source/examples/usm-system-shared.cpp @@ -0,0 +1,34 @@ +// SPDX-FileCopyrightText: 2026 The Khronos Group Inc. +// +// SPDX-License-Identifier: Apache-2.0 + +#include + +#include +#include + +int main() { + // Create a default queue to enqueue work to the default device + sycl::queue myQueue; + + // Allocate normal system memory - let the Linux kernel do the work! + // You can also use a std::vector, etc. + int *data = (int*) malloc(1024); + + myQueue.parallel_for(1024, [=](sycl::id<1> idx) { + // Initialize each buffer element with its own rank number starting at 0 + data[idx] = idx; + }); // End of the kernel function + + // Explicitly wait for kernel execution since there is no accessor involved + myQueue.wait(); + + // Print result + for (int i = 0; i < 1024; i++) + std::cout << "data[" << i << "] = " << data[i] << std::endl; + + // Clean up with normal system free + free(data); + + return 0; +} diff --git a/source/iface/usm_basic_concept.rst b/source/iface/usm_basic_concept.rst index 4d143d5c..c83e010c 100644 --- a/source/iface/usm_basic_concept.rst +++ b/source/iface/usm_basic_concept.rst @@ -1,5 +1,5 @@ .. - Copyright 2024 The Khronos Group Inc. + Copyright 2026 The Khronos Group Inc. SPDX-License-Identifier: CC-BY-4.0 .. _usm_basic_concept: @@ -170,7 +170,7 @@ supports ``sycl::aspect::usm_device_allocations``. .. rubric:: Example -See `usm-example-2`_. +See `usm-example-3`_. Shared allocations @@ -239,7 +239,7 @@ shared allocations can be queried through the aspect .. rubric:: Example -See `usm-example-1`_. +See `usm-example-2`_. System allocations @@ -260,15 +260,17 @@ through ``sycl::aspect::usm_system_allocations``. must still be allocated using their respective USM functions in order to guarantee their behavior. +See `usm-example-1`_. + .. _usm-example-1: ========= Example 1 ========= -Example of how shared memory can be used between host and device: +Example of how system shared memory can be used between host and device: -.. literalinclude:: /examples/usm-shared.cpp +.. literalinclude:: /examples/usm-system-shared.cpp :lines: 5- :linenos: @@ -278,6 +280,18 @@ Example of how shared memory can be used between host and device: Example 2 ========= +Example of how shared memory can be used between host and device: + +.. literalinclude:: /examples/usm-shared.cpp + :lines: 5- + :linenos: + +.. _usm-example-3: + +========= +Example 3 +========= + Example of using less capable device memory, which requires an explicit copy between the device and the host: From 6c61a46ea7ef5286f2109ca99657fa01884e2b0a Mon Sep 17 00:00:00 2001 From: Mark Abraham Date: Thu, 19 Mar 2026 15:30:37 +0100 Subject: [PATCH 2/2] Keep linter happy --- source/examples/usm-system-shared.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/source/examples/usm-system-shared.cpp b/source/examples/usm-system-shared.cpp index 6e55d826..be2b6109 100644 --- a/source/examples/usm-system-shared.cpp +++ b/source/examples/usm-system-shared.cpp @@ -13,7 +13,7 @@ int main() { // Allocate normal system memory - let the Linux kernel do the work! // You can also use a std::vector, etc. - int *data = (int*) malloc(1024); + int *data = (int *) malloc(1024); myQueue.parallel_for(1024, [=](sycl::id<1> idx) { // Initialize each buffer element with its own rank number starting at 0