CUDA编程（十二）使用同一内存编程

grace 1314

已于 2023-11-02 10:02:30 修改

阅读量320

点赞数

文章标签： c++ c语言

于 2023-10-28 18:52:18 首次发布

本文链接：https://blog.csdn.net/qq_43616651/article/details/134094696

版权

统一内存（unifiedmemory）编程。在这种新的编程模型下，将不再需要手动地在主机与设备间传输数据。统一内存编程模型由CUDA 6 引入，从开普勒架构开始就可用，而且其易用程度和性能都随着GPU架构的更新有所提高。然而，到目前为止，统一内存的某些功能在Windows操作系统中依然受到限制。

缓存：是位于处理器（CPU）和主内存之间的高速存储器，用于加速数据访问，存储了频繁使用的数据，临时存储，速度快，与处理器速度匹配，有很多级别，L1，L2，L3。L1最接近处理器，L3最远。

内存：计算机系统主要存储设备，用于存储正在运行的程序和数据，内存容易丢失，内存的容量比缓存大很多，内存的访问速度比硬盘等存储设备块，但仍然比缓存慢。

统一内存简介

统一内存是一种逻辑上的概念，它既不是显存，也不是主机的内存，而是一种系统中的任何处理器（CPU或GPU）都可以访问，并能保证一致性的虚拟存储器。这种虚拟存储器是通过CPU和GPU各自内部集成的内存管理单元（memory managementunit）实现的。在某种程度上，可以将一个系统中某个处理器的内存看成整个统一内存的超级大缓存。

统一内存是CUDA编程模型的一个组成部分，首次引入于CUDA 6.0，它定义了一个托管内存空间，在这个空间中，所有处理器都能看到一个具有共同地址空间的单一一致内存图像。

处理器指的是具有专用内存管理单元（MMU）的任何独立执行单元。这包括任何类型和架构的CPU和GPU。

在统一内存之前，还有一种零复制内存（zero-copymemory）。它们都提供了一种统一的能被CPU和GPU都访问到的存储器，但零复制内存只是用主机内存作为存储介质，而统一内存则能将数据放在一个最合适的地方（可以是主机，也可以是设备）

另外，统一内存编程在很大程度上涉及多GPU编程。

使用统一内存对硬件有较高的要求：

• 对于所有的功能，GPU的架构都必须不低于开普勒架构，主机应用程序必须为64位的。

• 对于一些较新的功能（在适当的时候会具体指出），至少需要帕斯卡架构的GPU，而且主机要用Linux系统。也就是说，不管用什么GPU，在Windows系统中都只能用第一代统一内存的功能。

• 在具有IBMPower9和NVLink的系统中，伏特架构的GPU支持设备访问任何主机内存，包括用malloc分配的动态数组和在栈上分配的静态数组。这是非常诱人的特征，但普通用户（包括作者）可能没有这样的硬件资源。

这里如果后续不能使用windows去做的话，可能会转到linux实现。

下面是使用统一内存可能带来的好处：

1. 统一内存使CUDA编程更加简单。使用统一内存，将不再需要手动将数据在主机与设备之间传输，也不需要针对同一组数据定义两个指针，并分别分配主机和设备内存。 对于某个统一内存变量，可以直接从GPU或者CPU中进行访问。

2. 可能会提供比手工移动数据提供更好的性能。底层的统一内存实现，可能会自动将一部分数据放置到离某个存储器更近的位置（如部分放置到某卡的显存中，部分放置到内存中），这种自动的就近数据存放，有可能提升性能。

3. 允许GPU在使用了统一内存的情况下，进行超量分配。超出GPU内存额度的部分可能存放在主机上。这可能是使用统一内存时最大的好处，因为一般来说CPU的内存可以更多，但处理速度较低，而GPU虽然处理速度较高，但内存（显存）数量有限 （参看第1章中所列的数据）。该功能要求帕斯卡架构或更高的架构及Linux操作系统。

底层系统在不需要显式内存复制调用的情况下管理CUDA程序中的数据访问和局部性。这为GPU编程带来两个主要好处：

通过在系统中一致地统一内存空间，使所有GPU和CPU之间的内存具有一致性，为CUDA程序员提供更紧密、更直接的语言集成，简化了GPU编程。
通过透明地将数据迁移到正在使用的处理器，最大限度地提高了数据访问速度。

简而言之，统一内存消除了通过cudaMemcpy*()例程显式移动数据的需求，而不会因将所有数据放入零拷贝内存而产生性能损失。当然，数据仍然会移动，因此程序的运行时间通常不会减少；统一内存使编写更简单和易于维护的代码成为可能。

统一内存提供了一个“单一指针指向数据”的模型，从概念上类似于CUDA的零拷贝内存。两者之间的一个关键区别在于，在零拷贝内存分配中，内存的物理位置在CPU系统内存中被固定，因此根据从何处访问内存，程序可能具有快速或较慢的访问速度。另一方面，统一内存将内存与执行空间解耦，因此所有数据访问都很快。

统一内存这个术语描述了一个系统，为各种程序提供内存管理服务，从针对运行时API的程序到使用虚拟ISA（PTX）的程序。该系统的一部分定义了选择加入统一内存服务的托管内存空间。

托管内存与特定设备分配（如使用cudaMalloc()例程创建的分配）是可互操作和可互换的；所有对设备内存有效的CUDA操作在托管内存上也有效；主机部分的程序能够引用和访问这些内存。

统一内存不支持连接到Tegra的离散GPU。

统一内存的基本使用方法

统一内存在设备中是当作全局内存使用的，而且必须在主机端定义或分配内存，而不能在设备端（核函数和__device__函数）定义或分配内存。例如，在核函数中由malloc分配的堆内存不属于统一内存，从而无法被CPU访问。

动态统一内存

先来看看一个例子；

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include "error.cuh" 
#include <math.h>

const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;
void __global__ add(const double* x, const double* y, double* z);
void check(const double* z, const int N);

int main(void)
{
    const int N = 1000000;
    const int M = sizeof(double) * N;
    double* x, * y, * z;
    CHECK(cudaMallocManaged((void**)&x, M));
    CHECK(cudaMallocManaged((void**)&y, M));
    CHECK(cudaMallocManaged((void**)&z, M));

    for (int n = 0; n < N; ++n)
    {
        x[n] = a;
        y[n] = b;
    }

    const int block_size = 128;
    const int grid_size = N / block_size;
    add << <grid_size, block_size >> > (x, y, z);

    CHECK(cudaDeviceSynchronize());
    check(z, N);

    CHECK(cudaFree(x));
    CHECK(cudaFree(y));
    CHECK(cudaFree(z));
    return 0;
}

void __global__ add(const double* x, const double* y, double* z)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    z[n] = x[n] + y[n];
}

void check(const double* z, const int N)
{
    bool has_error = false;
    for (int n = 0; n < N; ++n)
    {
        if (fabs(z[n] - c) > EPSILON)
        {
            has_error = true;
        }
    }
    printf("%s\n", has_error ? "Has errors" : "No errors");
}

还是那个老程序，数组相加，不同的是，只定义了3个数据指针，x,y,z，使用的也是一个新的cuda函数：cudaMallocManaged（）：

cudaMallocManaged() 是CUDA（Compute Unified Device Architecture）编程模型中的一个函数，用于分配统一内存（Unified Memory）。统一内存是一种内存管理技术，它允许将主机（CPU）和设备（GPU）之间的内存分配和访问变得更加无缝和简单。

cudaMallocManaged() 函数用于在GPU上分配一块统一内存，并将其分配的地址返回给调用者。分配的内存可以由主机和设备代码访问，而不需要显式地进行数据传输。这使得在GPU上执行的代码可以方便地访问主机分配的内存，反之亦然。

cudaMallocManaged(void** devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal);

devPtr：一个指向指针的指针，用于接收分配的内存地址。
size：要分配的内存大小（以字节为单位）。
flags：可选参数，用于控制内存分配的行为。默认情况下，可以使用 cudaMemAttachGlobal，表示内存可被多个GPU设备访问。

相比 cudaMalloc 函数，该函数多了一个可选参数 flags。该参数的默认值是 cudaMemAttachGlobal。如果取默认值，代表分配的全局内存可由任何设备通过任何 CUDA 流访问。flags 另一个可取的值是 cudaMemAttachHost，但我们不讨论这种情形。统一内存的释放依然是用之前用过的 cudaFree ()函数。值得强调的是，只能在主机端使用该函数分配统一内存，而不能在核函数中使用该函数。

for循环中为变量赋值，这是主机访问，后面调用核函数是设备对统一内存的访问。

从核函数的角度来看，统一内存和普通设备内存在使用上没有任何区别。也就是说，在将一个程序从不使用统一内存的版本改为使用统一内存的版本（或者反过来）时，不需要对核函数进行修改。而且，可以一点一点地将非统一内存改为统一内存（或者反过来），即同一个程序中可以同时使用统一内存和非统一内存。

对于开普勒架构和麦克斯韦架构（第一代统一内存）来说，主机与设备不能并发地访问统一内存。又因为核函数的调用是异步的，故在调用任何核函数之后，不能紧接着从 CPU 访问任何统一内存变量（不管核函数是否访问了 CPU 将要访问的某段统一内存地址），必须在中间加上一个同步函数，以确保核函数对统一内存的访问已经结束，如 cudaDeviceSynchronize。

目前不管用什么 GPU，在 Windows 系统中都只能用第一代统一内存的功能。

与之前的add.cu程序比较，可以看出在使用统一内存后，程序确实简化了许多，不需要再针对同一组数据定义两个数组（一个在主机，一个在设备），而且不需要显式地进行主机与设备间的数据传输。至于程序的性能，很难通过这种简单的程序进行测试。如果只是针对核函数来说，可以说两个版本的核函数具有同样的性能。

静态统一内存：

正如 GPU 中的全局内存除可以动态分配外，还可以静态分配，统一内存也可以静态地分配。要定义静态统一内存，只需要在修饰符 __device__ 的基础上再加上修饰符 __managed__ 即可。注意，这样的变量是在任何函数外部定义的，可见范围是所在源文件（更准确地说是所在翻译单元）。

#include <stdio.h>

__device__ __managed__ int ret[1000];

__global__ void AplusB(int a, int b)
{
    ret[threadIdx.x] = a + b + threadIdx.x;
}

int main()
{
     AplusB<<<1, 1000>>>(10, 100);
     cudaDeviceSynchronize(); // 第一代统一内存或者 Windows 系统需要
     for(int i = 0; i < 1000; i++)
     {
         printf("%d: A+B = %d\n", i, ret[i]);
     }
    return 0;
}

使用统一内存申请超量的内存：目前还不能在win上测试，就不细说了。

使用统一内存的一个好处是在适当的时候可以超量申请设备内存。

程序可以通过两种方式分配托管内存：使用cudaMallocManaged()例程，它在语义上类似于cudaMalloc()；或者通过定义全局的__managed__变量，它在语义上类似于__device__变量。

在支持计算能力为6.x及更高版本的设备的平台上，统一内存将使应用程序能够使用默认系统分配器分配和共享数据。这允许GPU访问整个系统虚拟内存，而无需使用特殊分配器。有关更多详细信息，请参阅系统分配器。

以下代码示例说明了使用托管内存可以改变主机代码编写方式的方式。首先，演示了没有使用统一内存的简单程序：

__global__ void AplusB(int *ret, int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    int *ret;
    cudaMalloc(&ret, 1000 * sizeof(int));
    AplusB<<< 1, 1000 >>>(ret, 10, 100);
    int *host_ret = (int *)malloc(1000 * sizeof(int));
    cudaMemcpy(host_ret, ret, 1000 * sizeof(int), cudaMemcpyDefault);
    for(int i = 0; i < 1000; i++)
        printf("%d: A+B = %d\n", i, host_ret[i]);
    free(host_ret);
    cudaFree(ret);
    return 0;
}

这个第一个示例在GPU上将两个数字与每个线程的ID相结合，并将结果值返回到一个数组中。如果没有使用托管内存，就需要在主机端和设备端都为返回值分配存储空间（如示例中的host_ret和ret），并且需要使用cudaMemcpy()来显式复制数据。

与这个程序的统一内存版本相比，你可以直接从主机端访问GPU上的数据。请注意cudaMallocManaged()例程，它返回一个指针，可以在主机和设备代码中都有效。这允许使用ret而无需单独的host_ret复制，大大简化了程序的编写并减小了程序的大小。

__global__ void AplusB(int *ret, int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    int *ret;
    cudaMallocManaged(&ret, 1000 * sizeof(int));
    AplusB<<< 1, 1000 >>>(ret, 10, 100);
    cudaDeviceSynchronize();
    for(int i = 0; i < 1000; i++)
        printf("%d: A+B = %d\n", i, ret[i]);
    cudaFree(ret);
    return 0;
}

最后，语言集成允许直接引用GPU声明的__managed__变量，并在使用全局变量时进一步简化程序

__device__ __managed__ int ret[1000];
__global__ void AplusB(int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    AplusB<<< 1, 1000 >>>(10, 100);
    cudaDeviceSynchronize();
    for(int i = 0; i < 1000; i++)
        printf("%d: A+B = %d\n", i, ret[i]);
    return 0;
}

请注意，在没有使用托管内存的示例中，使用了同步的cudaMemcpy()例程来同步内核（即等待它运行完成）并将数据传输到主机。统一内存示例不调用cudaMemcpy()，因此需要在主机程序安全使用来自GPU的输出之前，显式调用cudaDeviceSynchronize()来进行同步。

优化使用统一内存的程序

为了在使用统一内存时获得较高性能，需要避免缺页异常、保持数据的局部性（让相关数据尽量靠近对应的处理器）但避免内存抖动（即频繁地在不同的处理器之间传输数据）。 CUDA的统一内存机制可以部分地自动做到这些，但很多情况下还是需要手动地给编译器一些提示（hints），如使用CUDA运行时函数cudaMemAdvise和cudaMemPrefetchAsync。

cudaMemAdvise 是 CUDA（Compute Unified Device Architecture）编程中用于管理内存访问模式的函数之一。通过使用 cudaMemAdvise，开发者可以提供关于内存的建议，以优化数据迁移和访问性能。这些建议通常用于告诉CUDA运行时如何处理内存，以便更好地适应应用程序的访问模式。

cudaError_t cudaMemAdvise(const void* devPtr, size_t count, cudaMemoryAdvise advice, int device);

devPtr：要提供建议的设备指针。
count：内存块的元素数目。
advice：建议类型，指定如何处理内存。可能的值包括 cudaMemAdviseSetReadMostly、cudaMemAdviseSetPreferredLocation、cudaMemAdviseSetAccessedBy 等。
device：可选参数，指定设备的ID，用于多GPU系统，表示建议将应用于哪个设备。如果只有一个GPU，可以将其设为 -1。

不同的 cudaMemoryAdvise 枚举值表示不同的内存访问建议，这些建议包括但不限于：

cudaMemAdviseSetReadMostly：建议数据主要用于读取，可提高读取性能。
cudaMemAdviseSetPreferredLocation：建议数据的首选位置，以便更快地访问。
cudaMemAdviseSetAccessedBy：建议哪些设备（GPU）将访问数据，以便优化数据的共享。

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：这个建议将数据的首选位置设置为属于device的内存。将device的值设置为cudaCpuDeviceId将首选位置设置为CPU内存。设置首选位置不会立即导致数据迁移到该位置。相反，它在内存区域发生故障时引导迁移策略。如果数据已经位于其首选位置，并且故障的处理器可以建立映射而不需要迁移数据，那么迁移将被避免。另一方面，如果数据不在其首选位置，或者无法建立直接映射，那么它将被迁移到访问它的处理器。重要的是要注意，设置首选位置不会阻止使用cudaMemPrefetchAsync进行数据预取。
cudaMemAdviseSetAccessedBy：这个建议意味着数据将被device访问。这不会导致数据迁移，也不会影响数据本身的位置。相反，它会导致数据始终在指定处理器的页表中进行映射，只要数据的位置允许建立映射。如果出于任何原因数据被迁移，映射将相应地更新。这个建议在数据局部性不重要但要避免故障的情况下很有用。例如，考虑一个启用了对等访问的多个GPU的系统，其中位于一个GPU上的数据偶尔被其他GPU访问。在这种情况下，将数据迁移到其他GPU可能并不那么重要，因为访问不频繁，而迁移的开销可能过高。但是，避免故障仍然有助于提高性能，因此提前设置映射是有用的。请注意，当CPU访问这些数据时，数据可能会迁移到CPU内存，因为CPU不能直接访问GPU内存。为此数据设置了cudaMemAdviceSetAccessedBy标志的任何GPU现在将其映射更新为指向CPU内存中的页面。

查询使用属性

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

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

这个函数查询从devPtr开始，大小为count字节的内存范围的属性。内存范围必须是通过cudaMallocManaged分配的托管内存或通过__managed__变量声明的。可以查询以下属性：

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