CUDA C编程权威指南 第四章 全局内存

基础

1. 一般内存的设计:寄存器->缓存->主存->磁盘存储器
2. GPU内存设计

修饰符	变量名	存储器	作用域	生命周期
float var	寄存器	线程	线程
float var[100]	本地	线程	线程
__shared__	float var +	共享	块	块
__device__	float var +	全局	全局	应用程序
__constant__	float var +	常量	全局	应用程序

全局变量声明__device__

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

__device__ float devData;

__global__ void checkGlobalVariable(){
    printf("Device: the value of the global variable is %f\n", devData);
    devData += 2.0f;
}

int main(void)
{
    float value = 3.14f;
    cudaMemcpyToSymbol(devData, &value, sizeof(float));
    printf("Host:   copied %f to the global variable\n", value);

    checkGlobalVariable<<<1, 1>>>();
    cudaMemcpyFromSymbol(&value, devData, sizeof(float));
    printf("Host:   the value changed by the kernel to %f\n", value);
    cudaDeviceReset();
    return EXIT_SUCCESS;
}

cudaMemcpyToSymbol(devData, &value, sizeof(float)); 将host内存拷贝到device
cudaMemcpyFromSymbol(&value, devData, sizeof(float));将device内存拷贝到device

不能使用cudaMemcpy(&devData,&value,siezeof(float),cudaMemcopyHostToDevice);,因为不能使用"&",&devData在device端,而不在host端

通过cudaError_t cudaGetSymbolAddress(void**,devPtr, const void* symbol);获得全局变量的地址:

float *dftr = NULL;
cudaGetSymbolAddress((void**)&dptr, devData);
cudaMemcpy(dptr,&value,sizeof(float), cudaMemcpyHostToDevice);

内存管理

内存分配与释放

分配和释放内存耗时较高
cudaError_t cudaMalloc(void** devPtr, size_t count)来分配全局内存
cudaError_t cudaMemset(void *devPtr, int value, size_t count);将host内存拷贝到device
cuda Error_t cudaFree(void *devPtr);释放内存

内存传输

cudaError_t cudaMemcpy(void *dst, const void *src,size_t count, enum cudaMemcpyKind kind);从host的src赋值count字节到device的dst

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

int main(int argc, char **argv)
{

    unsigned int isize = 1 << 22;
    unsigned int nbytes = isize * sizeof(float);

    cudaDeviceProp deviceProp;
    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));

    float *h_a = (float *)malloc(nbytes);
    float *d_a;
    cudaMalloc((float **)&d_a, nbytes);

    for(unsigned int i = 0; i < isize; i++) h_a[i] = 0.5f;

    cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);

    cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost);

    cudaFree(d_a);
    free(h_a);

    cudaDeviceReset();
    return EXIT_SUCCESS;
}

固定内存

cudaError_t cudaMallocHost(void **devPtr, size_t count);在host分配count字节内存,页面锁定,且device可以访问,比可分页内存有更高的带宽

固定内存的分配和释放成本更高,但是它为大规模数据传输提供了更高的传输吞吐量。

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


int main(int argc, char **argv){

    // memory size
    unsigned int isize = 1 << 22;
    unsigned int nbytes = isize * sizeof(float);

    // host 固定内存(pin memory)
    float *h_a;
    cudaMallocHost ((float **)&h_a, nbytes);
    // 分配设备内存
    float *d_a;
    cudaMalloc((float **)&d_a, nbytes);

    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
    cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);

    // transfer data from the device to the host
    cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost);

    // free memory
    cudaFree(d_a);
    cudaFreeHost(h_a);//固定内存必须通过cudaFreeHost释放
    cudaDeviceReset();
    return EXIT_SUCCESS;
}

零拷贝内存

1. 一般来说host不能访问device内存,device不能访问host内存
2. host和device都可以访问零拷贝内存
3. 零拷贝内存是固定(不可分页)内存
   零拷贝内存作用:
4. 当device内存不足时可以利用host内存
5. 避免主机和设备的显示数据传输
6. 提高PCIe传输率

cuda Error_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);分配count字节host内存,必须使用cudaFreeHost函数释放

flags	描述
cudaHostAllocDefault	使cudaHostAlloc函数的行为与cudaMallocHost函数一致
cudaHostAllocPortable	返回能被所有CUDA上下文使用的固定内存，而不仅是执行内存分配的那一个
cudaHostAllocWriteCombined	返回写结合内存，该内存可以在某些系统配置上通过PCIe总线上更快地传输
cudaHostAllocMapped	该标志返回，可以实现主机写入和设备读取被映射到设备地址空间中的主机内存

通过cudaError_t cudaHostGetDevicePointer(void **pDevice , void *pHost, unsigned int flags);获得映射到固定内存的设备指针
该指针可以在device上被引用得到固定主机内存.
进行频繁读写的时候,零拷贝性能将显著降低,因为要经过PCIe总线来传递数据.
在集成架构(集显),CPU和GPU集成在一个芯片上,并且在物理地址上共享主存,零拷贝内存性能可能更好
在离散架构(独显),需要经过PCIe总线,零拷贝只有在特殊情况下才有优势

代码

统一寻址

统一虚拟寻址(UVA)

#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 data size of vectors
    int ipower = 10;
    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;
    cudaMalloc((float**)&d_A, nBytes);
    cudaMalloc((float**)&d_B, nBytes);
    cudaMalloc((float**)&d_C, nBytes);

    // transfer data from host to device
    cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
    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
    cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);

    // check device results
    checkResult(hostRef, gpuRef, nElem);

    // free device global memory
    cudaFree(d_A);
    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
    cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped);
    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
    cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0);
    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);//(h+A,H_B,d_C,nElem)才对吧

    // copy kernel result back to host side
    cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);

    // check device results
    checkResult(hostRef, gpuRef, nElem);

    // free  memory
    cudaFree(d_C);
    cudaFreeHost(h_A);
    cudaFreeHost(h_B);

    free(hostRef);
    free(gpuRef);

    // reset device
    cudaDeviceReset();
    return EXIT_SUCCESS;
}

内存访问模式

1. 核函数的内存请求通常是在DRAM设备和片上内存间以128字节或32字节内存事务来实现的。
2. 如果这两级缓存都被用到，那么内存访问是由一个128字节的内存事务实现的。如果只使用了二级缓存，那么这个内存访问是由一个32字节的内存事务实现
3. 一行一级缓存是128个字节，它映射到设备内存中一个128字节的对齐段。如果线程束中的每个线程请求一个4字节的值，那么每次请求就会获取128字节的数据，这恰好与缓存行和设备内存段的大小相契合。
4. 优化应用程序时，你需要注意设备内存访问的两个特性：对齐内存访问-·合并内存访问
5. 1. 对齐与合并内存访问,warp中所有的thread请求地址都在128B的缓存行范围内,完成内存加载只需要一个128B的事务,总线的使用率为100%

2. 访问时对齐的,引用的地址不是连续的thread ID,而是128B范围内的随机值.warp 中的Thread请求地址,仍然在一个缓存行范围内,所以只需要一个128B的事务来完成这一内存加载操作.总线利用率仍是100%.

3. warp请求32个连续4个字节的非对齐数据元素.在全局内存中warp的thread请求落在2个128字节的范围内.因此有两个128字节的事务来加载内存操作,总线利利用率为50%,并且在两个事务中加载的字节有一半是未使用的.

4. warp所有thread请求相同的地址,因为被引用的字节落在一个缓存范围内,所以只需要请求一个内存事务,但是总线的利用率非常低,如果加载的是4字节的,那么总线利用率就是4/128=3.125%

5. 最坏的情况:warp的thread分散于全局内存中的32个4字节地址.,地址要占用N个缓存(0<N<32),完成一次内存加载需要申请N次内存事务

结构体数组和数组结构体

SoA模式存储数据充分利用了GPU的内存带宽,GPU上SoA布局提供给了合并内存访问,对全局内存实现更高效的利用.

核函数可达到的带宽