CUDA全局内存(CUDA C编程权威指南)

grace 1314

已于 2023-12-02 13:51:30 修改

阅读量185

点赞数

文章标签： c语言

于 2023-11-30 14:29:08 首次发布

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

版权

1.内存访问的局部性

时间局部性：如果一个数据位置被引用，那么该数据在较短的时间周期内很可能再次被引用，随着时间的流逝，该数据被引用的可能性逐渐降低。

内存局部性：如果一个内存位置被引用，那么附近的位置也可能会被引用。

2.内存层次

寄存器缓存主存磁盘存储器

速度从左到右：快 -》慢

大小从左到右：小 -》大

3.CUDA内存模型：

存储器可以分为可编程的和不可编程的。

CUDA内存模型中可编程内存的类型：寄存器，共享内存，本地内存，常量内存，纹理内存，全局内存。

本地内存：一个线程具有，可读可写。

共享内存：一个线程块具有，块内线程均可访问，可读可写。

全局内存：一个网格具有，所有网格内线程均可访问，可读可写。

常量内存和纹理内存：所有网格内线程均可访问，只读。

纹理内存为各种数据布局提供了不同的寻址模式和滤波模式。

寄存器：GPU上运行速度最快的内存空间。核函数中声明的没有其他修饰符的自变量，通常存储在寄存器中，线程具有，可读可写。

本地内存，核函数中的变量，寄存器中放不下的或者不满足寄存器限定条件要求的将存放至本地内存。

共享内存，核函数中使用__shared__关键字修饰的变量存放在共享内存。与本地内存相比，宽带高，延迟低。类似于CPU一级缓存，但共享内存是可编程的。必须小心不能过度使用共享内存，容易限制活跃线程的数量，生命周期，与线程块同周期，线程块执行结束，其分配的共享内存将被释放并重新分配给其他块。

共享内存是线程之间相互通信的基本方式。一个块内的线程通过使用共享内存中的数据可以相互合作。访问共享内存必须同步。即使用void __syncthreads();即为线程块里的线程设立障碍点，同一个块里的线程必须在其他线程被运行执行前达到该处。可以避免潜在的数据冲突。

数据冲突：当一组未排序的多重访问通过不同线程访问相同的内存地址时，这些访问中至少有一个时可惜的，这是就会出现数据冲突。

在CUDA编程中,数据冲突主要指不同线程访问同一资源(如共享内存或全局内存位置)时可能出现的问题。

主要的CUDA数据冲突有:

读后写冲突(Read After Write hazard):多个线程对同一全局内存位置进行读写操作,后线程读到的数据可能不是预期值。
写后写冲突(Write After Write hazard):多个线程对同一内存位置进行写操作,后线程写的值可能会覆盖前线程写的数据。
读后读冲突(Read After Read hazard):多个线程重复读同一内存位置,但读指令的顺序不同,可能出现不同结果。
资源竞争(Resource Conflict):多个线程试图同时访问共享资源如锁或信号量时会产生竞争。

防止数据冲突的手段包括:

使用共享内存或全局锁进行同步。
对数据采用只读或只写策略避免读写操作冲突。
调整线程块、warp或线程的执行顺序避免重复读写同一位置。
将读取的指令提前或读后运算延迟避免时序问题。
避免直接读写数组下标未知的全局内存位置。

常量内存：使用__constant__来修饰常量变量，必须在全局空间内和核函数之外进行声明，大小只有64KB，常量内存时静态声明的，并对统一编译单元中所有的核函数可见。核函数只能从常量内存中读取数据。

常量内存必须在主机端使用

cudaMemcpyToSymbol(const void* symbol,const void* src,size_t count)

函数进行初始化,src 到symbol。

所有线程从相同的内存地址中读取数据，常量内存表现良好，比如数学公式中的系数。每当从一个常量内存中读取一次数据都会广播给线程束里的所有线程。

纹理内存：位于设备内存中，通过一种只读缓存访问全局内存。纹理内存是对二维空间的局部性优化。线程束里使用纹理内存访问二维数据的线程可以达到最优性能。

全局内存：GPU中最大，延迟最高，最常使用的内存。可以在任何SM设备上被访问到，贯穿应用程序的整个声明周期。

使用__device__静态声明一个变量。

使用cudaMalloc分配全局内存，使用cudaFree释放全局内存。然后指向全局内存的指针会作为参数传递给核函数。由于线程的执行不能跨线程块同步，所以不同线程块里的多个线程并发修改全局内存的同一位置可能会出现问题。最好是读。写会出错。

GPU缓存：

GPU上有4中缓存，一级缓存，二级缓存，只读常量缓存，只读纹理缓存。

一级缓存，每个SM都有一个一级缓存。

二级缓存，所有的SM共享一个二级缓存。

一级缓存和二级缓存都被用来在存储本地内存和全局内存中的数据，也包括寄存器溢出的部分。

每个SM上只用一个只读常量缓存和只读纹理缓存。

比如在orb的cuda版本中orgd.cu文件：

https://github.com/Accustomer/CUDA-ORB/blob/main/orbd.cu

4.静态全局内存

上代码：代码来自https://media.wiley.com/product_ancillary/29/11187393/DOWNLOAD/CodeSamples.zip

#include <cuda_runtime.h>
#include <stdio.h>

/*
 * An example of using a statically declared global variable (devData) to store
 * a floating-point value on the device.
 */

__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;
}

主机代码不能直接访问设备变量，设备代码不能直接访问主机变量

可以使用cudaGetSymbolAddress(void** devPtr,const void* symbol);获取全局变量的地址。

cudaGetSymbolAddress是一个CUDA的运行时API函数,它用于获取设备符号在设备内存中的地址。

功能:

symbol参数指定需要获取地址的设备符号名称。它必须是一个在设备端使用__device__ __constant__等修饰符定义的全局变量或者静态变量。

devPtr参数作为输出参数,函数执行后它指向的内存地址就是符号symbol在设备内存中的地址。

返回值为cudaSuccess表示成功获取地址,否则为错误码。示例：

__device__ float d_data; 

void* d_ptr;
cudaGetSymbolAddress(&d_ptr, d_data);

// d_ptr now points to address of d_data in device memory

这个函数主要用于获取在主机端通过指针直接访问设备内存中的符号数据的地址。

结合cudaMemcpyToSymbol预先在设备符号位置初始化数据,就可以实现主机直接读写设备符号数据的功能。或者cudaMemcpy。

有一个例外，可以直接从主机引用GPU内存：CUDA固定内存。主机代码和设备代码都可以通过简单的指针引用直接访问固定内存。

5.内存管理

CUDA编程的内存管理需要程序员显式的管理主机和设备之间的数据移动。

核函数在设备内存系统空间中运行。

内存分配和释放：

cudaMalloc（分配全局内存），cudaMemset（初始化或者填充），cudaFree（释放内存空间）

其中cudaMalloc与cudaFree是必须成对出现的。

内存传输：

cudaMemcpy（内存传输），具有4种类型，HostToHost，HostToDevice，DeviceToHost，DeviceToDevice。主机，设备。

#include <cuda_runtime.h>
#include <stdio.h>

/*
 * An example of using CUDA's memory copy API to transfer data to and from the
 * device. In this case, cudaMalloc is used to allocate memory on the GPU and
 * cudaMemcpy is used to transfer the contents of host memory to an array
 * allocated using cudaMalloc.
 */

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;
}

cuda编程的一个基本原则尽可能减少主机与设备之间传输。

一般分配的主机内存默认是可分页的，即因页面错误导致的操作，该操作按照操作系统的要求将主机虚拟内存上的数据移动到不同位置。GPU不能在可分页主机内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数据时，它无法控制。

所以固定内存或者锁页内存应运而生，将主机数据复制到固定内存中，然后从固定内存传输数据给设备内存。

可以使用cudaMallocHost直接分配主机锁页内存/固定主机内存。这些内存时页面锁定的并且对设备来说可访问的，由于固定内存能被设备直接访问，所以它能用比可分页内存高得多的带宽进行读写。

但是分配过多的固定内存可能会降低主机系统的性能，因为它减少了用于存储虚拟内存数据的可分页内存的数量，其中分页内存对主机系统可用的。

可以使用cudaFreeHost对锁页内存进行释放。

#include <cuda_runtime.h>
#include <stdio.h>

/*
 * An example of using CUDA's memory copy API to transfer data to and from the
 * device. In this case, cudaMalloc is used to allocate memory on the GPU and
 * cudaMemcpy is used to transfer the contents of host memory to an array
 * allocated using cudaMalloc. Host memory is allocated using cudaMallocHost to
 * create a page-locked host array.
 */

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;
}

零拷贝内存：通常来说，主机不能直接访问设备变量，同时设备也不能直接访问主机变量。担忧一个例外，零拷贝内存。主机个设备都可以访问零拷贝内存。GPU设备可以直接访问零拷贝内存。

在CUDA核函数使用零拷贝内存的好处：当设备内存不足时可以利用主机内存，避免主机和设备间的显示数据传输，提高PCIe传输率。

当使用零拷贝内存来共享主机和设备间的数据时，必须同步主机和设备间的内存访问，同时更改主机和设备的零拷贝内存中的数据将导致不可预知的后果。

零拷贝内存是固定（不可分页）内存，该内存映射在设备地址空间中。

可以通过 cudaHostAlloc创建一个固定映射。分配了一定大小的主机内存，该内存是页面锁定，并且设备可访问。同时也必须使用cudaFreeHost进行释放。

cudaHostAlloc 函数用于内存分配。这里简要介绍它的功能和用法：

目的：cudaHostAlloc 用于分配页锁定的主机内存（也称为固定内存）。这种类型的内存分配对于主机（CPU）和设备（GPU）之间的异步内存传输非常有利，因为GPU可以直接访问它，无需将数据复制到可分页内存中的缓冲区。

cudaError_t cudaHostAlloc(void **pHost, size_t size, unsigned int flags);

pHost：指向分配的内存的指针的地址。
size：要分配的内存大小，以字节为单位。
flags：控制分配类型的标志，例如 cudaHostAllocDefault、cudaHostAllocPortable、cudaHostAllocMapped 等。

用法：

调用 cudaHostAlloc 分配内存。
使用分配的内存进行数据处理或准备数据以供 GPU 使用。
使用完毕后，使用 cudaFreeHost 来释放内存。

可以使用cudaHostGetDevicePointer获取映射到固定内存的设备指针。其返回一个设备指针，该指针可以在设备上被引用来访问映射到固定主机内存。

如果设备不支持映射得到的固定内存，该函数失效。flag将留作以后使用。

零拷贝技术在CUDA高版本中没有彻底消失,但确实发生了一些变化:

在CUDA 6.0版本之前,主要使用cudaHostAlloc分配主机可见内存实现零拷贝。
从CUDA 6.0开始,引入了新的managed memory概念。通过cudaMallocManaged可以在CPU/GPU间共享内存,替代cudaHostAlloc实现零拷贝。
新的managed memory相比cudaHostAlloc,分配位置可以控制,且支持Unified Virtual Address空间。但在某些情况下性能可能不如直接使用cudaHostAlloc。
后续CUDA 7/8/9等版本主流API变更为使用managed memory实现零拷贝,cudaHostAlloc支持依旧存在但不再主推。
高版本CUDA中也引入了直接从GPU中读取主机内存的能力,如cudaMemPrefetchAsync等,可以实现更高效的主机读取。
同时一些第三方库如hipify也提供类似cudaHostAlloc的能力兼容低版本代码。

所以总体来说,零拷贝技术目前主流实现方式是通过managed memory,cudaHostAlloc作为可选方案依然存在。
但推荐使用统一的managed memory方式进行CPU/GPU数据共享。只是实现细节上有一些变化。

如果只是共享主机和设备端的少量数据，零拷贝内存是一个不错选择，如果数据过大，零拷贝内存会导致性能显著下降。

统一虚拟寻址（UVA）：有了UVA，主机内存和设备内存可以共享同一个虚拟地址空间。

在UVA之前，需要管理那些指针指向主机内存，那些指针指向设备内存。有了UVA，由指针指向的内存空间对应用程序来说都是透明的。

通过UVA，由cudaHostAlloc分配的固定主机内存具有相同的主机和设备指针。所以可以直接将返回的指针直接传递给核函数。不需要获取设备指针或管理物理数据完全相同的两个指针。

小结一下：

最基础的数据传输：

申请主机内存（主机指针），申请设备内存（设备指针），将主机内存数据移动至设备内存。

// malloc host memory
    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float *)malloc(nBytes);
    h_B     = (float *)malloc(nBytes);
    h_c = (float *)malloc(nBytes);

// initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, 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));

// 进行接下来的核函数计算

// 将结果传回主机

零拷贝数据传输：

申请主机内存（主机直至），得到指向主机内存的设备指针

// 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));

//剩余同上

可以看到省区申请设备内存一过程。

统一虚拟寻址：

申请主机内存（主机指针），直接将主机指针当作设备指针。

// 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);


    sumArrays<<<grid, block>>>(h_A, h_B, d_C, nElem);

当然这里结果d_c也得申请内存空间。

统一内存寻址：同意内存中创建了一个托管内存池，内存池已分配的空间可以用相同的内存地址即指针，在CPU和GPU上进行访问。底层系统在统一内存空间中自动在主机和设备之间进行数据传输。

统一内存寻址依赖与UVA的支持，但而这不同，UVA为系统的所有处理器提供了一个单一的虚拟内存地址空间。但是，UVA不会自动将数据从一个物理位置转移到另一个位置，这是统一内存寻址的特有功能。

统一内存寻址提供了一个“单指针到数据”模型，在概念上类似于零拷贝内存，但零拷贝内存是在主机内存中进行分配，因此受到在PCIe总线上访问零拷贝内存的一个像，核函数性能具有高延迟。

而虚拟内存寻址将内存和执行空间分离，因此可以根据需要将数据透明地传输到主机或者设备上，以提升局部性和性能。

托管内存指的是由底层系统自动分配的统一内存，与特定于设备的分配内存可以互操作，如它们的创建都使用cudaMalloc程序。因此，可以在核函数中使用两种类型的内存，由系统控制的托管内存，以及由应用程序明确分配和调用的未托管内存。所有在设备内存有效的CUDA操作，同样也适用于托管内存。其主要区别是主机也能够引用和访问托管内存。

静态分配与动态分配：

静态分配和动态分配是针对内存的两种管理方式,主要区别如下:

静态分配(Static Allocation):

编译期间就确定内存块的大小。

分配的内存空间大小和生存周期与整个程序挂钩。

典型形式有全局和静态变量。

优点:效率高,不需要重复分配。缺点:大小固定,不便于动态调整。

动态分配(Dynamic Allocation):

运行期间根据实际需要动态分配内存。

使用malloc/free、new/delete等函数在堆上申请和释放内存。

分配的内存大小和生存周期由程序在运行过程中决定。

优点:灵活性好,可以动态申请和释放内存。缺点:效率略低于静态分配。

区别在于:

静态分配大小在编译期间确定,动态分配可以在运行期间决定。

静态分配内存周期与全局变量一致,动态分配可以在任意时刻释放。

一般来说,大小固定的场景使用静态分配,需要动态调整的场景使用动态分配。两种方式各有优势。

托管内存可以被静态分配，也可以被动态分配。可以通过添加__managed__注释，静态声明一个设备变量作为托管变量。但这个操作只能在文件范围和全局范围内进行，该变量可以从主机和设备代码中直接被引用。

__device__ __managed__ int y

也可以使用cudaMallocManaged函数进行动态分配托管内存。

6.内存访问模式

对齐与合并访问：

全局内存是一个逻辑内存空间（即不是物理真实存在的，只是人们划分的），一般所有的应用程序数据最初存在与DRAM上，即物理设备中。核函数的内存请求通常是在DRAM设备和片上内存间以128字节或32字节的内存事务来是实现的。

访问通过一级缓存，会变为128字节的访问，只通二级缓存则变为32字节内存访问。

grace 1314

关注

0
点赞
踩
2

收藏

觉得还不错? 一键收藏
0
评论
CUDA全局内存(CUDA C编程权威指南)

访问共享内存必须同步。常量内存：使用__constant__来修饰常量变量，必须在全局空间内和核函数之外进行声明，大小只有64KB，常量内存时静态声明的，并对统一编译单元中所有的核函数可见。这种类型的内存分配对于主机（CPU）和设备（GPU）之间的异步内存传输非常有利，因为GPU可以直接访问它，无需将数据复制到可分页内存中的缓冲区。统一内存寻址提供了一个“单指针到数据”模型，在概念上类似于零拷贝内存，但零拷贝内存是在主机内存中进行分配，因此受到在PCIe总线上访问零拷贝内存的一个像，核函数性能具有高延迟。
复制链接

扫一扫