http://blog.csdn.net/gamesdev/article/category/1778017
处理DATA_SIZE = 1048576个随机数（int）数据（4M）的平方和。
#define DATA_SIZE 1048576
#define THREAD_NUM 256 如果设置了多线程，线程数量是256
#define BLOCK_NUM 32 如果设置了多线程块，就是32
VERSION 1
要点：1. 用全局变量g_Data[DATA_SIZE]存放原始数据，cudaMalloc分配显存空间，再将g_Data复制过来
2. 只用一个块，一个线程计算
3. 结果也只用一个int指针保存，从cudaMalloc分配的显存空间复制回到局部变量result
===============================================================
__global__ static void kernerl_SquareSum( int* pOut, int *pIn, size_t* pDataSize )
{
for ( size_t i = 0; i <*pDataSize; ++i )
{
*pOut += pIn[i] * pIn[i];
}
}
===============================================================
VERSION2
要点：1. 4M的数据不再分配在全局变量区域，使用CUDA内存锁页存储器分配host端，锁页存储器是一种不受分页影响的存储器，由于这一特性，使得读取的速度非常快
2. 只用1个块，每个块中各有THREAD_NUM个线程在运行，每个线程串行处理DATA_SIZE/THREAD_NUM个数据
3. 结果用一个result[THREAD_NUM]，传回内存锁页存储器分配的host端后，将THREAD_NUM个线程的结果在CPU中累加
===============================================================
__global__ static void kernerl_SquareSum( int* pOut, int *pIn, size_t* pDataSize)
{
const size_t computeSize = *pDataSize/THREAD_NUM;
const size_t thID = threadIdx.x;
for(size_t i = thID *computeSize; i < ( thID + 1 ) * computeSize; ++i )
{
pOut[thID] += pIn[ i] * pIn[ i];
}
/*
for ( size_t i = 0; i < computeSize; ++i )
{
pOut[thID] += pIn[thID*computeSize + i] * pIn[thID*computeSize + i];
}
*/
}
===============================================================
VERSION3
要点： 在CUDA对线程进行调度时，通常将时间片按照线程的ID来平均分配给每一线程，尽管时间片的周期非常小，系统依然按照“线程0、线程1、线程2……”进行调度。而如果让每一个线程做到平均跨越式的访问，那么整体看来就像存储器（显存）被顺序访问一样。因此我们可以让每一个线程间隔THREAD_NUM个字节进行访问，那么第一轮所有的线程都会访问[0,THREAD_NUM-1]字节的数据，第二轮所有的线程都会访问[THREAD_NUM，2 * THREAD_NUM]字节的数据。这样做到了顺序访问。
for ( size_t i = 0; i < computeSize; ++i )
{
pOut[thID] += pIn[thID*computeSize + i] * pIn[thID*computeSize + i];
}
时间片周期非常小的周期内：这段代码的访问顺序是 线程0访问pIn[0]->线程1访问pIn[computeSize]->线程2访问pIn[2*computeSize].....存储器是顺序存储的但是却没有被顺序访问。
for(size_t i = thID ; i < *pDataSize;i = i+THREAD_NUM )
{
pOut[thID] += pIn[ i] * pIn[ i];
}
这段代码的的访问顺序是 线程0访问pIn[0]->线程1访问pIn[1]->线程2访问pIn[2].....做到顺序访问。
===============================================================
__global__ static void kernerl_SquareSum( int* pOut, int *pIn, size_t* pDataSize)
{
const size_t thID = threadIdx.x;
for(size_t i = thID ; i < *pDataSize;i = i+THREAD_NUM )
{
pOut[thID] += pIn[ i] * pIn[ i];
}
}
===============================================================
VERSION4
要点： 1. 做到了"顺序"访问的同时，划出更多的块，让一个SM多个块并行，多个SM并行。
2. 用一个BLOCK_NUM*THREAD_NUM的矩阵来记录每个线程的结果。
===============================================================
__global__ static void kernerl_SquareSum( int* pOut, int *pIn, size_t* pDataSize)
{
const size_t thID = threadIdx.x;
const size_t bloID = blockIdx.x;
for(size_t i = bloID*THREAD_NUM + thID ; i < *pDataSize;i = i+BLOCK_NUM*THREAD_NUM )
{
pOut[bloID*THREAD_NUM + thID] += pIn[ i] * pIn[ i];
}
}
===============================================================
VERSION5
要点：1.在CPU上需要将所有线程的结果串行求和：
for ( int i = 0; i < BLOCK_NUM*THREAD_NUM; ++i )
{
result += pResult[i];
}
这是很大一部分的负担。这次主要精在如何将这一部分代码放在GPU上运行。也就是每个线程块的结果用result[BLOCK_NUM]保存，传回内存锁页存储器分配的host端后，将BLOCK_NUM个线程的结果在CPU中累加。
2.引入共享内存。共享存储器位于每个多处理器内，是片上存储器，它的作用域在一个块（BLOCK）内。于是在块内进行各个线程的结果求和，比较符合我们的期望，也就是说，线程执行y=X^2这样的操作，将结果放在共享存储器中，再选择一个线程执行z=∑y的操作，最后将w=∑z交给CPU来执行。
结果：受限于GPU每个线程的延迟 带宽变小
===============================================================
__global__ static void kernerl_SquareSum( int* pOut, int *pIn, size_t* pDataSize)
{
const size_t thID = threadIdx.x;
const size_t bloID = blockIdx.x;
__shared__ int sharedData[THREAD_NUM];
for(size_t i = bloID*THREAD_NUM + thID ; i < *pDataSize;i = i+BLOCK_NUM*THREAD_NUM )
{
sharedData[thID] += pIn[i] * pIn[ i];
}
__syncthreads( );
if ( thID == 0 )// 由0号线程完成数据的累加
{
pOut[bloID] = 0;// 先初始化为
for ( size_t i = 1; i <THREAD_NUM; i++ )
{
pOut[bloID] += sharedData[i];
}
}
}
===============================================================
VERSION6
要点：上次的程序是使用了CUDA的共享存储器进行累加的运算，不过这里面有一个比较明显的资源浪费现象：
最后求和的时候通过__syncthreads()函数对一个块（BLOCK）中的所有线程进行同步，但是仅仅使用了一个线程进行求和累加的运算，其余的线程无事可做。为了防止这样的现象发生，我们需要借助并行缩减算法中比较常见的缩减树算法。
int offset = 1; // 记录每轮增倍的步距
int mask = 1; // 选择合适的线程
while ( offset < THREAD_NUM )
{
if ( ( tID & mask ) == 0 ) //2的倍数，4的倍数，8的倍数
{
sharedData[tID] += sharedData[tID + offset];
}
offset += offset; // 左移一位
mask = offset + mask; // 掩码多一位二进制位
__syncthreads( );
}
===============================================================
__global__ static void kernerl_SquareSum( int* pOut, int *pIn, size_t* pDataSize)
{
const size_t thID = threadIdx.x;
const size_t bloID = blockIdx.x;
int offset = 1;
int mask = 1;
__shared__ int sharedData[THREAD_NUM];
for(size_t i = bloID*THREAD_NUM + thID ; i < *pDataSize;i = i+BLOCK_NUM*THREAD_NUM )
{
sharedData[thID] += pIn[i] * pIn[ i];
}
__syncthreads( );
while ( offset < THREAD_NUM )
{
if ( ( thID & mask ) == 0 )
{
sharedData[thID] += sharedData[thID + offset];
}
offset += offset; // 左移一位
mask = offset + mask; // 掩码多一位二进制位
__syncthreads( );
}
if ( thID == 0 )// 如果线程ID为，那么计算结果
{
pOut[bloID] = sharedData[0];
}
}
===============================================================
VERSION7
要点：1.在使用共享存储器的时候要适当地避免bank冲突，。
2.一般来说，使用流程控制产生的指令比顺序执行所产生的指令要多，如果尝试将两个执行效果等价但一个使用了流程控制另一个只是顺序执行的程序代码所生成的汇编代码相比较，那么结果表明使用流程控制控制所产生的汇编代码会更多。
此外，过多地使用流程控制将会产生上述对warp内的线程产生阻塞，我们应当适当地规避。于是我们这次将内核代码中的流程控制改写一下，希望能够产生更少的指令以及更高的数据带宽。
===============================================================
__global__ static void kernerl_SquareSum( int* pOut, int *pIn, size_t* pDataSize)
{
const size_t thID = threadIdx.x;
const size_t bloID = blockIdx.x;
__shared__ int sharedData[THREAD_NUM];
for(size_t i = bloID*THREAD_NUM + thID ; i < *pDataSize;i = i+BLOCK_NUM*THREAD_NUM )
{
sharedData[thID] += pIn[i] * pIn[ i];
}
__syncthreads( );
if(thID<128) sharedData[thID] += sharedData[thID+128];
__syncthreads( );
if ( thID < 64 ) sharedData[thID] += sharedData[thID + 64];
__syncthreads( );
if ( thID < 32 ) sharedData[thID] += sharedData[thID + 32];
//__syncthreads( ); 为什么会注释掉呢？因为每一个warp包含了32个线程，而warp是线程同时执行的最小单位，因此我们确定在一个warp中的线程没有必要进行同步。
if ( thID < 16 ) sharedData[thID] += sharedData[thID + 16];
//__syncthreads( );
if ( thID < 8 ) sharedData[thID]+= sharedData[thID + 8];
//__syncthreads( );
if ( thID < 4 ) sharedData[thID]+= sharedData[thID + 4];
//__syncthreads( );
if ( thID < 2 ) sharedData[thID]+= sharedData[thID + 2];
//__syncthreads( );
if ( thID < 1 ) sharedData[thID]+= sharedData[thID + 1];
if ( thID == 0 )// 如果线程ID为0，那么计算结果
{
pOut[bloID] = sharedData[0];
}
}
===============================================================
总结：
模仿CUDA程序优化（一）~（九），我自己实现了一些自己能理解的，其中有两个“在使用共享存储器的时候要适当地避免bank冲突”还有最后用PTX的指令代替乘法指令*没有实现。