cuda编程知识普及

本帖经过多方整理,大多来自各路书籍《GPGPU编程技术》《cuda高性能》
 
1 gridblock都可以用三元向量来表示:
 
  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关键字
 
  使用这个关键字定义数组,设备会知道这个数组随时都会改变,就会自动重新读取数组(但是不能保证线程间读取的数据一致)
 
 
posted @ 2013-08-29 10:49  xingoo  阅读(5353)  评论(0编辑  收藏  举报