数据预取：

数据预取意味着将数据迁移到处理器的内存，并在处理器开始访问该数据之前将其映射到该处理器的页表中。 数据预取的目的是避免错误，同时建立数据局部性。 这对于在任何给定时间主要从单个处理器访问数据的应用程序都是最有价值的。 随着访问处理器在应用程序的生命周期内发生变化，可以相应地预取数据以跟踪应用程序的执行流程。 由于工作是在CUDA中的流中启动的，因此预计数据预取也是流式操作，如以下API所示：

cudaError_t cudaMemPrefetchAsync(const void *devPtr,    size_t count,    int dstDevice,    cudaStream_t stream);

其中由devPtr指针指定的内存区域和计数字节数（ptr向下舍入到最接近的页边界并向上舍入到最接近的页边界的计数）将通过在流中排列迁移操作而迁移到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遵循流排序语义，即直到流中的所有先前操作都完成后才开始迁移，并且在迁移完成之前，流中的任何后续操作都不会开始。数据使用提示：当多个处理器需要同时访问相同的数据时，数据预取本身是不够的。 在这种情况下，应用程序提供有关如何实际使用数据的提示很有用。 以下咨询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 cudaMallocManageddataPtr = (char *)malloc(dataSize);// Set the advice on the memory regioncudaMemAdvise(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：这个建议意味着数据将被设备访问。这不会导致数据迁移，并且不会影响数据本身的位置。相反，只要数据的位置允许建立映射，它就会始终将数据映射到指定处理器的页表中。如果数据因任何原因而被迁移，则映射会相应更新。这种建议在数据局部性不重要的情况下非常有用，但避免出现故障。例如，考虑一个包含多个启用了对等访问的GPU的系统，其中位于一个GPU上的数据偶尔会被其他GPU访问。在这种情况下，将数据迁移到其他GPU并不重要，因为访问不频繁，迁移开销可能过高。但是预防故障仍然有助于提高性能，因此预先设置映射是有用的。请注意，在CPU访问此数据时，可能会将数据迁移到CPU内存，因为CPU无法直接访问GPU内存。任何为此数据设置了cudaMemAdviceSetAccessedBy标志的GPU现在都将其映射更新为指向CPU内存中的页面

也可以使用以下某个值取消设置每个建议：cudaMemAdviseUnsetReadMostly，cudaMemAdviseUnsetPreferredLocation和cudaMemAdviseUnsetAccessedBy。

查询使用属性:程序可以通过使用以下API查询通过cudaMemAdvise或cudaMemPrefetchAsync分配的内存范围属性：

cudaMemRangeGetAttribute(void *data,    size_t dataSize,    enum cudaMemRangeAttribute attribute,    const void *devPtr,    size_t count);

该函数查询以devPtr开始的内存范围的属性，其大小为计数字节。 内存范围必须指向通过cudaMallocManaged分配或通过__managed__变量声明的托管内存。 有可能查询

以下属性：

• cudaMemRangeAttributeReadMostly：如果给定内存范围内的所有页面都启用了重复数据消除，则返回的结果为1，否则为0。
• cudaMemRangeAttributePreferredLocation：如果内存范围内的所有页面都返回，则返回的结果将是GPU设备ID或cudaCpuDeviceId
  有相应的处理器作为他们的首选位置，否则将返回cudaInvalidDeviceId。 应用程序可以使用此查询API根据托管指针的首选位置属性决定是通过CPU还是GPU来分段数据。 请注意，查询时页面在内存范围内的实际位置可能与首选位置不同。
• cudaMemRangeAttributeAccessedBy：将返回那些建议为该内存范围设置的设备列表。
• cudaMemRangeAttributeLastPrefetchLocation：将返回使用cudaMemPrefetchAsync显式预取显示内存范围中的所有页面的最后一个位置。 请注意，这只是返回应用程序请求预取内存范围的最后一个位置。 它没有提供关于到那个位置的预取操作是否已经完成或者甚至开始的指示。

另外，可以使用相应的cudaMemRangeGetAttributes函数来查询多个属性。