CUDA全局内存模型

原创

已于 2024-11-24 09:58:19 修改 · 1.1k 阅读

25 ·

CC 4.0 BY-SA版权

文章标签：

#linux #gpu算力 #c++

于 2024-11-06 21:40:53 首次发布

CUDA内存模型

1. CUDA内存模型概述
2. 内存管理
3. 内存访问模式
5. 使用统一内存的矩阵加法

1. CUDA内存模型概述

1.1 CUDA内存模型

CUDA内存模型提出了多种可编程内存的模型

寄存器
共享内存
本地内存
常量内存
纹理内存
全局内存

1.2 本地内存

核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中。

1.3 共享内存

核函数中使用如下修饰符修饰的变量存放在共享内存中：
__shared__
访问共享内存必须同步使用如下调用
void __syncthreads();
SM中的一级缓存和共享内存都是用64KB的片上内存，通过静态划分，在运行时可以通过如下指令进行动态分配：
cudaError_t cudaFuncSetCacheConfig(const void* func, cudaFuncCache cacheConfig);
cudaFuncCache 是一个枚举类型，它定义了可能的缓存配置选项。这些选项包括：

cudaFuncCachePreferNone：不偏好共享内存或L1缓存，使用GPU默认设置。
cudaFuncCachePreferShared：偏好共享内存，即增加共享内存的大小，减小L1缓存的大小。
cudaFuncCachePreferL1：偏好L1缓存，即增加L1缓存的大小，减小共享内存的大小。
cudaFuncCachePreferEqual：平衡共享内存和L1缓存的大小。

1.4 常量内存

常量内存驻留在设备内存中，每个SM专用的常量缓存中缓存，使用如下修饰符来修饰：
__constant__
核函数只能从常量内存中读取数据。因此，常量内存必须在主机端使用下面的函数来初始化：
cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count, size_t offset = 0, cudaMemcpyKind kind = cudaMemcpyHostToDevice);

1.5 纹理内存

纹理内存驻留在设备内存中，并在每个SM的只读缓存中缓存。纹理内存是对二维空间局部性的优化，所以线程束里使用纹理内存访问二维数据的线程可以达到最优性能。

1.6 全局内存

一个全局内存变量可以被静态声明或动态声明。使用如下修饰符：
__device__
全局内存常驻于设备内存中，可通过32字节、64字节或128字节的内存事务进行访问。这些内存事务必须自然对齐。

1.7 静态全局内存

下面的代码说明了如何静态声明一个全局变量:

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>
__device__ float devData;
__global__ void checkGlobalVariable()
{
   
   
    // display the original value
    printf("Device: the value of the global variable is %f\n", devData);
    // alter the value
    devData += 2.0f;
}
int main(void)
{
   
   
    // initialize the global variable
    float value = 3.14f;
    CHECK(cudaMemcpyToSymbol(devData, &value, sizeof(float)));
    printf("Host:   copied %f to the global variable\n", value);
    // invoke the kernel
    checkGlobalVariable<<<1, 1>>>();
    // copy the global variable back to the host
    CHECK(cudaMemcpyFromSymbol(&value, devData, sizeof(float)));
    printf("Host:   the value changed by the kernel to %f\n", value);
    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

输出打印如下：

$ ./test
Host:   copied 3.140000 to the global variable
Device: the value of the global variable is 3.140000
Host:   the value changed by the kernel to 5.140000

2. 内存管理

本节如何使用CUDA函数来显示地管理内存和数据移动

分配和释放设备内存
在主机和设备之间传输数据

2.1 内存分配和释放

在主机上使用下列函数分配全局内存：
cudaError_t cudaMalloc(void** devPtr, size_t size);
使用下列函数将其初始化：
cudaError_t cudaMemset(void *devPtr, int value, size_t count);
一旦应用程序不再使用已分配的全局内存，使用如下代码释放该内存空间：
cudaError_t cudaFree(void* devPtr);

2.2 内存传输

一旦分配了全局内存，可以使用下列函数从主机向设备传输数据：
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);
下面使用一个例子:

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>
int main(int argc, char **argv)
{
   
   
    // set up device
    int dev = 0;
    CHECK(cudaSetDevice(dev));
    // memory size
    unsigned int isize = 1 << 22;
    unsigned int nbytes = isize * sizeof(float);
    // get device information
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("%s starting at ", argv[0]);
    printf("device %d: %s memory size %d nbyte %5.2fMB\n", dev,
           deviceProp.name, isize, nbytes / (1024.0f * 1024.0f));
    // allocate the host memory
    float *h_a = (float *)malloc(nbytes);
    // allocate the device memory
    float *d_a;
    CHECK(cudaMalloc((float **)&d_a, nbytes));
    // initialize the host memory
    for(unsigned int i = 0; i < isize; i++) h_a[i] = 0.5f;
    // transfer data from the host to the device
    CHECK(cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice));
    // transfer data from the device to the host
    CHECK(cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost));
    // free memory
    CHECK(cudaFree(d_a));
    free(h_a);
    // reset device
    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

2.3 固定内存

分配的主机内存默认是可分页。GPU不能在可分页主机内存上安全地访问数据，当主机操作系统在物理位置上移动该数据，它无法控制。当可分页主机内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，从固定内存传输数据给设备内存。

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>
int main(int argc, char **argv)
{
   
   
    // set up device
    int dev = 0;
    CHECK(cudaSetDevice(dev));
    // memory size
    unsigned int isize = 1 << 22;
    unsigned int nbytes = isize * sizeof(float);
    // get device information
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    if (!deviceProp.canMapHostMemory)
    {
   
   
        printf("Device %d does not support mapping CPU host memory!\n", dev);
        CHECK(cudaDeviceReset());
        exit(EXIT_SUCCESS);
    }
    printf("%s starting at ", argv[0]);
    printf("device %d: %s memory size %d nbyte %5.2fMB canMap %d\n", dev,
           deviceProp.name, isize, nbytes / (1024.0f * 1024.0f),
           deviceProp.canMapHostMemory);
    // allocate pinned host memory
    float *h_a;
    CHECK(cudaMallocHost ((float **)&h_a, nbytes));
    // allocate device memory
    float *d_a;
    CHECK(cudaMalloc((float **)&d_a, nbytes));
    // initialize host memory
    memset(h_a, 0, nbytes);
    for (int i = 0; i < isize; i++) h_a[i] = 100.10f;
    // transfer data from the host to the device
    CHECK(cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice));
    // transfer data from the device to the host
    CHECK(cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost));
    // free memory
    CHECK(cudaFree(d_a));
    CHECK(cudaFreeHost(h_a));
    // reset device
    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

2.4 零拷贝内存

通常来说主机不能直接访问设备变量，同时设备也不能直接访问主机变量。零拷贝内存是固定内存，该内存映射到设备地址空间中。

当设备内存不足时可利用的主机内存
避免主机和设备间的显式数据传输
提高PCIe传输率

通过下列函数创建一个到固定内存的映射：
__host__ cudaError_t cudaHostAlloc(void** pHost, size_t size, unsigned int flags);
为了测试零拷贝内存读写的性能。给数组A和B分配零拷贝内存，并在设备内存上为数组C分配空间。

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>
void checkResult(float *hostRef, float *gpuRef, const int N)
{
   
   
    double epsilon = 1.0E-8;
    for (int i = 0; i < N; i++)
    {
   
   
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
   
   
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i],
                    gpuRef[i], i);
            break;
        }
    }
    return;
}
void initialData(float *ip, int size)
{
   
   
    int i;
    for (i = 0; i < size; i++)
    {
   
   
        ip[i] = (float)( rand() & 0xFF ) / 10.0f;
    }
    return;
}
void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
   
   
    for (int idx = 0; idx < N; idx++)
    {
   
   
        C[idx] = A[idx] + B[idx];
    }
}
__global__ void sumArrays(float *A, float *B, float *C, const int N)
{
   
   
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) C[i] = A[i] + B[i];
}
__global__ void sumArraysZeroCopy(float *A, float *B, float *C, const int N)
{
   
   
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) C[i] = A[i] + B[i];
}
int main(int argc, char **argv)
{
   
   
    // set up device
    int dev = 0;
    CHECK(cudaSetDevice(dev));
    // get device properties
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    // check if support mapped memory
    if (!deviceProp.canMapHostMemory)
    {
   
   
        printf("Device %d does not support mapping CPU host memory!\n", dev);
        CHECK(cudaDeviceReset());
        exit(EXIT_SUCCESS);
    }
    printf("Using Device %d: %s ", dev, deviceProp.name);
    // set up data size of vectors
    int ipower = 10;
    if (argc > 1) ipower = atoi(argv[1]);
    int nElem = 1 << ipower;
    size_t nBytes = nElem * sizeof(float);
    if (ipower < 18)
    {
   
   
        printf("Vector size %d power %d  nbytes  %3.0f KB\n", nElem, ipower,
               (float)nBytes / (1024.0f));
    }
    else
    {
   
   
        printf("Vector size %d power %d  nbytes  %3.0f MB\n", nElem, ipower,
               (float)nBytes / (1024.0f * 1024.0f));
    }
    // part 1: using device memory
    // malloc host memory
    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float *)malloc(nBytes);
    h_B     = (float *)malloc(nBytes);
    hostRef = (float *)malloc(nBytes);
    gpuRef  = (float *)malloc(nBytes);
    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);
    // add vector at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);
    // malloc device global memory
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((float**)&d_A, nBytes));
    CHECK(cudaMalloc((float**)&d_B, nBytes));
    CHECK(cudaMalloc((float**)&d_C, nBytes));
    // transfer data from host to device
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));
    // set up execution configuration
    int iLen = 512;
    dim3 block (iLen);
    dim3 grid  ((nElem + block.x - 1) / block.x);
    sumArrays<<<grid, block>>>(d_A, d_B, d_C, nElem);
    // copy kernel result back to host side
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    // check device results
    checkResult(hostRef, gpuRef, nElem);
    // free device global memory
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_B));
    // free host memory
    free(h_A);
    free(h_B);
    // part 2: using zerocopy memory for array A and B
    // allocate zerocpy memory
    CHECK(cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped));
    CHECK(cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped));
    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);
    // pass the pointer to device
    CHECK(cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0));
    CHECK(cudaHostGetDevicePointer((void **)&d_B, (void *)h_B, 0));
    // add at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);
    // execute kernel with zero copy memory
    sumArraysZeroCopy<<<grid, block>>>(d_A, d_B, d_C, nElem);
    // copy kernel result back to host side
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    // check device results
    checkResult(hostRef, gpuRef, nElem);
    // free  memory
    CHECK(cudaFree(d_C));
    CHECK(cudaFreeHost(h_A));
    CHECK(cudaFreeHost(h_B));
    free(hostRef);
    free(gpuRef);
    // reset device
    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

输出结果如下：

 nvprof ./test
==5127== NVPROF is profiling process 5127, command: ./test
Using Device 0: NVIDIA GeForce MX250 Vector size 1024 power 10  nbytes    4 KB
==5127== Profiling application: ./test
==5127== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   52.38%  7.7440us         1  7.7440us  7.7440us  7.7440us  sumArraysZeroCopy(float*, float*, float*, int)
                   19.91%  2.9440us         1  2.9440us  2.9440us  2.9440us  sumArrays(float*, float*, float*, int)
                   19.91%  2.9430us         2  1.4710us  1.4710us  1.4720us  [CUDA memcpy DtoH]
                    7.79%  1.1520us         2     576ns     576ns     576ns  [CUDA memcpy HtoD]

从结果看出，共享主机和设备端少量数据，零拷贝内存可以是个不错的选择。

2.5 统一虚拟寻址

计算能力为2.0及以上版本的设备支持一种特殊的寻址方式，称为统一虚拟寻址（UVA）。

3. 内存访问模式

为了读写数据时达到最佳的性能，内存访问操作必须满足一定的条件。

3.1 对齐与合并访问

核函数的内存请求通常在DRAM设备和片上内存间以128字节或32字节内存事务来实现的。
所有对全局内存的访问都会通过二级缓存，也有许多访问会通过一级缓存，取决于访问类型和GPU架构。这两级缓存都被用到，那么内存访问是由一个128字节的内存事务实现的。如果使用了二级缓存，那么这个内存访问由一个32字节的内存事务实现的。
优化应用程序时，注意内存访问的两个特性：

对齐内存访问
合并内存访问

3.2 全局内存读取

数据通过以下3种缓存/缓冲路径进行传输

一级和二级缓存
常量缓存
只读缓存

一/二级缓存时默认路径。想要通过其它两种路径传递数据需要应用程序显式说明。以下标志通知编译器禁用一级缓存：
-Xptxas -dlcm=cg
一级缓存可以使用下列标识符直接启用：
-Xptxas -dlcm=ca

3.2.1 缓存加载

缓存加载操作经过一级缓存，在粒度为128字节的一级缓存行上由设备内存事务进行传输。缓存加载可以分为对齐/非对齐及合并/非合并。
CPU一级缓存和GPU一级缓存之间的差异：
CPU一级缓存优化了时间和空间的局部性。CPU一级缓存是专为空间局部而不是为时间局部性优化的。频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。

3.2.2 没有缓存的加载

没有缓存的加载不经过一级缓存，它在内存段的粒度上（32个字节）而非缓存池的粒度（128个字节）执行。这是更细粒度的加载，可以为非对其或非合并的内存访问带来更好的总线利用率。

3.2.3 非对齐读取的示例

为了说明核函数中非对齐访问对性能的影响，修改向量加法代码进行修改，去掉所有的内存加载操作，来指定一个偏移量。

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>
void checkResult(float *hostRef, float *gpuRef, const int N)
{
   
   
    double epsilon = 1.0E-8;
    bool match = 1;
    for (int i = 0; i < N; i++)
    {
   
   
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
   
   
            match = 0;
            printf("different on %dth element: host %f gpu %f\n", i, hostRef[i],
                    gpuRef[i]);
            break;
        }
    }
    if (!match)  printf("Arrays do not match.\n\n");
}
void initialData(float *ip,  int size)
{
   
   
    for (int i = 0; i < size; i++)
    {
   
   
        ip[i] = (float)( rand() & 0xFF ) / 100.0f;
    }
    return;
}
void sumArraysOnHost(float *A, float *B, float *C, const int n, int offset)
{
   
   
    for (int idx = offset, k = 0; idx < n; idx++, k++)
    {
   
   
        C[k] = A[idx] + B[idx];
    }
}
__global__ void warmup(float *A, float *B, float *C, const int n, int offset)
{
   
   
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;
    if (k < n) C[i] = A[k] + B[k];
}
__global__ void readOffset(float *A, float *B, float *C, const int n,
                           int offset)
{
   
   
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;
    if (k < n) C[i] = A[k] + B[k];
}
int main(int argc, char **argv)
{
   
   
    // set up device
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("%s starting reduction at ", argv[0]);
    printf("device %d: %s ", dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));
    // set up array size
    int nElem = 1 << 20; // total number of elements to reduce
    printf(" with array size %d\n", nElem);
    size_t nBytes = nElem * sizeof(float);
    // set up offset for summary
    int blocksize = 512;
    int offset = 0;
    if (argc > 1) offset    = atoi(argv[1]);
    if (argc > 2) blocksize = atoi(argv[2]);
    // execution configuration
    dim3 block (blocksize, 1);
    dim3 grid  ((nElem + block.x - 1) / block.x, 1);
    // allocate host memory
    float *h_A = (float *)malloc(nBytes);
    float *h_B = (float *)malloc(nBytes);
    float *hostRef = (float *)malloc(nBytes);
    float *gpuRef  = (float *)malloc(nBytes);
    //  initialize host array
    initialData(h_A, nElem);
    memcpy(h_B, h_A, nBytes);
    //  summary at host side
    sumArraysOnHost(h_A, h_B, hostRef, nElem, offset);
    // allocate device memory
    float *d_A, *d_B, *d_C