GPU结构

每个GPU中有若干个GPC（图形计算簇），每个GPC中又有若干个TPC（Texture Processing cluster), 每个TPC中有若干个SM（Stream Multiprocessor),每个SM中有若干个block，在block中有各种各样的核和寄存器。

在GPU编程中，通常把cpu称作host，把gpu称作device。

计算过程：

  • 数据从cpu拷贝到gpu显存中
  • 加载GPU程序并运行
  • 将运行完后的数据又送回CPU中

CUDA编程

基本结构：

malloc();
cudamalloc(); //分配显存
cudaMemcpy(): //将内存中的内容放到显存中
kernelName<<<gridSize, blockSize>>>(input_params);//计算过程，input_params是运行的函数

cudaMemcpy();
cudaFree();

例如：

__global__ void add(int *a, int *b, int *c)
{
    *c = *a + *b;
}

int main(void) 
{
    int a, b, c;// host copies of a, b, c
    int *d_a, *d_b, *d_c;// device copies of a, b, c
    in tsize = sizeof(int);// Allocate space for device copies of a, b, c，
    
    分配内存
    cudaMalloc((void**)&d_a, size);
    cudaMalloc((void**)&d_b, size);
    cudaMalloc((void**)&d_c, size);// Setup input values，
    
    设置输入数据
    a = 2;
    b = 7;
    
    // Copy inputs to device，拷贝数据到GPU的显存
    cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);
    
    // Launch add() kernel on GPU，执行kernel
    add<<<1,1>>>(d_a, d_b, d_c);
    
    // Copy result back to host，数据拷贝回host内存
    cudaMemcpy(&c, d_c, size, cudaMemcpyDeviceToHost);
    
    // Cleanup，释放内存cudaFree(d_a); 
    cudaFree(d_b); cudaFree(d_c);return 0;
}

函数前的标识：

  • __global__: 表示在device上执行，从host中调用。它的返回类型必须是void，并且它不会等待函数执行完主线程就会进行下一步。
  • __device__: 在device上执行，并且从device中调用
  • __host__: host执行，host调用，默认值

函数中变量类型：

  • __device__: grid中所有线程都可以访问
  • __constant__: 常量内存
  • __shared__: 只有同一个block中的线程才可以访问
  • <<>>: 调用的基本结构为 grid(网格）->block(块)->thread(线程).例如<<<1, 1>>>代表这个网格中有一个块，这个块中有一个线程。

grid和thread都是dim3型变量，dim3可以看成包含三个整数的结构体。我们可以自由的定义grid或thread为一维、二维或三维,他们的默认值都是1。例如：

dim3 block(3, 3);
add<<<1, block>>>(...);
表示有一个块，每块中有(3, 3)个线程

上面的例子可以变成
add<<<N, 1>>>(d_a, d_b, d_c);

__global__ void add(int *a, int *b, int *c)
{
    c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

每一个线程都可以用(blockIdx, threadIdx)进行标识，这两个变量都是dim3类型的，因此上面的例子中就可以使用blockIdx.x表示这是第几个block.

变量：

  • gridDim: grid维度,dim3型
  • blockDim: block维度
  • blockIdx:
  • threadIdx:
  • warpSize: 大小固定为32

每个grid都是在一个SM上的，而SM空间有限，因此每次可以使用的线程是有上限的，一般是1024.

内存分配：

内存分为两种

  • 可分配内存： malloc()/new() 操作所得到的内存，这种内存由cpu进行管理，可能被分页机制换下

  • 分页锁定内存： 这种内存受gpu控制，不受cpu调度影响，可以使用dma直接读取

    • cudaMallocHost(void **ptr, size_t size)
    • cudaHostAlloc(void **pHost, size_t size, unsigned int flags):
      • cudaHostAllocPortable：多个CPU线程都可访问
      • cudaHostAllocWriteCombined：数据传输快，但只允许host写（读速度慢)
      • cudaHostAllocMapped：CPU/GPU都可访，属于内存映射
    • cudaHostRegister( void* ptr, size_tsize, unsigned intflags ): 可分页内存转换为分页锁定内存

    显存分配

    • cudaMalloc()： 一维数组，数据连续存储
    • cudaMallocPitch()： 二维数组，会自动对齐，可能会浪费部分内存空间
    • cudaMalloc3D()：三维数组

统一寻址

统一寻址指的是系统自动在host和device之间搬运数据，同时保证host和device都可访问，并且需要进行显式同步。（因为系统会在运行函数时同时进行下一步）

例如：

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

int main()
{
    int *ret;
    cudaMallocManaged(&ret, 1000 * sizeof(int));//自动进行显存内存分配及数据搬运
    add<<<1, 1000>>>(ret, 10, 100);
    cudaDeviceSynchronized();//进行同步
    for(int i=0; i<1000; i++)
    {
        printf("%d", ret[i]);
    }
    cudaFree(ret);
    return 0;
}