矩阵赋值CUDA实现

程序源码如下：

#include<stdlib.h>//系统头文件
#include<stdio.h>
#include<string.h>
#include<math.h>
#include<cutil.h>//项目头文件

/************矩阵赋值********/
__global__ void testKernel(float* g_idata,float * g_odata)
{
 //shared memory,extern表示大小由host端的Ns参数确定
 extern __shared__ float sdata[];

 const unsigned int bid=blockIdx.x;//线程序 所在的bloc的索引号
 const unsigned int tid_in_block=threadIdx.x;//线程在block中的位置
 const unsigned int tid_in_grid=blockDim.x*blockIdx.x+threadIdx.x;
 //按行划分任务时，线程在整个grid中的位置

 //将数据从global读入shared memory,读入数据后进行一次同步，
 //保证计算时所有数据均已到位
 sdata[tid_in_block]=g_idata[tid_in_grid];
 __syncthreads();

 //计算，需要进行同步，确保要写入的数据已经被更新
 sdata[tid_in_block]*=(float)bid;
 __syncthreads();
 //将shred memory 中的数据写到global memory
 g_odata[tid_in_grid]=sdata[tid_in_block];
}

//函数声明
void runTest(int argc,char** argv);
//主函数
int main(int argc,char** argv)
{
 runTest(argc,argv);
 CUT_EXIT(argc,argv);//退出CUDA
}

void runTest(int argc,char** argv)
{
 CUT_DEVICE_INIT(argc,argv);//启动CUDA
 unsigned int i,j;

 //为数据分配的存储器大小，这里用每一个线程计算一个单精度浮点数
 unsigned int num_blocks=4;
 unsigned int num_threads=4;//
 unsigned int mem_size=sizeof(float)* num_threads*num_blocks;//
 
 //在host端分配显存，d_表示device端
 float* h_idata=(float*)malloc(mem_size);
 float* h_odata=(float*)malloc(mem_size);
 //在device端分配显存，d_表示device端
 float* d_idata;
 CUDA_SAFE_CALL(cudaMalloc((void**)&d_idata,mem_size));
 float* d_odata;
 CUDA_SAFE_CALL(cudaMalloc((void**)&d_odata,mem_size));

 //初始化内在中的值
 for(i=0;i<num_threads*num_blocks;i++)
  h_idata[i]=1.0f;
 //将内存中的输入数据读入显存，这样就完成了主机对设备的数据写入
 CUDA_SAFE_CALL(cudaMemcpy(d_idata,h_idata,mem_size,cudaMemcpyHostToDevice));

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

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

 //检查GPU是否正常运行
 CUT_CHECK_ERROR("Kernel execution failed");

 //将结果从显存写入内存
 CUDA_SAFE_CALL(cudaMemcpy(h_odata,d_odata,mem_size,cudaMemcpyDeviceToHost));

 //打印结果
 for(i=0;i<num_blocks;i++)
 {
  for(j=0;j<num_threads;j++)
  {
   printf("%5.0f",h_odata[i*num_threads+j]);
  }
  printf("\n");
 }

 //释放存储器
 free(h_idata);
 free(h_odata);
 CUDA_SAFE_CALL(cudaFree(d_idata));
 CUDA_SAFE_CALL(cudaFree(d_odata));
}
函数解释：

1.CUT_DEVICE_INIT(argc,argv):启动CUDA宏函数

2.CUT_EXIT(argc,argv):退出CUDA宏函数

3.CUDA_SAFE_CALL()：调用cutil库操作需要的宏函数，调用后的返回值为cudaerr型

4.CUT_CHECK_ERROR():宏函数可以接受最近一次的cudaerror_t异常，如果发生异常，将输出 相应的错误类型，cudaerror_t定义在toolkit里的\include\driver_types.h

5.CUDA_SAFE_CALL(cudaMalloc(size))：在显存global memory上分配大小为size字节的空间。

6.CUDA_SAFE_CALL(cudaFree())：在释放显存global memory上分配空间。

7.CUDA_SAFE_CAL(cudaMemcpy()):用于拷贝存储器中的数据，其中第一个参数是指向目的的指针，第二上参数是指向源的指针，第三个参数是需要拷贝的字节数，第四个参数是拷贝操作的类型。拷贝操作类型共有三种：

       cudaMemcpyDeviceToHost将显存中的数据拷贝到内存中。

       cudaMemcpyHostToDevice将内存中的数据拷贝到显存中。

       cudaMemcpyDeviceToDevice将global memory中的数据拷贝到同一CUDA上下文的global memory的另一区域中。