NVIDIA CUDA 开发者可以利用多种工具和库来简化开发与部署，使用户能够专注于应用程序的“内容”和“方式”。
多进程服务 (MPS) 就是一个例子，用户可以通过跨进程共享 GPU 资源来提升 GPU 利用率。重要的是，这一过程可以透明地完成，应用程序无需了解 MPS，也无需修改代码。
MLOPart 简介
NVIDIA Blackwell GPU 提供高带宽，适合训练当前的大语言模型。然而，在某些情况下，应用程序无法充分利用 Blackwell 的全部带宽，且对延迟更为敏感。
内存局部性优化分区 (MLOPart) 设备源自 NVIDIA CUDA 架构的 GPU，并针对低延迟性能进行了优化。作为 CUDA MPS 的一项功能，MLOPart 支持多 GPU 感知的应用程序访问 MLOPart 设备。
在现实世界中，判断应用程序是否受延迟限制或带宽限制并不总是容易。MLOPart 旨在通过 MPS 控制器实现启用与禁用，无需重写应用程序。开发者可进行简单的 A/B 测试，以确认应用程序是否从 MLOPart 中受益。
MLOPart 设备枚举
MLOPart 的定义是，启用后，支持 MLOPart 的设备会显示为多个独立的 CUDA 设备，各自拥有独立的计算和内存资源。从这一角度来看，它与 NVIDIA 的多实例 GPU（MIG）类似。稍后我们将在本文中对 MLOPart 与 MIG 进行比较。
MLOPart 基于 GPU 的底层架构创建 CUDA 设备。在可能的情况下，CUDA 设备会沿着可能对内存延迟产生负面影响的边界进行划分，边界的每一侧均包含代表 MLOPart 设备的内存和计算资源。对于 Blackwell 而言，划分是沿着子边界进行的。
如果 GPU 不具备此类边界，则不会创建 MLOPart 设备，GPU 将正常呈现给 CUDA 应用程序。NVIDIA DGX B200 和 NVIDIA B300 的每个 GPU 最多可支持两台 MLOPart 设备。该数量可能随未来架构的更新而变化，因此建议开发者不要将 GPU 支持的 MLOPart 设备数量进行硬编码。
MLOPart 设备的功能和特性
MLOPart 设备与底层设备相似，但存在一些值得注意的例外。尽管原则上开发者无需重写应用程序即可使用 MLOPart 设备，但仍需注意，它们并不具备底层设备的全部功能和特性。
与底层设备共享的功能及特性包括：
计算能力
MLOPart 设备具有相同的计算能力，并且能够执行与底层设备相同的 GPU 二进制文件。例如，支持计算能力 10.0 的 MLOPart 设备即具备计算能力 10.0。
点对点能力
MLOPart 设备将能够与底层设备保持相同的点对点通信能力。例如，若两个物理设备通过 NVIDIA NVLink 相连，则由这两个底层设备生成的任何 MLOPart 设备也将通过 NVLink 实现互联。
属于同一底层设备的 MLOPart 设备之间不适用此规则。在此情况下，它们仍可进行点对点通信，但无需依赖 NVLink 或 PCIe 等点对点通信方式。
当对端设备为同一底层设备的 MLOPart 设备时，相较于通过其他方式连接的对端设备，其延迟更低，对端带宽更高。
PCI ID
MLOPart 设备与底层设备共享相同的 PCI ID（bus.device.domain）。
与底层设备不同的功能和特性包含以下内容。
流式传输多处理器数量
每个 MLOPart 设备的流式多处理器 (SM) 数量将少于底层设备。此外，在共享同一底层设备的所有 MLOPart 设备中，SM 的总量也可能少于底层设备中的 SM 总量。
属于同一底层设备的 MLOPart 设备之间 SM 数量相同，且在相同型号的 NVIDIA GPU 中保持一致。
例如，若启用 MLOPart，配备 8 个 Blackwell GPU（通常每个 GPU 含 148 个 SM）的 NVIDIA HGX B200 系统，将生成 16 台 MLOPart 设备，每台配备 70 个 SM。
可用内存
MLOPart 设备对底层设备的总显存进行分区，并仅从该分区中分配内存，但 CUDA 托管内存分配除外。每个 MLOPart 设备的内存容量均小于底层设备。同一底层设备下的各个 MLOPart 设备具有相同的总显存大小。
在当前版本的 MLOPart 中，一个 MLOPart 设备上分配的内存可能会影响同一底层设备上另一个 MLOPart 设备通过 cuMemGetInfo 和 cudaMemGetInfo 报告的可用内存，即使这些设备具有独立的分区。未来的驱动程序将在 MLOPart 设备之间实现更严格的内存隔离。
虚拟地址空间
同一底层设备上的 MLOPart 设备共享一个虚拟地址空间。这意味着，在某个 MLOPart 设备上分配的内存若发生缓冲区溢出，可能会影响同一进程中在其他 MLOPart 设备上分配的内存，造成数据损坏。
通用唯一标识符
每个 MLOPart 设备都具有一个通用唯一标识符 (UUID)，可通过 CUDA API 进行查询。该标识符可用于唯一识别 MLOPart 设备，并利用 CUDA_VISIBLE_DEVICES 对可用的 CUDA 设备进行筛选。
使用 MLOPart 进行部署
与其他 CUDA MPS 功能类似，用户可通过 MPS 控制器命令来控制其行为。
使用 start_server 命令启动 MPS 服务器。在 CUDA 13.1 中，我们为此命令引入了 -mlopart 选项，使用户能够启动支持 MLOPart 的 MPS 服务器，从而创建相应的 MPS 客户端。由于该操作在每台服务器上独立进行，不同用户可根据自身需求配置不同的 MLOPart 设置。
在 CUDA 13.0 中，我们引入了 device_query MPS 控制器命令，用于提供 MPS 枚举的 CUDA 设备的相关信息。服务器创建后，可使用 device_query 获取将向该服务器客户端公开的设备信息，例如设备名称、设备序数和 UUID。
$ echo device_query | nvidia-cuda-mps-control
Default
Device Ordinal PCI IDs UUID Name Attributes
0 0000:1b.00.00 GPU-ebebf640-14d4-de34-f16e-a5e7da272ac4 NVIDIA B200
1 0000:43.00.00 GPU-6d3a75da-dd2e-173e-e797-c0b8ed47a100 NVIDIA B200
2 0000:52.00.00 GPU-a517c26e-0f2f-945a-1672-ea75149f54d6 NVIDIA B200
3 0000:61.00.00 GPU-999b1bd5-82d8-3db2-e2ec-fdae5d1103b1 NVIDIA B200
4 0000:9d.00.00 GPU-b5830513-614b-38ac-b177-5cc2f850ea3d NVIDIA B200
5 0000:c3.00.00 GPU-05f3779e-bfa6-f9c8-256f-6cee98b8871d NVIDIA B200
6 0000:d1.00.00 GPU-2facdb95-1af2-26e3-2c9d-e02f4651675d NVIDIA B200
7 0000:df.00.00 GPU-7e555b40-ffe0-e066-4db3-4ddd96344f0d NVIDIA B200
Server 14056
Device Ordinal PCI IDs UUID Name Attributes
N/A 0000:1b.00.00 GPU-ebebf640-14d4-de34-f16e-a5e7da272ac4 NVIDIA B200 M
0 0000:1b.00.00 GPU-1bd9c0d8-c86a-5a37-acee-411ebcef5fd0 NVIDIA B200 MLOPart 0 MD
1 0000:1b.00.00 GPU-58e7f54c-f60f-56b7-a4c4-b3fb418fde3e NVIDIA B200 MLOPart 1 MD
N/A 0000:43.00.00 GPU-6d3a75da-dd2e-173e-e797-c0b8ed47a100 NVIDIA B200 M
2 0000:43.00.00 GPU-68fb01e9-499c-56d4-b768-8fca70a5ddff NVIDIA B200 MLOPart 0 MD
3 0000:43.00.00 GPU-6cf0c4ea-3a05-52b1-aec6-63acf60df19b NVIDIA B200 MLOPart 1 MD
N/A 0000:52.00.00 GPU-a517c26e-0f2f-945a-1672-ea75149f54d6 NVIDIA B200 M
4 0000:52.00.00 GPU-dd670b14-ca31-5dfd-a49b-7220701f4fc6 NVIDIA B200 MLOPart 0 MD
5 0000:52.00.00 GPU-d7433996-1714-5baa-9812-22cecdc792d3 NVIDIA B200 MLOPart 1 MD
N/A 0000:61.00.00 GPU-999b1bd5-82d8-3db2-e2ec-fdae5d1103b1 NVIDIA B200 M
6 0000:61.00.00 GPU-cff5ab0b-a509-54c8-a9c0-c5ebe3fbd3a0 NVIDIA B200 MLOPart 0 MD
7 0000:61.00.00 GPU-7933cfe7-5139-50d8-ad90-0f7f1ddba559 NVIDIA B200 MLOPart 1 MD
N/A 0000:9d.00.00 GPU-b5830513-614b-38ac-b177-5cc2f850ea3d NVIDIA B200 M
8 0000:9d.00.00 GPU-f973284b-7385-576b-80d7-3ea083bcea94 NVIDIA B200 MLOPart 0 MD
9 0000:9d.00.00 GPU-668e4145-b221-5495-a3fe-a5cdc0e6f6eb NVIDIA B200 MLOPart 1 MD
N/A 0000:c3.00.00 GPU-05f3779e-bfa6-f9c8-256f-6cee98b8871d NVIDIA B200 M
10 0000:c3.00.00 GPU-53858feb-87eb-5963-8d47-6fbf4b24cd4a NVIDIA B200 MLOPart 0 MD
11 0000:c3.00.00 GPU-700b029a-be98-5d13-9a4e-5e8e21386e34 NVIDIA B200 MLOPart 1 MD
N/A 0000:d1.00.00 GPU-2facdb95-1af2-26e3-2c9d-e02f4651675d NVIDIA B200 M
12 0000:d1.00.00 GPU-563db4f2-f70a-564d-aa4a-dbd52d6dfc0b NVIDIA B200 MLOPart 0 MD
13 0000:d1.00.00 GPU-b643e07a-6eda-5cd8-bdde-1788590d0b4b NVIDIA B200 MLOPart 1 MD
N/A 0000:df.00.00 GPU-7e555b40-ffe0-e066-4db3-4ddd96344f0d NVIDIA B200 M
14 0000:df.00.00 GPU-f8f5b46d-7774-57a1-97d2-88f23c3457f0 NVIDIA B200 MLOPart 0 MD
15 0000:df.00.00 GPU-46d7f9b7-0303-5432-b50a-16381f37e365 NVIDIA B200 MLOPart 1 MD
启用 MLOPart 后，device_query 会在其衍生设备下方显示 MLOPart 设备。此方法是确定应用程序启动时用于 CUDA_VISIBLE_DEVICES 的UUID 值的推荐方式。由于 CUDA 列举的设备数量多于系统实际存在的设备，因此在设备枚举过程中存在模糊性。
请注意，MLOPart 设备仅存在于 MPS 和 CUDA 环境中。nvidia-smi 未提供有关 MLOPart 设备的信息。
随后，扩展了 ps MPS 控制器命令，以显示进程是否正在使用 MLOPart 设备。
$ while1 -a &
[1] 52845
$ echo ps | nvidia-cuda-mps-control
PID ID SERVER DEVICE NAMESPACE COMMAND ATTRIBUTES
52845 1 52837 GPU-b13add01-c28c 4026531836 while1 MD
使用中的 MLOPart
现在，我们来探讨 MLOPart 对内存延迟和带宽的影响。
延迟
作为示例，我们来看看 MLOPart 如何通过在循环中执行一些原子操作的简单内核来影响内存延迟。
首先，我们定义核函数和辅助程序：
#include <cuda_runtime.h>
#include <vector>
#include <cstdio>
// Helper macro to check for CUDA errors
#define CUDA_CHECK_FAILURE(x) \
if (cudaSuccess != (cudaError_t)x)\
{\
const char* errName = cudaGetErrorName(x);\
const char* errStr = cudaGetErrorString(x);\
printf("%s:%d - %s: %s\n", __FILE__, __LINE__, errName, errStr);\
exit(EXIT_FAILURE);\
}
// Device memory variable to use to prevent the compiler from optimizing away the memory access
__device__ volatile int dummy;
// Trivial kernel to touch the memory so we can measure latency
__global__ void accessMemoryHighLatency(int *startAddress, size_t memorySizeInBytes) {
for (int i = 0 ; i < memorySizeInBytes / sizeof(int) ; ++i) {
dummy = atomicAdd(&startAddress[i], 1);
}
}
原子操作对延迟敏感，因此可以清晰地衡量使用 MLOPart 与不使用 MLOPart 之间的差异。以下函数利用 CUDA 事件来测量内核 accessMemoryHighLatency 的执行时间。
// Function to launch the kernel and measure the runtime using CUDA events
float measureKernelRuntime(int *memoryDevPtr, size_t memorySizeInBytes, int numBlocks, int numThreads) {
cudaEvent_t start = NULL, stop = NULL;
float time = 0;
CUDA_CHECK_FAILURE(cudaEventCreate(&start));
CUDA_CHECK_FAILURE(cudaEventCreate(&stop));
CUDA_CHECK_FAILURE(cudaEventRecord(start, 0));
accessMemoryHighLatency<<<numBlocks, numThreads>>>(memoryDevPtr, memorySizeInBytes);
CUDA_CHECK_FAILURE(cudaPeekAtLastError());
CUDA_CHECK_FAILURE(cudaEventRecord(stop, 0));
CUDA_CHECK_FAILURE(cudaEventSynchronize(stop));
CUDA_CHECK_FAILURE(cudaEventElapsedTime(&time, start, stop));
CUDA_CHECK_FAILURE(cudaEventDestroy(start));
CUDA_CHECK_FAILURE(cudaEventDestroy(stop));
return time;
}
最后，我们可以通过构建一个简单的多 GPU 感知程序来整合这一切。
int main(int argc, char *argv[]) {
size_t memorySizeInBytes = 32 * 1024 * 1024; // 32 MB
int numBlocks = 32;
int numThreads = 1;
int numDevices = 0;
float totalTime = 0;
CUDA_CHECK_FAILURE(cudaGetDeviceCount(&numDevices));
// Measure the runtime for each device
for (int i = 0; i < numDevices; i++) {
// Set the current device
CUDA_CHECK_FAILURE(cudaSetDevice(i));
// Allocate memory on the device
int *memoryDevPtr;
CUDA_CHECK_FAILURE(cudaMalloc(&memoryDevPtr, memorySizeInBytes));
// Measure the runtime
float time = measureKernelRuntime(memoryDevPtr, memorySizeInBytes, numBlocks, numThreads);
totalTime += time;
printf("Device %d - Total time: %f milliseconds\n", i, time);
// Free the memory
CUDA_CHECK_FAILURE(cudaFree(memoryDevPtr));
}
printf("Average time: %f milliseconds\n", totalTime / numDevices);
return EXIT_SUCCESS;
}
我们将此文件命名为 atomic_memory_access.cu，并使用 nvcc atomic_memory_access.cu -arch=sm_100 -o atomic_memory_access. 进行编译
要建立基准，我们使用 MPS 运行示例，但不采用 MLOPart。
$ nvidia-cuda-mps-control -d
# Optional step of explicitly creating an MPS server. This is also done implicitly when we launch a CUDA application while MPS is active.
$ echo start_server -uid $UID | nvidia-cuda-mps-control
$ ./atomic_memory_access
Device 0 - Total time: 2320.550537 milliseconds
Device 1 - Total time: 2323.710938 milliseconds
Device 2 - Total time: 2334.533447 milliseconds
Device 3 - Total time: 2304.551025 milliseconds
Device 4 - Total time: 2304.328125 milliseconds
Device 5 - Total time: 2316.102295 milliseconds
Device 6 - Total time: 2306.165283 milliseconds
Device 7 - Total time: 2306.362061 milliseconds
Average time: 2314.537842 milliseconds
在这里，我们看到每个设备的平均时间大约为 2300 毫秒。接下来，让我们启用 MLOPart 并重新运行。
# Quit the MPS controller to cleanup the previous server.
$ echo quit | nvidia-cuda-mps-control
# Now repeat the above steps, with MLOPart enabled.
$ nvidia-cuda-mps-control -d
# Note that we must explicitly start the server with "-mlopart".
$ echo start_server -uid $UID -mlopart | nvidia-cuda-mps-control
$ ./atomic_memory_access
Device 0 - Total time: 1500.194946 milliseconds
Device 1 - Total time: 1475.914062 milliseconds
Device 2 - Total time: 1479.729492 milliseconds
Device 3 - Total time: 1480.196045 milliseconds
Device 4 - Total time: 1478.959106 milliseconds
Device 5 - Total time: 1490.808716 milliseconds
Device 6 - Total time: 1468.943237 milliseconds
Device 7 - Total time: 1479.297241 milliseconds
Device 8 - Total time: 1467.947632 milliseconds
Device 9 - Total time: 1476.900757 milliseconds
Device 10 - Total time: 1477.081421 milliseconds
Device 11 - Total time: 1490.295044 milliseconds
Device 12 - Total time: 1484.558594 milliseconds
Device 13 - Total time: 1481.660156 milliseconds
Device 14 - Total time: 1476.067383 milliseconds
Device 15 - Total time: 1484.143921 milliseconds
Average time: 1480.793457 milliseconds
在本示例中，我们看到使用 MLOPart 时，每个设备的执行时间明显缩短。虽然这是一个精心设计的示例，但在决定如何部署特定应用程序时，仍需权衡是否采用 MLOPart。
带宽
鉴于 MLOPart 设备的内存少于完整设备，其 DRAM 带宽也相应低于未使用 MLOPart 的设备。
与必须通过 NVLink 或 PCIe 进行通信的设备相比，MLOPart 设备在相同底层 GPU 上的 MLOPart 实例之间具有更优的点对点带宽。
我们来看看在相同底层设备与不同底层设备上的 MLOPart 设备之间进行双向 P2P 带宽测试的（部分）结果：
$ ./nvbandwidth -t device_to_device_memcpy_read_ce
...
Running device_to_device_memcpy_read_ce.
memcpy CE GPU(row) -> GPU(column) bandwidth (GB/s)
0 1 2 3 4
0 N/A 2352.76 766.82 743.46 767.51
1 2402.78 N/A 765.86 744.04 767.03
2 767.23 744.30 N/A 2349.54 766.00
3 767.37 743.91 2372.91 N/A 767.30
4 766.75 743.52 766.89 743.97 N/A
在上述示例中，设备 0 和 1 位于同一底层 GPU 上，设备 2 和 3 也位于同一底层 GPU 上。
对于 B200，在启动 cuMemcpyAsync 等操作时，对等通信通常会使用 NVLink。若这些 B200 对等端位于同一 B200 芯片上的 MLOPart 设备，则可改用速度更快的 NV-HBI。
使用 MLOPart 时的注意事项
如前所述，使用 MLOPart 意味着在延迟与带宽之间做出取舍，优先选择更低的延迟而非更高的带宽。这并非采用 MLOPart 时唯一需要权衡的因素。
通过 CUDA_VISIBLE_DEVICES 进行设备过滤
可用于 MPS 服务器和客户端的设备可通过 CUDA_VISIBLE_DEVICES 环境变量进行过滤和/或重新映射，该操作通常基于设备序数完成。在使用 MPS 时，若控制器与服务器/客户端均采用相同的 CUDA_VISIBLE_DEVICES 值，则可能在不考虑重映射的情况下引发错误。
例如，给定一个具有 8 个 CUDA 设备的系统，则可以初始化 MPS 控制器以过滤掉奇数设备 (CUDA_VISIBLE_DEVICES=0,2,4,6)。在此场景中，MPS 服务器和客户端仅能看见 4 个 CUDA 设备，即使未使用 CUDA_VISIBLE_DEVICES 也是如此。若对 CUDA_VISIBLE_DEVICES 使用相同的值将导致失败，因为我们只能看见设备 0-3。因此，建议使用明确无误的 UUID。
启用 MLOPart 后，还需注意其他不一致的情况。MPS 控制器与启用 MLOPart 的 MPS 服务器/客户端所见设备的 UUID 并不相同。使用 CUDA_VISIBLE_DEVICES 时，建议在启用 MLOPart 的 MPS 服务器启动后执行 device_query 命令，以确定 MPS 客户端可用的 UUID。
计算资源减少
启用 MLOPart 后，MLOPart 设备可能会禁用部分 SM。内存延迟降低所带来的性能提升与计算资源减少所导致的性能损失之间需要进行权衡。应针对每个应用程序对这些指标进行综合评估。
托管显存
受管理内存无法从 MLOPart 中受益。由于 MLOPart 需要创建 GPU 显存以实现低延迟分配，而使用受管理内存无法满足这一需求，因此尝试使用受管理内存将像往常一样正常运行，且仍可通过受管理内存 API 创建分配，但预计无法获得性能提升。
访问修改器
借助 cuMemSetAccess API，程序员可为 CUDA 分配指定访问属性。当对 MLOPart 设备使用此 API 时，所应用的属性将遵循同一底层 GPU 上所有 MLOPart 设备中限制较宽松的设置。例如，若对一个 MLOPart 设备将缓冲区设为只读，而对另一个 MLOPart 设备将缓冲区设为读写（默认），则两个 MLOPart 设备均会具有读写访问权限，直至两者均更新为限制更严格的访问类型。
x86 要求
MLOPart 目前仅支持 x86 平台。后续版本将提供对 ARM 平台的支持。
与 MIG 的比较
与 MLOPart 类似，MIG 可用于从单个 GPU 创建多个 CUDA 设备。某些 MIG 配置还能在降低延迟的同时牺牲部分带宽，且无需修改代码。
| 主题 | MIG | MLOPart / MPS |
| 权限 | 配置需要超级用户权限 | 不需要超级用户权限 |
| 范围 | 系统范围设置 | 每个用户 / 每个服务器设置 |
| 内存隔离 | 在 MIG GPU 实例之间执行严格的内存隔离 | 来自一个 MLOPart 设备的内存可能会损坏同一 GPU 上的另一个设备 |
| 性能隔离 | 在 MIG 计算实例之间执行严格的性能隔离 | Perf |
如需详细了解 MLOPart、CUDA MPS 以及如何进一步提升 GPU 利用率，请查阅 MPS 文档。
致谢：感谢以下 NVIDIA 贡献者：Alfred Barnat、Ehren Bendler、Alicia Hu、Balint Joo、Ze Long、Yashwant Marathe、Vance Miller、Kyrylo Perelygen、Will Pierce、Yifan Yang
