Clarification on usage of ur `USMImportExp` and sycl `prepare_for_device_copy`

[nagic0]

Hi,

I am working on Unified Memory Architecture (UMA) platforms (e.g., Intel iGPUs) and trying to import an existing CPU buffer into the GPU as a USM buffer to be accessed directly by GPU kernels, like cudaHostRegister.

Background & Current Approach

My initial attempt using Level Zero directly was successful. I used zeMemAllocShared with the ZEX_HOST_MEM_ALLOC_FLAG_USE_HOST_PTR flag, which correctly achieved the desired zero-copy behavior for kernel execution.

However, since my upper-level application is built on the SYCL -> Unified Runtime -> Level Zero stack, I investigated how this is handled in UR. I noticed that sycl::ext::oneapi::experimental::prepare_for_device_copy eventually calls urUSMImportExp. Then I referred to the documentation here: (https://github.com/oneapi-src/unified-runtime/blob/main/scripts/core/EXP-USM-IMPORT-RELEASE.rst)

An Attempt

It appears my previous observation and assumption was incorrect. This might not be working as I expected.

I created a test case for USM shared memory to verify this. The logic is simple: verify that a buffer initialized by the CPU (data[i] = i) is correctly read by the GPU (checking if data[i] == i).

In the attached reproduction code (around line 50), I implemented three different allocation/import strategies for comparison:

• 1. Standard USM: Directly allocating a USM buffer using malloc_shared.
• 2. L0 Import (Manual): Allocating via standard CPU malloc and importing to GPU using Level Zero's zeMemAllocShared (with ZEX_HOST_MEM_ALLOC_FLAG_USE_HOST_PTR).
• 3. SYCL Import (Experimental): Allocating via standard CPU malloc and importing using sycl::ext::oneapi::experimental::prepare_for_device_copy (which wraps urUSMImportExp).

Results:

• Methods 1 and 2 work correctly: The GPU sees the correct data.
• Method 3 fails: The GPU reads all zeros.

It seems that while zeMemAllocShared handles the mapping correctly for compute access, prepare_for_device_copy (via urUSMImportExp) does not seem to expose the CPU-initialized data to the kernel in the same way on this platform.

Question

The documentation states that urUSMImportExp is designed for Data Transfer (e.g., optimizing memcpy operations in SYCL).
Since I am using the imported buffer for GPU Computation (accessing it directly within a kernel) rather than just for data movement.

Is this usage considered correct and spec-compliant?

Any clarification and suggestion would be appreciated. Thanks!

Code

#include <cstdlib>
#include <sycl/sycl.hpp>
#include <iostream>
#include <vector>
#include <level_zero/ze_api.h>

using namespace sycl;

constexpr size_t GB = 1024L * 1024L * 1024L;
constexpr size_t bytes = 2L * GB; // Buffer size
constexpr size_t N = bytes / sizeof(int);

constexpr uint32_t ZEX_HOST_MEM_ALLOC_FLAG_USE_HOST_PTR = ZE_BIT(30);

#define L0_SAFE_CALL(call)                                  \
    {                                                       \
        ze_result_t status = (call);                        \
        if (status != ZE_RESULT_SUCCESS) {                  \
            std::cerr << "Level Zero call failed: "        \
                      << #call << " returned " << status    \
                      << " at " << __FILE__ << ":"         \
                      << __LINE__ << std::endl;             \
            std::exit(1);                                  \
        }                                                   \
    }

void import_by_ze(sycl::queue& q, void* ptr, size_t size) {
    ze_context_handle_t ze_context = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(q.get_context());
    ze_device_handle_t ze_device = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(q.get_device());

    L0_SAFE_CALL(zeContextGetStatus(ze_context));
    L0_SAFE_CALL(zeDeviceGetStatus(ze_device));

    ze_device_mem_alloc_desc_t usm_device_desc = {};
    ze_host_mem_alloc_desc_t usm_host_desc = {};
    usm_host_desc.flags = ZEX_HOST_MEM_ALLOC_FLAG_USE_HOST_PTR;
    L0_SAFE_CALL(zeMemAllocShared(ze_context, &usm_device_desc, &usm_host_desc, size, 4096, ze_device, &ptr));
}

int main() {
    try {
        queue q(gpu_selector_v);
        std::cout << "Running on Device: " 
                  << q.get_device().get_info<info::device::name>() << std::endl;

        std::cout << "Allocating " << bytes / (1024.0 * 1024.0 * 1024.0) 
                  << " GB of Unified Shared Memory..." << std::endl;


        /*
        Three options to allocate USM memory:
        1. USM malloc_shared
        2. CPU alloc + import_by_ze
        3. CPU alloc + prepare_for_device_copy (Not working)
        */

        // (1) 
        // int* data = malloc_shared<int>(N, q);

        // (2)
        // int* data = static_cast<int*>(std::aligned_alloc(4096, bytes));
        // import_by_ze(q, data, bytes);

        // (3)
        int* data = static_cast<int*>(std::aligned_alloc(4096, bytes));
        sycl::ext::oneapi::experimental::prepare_for_device_copy(data, N * sizeof(int), q);
        
        int* data_out = malloc_shared<int>(N, q);

        int* err_count = malloc_shared<int>(1, q);
        *err_count = 0;

        if (data == nullptr || err_count == nullptr) {
            std::cerr << "Memory allocation failed!" << std::endl;
            return 1;
        }

        std::cout << "Host (CPU): Filling data pattern..." << std::endl;
        for (size_t i = 0; i < N; ++i) {
            data[i] = static_cast<int>(i);
        }

        std::cout << "Device (GPU): Verifying data pattern..." << std::endl;

        q.parallel_for(range<1>(N), [=](id<1> idx) {
            size_t i = idx[0];
            int val = data[i];
            int expected = static_cast<int>(i);

            if (val != expected) {
                auto err_ref = atomic_ref<int, memory_order::relaxed, memory_scope::device, access::address_space::global_space>(*err_count);
                err_ref.fetch_add(1);
            }
            data_out[i] = val;
        }).wait();

        if (*err_count == 0) {
            std::cout << "SUCCESS: All " << N << " elements matched!" << std::endl;
        } else {
            std::cout << "FAILURE: Found " << *err_count << " errors!" << std::endl;
        }

        for (size_t i = 0; i < 10; ++i) {
            std::cout << "data_out[" << i << "] = " << data_out[i] << std::endl;
        }

        // free(data, q);
        free(err_count, q);

    } catch (exception const& e) {
        std::cerr << "SYCL Exception: " << e.what() << std::endl;
        return 1;
    }

    return 0;
}

Compile commands:

icpx -fsycl -O3 usm_check.cpp -o usm_check -lze_loader

Platform

OS: Ubuntu 24.04.2 LTS (x86_64)
GCC version: (Ubuntu 14.2.0-4ubuntu2~24.04) 14.2.0
Clang version: Could not collect
CMake version: version 4.1.0
Libc version: glibc-2.39

Intel GPU driver version:
* intel-opencl-icd:	25.35.35096.9-1~24.04~ppa3
* libze1:	1.22.5-1~24.04~ppa1

Intel GPU models detected:
* [0] _XpuDeviceProperties(name='Intel(R) Arc(TM) Graphics', platform_name='Intel(R) oneAPI Unified Runtime over Level-Zero', type='gpu', driver_version='1.6.35096+9', total_memory=89670MB, max_compute_units=128, gpu_eu_count=128, gpu_subslice_count=8, max_work_group_size=1024, max_num_sub_groups=128, sub_group_sizes=[8 16 32], has_fp16=1, has_fp64=1, has_atomic64=1)

SYCL version: 2025.2 (20250200)

CPU:
Architecture:                            x86_64
CPU op-mode(s):                          32-bit, 64-bit
Address sizes:                           46 bits physical, 48 bits virtual
Byte Order:                              Little Endian
CPU(s):                                  16
On-line CPU(s) list:                     0-15
Vendor ID:                               GenuineIntel
Model name:                              Intel(R) Core(TM) Ultra 9 285H
CPU family:                              6
Model:                                   197
Thread(s) per core:                      1
Core(s) per socket:                      16
Socket(s):                               1
Stepping:                                2
CPU(s) scaling MHz:                      64%
CPU max MHz:                             5400.0000
CPU min MHz:                             400.0000
BogoMIPS:                                7372.80
Flags:                                   fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf tsc_known_freq pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault epb ssbd ibrs ibpb stibp ibrs_enhanced tpr_shadow flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid rdseed adx smap clflushopt clwb intel_pt sha_ni xsaveopt xsavec xgetbv1 xsaves split_lock_detect user_shstk avx_vnni lam wbnoinvd dtherm ida arat pln pts hwp hwp_notify hwp_act_window hwp_epp hwp_pkg_req hfi vnmi umip pku ospke waitpkg gfni vaes vpclmulqdq rdpid bus_lock_detect movdiri movdir64b fsrm md_clear serialize pconfig arch_lbr ibt flush_l1d arch_capabilities
Virtualization:                          VT-x
L1d cache:                               480 KiB (12 instances)
L1i cache:                               768 KiB (12 instances)
L2 cache:                                28 MiB (9 instances)
L3 cache:                                24 MiB (1 instance)
NUMA node(s):                            1
NUMA node0 CPU(s):                       0-15
Vulnerability Gather data sampling:      Not affected
Vulnerability Ghostwrite:                Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit:             Not affected
Vulnerability L1tf:                      Not affected
Vulnerability Mds:                       Not affected
Vulnerability Meltdown:                  Not affected
Vulnerability Mmio stale data:           Not affected
Vulnerability Old microcode:             Not affected
Vulnerability Reg file data sampling:    Not affected
Vulnerability Retbleed:                  Not affected
Vulnerability Spec rstack overflow:      Not affected
Vulnerability Spec store bypass:         Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:                Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:                Mitigation; Enhanced / Automatic IBRS; IBPB conditional; PBRSB-eIBRS Not affected; BHI BHI_DIS_S
Vulnerability Srbds:                     Not affected
Vulnerability Tsa:                       Not affected
Vulnerability Tsx async abort:           Not affected
Vulnerability Vmscape:                   Mitigation; IBPB before exit to userspace

[nagic0]

More information from compute-runtime:

prepare_for_device_copy uses zexDriverImportExternalPointer L0 API and it manages the external buffer with hostPointerManager->createHostPointerMultiAllocation (here).

And zeMemAllocShared with use host ptr flag uses standard USM allocation path with this->driverHandle->svmAllocsManager->createSharedUnifiedMemoryAllocation and insertSVMAlloc(retPtr, allocData); (here)