#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); }
时间: 2024-11-10 00:59:27