Is this page helpful?

Usage with CUDA APIs#

This chapter describes how to detect an iGPU, choose a CUDA allocation path, query memory budgets, and respond to trim callbacks on Windows.

iGPU Detection#

Do not assume that only ARM platforms have an iGPU and therefore unified memory. Windows on ARM supports x86_64 emulation, so an application that is not fully ported to ARM should still detect and account for an iGPU. The emulation layer translates CPU instructions only. GPU device code runs natively on the iGPU and is subject to the same unified memory architecture and budgeting constraints as a native ARM64 application. To detect an iGPU in CUDA, call cuDeviceGetAttribute and check CU_DEVICE_ATTRIBUTE_INTEGRATED. If CU_DEVICE_ATTRIBUTE_INTEGRATED is set to 1, the application is running on a Windows iGPU system. To detect an iGPU in DirectX, call CheckFeatureSupport with D3D12_FEATURE_ARCHITECTURE1 and check the archInfo.UMA and archInfo.CacheCoherentUMA flags.

if (FAILED(D3DDev_0001->CheckFeatureSupport(D3D12_FEATURE_ARCHITECTURE1, &archInfo, sizeof(archInfo)))) {
    printf("ERROR: failed to retrieve D3D12_FEATURE_ARCHITECTURE1\n");
} else {
    if (archInfo.UMA) {
        printf("\tarchInfo.UMA == TRUE\n");
    } else {
        printf("\tarchInfo.UMA == FALSE\n");
    }

    if (archInfo.CacheCoherentUMA) {
        printf("\tarchInfo.CacheCoherentUMA == TRUE\n");
    } else {
        printf("\tarchInfo.CacheCoherentUMA == FALSE\n");
    }
}

Allocating Memory on iGPU#

On RTX Spark, the available GPU memory is split into two primary segments:

  • dedicated, a carveout from system memory accessible only by the GPU, and

  • shared, backed by system memory and accessible by both the CPU and the GPU.

Unlike Windows systems with a dedicated GPU, where shared system memory is classified as non-local, the RTX Spark iGPU treats shared memory as local to the GPU. The segment that contains an allocation affects page size, reporting, and budgeting.

The following sections describe three recommended allocation paths.

Avoid using cudaMallocManaged. While technically supported, it relies on a compatibility path that can lead to performance degradation.

cudaMalloc / cudaMallocAsync — GPU-Exclusive Allocations#

Use cuMemAlloc(Async), cudaMalloc(Async), or cuMemCreate for memory that is accessed exclusively (or predominantly) by the GPU. These allocations first land in the dedicated memory segment and allow the driver to apply internal alignment, page-size, and mapping optimizations. When the dedicated memory segment is exhausted, new allocations spill to the shared segment. Allocations in the shared segment are backed by smaller pages than allocations in the dedicated segment, so access can be slower. Allocation placement is non-configurable—the driver determines whether memory resides in the dedicated or shared segment.

cudaMallocHost — CPU/GPU Shared Allocations#

Use cudaMallocHost for CPU/GPU coherent memory that both processors access. These allocations land in the shared memory segment. Get the GPU-side pointer by calling cudaHostGetDevicePointer. The GPU can access the memory without an explicit copy.

Note

Allocations made with cudaMallocHost are not L2-cacheable on the iGPU, so they are useful for copy and streaming due to a higher memory bandwidth. For L2-cacheable memory, use VirtualAlloc2 + cudaHostRegister / cuMemHostRegister.

VirtualAlloc2 + cudaHostRegister / cuMemHostRegister — Shared CPU/GPU Access#

For GPU-L2-cacheable CPU/GPU coherent memory, use the CUDA cudaHostRegister or cuMemHostRegister APIs to share buffers between the CPU and GPU. Allocate the host memory with VirtualAlloc2 and align the base address and buffer size by 2MB for large buffers. Obtain the GPU-side pointer using cudaHostGetDevicePointer.

../_images/UMAShareDataPointers.png

Virtual Address Space pointers for unified memory on RTX Spark#

Note

Limit CPU/GPU shared allocations to what the workload requires. Use cudaHostRegister / cuMemHostRegister only for memory that both the CPU and GPU must access. Allocating more shared memory than required reduces the budget available to the rest of the system.

Recommendation#

Use CPU/GPU coherent allocations only when necessary, and prefer cudaMallocHost for those allocations unless you need L2-cacheable memory. For GPU-exclusive memory, use cudaMalloc(Async) or cuMemAlloc(Async) so the allocation can land in the dedicated segment. One use case for CPU/GPU coherent allocations is input and output buffers for AI inference. Keep intermediate buffers and model weights GPU-exclusive.

Memory Reporting#

Windows exposes GPU memory information through several APIs. Each has different visibility into the dedicated and shared memory segments.

cudaMemGetInfo#

cudaMemGetInfo reports total and free memory for the dedicated and shared memory budget. Therefore, the free memory reported by cudaMemGetInfo only indicates how much memory can be allocated using an API such as cudaMalloc. CPU/GPU coherent allocations have less available memory because they are limited to the shared memory segment.

NVML — nvmlDeviceGetMemoryInfo_v2#

nvmlDeviceGetMemoryInfo_v2 reports dedicated (device) memory. On UMA, that value is the dedicated carveout. It does not report the shared memory segment.

DirectX — QueryVideoMemoryInfo#

On UMA systems, Windows folds both dedicated and shared memory into the local segment. The non-local segment (DXGI_MEMORY_SEGMENT_GROUP_NON_LOCAL) reports zero for budget and usage. Do not assume shared memory is tracked under non-local.

DXGI_QUERY_VIDEO_MEMORY_INFO local{};
adapter->QueryVideoMemoryInfo(0, DXGI_MEMORY_SEGMENT_GROUP_LOCAL, &local);
// local.Budget        - total available memory (dedicated + shared)
// local.CurrentUsage  - current usage across both segments

To distinguish between the dedicated carveout and shared system memory, query the adapter descriptor:

DXGI_ADAPTER_DESC1 desc{};
adapter->GetDesc1(&desc);
// desc.DedicatedVideoMemory  - size of the dedicated carveout (VRAM)
// desc.SharedSystemMemory    - size of the shared system memory pool

Use desc.DedicatedVideoMemory to read the carveout size. Use desc.SharedSystemMemory to read the shared memory pool size. Use local.Budget and local.CurrentUsage to track combined available and used memory at runtime.

Recommendation#

On Windows iGPU, use DirectX QueryVideoMemoryInfo when you need budget and usage for the combined local segment, and use the adapter descriptor when you need dedicated and shared sizes separately. Use cudaMemGetInfo for a single CUDA view of the dedicated and shared budget. If you need dedicated memory only, NVML is sufficient.

Most NVIDIA SDKs (TensorRT, TensorRT RTX, cuDNN, cuBLAS) do not query GPU memory themselves; they either allocate and error out when memory is insufficient, or expect the caller to supply workspace memory. PyTorch exposes both cudaMemGetInfo and NVML as query paths — be aware of the limitations described above when interpreting their output.

Memory Budgeting and Trim Callbacks#

Windows limits the GPU budget based on the total available system memory, because exposing too much system memory to the GPU can starve the rest of the system. The available GPU budget consists of dedicated (carveout) memory and shared system memory.

For both dedicated and shared segments, avoid allocating all available memory. The memory available to the CPU can be less than the memory available to the GPU. Allocating the full GPU budget can leave too little host memory and can make the system unresponsive.

Trim Callbacks#

Instead of polling memory at startup, subscribe to asynchronous budget-change notifications using cudaDeviceRegisterAsyncNotification and listen for cudaAsyncNotificationType::cudaAsyncNotificationTypeOverBudget.

Trim callbacks are fired whenever there is a relevant budget change in either the dedicated or the shared memory segment. There is no single fixed threshold that triggers the callback. The system considers overall memory pressure, other applications’ demands, and the current usage relative to the budget. Do not hard-code assumptions about when the callback will fire.

When a trim callback is received, the application should:

  1. Query current memory state: Use QueryVideoMemoryInfo (or NVML for dedicated memory) to determine the updated budgets and current usage for both segments.

  2. React to the budget change: Free or migrate allocations to bring usage back within the reported budgets. Prioritize releasing memory from the over-budget segment.

  3. Continue listening: Callbacks might be sent repeatedly while the application remains over budget.

Use trim callbacks on dGPU devices because CUDA can swap GPU allocations to shared memory and oversubscribe VRAM. Swapping GPU allocations to shared memory can increase latency compared with keeping those allocations in VRAM. Either keep allocations within VRAM or inform the user that oversubscription is in use.

Do not query available memory, allocate that amount, and rely on trim callbacks to recover. That pattern is common on iGPU, dGPU, and eGPU systems, and it is risky: the initial query can include memory that is allocated but not currently used by the GPU. Allocating to that reported budget can exceed the memory that is actually free and evict other allocations.