Is this page helpful?

4.8. Locality Domains#

A locality domain is a portion of a GPU that contains streaming multiprocessors (SMs) and device memory. An application can allocate device memory in a locality domain and create a green context with SM resources in the same locality domain. Co-locating computation near the memory it accesses can improve performance on devices that have more than one locality domain. The compute localization API is described in Localizing Compute Resources and memory localization in Allocating Localized Device Memory.

CUDA identifies the locality domains of a device by zero-based locality domain IDs. A locality domain ID is meaningful only together with its device ID. The same ID is used when localizing SM resources and memory, allowing an application to request matching placement for both.

4.8.1. Discovering Locality Domains#

Use cudaDeviceGetAttribute with the following cudaDeviceAttr values:

  • cudaDevAttrLocalityDomainCount returns the number of locality domains on the device. Valid locality domain IDs are in the range [0, cudaDevAttrLocalityDomainCount - 1].

  • cudaDevAttrLocalityDomainMultiprocessorCount returns the stable number of SMs available in each locality domain. This is the number of SMs that each locality domain on the device is guaranteed to contain.

Applications should query these attributes rather than assume a particular topology. The number of locality domains available on a device is affected by the system configuration. The same device type may be presented with a different number of locality domains under different scenarios (Section 4.8.5.3). All devices, including those with a single locality domain available, can use the locality domain APIs. On devices with a single locality domain, these APIs are the same as targeting the full device. Use the device attributes to programmatically query the locality domain properties of the device:

int localityDomainCount = 0;
int smsPerLocalityDomain = 0;

cudaDeviceGetAttribute(&localityDomainCount, cudaDevAttrLocalityDomainCount, device);
cudaDeviceGetAttribute(&smsPerLocalityDomain, cudaDevAttrLocalityDomainMultiprocessorCount, device);

4.8.2. Localizing Compute Resources#

To localize SM resources to a specific locality domain, use the CUDA green context API. SMs placement is expressed through the green context resource APIs described in Section 4.7. The main steps to construct a localized SM partition are:

  1. Obtain the device’s cudaDevResourceTypeSm resource with cudaDeviceGetDevResource.

  2. Describe a localized SM group with cudaDevSmResourceGroupParams (details).

  3. Pass the group parameters to cudaDevSmResourceSplit.

  4. Generate a resource descriptor with cudaDevResourceGenerateDesc, create a green context with cudaGreenCtxCreate, and finally create a stream via cudaExecutionCtxStreamCreate.

  5. Launch kernels on the stream.

Typical green context creation is described in Section 4.7.2. The additions necessary to localize SM resources are illustrated below:

To describe a localized SM group, set cudaDevSmResourceGroupLocalityDomainId in flags and set the ID in localityDomainId. Valid IDs are between 0 and cudaDevAttrLocalityDomainCount - 1.

cudaDevResource smResource;
cudaDeviceGetDevResource(device, &smResource, cudaDevResourceTypeSm);

{
    cudaDevSmResourceGroupParams params = {};
    // all fields in params must be zero-initialized
    params.flags = cudaDevSmResourceGroupLocalityDomainId;
    params.localityDomainId = localityDomainId;

    cudaDevResource result;
    cudaDevSmResourceSplit(&result, 1, &smResource, nullptr, 0, &params);
}

4.8.2.1. Discovery Mode#

Use the discovery mode of cudaDevSmResourceSplit when the number of SMs that satisfy a set of constraints is not known in advance. Enable discovery mode by setting smCount to zero. The API writes the discovered count to params.smCount. Passing nullptr for result performs a dry run that discovers the count without creating a device resource:

cudaDevSmResourceGroupParams params {};
params.flags = cudaDevSmResourceGroupLocalityDomainId;
params.localityDomainId = localityDomainId;
params.smCount = 0;

cudaDevSmResourceSplit(nullptr, 1, &smResource, nullptr, 0, &params);
unsigned int availableSms = params.smCount;

If result is non-null, the same call both discovers the count and returns a resource containing that many SMs.

4.8.2.2. Supplemental SMs#

A device can associate additional, supplemental SMs with a locality domain in excess of the SMs included in cudaDevAttrLocalityDomainMultiprocessorCount. A locality domain’s supplemental SMs provide the same memory-proximity benefit as stable SMs, but their number can vary between locality domains and devices. The actual set of SMs associated with a locality domain is the union of its stable and supplemental SMs.

To discover the total number of stable and supplemental SMs associated with a locality domain, include cudaDevSmResourceGroupLocalityDomainSupplementalSms among the discovery constraints:

cudaDevSmResourceGroupParams supplementalParams {};
supplementalParams.flags =
    cudaDevSmResourceGroupLocalityDomainId |
    cudaDevSmResourceGroupLocalityDomainSupplementalSms;
supplementalParams.localityDomainId = localityDomainId;
supplementalParams.smCount = 0;

cudaDevSmResourceSplit(nullptr, 1, &smResource, nullptr, 0,
                       &supplementalParams);
unsigned int totalLocalizedSms = supplementalParams.smCount;
unsigned int supplementalSms = totalLocalizedSms - smsPerLocalityDomain;

4.8.2.3. SM Split Constraints#

Multiple constraints may be specified in params in combination with requesting a locality domain. For example:

Set smCount to a positive value to request that many SMs, or to zero to use discovery mode as described above. Without supplemental SMs or backfill, a request larger than cudaDevAttrLocalityDomainMultiprocessorCount fails. To include supplemental SMs, add the cudaDevSmResourceGroupLocalityDomainSupplementalSms flag.

When coscheduledSmCount is 0 the API does not restrict SMs by coscheduled capability and the maximum number of localized SMs are available. The cluster capability of the partition can be dynamically queried later with cudaOccupancyMaxPotentialClusterSize. When coscheduledSmCount is nonzero, the API will select only localized SMs satisfying the requested coscheduledSmCount.

To select SMs from any locality domain in the input resource, set localityDomainId to cudaDevSmResourceGroupLocalityDomainIdAny while keeping cudaDevSmResourceGroupLocalityDomainId in flags. This value allows the split to draw from the union of every available locality domain while excluding supplemental SMs.

group.flags = cudaDevSmResourceGroupLocalityDomainId;
group.localityDomainId = cudaDevSmResourceGroupLocalityDomainIdAny;

Any unused SMs not returned in the result are returned in the API’s remainder parameter.

4.8.2.4. Backfill#

The cudaDevSmResourceGroupBackfill flag relaxes the split constraints so CUDA can fill a group that cannot otherwise reach the requested smCount. When splitting by locality domain, backfill behaves as follows:

  • Add cudaDevSmResourceGroupBackfill to fill the group first with non-supplemental SMs from the requested locality domain, then supplemental SMs from that domain, and finally SMs from other locality domains. The call fails with cudaErrorInvalidResourceConfiguration if no SM can be found in the requested locality domain.

4.8.2.5. Inspecting Split Results#

The result’s cudaDevResource.sm.flags and cudaDevResource.sm.localityDomainId report the resource constraints. localityDomainId identifies the requested locality domain when cudaDevSmResourceGroupLocalityDomainId is present in flags, otherwise its value is not meaningful. A resource created with backfill can contain SMs from other locality domains, so localityDomainId only reflects the inital requested locality domain, not every locality domain present in such a resource.

When combining SM resources in one descriptor, the resources must have matching constraints, including flags and localityDomainId values.

4.8.2.6. Creating Localized Streams#

The following code queries the number of locality domains, and constructs one stream per locality domain.

std::vector<cudaExecutionContext_t> contexts(localityDomainCount);
std::vector<cudaStream_t> streams(localityDomainCount);

{
    std::vector<cudaDevSmResourceGroupParams> params(localityDomainCount);
    std::vector<cudaDevResource> result(localityDomainCount);

    for (int i = 0; i < localityDomainCount; i++) {
        params[i] = {};
        params[i].flags = cudaDevSmResourceGroupLocalityDomainId;
        params[i].localityDomainId = i;
    }

    // A single split produces non-overlapping resources for all domains.
    cudaDevSmResourceSplit(result.data(), localityDomainCount, &smResource, nullptr, 0, params.data());

    for (int i = 0; i < localityDomainCount; i++) {
        cudaDevResourceDesc_t desc;
        cudaDevResourceGenerateDesc(&desc, &result[i], 1);
        cudaGreenCtxCreate(&contexts[i], desc, device, 0);
        cudaExecutionCtxStreamCreate(&streams[i], contexts[i], cudaStreamDefault, 0);
    }
}

Each stream in streams executes on SMs from the locality domain with the corresponding index. Destroy the streams with cudaStreamDestroy and then destroy their contexts with cudaExecutionCtxDestroy when they are no longer needed.

4.8.3. Allocating Localized Device Memory#

CUDA represents a localized memory location with a cudaMemLocation whose type is cudaMemLocationTypeDeviceLocalityDomain. For this location type, initialize both members of cudaMemLocation.localized:

  • deviceId identifies the GPU.

  • localityDomainId identifies a locality domain on that GPU.

The runtime API uses this location with stream-ordered memory pools (Section 4.3). The driver API additionally uses the corresponding CUmemLocation with virtual memory management allocations (Section 4.18).

Allocations created through other mechanisms, for example cudaMalloc or unified memory, can not be localized. MPS (Multi-Process Service) in MLOPart mode can be used to transparently localize applications that rely on cudaMalloc.

4.8.3.1. Creating a Localized Memory Pool#

To create a localized memory pool, set the location in cudaMemPoolProps and call cudaMemPoolCreate:

cudaMemPool_t pool {};
cudaMemPoolProps props {};

props.allocType = cudaMemAllocationTypePinned;
props.location.type = cudaMemLocationTypeDeviceLocalityDomain;
props.location.localized.deviceId = device;
props.location.localized.localityDomainId = localityDomainId;

cudaMemPoolCreate(&pool, &props);

Allocate from this pool with cudaMallocFromPoolAsync and release the allocation with cudaFreeAsync. The location can also be passed to cudaMemGetDefaultMemPool, cudaMemGetMemPool, and cudaMemSetMemPool when selecting the default or current pinned-memory pool for that locality domain. For the general memory-pool programming model, see Memory Pools.

4.8.3.2. Creating a Localized VMM Allocation#

The driver API can place physical memory created by cuMemCreate in a locality domain. Set CUmemAllocationProp::location.type to CU_MEM_LOCATION_TYPE_DEVICE_LOCALITY_DOMAIN and initialize location.localized.deviceId and location.localized.localityDomainId. The resulting physical allocation is reserved, mapped, and granted access in the same way as other VMM allocations. Access control set via cuMemSetAccess operates at full device granularity, not individual locality domains. See the localized VMM allocation section for the complete allocation sequence and granularity requirements.

Managed memory can not be localized: Requesting cudaMemAllocationTypeManaged with cudaMemLocationTypeDeviceLocalityDomain will return CUDA_ERROR_NOT_SUPPORTED. Similarly, cudaMemPrefetchAsync and cudaMemAdvise do not accept a device locality domain as a location and will return CUDA_ERROR_NOT_SUPPORTED.

See Locality Domains and Unified Memory for the relationship between locality domains and Unified Memory.

4.8.4. Querying Memory Placement#

Use cudaMemPoolGetAttribute with cudaMemPoolAttrLocalityDomainId to query a pool. It returns the locality domain ID for a localized pool and -1 for a pool that is not localized.

Use cudaPointerGetAttributes to query an allocation through its pointer. The returned cudaPointerAttributes.localityDomainOrdinal field contains the allocation’s locality domain ordinal, or -1 when the allocation is not localized:

cudaPointerAttributes attributes {};
cudaPointerGetAttributes(&attributes, ptr);
int localityDomainId = attributes.localityDomainOrdinal;

4.8.5. Programming Guidance#

To localize a workload, launch kernels using compute resources from the same locality domain as their memory allocations.

Localizing only compute or only memory resources typically does not improve performance.

If a kernel must use memory from several locality domains, measure whether compute localization is beneficial.

4.8.5.1. Localized Kernels With Fork/Join Pattern#

The fork/join pattern uses events and stream waits to express a parallel region. The application creates one stream from each localized green context by using cudaExecutionCtxStreamCreate. A separate main stream coordinates those localized streams. Create one fork event for the main stream and one join event per localized stream with cudaEventCreateWithFlags; timing should be disabled when the events are used only for ordering.

To fork, record one event on the parent stream and make every localized stream wait for that event. Work previously submitted to the parent is therefore ordered before the localized work. Submit independent work to the localized streams, using memory from their corresponding locality domains.

To join, record one event on each localized stream after its work and make the parent stream wait for all of those events. Work subsequently submitted to the parent is then ordered after the localized work.

produceInputs<<<grid, block, 0, parentStream>>>(...);
cudaEventRecord(forkEvent, parentStream);

// localizedStreams[i] belongs to the green context for locality domain i.
for (int i = 0; i < localityDomainCount; i++) {
    cudaStreamWaitEvent(localizedStreams[i], forkEvent);
    processPartition<<<grid, block, 0, localizedStreams[i]>>>(...);
    cudaEventRecord(joinEvents[i], localizedStreams[i]);
}

for (int i = 0; i < localityDomainCount; i++) {
    cudaStreamWaitEvent(parentStream, joinEvents[i]);
}
consumeResults<<<grid, block, 0, parentStream>>>(...);

4.8.5.2. CUDA Graphs#

CUDA Graphs (Section 4.2) support localization through stream capture and manual graph construction. For some applications, CUDA graphs may offer improved performance over stream-based programming.

4.8.5.2.1. Stream Capture#

A fork/join pattern can be captured to a CUDA graph. Begin capture on the parent stream, record the fork event, and wait on the fork event in each localized stream. Waiting on the fork event enrolls the localized streams in the same capture graph as the parent stream. Operations submitted to the localized streams become parallel graph branches. After launching all work on the localized streams, record the join event and wait on all the join events in the parent stream. These waits reconnect all branches to the parent stream. Every participating stream must be joined back to the origin stream before cudaStreamEndCapture is called.

cudaStreamBeginCapture(parentStream, cudaStreamCaptureModeGlobal);

// Record forkEvent, submit work to each localized stream, record the
// joinEvents, and make parentStream wait as in the fork/join example above.

cudaGraph_t graph;
cudaStreamEndCapture(parentStream, &graph);

The locality restriction of each child stream is retained by the captured kernel nodes. Graph replay uses the original localized SMs even when the executable graph is launched on a different stream. The graph launch stream supplies ordering but does not replace the execution resources recorded in the nodes.

4.8.5.2.2. Graph Construction#

The same graph topology can be constructed explicitly. Add one kernel node for each locality domain, using the corresponding localized execution context in cudaKernelNodeParamsV2.ctx.

cudaGraph_t graph = nullptr;
cudaGraphCreate(&graph, 0);

std::vector<cudaGraphNode_t> localizedNodes(localityDomainCount);
for (int i = 0; i < localityDomainCount; i++) {
    cudaGraphNodeParams params = { CU_GRAPH_NODE_TYPE_KERNEL };
    params.kernel.kern = kernel;
    params.kernel.gridDim = grid;
    params.kernel.blockDim = block;
    params.kernel.kernelParams = kernelArgs[i];
    params.kernel.ctx = localizedContexts[i];

    cudaGraphAddNode(&localizedNodes[i], graph, nullptr, nullptr, 0, &params);
}

4.8.5.3. Compatibility Notes#

vGPU, MIG, and devices in NUMA mode do not support localization and will report only 1 locality domain.

CUDA on Windows does not support localization and will report only 1 locality domain.

CUDA Checkpoint does not guarantee that localized green contexts retain their localization properties when restored on a different device. The application may exhibit reduced performance after a restore.

MPS (Multi-Process Service) in MLOPart mode presents each MLOPart device as a CUDA device with a single locality domain. MPS static partitioning treats the full device as a single locality domain.