cuda编程知识普及 - 码农教程

本帖经过多方整理，大多来自各路书籍《GPGPU编程技术》《cuda高性能》

1 grid 和 block都可以用三元向量来表示：

grid的数组元素是block

　　block的数组元素是grid

但是1.x计算能力的核心，grid的第三元必须为1.block的X和Y索引最大尺寸为512

2 通过__launch_bounds__(maxBlockSize,minBlocksPerMp)来限制每个block中最大的线程数，及每个多处理器上最少被激活的block数

3 SM streaming multiprocessor 多流处理器

SP scalar processor cores 标量处理核心

一个Block中的所有线程在一个多处理器上面并发执行。当这个Block的所有线程执行完后，再激活其他等待的Block.一个多处理器上也可以执行多个block。但是一个block却不能拆分为多个处理器上面执行

对于同一个Block里面的线程：

1 同一个Block里的线程可以被同步

2 可以共同访问多处理器里的共享存储器

到2.x为止，多处理器执行任务时，以32个并行线程为单位，称为一个wrap。

当以个block到来的时候，会被分成线程号连续的多个wrap，然后多处理器上的SIMT控制器以wrap为单位控制调度线程。所以block中的线程数要是以32的整数倍来设计，就不会出现空闲的SP。组织WARP的时候，从线程号最小的开始

4 各个存储器存储位置及作用

5 寄存器放在SP中，如果溢出，会被放在设备处理器上面，发生严重滞后，影响性能。

1.0　　 4KB 
2.0　　 16kb

共享存储器位于SM中，大约两个时钟周期读写4B，静态分配 __shared__ int shared[16];

1.0　　 16KB 
2.0　　 48kb

6 共享存储器，是以4个字节为单位的16个存储器组

　　bank冲突：半个warp中的多线程访问的数组元素处于同一个bank时，访问串行化，发生冲突

　　避免冲突：最多的数据类型是int、float等占用4个字节的类型

7线程设计

float shared=data[base+tid];
base访问的起始元素下标 tid线程号

如果要是char类型，每个元素占1个字节，就会冲突

float shared = data[base+4*tid];

8 共享存储器广播访问：半个warp线程都访问一个数据

9 补白策略

    shared[tid]=global[tid];
 
    int number = shared[tid*16];
    int nRow = tid/16;
    int nColumn = tid%16;
    shared[nColumn*17+nRow] = global[tid];
 
    int number = shared[17*tid];

10 一次性访问全局存储器：数据的起始地址应为每个线程访问数据大小的16倍的整数倍

11 主机锁页存储器：cudaHostMalloc()分配。

不参与操作系统分页管理的存储空间，访问锁页文件不会耗费主机内存分页管理方面的开销。不会被操作系统放到硬盘的页面文件中，因此比访问普通的主机存储器更快。

12 计算能力2.x的GPU上面，每个SM有独立的一级缓存，有唯一的二级缓存

13 异步并发：

主机上的计算、

设备上的计算、

主机到设备上的传输、

设备到主机上的传输共同执行

14 设备存储器 类型是DRAM，动态随机存储器。使用它最高效的方式就是顺序读取。为了保证顺序：

__global__ static void sumof(int *pnNumber,int* pnResult,clock_t* pclock_tTime){
    const int tid = threadIdx.x;
    int nSum = 0;
    int i;
    clock_t clock_tStart;
    if(tid == 0) clock_tStart = clock();
 
    for(i = tid;i<DATA_SIZE;i+=THREAD_NUM){
        nSum += pnNumber[i]*pnNumber[i];
}
 
    pnResult[tid] = nSum;
    if(tid == 0)
        *pclock_tTime = clock()-clock_tStart;
}

每个block 在1.x的计算能力的GPU下，最多只有512的线程数

__global__ static void sumof(int *pnNumber,int* pnResult,clock_t* pclock_tTime){
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    int nSum = 0;
    int i;
    clock_t clock_tStart;
    if(tid == 0) pclock_tTime[bid] = clock();
 
    for(i = bid*THREAD_NUM+tid;i<DATA_SIZE;i+=BLOCK_NUM*THREAD_NUM){
        nSum += pnNumber[i]*pnNumber[i];
}
 
    pnResult[bid*THREAD_NUM+tid] = nSum;
 
    if(tid == 0)
        *pclock_tTime[bid+BLOCK_NUM] = clock();
}

15 用缩减树避免bank冲突：

　　bank冲突指的是，一个warp内的线程同时访问一个bank列，导致串行读取数据

    noffset = THREAD_NUM/2;
    while(noffset > 0){
        if(tid < offset)
            nshared[tid] += nshared[tid+noffset];
    }
    noffset >>= 1;
 
    __syncthreads();

16 CPU有强大的分支预测、程序堆栈、循环优化等针对控制采取的复杂逻辑。

GPU相对简单，适合处理顺序的，单一的，少循环，少跳转的语句。

17 #progma unroll 5下面的程序循环5次

18 cuda中的同步

1》__syncthreads()同步

　　同一个warp内的线程总是被一同激活且一同被分配任务，因此不需要同步。因此最好把需要同步的线程放在同一个warp内，这样就减少了__syncthreads()的指令

2》__threadfence() __threadfence_block()同步

　　前者针对grid的所有线程，后者针对block内的所有线程。告知线程，全局存储器或共享存储器已经被改变

3》cudaThreadSynchronize() 主机与设备间的同步

　　在主机程序里同步线程。该函数以上的设备线程完成后，控制权才交给cpu

4》volatile关键字