本帖经过多方整理，大多来自各路书籍《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关键字
使用这个关键字定义数组，设备会知道这个数组随时都会改变，就会自动重新读取数组（但是不能保证线程间读取的数据一致）
