题目:给一个二维数组赋值。
分析:主机端代码完成的主要功能:
- 启动CUDA,使用多卡时应加上设备号,或使用cudaSetDevice()设置GPU设备。
- 为输入数据分配内存空间
- 初始化输入数据
- 为GPU分配显存,用于存放输入数据
- 将内存中的输入数据拷贝到显存。
- 为GPU分配显存,用于存放输出数据。
- 调用device端的Kernel进行计算,将结果写到显存中的对应区域。
- 为CPU分配内存,用于存放GPU传回的输出数据
- 将显存中的结果回读到内存。
- 使用CPU对传回的数据进行其他的处理。
- 释放内存和显存空间。
- 退出CUDA
设备端代码要完成的任务:
- 从显存读取数据到GPU片内。
- 对数据进行处理。
- 将处理后的数据写回显存。
实现代码:
#include <stdlib.h> //系统头文件
#include <stdio.h>
#include <string.h>
#include <math.h>
// 核函数,GPU端代码
#ifndef _EXAMPLE_1_KERNEL_H_
#define _EXAMPLE_1_KERNEL_H_
//#include <stdio.h> //在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__ void
testKernel( 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);
}