CUDA统一内存知识点总结

精选原创

wx5f184b1820e35 2024-01-06 23:20:15 ©著作权

©著作权归作者所有：来自51CTO博客作者wx5f184b1820e35的原创作品，请联系作者获取转载授权，否则将追究法律责任

背景及概念

在典型的个人计算机或集群节点中，CPU和GPU的内存物理上是分离的，通过PCI-Express总线连接。在CUDA 6之前，程序员必须将共享数据分配到两个不同的内存中，并显式复制，给CUDA程序带来了繁琐的复杂性。 CUDA6.0于2014年发布，Unified Memory（UM）就是本次发布引入的。

UM允许开发者在编写CUDA程序时使用一致的内存地址空间，而不需要手动管理CPU和GPU之间的数据传输。这一特性在很大程度上简化了GPU编程，并提高了开发效率。 UM 的引入使得数据在CPU和GPU之间的传输更为透明，由运行时系统负责在需要时自动进行数据迁移。

统一内存创建了一个在 CPU 和 GPU 之间共享的托管内存池，弥合了 CPU-GPU 鸿沟。 CPU 和 GPU 都可以使用cudaMallocManaged分配的指针访问托管内存（即相同的数据），无需为不同的设备维护不同的指针。关键是系统会根据需要动态地自动在CPU和GPU之间迁移统一内存中分配的数据，以确保数据在使用时位于合适的设备上。

简单应用-分配UM

cudaMallocManaged 是用于在CUDA程序中分配Unified Memory的函数。这个函数会返回一个指针，可以在CPU和GPU上使用，而且内存是在共享的虚拟地址空间中分配的。

函数定义

<code>cudaError_t cudaMallocManaged(void** devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal);</code>

devPtr：一个指向指针的指针，用于存储分配的内存块的起始地址。在这个指针中，你可以在 CPU 和 GPU 上访问分配的内存。
size：要分配的内存块的字节数。
flags：附加标志，用于指定内存的使用方式。在大多数情况下，使用默认值 cudaMemAttachGlobal 即可。
返回值：返回值是一个 cudaError_t 类型的错误码，用于指示函数是否成功执行。如果返回值是 cudaSuccess，则说明分配成功。

函数的简单使用以下是一个简单的使用cudaMallocManaged 分配 Unified Memory的样例：

#include <iostream>

__global__ void kernel(int *data) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    data[tid] += 100; // 在GPU上修改Unified Memory的数据
}

int main() {
    const int N = 10;
    int *data;

    // 使用 cudaMallocManaged 分配 Unified Memory
    cudaMallocManaged(&data, N * sizeof(int));

    // 在 CPU 上访问和修改数据
    for (int i = 0; i < N; ++i) {
        data[i] = i;
    }

    // 在 GPU 上调用核函数，修改数据
    // 请注意，由于 UM，data 指针可以在 CPU 和 GPU 上使用
    // 所以不需要显式的数据传输
    kernel<<<1, N>>>(data);
    cudaDeviceSynchronize(); // 等待 GPU 执行完成

    // 再次在 CPU 上访问数据
    for (int i = 0; i < N; ++i) {
        std::cout << data[i] << " ";
    }

    // 释放分配的 Unified Memory
    cudaFree(data);

    return 0;
}

UM的数据迁移

自动数据迁移：当 CPU 试图访问 GPU 分配的 UM 内存时，或者当 GPU 试图访问 CPU 分配的 UM 内存时，CUDA 运行时系统会自动触发数据的迁移。

延迟和异步操作：

数据迁移可能引入一定的延迟，因为系统需要判断何时以及如何进行迁移。在某些情况下，这可能导致性能下降。为了优化性能，可以使用异步操作，如异步内存迁移和异步执行核函数。异步操作允许 CPU 或 GPU 在等待数据迁移的同时执行其他操作，提高整体并行性异步内存迁移：

使用 CUDA 流（stream）可以实现异步内存迁移。CUDA 流是一种在 GPU 上并发执行操作的机制每个流都代表了一条独立的指令序列，可以在 GPU 上并发执行。通过将内存迁移与核函数执行异步化，可以减小对性能的影响。

#include <iostream>

__global__ void kernel(int* data, int val) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    data[tid] += val;
}

int main() {
    const int N = 256;
    int *data;

    // 使用 cudaMallocManaged 分配 Unified Memory
    cudaMallocManaged(&data, N * sizeof(int));

    // 在 CPU 上访问和初始化数据
    for (int i = 0; i < N; ++i) {
        data[i] = i;
    }

    // 在 GPU 上调用核函数，异步执行
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    kernel<<<1, N, 0, stream>>>(data, 10);

    // 在流上启动异步内存迁移
    cudaMemcpyAsync(data + N/2, data, N/2 * sizeof(int), cudaMemcpyDeviceToDevice, stream);

    // 在流上启动异步核函数执行
    kernel<<<1, N/2, 0, stream>>>(data + N/2, 20);

    cudaStreamSynchronize(stream); // 等待流执行完成

    // 在 CPU 上访问修改后的数据
    for (int i = 0; i < N; ++i) {
        std::cout << data[i] << " ";
    }

    // 销毁流
    cudaStreamDestroy(stream);

    // 释放分配的 Unified Memory
    cudaFree(data);

    return 0;
}

数据迁移的方向：

数据可以从 CPU 到 GPU 迁移，也可以从 GPU 到 CPU 迁移。

内存访问亲和性

UM 允许指定内存的访问亲和性，即数据在被特定设备首次访问时将会被分配到该设备的内存。

这可以通过 cudaMemPrefetchAsync 来实现。

<code>cudaError_t cudaMemPrefetchAsync(const void* devPtr, size_t count, int dstDevice, cudaStream_t stream = 0);</code> devPtr：指向要迁移的内存块的指针。 count：要迁移的字节数。 dstDevice：目标设备的设备 ID。 stream：CUDA 流，表示异步执行的流。默认为0，表示默认流。

//示例代码：
#include <iostream>

int main() {
    const int N = 10;
    int *data;

    // 使用 cudaMallocManaged 分配 Unified Memory
    cudaMallocManaged(&data, N * sizeof(int));

    // 将数据迁移到 GPU 0 上的内存
    cudaMemPrefetchAsync(data, N * sizeof(int), 0);

    // 在 CPU 上访问和修改数据
    for (int i = 0; i < N; ++i) {
        data[i] = i;
    }

    // 在 GPU 0 上调用核函数，修改数据
    // 请注意，由于 UM，data 指针可以在 CPU 和 GPU 上使用
    // 所以不需要显式的数据传输
    kernel<<<1, N>>>(data);
    cudaDeviceSynchronize(); // 等待 GPU 执行完成

    // 释放分配的 Unified Memory
    cudaFree(data);

    return 0;
}

在这个示例中，cudaMemPrefetchAsync 函数将数据 data 预取到 GPU 0 上的内存，以提高对该设备的访问性能。

共享内存与常量内存

共享内存：共享内存是一种在线程块内部共享的内存，它可以被同一线程块中的所有线程访问。通过 shared 关键字定义共享内存，它在 GPU 上的每个线程块中都有一份拷贝。共享内存的使用可以显著提高访问速度，因为它位于 GPU 的多个线程之间的共享存储器中。常量内存：常量内存是一种只读内存，对所有线程块和线程都是全局唯一的。它对于那些在执行期间不变的数据非常有用。常量内存通常用于存储在内核函数执行期间不变的常量数据。内存共享策略：在 UM 中，内存的共享策略可以通过内存的作用域和生命周期来决定。全局内存是 UM 中的默认内存类型，所有设备都能访问，生命周期取决于分配和释放的时间点。通过共享内存，你可以在线程块中共享数据，提高访问速度。常量内存适用于在内核函数执行期间不变的常量数据。

#include <iostream>

__global__ void kernel(int* data) {
    // 在共享内存中分配数组，每个线程块有一份拷贝
    __shared__ int sharedData[256];

    // 在常量内存中定义常量
    __constant__ int constantValue = 42;

    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    // 使用共享内存
    sharedData[threadIdx.x] = data[tid];
    __syncthreads(); // 确保所有线程都已经写入共享内存

    // 使用常量内存
    int result = sharedData[threadIdx.x] + constantValue;
    data[tid] = result;
}

int main() {
    const int N = 256;
    int *data;

    // 使用 cudaMallocManaged 分配 Unified Memory
    cudaMallocManaged(&data, N * sizeof(int));

    // 在 CPU 上访问和初始化数据
    for (int i = 0; i < N; ++i) {
        data[i] = i;
    }

    // 在 GPU 上调用核函数，修改数据
    kernel<<<1, N>>>(data);
    cudaDeviceSynchronize(); // 等待 GPU 执行完成

    // 在 CPU 上访问修改后的数据
    for (int i = 0; i < N; ++i) {
        std::cout << data[i] << " ";
    }

    // 释放分配的 Unified Memory
    cudaFree(data);

    return 0;
}

共享内存在同一线程块中的所有线程之间共享

常量内存是全局唯一的，适用于不变的常量数据。