|
|
N.3. Performance Tuning
为了使用统一内存实现良好的性能，必须满足以下目标：
- 应避免错误：虽然可重放错误是启用更简单的编程模型的基础，但它们可能严重损害应用程序性能。故障处理可能需要几十微秒，因为它可能涉及 TLB 无效、数据迁移和页表更新。与此同时，应用程序某些部分的执行将停止，从而可能影响整体性能。
- 数据应该位于访问处理器的本地：如前所述，当数据位于访问它的处理器本地时，内存访问延迟和带宽明显更好。因此，应适当迁移数据以利用较低的延迟和较高的带宽。
- 应该防止内存抖动：如果数据被多个处理器频繁访问并且必须不断迁移以实现数据局部性，那么迁移的开销可能会超过局部性的好处。应尽可能防止内存抖动。如果无法预防，则必须进行适当的检测和解决。
为了达到与不使用统一内存相同的性能水平，应用程序必须引导统一内存驱动子系统避免上述陷阱。值得注意的是，统一内存驱动子系统可以检测常见的数据访问模式并自动实现其中一些目标，而无需应用程序参与。但是，当数据访问模式不明显时，来自应用程序的明确指导至关重要。 CUDA 8.0 引入了有用的 API，用于为运行时提供内存使用提示 (cudaMemAdvise()) 和显式预取 (cudaMemPrefetchAsync())。这些工具允许与显式内存复制和固定 API 相同的功能，而不会恢复到显式 GPU 内存分配的限制。
注意：Tegra 设备不支持 cudaMemPrefetchAsync()。
N.3.1. Data Prefetching
数据预取意味着将数据迁移到处理器的内存中，并在处理器开始访问该数据之前将其映射到该处理器的页表中。 数据预取的目的是在建立数据局部性的同时避免故障。 这对于在任何给定时间主要从单个处理器访问数据的应用程序来说是最有价值的。 由于访问处理器在应用程序的生命周期中发生变化，因此可以相应地预取数据以遵循应用程序的执行流程。 由于工作是在 CUDA 中的流中启动的，因此预计数据预取也是一种流操作，如以下 API 所示：
cudaError_t cudaMemPrefetchAsync(const void *devPtr,
size_t count,
int dstDevice,
cudaStream_t stream);其中由 devPtr 指针和 count 字节数指定的内存区域，ptr 向下舍入到最近的页面边界，count 向上舍入到最近的页面边界，通过在流中排队迁移操作迁移到 dstDevice。 为 dstDevice 传入 cudaCpuDeviceId 会导致数据迁移到 CPU 内存。 考虑下面的一个简单代码示例：
void foo(cudaStream_t s) {
char *data;
cudaMallocManaged(&data, N);
init_data(data, N); // execute on CPU
cudaMemPrefetchAsync(data, N, myGpuId, s); // prefetch to GPU
mykernel<<<..., s>>>(data, N, 1, compare); // execute on GPU
cudaMemPrefetchAsync(data, N, cudaCpuDeviceId, s); // prefetch to CPU
cudaStreamSynchronize(s);
use_data(data, N);
cudaFree(data);
}如果没有性能提示，内核 mykernel 将在首次访问数据时出错，这会产生额外的故障处理开销，并且通常会减慢应用程序的速度。 通过提前预取数据，可以避免页面错误并获得更好的性能。 此 API 遵循流排序语义，即迁移在流中的所有先前操作完成之前不会开始，并且流中的任何后续操作在迁移完成之前不会开始。
N.3.2. Data Usage Hints
当多个处理器需要同时访问相同的数据时，单独的数据预取是不够的。 在这种情况下，应用程序提供有关如何实际使用数据的提示很有用。 以下咨询 API 可用于指定数据使用情况：
cudaError_t cudaMemAdvise(const void *devPtr,
size_t count,
enum cudaMemoryAdvise advice,
int device);其中，为从 devPtr 地址开始的区域中包含的数据指定的通知和计数字节的长度，四舍五入到最近的页面边界，可以采用以下值：
- cudaMemAdviseSetReadMostly：这意味着数据大部分将被读取并且只是偶尔写入。 这允许驱动程序在处理器访问数据时在处理器内存中创建数据的只读拷贝。 同样，如果在此区域上调用 cudaMemPrefetchAsync，它将在目标处理器上创建数据的只读拷贝。 当处理器写入此数据时，相应页面的所有副本都将失效，但发生写入的拷贝除外。 此建议忽略设备参数。 该建议允许多个处理器以最大带宽同时访问相同的数据，如以下代码片段所示：
char *dataPtr;
size_t dataSize = 4096;
// Allocate memory using malloc or cudaMallocManaged
dataPtr = (char *)malloc(dataSize);
// Set the advice on the memory region
cudaMemAdvise(dataPtr, dataSize, cudaMemAdviseSetReadMostly, 0);
int outerLoopIter = 0;
while (outerLoopIter < maxOuterLoopIter) {
// The data is written to in the outer loop on the CPU
initializeData(dataPtr, dataSize);
// The data is made available to all GPUs by prefetching.
// Prefetching here causes read duplication of data instead
// of data migration
for (int device = 0; device < maxDevices; device++) {
cudaMemPrefetchAsync(dataPtr, dataSize, device, stream);
}
// The kernel only reads this data in the inner loop
int innerLoopIter = 0;
while (innerLoopIter < maxInnerLoopIter) {
kernel<<<32,32>>>((const char *)dataPtr);
innerLoopIter++;
}
outerLoopIter++;
}
- cudaMemAdviseSetPreferredLocation：此建议将数据的首选位置设置为属于设备的内存。传入设备的 cudaCpuDeviceId 值会将首选位置设置为 CPU 内存。设置首选位置不会导致数据立即迁移到该位置。相反，它会在该内存区域发生故障时指导迁移策略。如果数据已经在它的首选位置并且故障处理器可以建立映射而不需要迁移数据，那么迁移将被避免。另一方面，如果数据不在其首选位置，或者无法建立直接映射，那么它将被迁移到访问它的处理器。请务必注意，设置首选位置不会阻止使用 cudaMemPrefetchAsync 完成数据预取。
- cudaMemAdviseSetAccessedBy：这个advice意味着数据将被设备访问。这不会导致数据迁移，并且对数据本身的位置没有影响。相反，只要数据的位置允许建立映射，它就会使数据始终映射到指定处理器的页表中。如果数据因任何原因被迁移，映射会相应更新。此advice在数据局部性不重要但避免故障很重要的情况下很有用。例如，考虑一个包含多个启用对等访问的 GPU 的系统，其中位于一个 GPU 上的数据偶尔会被其他 GPU 访问。在这种情况下，将数据迁移到其他 GPU 并不那么重要，因为访问不频繁并且迁移的开销可能太高。但是防止故障仍然有助于提高性能，因此提前设置映射很有用。请注意，在 CPU 访问此数据时，由于 CPU 无法直接访问 GPU 内存，因此数据可能会迁移到 CPU 内存。任何为此数据设置了 cudaMemAdviceSetAccessedBy 标志的 GPU 现在都将更新其映射以指向 CPU 内存中的页面。
每个advice也可以使用以下值之一取消设置：cudaMemAdviseUnsetReadMostly、cudaMemAdviseUnsetPreferredLocation 和 cudaMemAdviseUnsetAccessedBy。 |
|
发表于 2023-1-18 12:43:57
