题目：给一个二维数组赋值。
分析：主机端代码完成的主要功能：
- 启动CUDA，使用多卡时应加上设备号，或使用cudaSetDevice（）设置GPU设备。
- 为输入数据分配内存空间
- 初始化输入数据
- 为GPU分配显存，用于存放输入数据
- 将内存中的输入数据拷贝到显存。
- 为GPU分配显存，用于存放输出数据。
- 调用device端的Kernel进行计算，将结果写到显存中的对应区域。
- 为CPU分配内存，用于存放GPU传回的输出数据
- 将显存中的结果回读到内存。
- 使用CPU对传回的数据进行其他的处理。
- 释放内存和显存空间。
- 退出CUDA
设备端代码要完成的任务：
- 从显存读取数据到GPU片内。
- 对数据进行处理。
- 将处理后的数据写回显存。
实现代码：
#include//系统头文件#include #include #include // 核函数，GPU端代码#ifndef _EXAMPLE_1_KERNEL_H_#define _EXAMPLE_1_KERNEL_H_//#include //在emu模式下包含这个头文件，以便输出一些中间结果来观察，在GPU实际运行时是不能使用的//! Simple test kernel for device functionality//! @param g_idata input data in global memory//! @param g_odata output data in global memory__global__ voidtestKernel( float* g_idata, float* g_odata) { // shared memory // extern表示大小由host端的Ns参数确定 extern __shared__ float sdata[]; // access thread id const unsigned int bid = blockIdx.x; //线程所在的block的索引号 const unsigned int tid_in_block = threadIdx.x; //线程在block中的位置 const unsigned int tid_in_grid = blockDim.x * blockIdx.x + threadIdx.x;//按行划分任务时，线程在整个grid中的位置 // 将数据从global memory读入shared memory sdata[tid_in_block] = g_idata[tid_in_grid]; //读入数据后进行一次同步，保证计算时所有数据均已到位 __syncthreads(); // 计算 sdata[tid_in_block] *= (float)bid;// sdata[tid_in_block] *= (float)tid_in_block;// sdata[tid_in_block] *= (float)tid_in_grid; //进行同步，确保要写入的数据已经被更新 __syncthreads(); // 将shared memory中的数据写到global memory g_odata[tid_in_grid] = sdata[tid_in_block];}#endif // #ifndef _EXAMPLE_1_KERNEL_H_// 函数声明void runTest( int argc, char** argv);// 主函数int main( int argc, char** argv) { runTest( argc, argv);}void runTest( int argc, char** argv) { unsigned int num_blocks = 4; //定义网格中的线程块数量 unsigned int num_threads = 4;//定义每个线程块中的线程数量 unsigned int mem_size = sizeof(float) * num_threads * num_blocks;//为数据分配的存储器大小，这里我们用每一个线程计算一个单精度浮点数。 // 在host端分配内存，h_表示host端，i表示input，o表示output //输入数据 float* h_idata = (float*) malloc( mem_size); //输出数据 float* h_odata = (float*) malloc( mem_size); // 在device端分配显存，d_表示device端 //显存中的输入数据 float* d_idata; cudaMalloc( (void**) &d_idata, mem_size); //显存中的输出数据 float* d_odata; cudaMalloc( (void**) &d_odata, mem_size); // 初始化内存中的值 for( unsigned int i = 0; i < num_threads * num_blocks; i++) { h_idata[i] = 1.0f; } // 将内存中的输入数据读入显存，这样就完成了主机对设备的数据写入 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); // 将结果从显存写入内存 cudaMemcpy( h_odata, d_odata, mem_size,cudaMemcpyDeviceToHost ); // 打印结果 for( unsigned int i = 0; i < num_blocks; i++) { for( unsigned int j = 0; j < num_threads; j++) { printf( "%5.0f", h_odata[ i * num_threads + j]); } printf("\n"); } // 释放存储器 free( h_idata); free( h_odata); cudaFree(d_idata); cudaFree(d_odata);}