4.7. 局部性域

4.7.1. 发现局部性域

使用cudaDeviceGetAttribute，并指定以下cudaDeviceAttr值：

• cudaDevAttrLocalityDomainCount返回设备上的局部性域数量。

• cudaDevAttrLocalityDomainMultiprocessorCount返回每个局部性域中的 SM 数量。

应用程序应查询这些属性，而不是假定某种特定的拓扑结构。设备上可用的局部性域数量受系统配置影响。在不同情形下，相同类型的设备可能呈现不同数量的局部性域（第 4.7.5.3 节）。所有设备，包括仅有一个可用局部性域的设备，都可以使用局部性域 API。在仅有一个局部性域的设备上，这些 API 与面向整个设备的操作相同。使用设备属性以编程方式查询设备的局部性域属性：

int localityDomainCount = 0;
int smsPerLocalityDomain = 0;

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

4.7.2. 计算资源局部化

要将 SM 资源局部化到特定的局部性域，请使用 CUDA 绿色上下文 API。SM 的放置通过第 4.6 节中介绍的绿色上下文资源 API 来表达。构建局部化 SM 分区的主要步骤如下：

1. 使用cudaDeviceGetDevResource获取设备的cudaDevResourceTypeSm资源。

2. 使用cudaDevSmResourceGroupParams描述局部化 SM 组（详细信息）。

3. 将组参数传递给cudaDevSmResourceSplit。

4. 使用cudaDevResourceGenerateDesc生成资源描述符，使用cudaGreenCtxCreate创建绿色上下文，最后通过cudaExecutionCtxStreamCreate创建流。

5. 在该流上启动核函数。

典型的绿色上下文创建过程在第 4.6.2 节中介绍。为使 SM 资源局部化而必须添加的内容如下所示：

要描述局部化 SM 组，请在flags中设置cudaDevSmResourceGroupLocalityDomainId，并在localityDomainId中设置 ID。有效 ID 介于0和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);
}

可以在请求局部性域的同时，在params中指定多个约束。例如：

smCount约束可以为0（发现模式）或正值。在发现模式下，API 在params.smCount中返回该局部性域中的 SM 数量。对于完整、未分区的设备资源，此值与cudaDevAttrLocalityDomainMultiprocessorCount相同。当smCount为正值时，API 最多分配cudaDevAttrLocalityDomainMultiprocessorCount所指定的数量。更大的smCount请求将失败。

当coscheduledSmCount为0时，API 不会根据协同调度能力限制 SM，因此可使用最大数量的局部化 SM。之后可以通过cudaOccupancyMaxPotentialClusterSize动态查询该分区的簇能力。当coscheduledSmCount非零时，API 将仅选择满足所请求的coscheduledSmCount的局部化 SM。

默认情况下，result中的所有 SM 都必须满足局部性域约束以及指定的任何附加约束。要放宽此行为，请将cudaDevSmResourceGroupBackfill添加到params.flags。启用回填后，API 首先从所请求的局部性域填充该组，然后使用不归属于任何局部性域的 SM，最后使用其他局部性域中的 SM。如果在所请求的局部性域中找不到任何 SM，该调用将失败并返回cudaErrorInvalidResourceConfiguration。

所有未在result中返回的未使用 SM 都会在 API 的remainder参数中返回。

result的cudaDevResource.sm.flags和cudaDevResource.sm.localityDomainId报告资源的放置情况。将 SM 资源组合到一个描述符中时，这些资源的约束必须匹配，包括flags和localityDomainId的值。

以下代码查询局部性域的数量，并为每个局部性域构建一个流。

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);
    }
}

streams中的每个流都在具有对应索引的局部性域中的 SM 上执行。不再需要这些流时，使用cudaStreamDestroy销毁流，然后使用cudaExecutionCtxDestroy销毁其上下文。

4.7.3. 分配局部化设备内存

CUDA 使用type为cudaMemLocationTypeDeviceLocalityDomain的cudaMemLocation来表示局部化内存位置。对于这种位置类型，请初始化cudaMemLocation.localized的两个成员：

• deviceId标识 GPU。

• localityDomainId标识该 GPU 上的一个局部性域。

运行时 API 将此位置用于流有序内存池（第 4.3 节）。驱动程序 API 还将相应的CUmemLocation用于虚拟内存管理分配（第 4.17 节）。

通过其他机制（例如cudaMalloc或统一内存）创建的分配无法局部化。MLOPart 模式下的MPS（多进程服务）可用于对依赖cudaMalloc的应用程序透明地进行局部化。

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);

4.7.4. 查询内存放置情况

使用cudaMemPoolGetAttribute并指定cudaMemPoolAttrLocalityDomainId来查询池。对于局部化池，它返回局部性域 ID；对于非局部化池，它返回-1。

使用cudaPointerGetAttributes通过指针查询分配。返回的cudaPointerAttributes.localityDomainOrdinal字段包含该分配的局部性域序号；如果该分配未局部化，则为-1：

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

4.7.5. 编程指导

要使工作负载局部化，请使用与核函数的内存分配位于同一局部性域的计算资源来启动核函数。

仅对计算资源或仅对内存资源进行局部化通常不会提高性能。

如果核函数必须使用来自多个局部性域的内存，请测量计算局部化是否有益。

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>>>(...);

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);

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);
}