矩阵棋盘划分的CUDA程序

/*使用二维的grid和block将4*4的矩阵块划分为4个2*2的子阵，即棋盘划分*/
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <cutil.h>
#define SDATA(index) CUT_BANK_CHECKER(sdata,index)

__global__ void testKernel(float* g_idata,float* g_odata,int width,int height)
{
 //shared memory没有extern 前缀，静态分配空间，因此[]中必须指定为数组分配的大小
 __shared__ float sdata[4];
 //计算线程索引
 unsigned int bid_in_grid=__mul24(blockIdx.y,gridDim.x)+blockIdx.x;
 unsigned int tid_in_block=__mul24(blockIdx.y,gridDim.x)+threadIdx.x;
 unsigned int tid_in_grid_x=__mul24(blockDim.x,blockIdx.x)+threadIdx.x;
 unsigned int tid_in_grid_y=__mul24(blockDim.y,blockIdx.y)+threadIdx.y;
 unsigned int tid_in_grid=__mul24(tid_in_grid_y,width)+tid_in_grid_x;
 //从global memory中读入数据。host模拟运行时使用，bank checker检查是否产生bank conflict
 SDATA(tid_in_block)=g_idata[tid_in_grid];
 __syncthreads();
 //进行运算
    SDATA(tid_in_block)=(float)bid_in_grid*SDATA(tid_in_block);
 __syncthreads();
 //将数据写回global memory
 g_odata[tid_in_grid]=SDATA(tid_in_block);
}
__host__ void runTest(int argc,char** argv)
{
 CUT_DEVICE_INIT(argc,argv);//启动CUDA

 //数据大小，这里用每一个线程计算计算一个单精度浮点数
 unsigned int mem_size=sizeof(float)*4*4;

 //在host端分配内存，注意这里分配的是pinned memory
 float* h_idata;
 CUDA_SAFE_CALL(cudaMallocHost((void**)&h_idata,mem_size));
 //初始化线程中的值
 for(unsigned int i=0;i<4;i++)
  for(unsigned int j=0;j<4;j++)
   h_idata[i*4+j]=2.0f;
 //在device端分配显存
 float* d_idata;
 CUDA_SAFE_CALL(cudaMalloc((void**)&d_idata,mem_size));
 //将内存中的值读入显存
 CUDA_SAFE_CALL(cudaMemcpy(d_idata,h_idata,mem_size,cudaMemcpyHostToDevice));
 //在device端分配显存，用于存储结果
    float* d_odata;
 CUDA_SAFE_CALL(cudaMalloc((void**)&d_odata,mem_size));

 //设置运行参数，即网格的形状和线程块的形状
 dim3 grid(2,2,1);
 dim3 threads(2,2,1);

 //****************运行内核函数，调用调和GPU进行运算
 testKernel<<<grid,threads>>>(d_idata,d_odata,4,4);

 //检查GPU 是否正常运行
 CUT_CHECK_ERROR("Kernel execution failed");
 //
 float* h_odata;
 CUDA_SAFE_CALL(cudaMallocHost((void**)&h_odata,mem_size));
 //
 CUDA_SAFE_CALL(cudaMemcpy(h_odata,d_odata,mem_size,cudaMemcpyDeviceToHost));
 //
 for(unsigned int i=0;i<4;i++)
 {
  for(unsigned int j=0;j<4;j++)
  {
   printf("%5.0f",h_odata[i*4+j]);
        }
  printf("\n");
 }
 //释放存储器
 CUDA_SAFE_CALL(cudaFreeHost(h_idata));
 CUDA_SAFE_CALL(cudaFreeHost(h_odata));
    CUDA_SAFE_CALL(cudaFree(d_idata));
 CUDA_SAFE_CALL(cudaFree(d_odata));
}
//主函数
int main(int argc,char** argv)
{
 runTest(argc,argv);
 CUT_EXIT(argc,argv);// 退出CUDA
}