矩阵赋值实例(matrixAssign) 题⽬:给⼀个⼆维数组赋值。
分析:主机端代码完成的主要功能:
1. 启动CUDA,使⽤多卡时应加上设备号,或使⽤cudaSetDevice()设置GPU设备。
2. 为输⼊数据分配内存空间
3. 初始化输⼊数据
4. 为GPU分配显存,⽤于存放输⼊数据
5. 将内存中的输⼊数据拷贝到显存。
6. 为GPU分配显存,⽤于存放输出数据。
7. 调⽤device端的Kernel进⾏计算,将结果写到显存中的对应区域。
8. 为CPU分配内存,⽤于存放GPU传回的输出数据
9. 将显存中的结果回读到内存。
10. 使⽤CPU对传回的数据进⾏其他的处理。
11. 释放内存和显存空间。
12. 退出CUDA
设备端代码要完成的任务:
1. 从显存读取数据到GPU⽚内。
2. 对数据进⾏处理。
3. 将处理后的数据写回显存。
实现代码:
#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 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++)
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);
}
版权声明:本站内容均来自互联网,仅供演示用,请勿用于商业和其他非法用途。如果侵犯了您的权益请与我们联系QQ:729038198,我们将在24小时内删除。
发表评论