CUDA

CUDA

Posted by infinityyf on April 10, 2020

基本使用

在CPU端定义函数，上传到GPU端定义

__global__ void helloWorldGPU(void) {
	std::cout << "hello" << std::endl;
//__global__从主机端调用，在设备端执行
//__device__从设备端调用，在设备端执行
//__host__从主机端调用，主机端执行
}


//使用该函数启动内核代码
helloWorldGPU<<<1,10>>>()//启动10个线程

cudaDeviceRest//使用该函数释放和清空当前设备持有的资源

  1. 使用步骤
    分配内存
    将CPU数据拷贝到GPU
    调用CUDA核函数
    拷贝回CPU
    释放GPU内存

  2. 内存管理
    使用cudaMalloc分配线性内存，使用cudaMemset,cudaMemcpy,cudaFree进行管理，这些函数将在GPU全局内存中分配。cudaMemcpy的调用混导致主机的阻塞

  3. 线程管理
    cuda使用两层的线程层次结构。一个内核启动所有的线程统称一个网格。同一个网格中的线程共享相同的全局内存空间。一个网格由多个线程块构成，一个线程块包含一组线程。CUDA在运行时会为每个线程分配坐标变量，blockIdx,threadIdx,通过这些变量将数据分配给线程。

这两个变量是基于uint3实现的，所以每个都有xyz三个分量。网格的维度核线程块的维度通过blockDim，gridDim来查询，同样也是包含三个分量。

一般网格回被组织成线程块的二维数组，而线程块为线程的三维数组。

  1. 启动核函数
    grid和block均可以自定义
    kernalname«<grid,block»>()
    grid：网格的维度
    block: 块的维度
    在核函数中使用__syncthreads()同步所有线程块中的线程

使用cudaDeviceSynchronize强制使主机端程序等待所有的核函数执行结束

  1. 在其他文件中调用
    用来被调用的CUDA函数中要加上extern “C” 的声明，并在cpp文件中进行声明（extern “C” int runtest(int *host_a, int *host_b, int *host_c);）后再调用

一些概念

  1. warp warp是基本的执行单元，包含32个并行thread，一个线程块的thread只能在一个SM上调度
  2. SP
    最基本的处理单元，也称为CUDA core
  3. SM
    多个SP组成SM
  4. 不同划分的threadid计算方法
    grid划分成1维，block划分为1维
    int threadId = blockIdx.x *blockDim.x + threadIdx.x;

grid划分成1维，block划分为2维
int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;

grid划分成1维，block划分为3维
int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z + threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;

grid划分成2维，block划分为1维   int blockId = blockIdx.y * gridDim.x + blockIdx.x;
int threadId = blockId * blockDim.x + threadIdx.x;

grid划分成2维，block划分为2维
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y)+ (threadIdx.y * blockDim.x) + threadIdx.x;

grid划分成2维，block划分为3维
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z) + (threadIdx.z * (blockDim.x * blockDim.y))+ (threadIdx.y * blockDim.x) + threadIdx.x;

grid划分成3维，block划分为1维
int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * blockDim.x + threadIdx.x;

grid划分成3维，block划分为2维   int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y)+ (threadIdx.y * blockDim.x) + threadIdx.x;

grid划分成3维，block划分为3维
int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;   int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)+ (threadIdx.z * (blockDim.x * blockDim.y))+ (threadIdx.y * blockDim.x) + threadIdx.x;

同步

块内通讯是通过共享内存进行，不同块之间不能通讯。通过__syncthreads()位于device_functions.h头文件中函数来同步以可block中所有的线程，而cudaThreadSynchronize()用于同步所有的线程。cudaDeviceSynchronize()用于同步CPU和GPU

共享内存

是GPU上可控的一级缓存

__shared__ float a_in[34]; //定义共享内存
extern __shared__ float a[];//动态分配共享内存
ddkernel<<<Grids, Blocks,smemSize>>>(paramenter);//调用核核函数的时候使用第三个参数指定大小

申请共享内存后，其内容在每一个用到的block被复制一遍，使得在每个block内，每一个thread都可以访问和操作这块内存，而无法访问其他block内的共享内存 共享内存的使用是优化CUDA程序的核心方法，通过对全局数据进行合并访问，让kernel内交错的内存需求去访问共享内存。共享内存是块内的，所以使用完之后需要使用__syncthreads()进行同步还有通过fence进行同步，分别是__threadfence_block()，__threadfence_grid()，__threadfence_system()

为了获得高带宽，共享内存分为32个相等的内存块，可以同时访问，wrap中有三种获取共享内存的模式

wrap

在硬件上并不是所有的thread都是并行的，wrap是SM的基本执行单元，一个wrap并行32个thread，若block所含的线程数不是wrap的整数倍，那么就会有一些thread是未激活的。

__restrict__修饰符

restrict修饰符是c语言中存在的一个语义限定修饰，例如在下面的实例中

void foo(const float* a,const float* b,float* c)
{
    c[0] = a[0] * b[0];
    c[1] = a[0] * b[0];
    c[2] = a[0] * b[0] * a[1];
    c[3] = a[0] * a[1];
    c[4] = a[0] * b[0];
    c[5] = b[0];
    ...
}
c和a可能指向同一个地址，提前写入可能会出现问题，通过使用restrict限定，编译器会进行优化
void foo(const float* __restrict__ a,const float* __restrict__ b,float* __restrict__ c)
{
    float t0 = a[0];
    float t1 = b[0];
    float t2 = t0 * t2;
    float t3 = a[1];
    c[0] = t2;
    c[1] = t2;
    c[4] = t2;
    c[2] = t2 * t3;
    c[3] = t0 * t3;
    c[5] = t1;
    ...
}
CATALOG
    FEATURED TAGS
    C++ CUDA