最终输出实例为:
0 0 1 1
0 0 1 1
2 2 3 3
2 2 3 3
#ifndef _EXAMPLE_2_KERNEL_H_ #define _EXAMPLE_2_KERNEL_H_ #define SDATA( index) __global__ void testKernel( float* g_idata, float* g_odata, int width, int height) { // shared memory // 没有extern前缀,静态分配空间,因此[ ]中必须指定为数组分配的大小 // __shared__ float sdata[4]; // 计算线程索引 unsigned int bid_in_grid = __mul24(blockIdx.y , gridDim.x) + blockIdx.x; unsigned int tid_in_block = __mul24(threadIdx.y , blockDim.x) + threadIdx.x; unsigned int tid_in_grid_x = __mul24(blockDim.x , blockIdx.x) + threadIdx.x; unsigned int tid_in_grid_y = __mul24(blockDim.y , blockIdx.y) + threadIdx.y; unsigned int tid_in_grid = __mul24(tid_in_grid_y , width)+ tid_in_grid_x; // 从global memory中读入数据 // host模拟运行时使用,bank checker检查是否产生bank conflict // SDATA( tid_in_block) = g_idata[ tid_in_grid]; g_odata[tid_in_grid]=bid_in_grid; __syncthreads(); // 进行运算 // SDATA( tid_in_block) = (float)bid_in_grid * SDATA( tid_in_block); // __syncthreads(); //将数据写回global memory // g_odata[ tid_in_grid] = SDATA( tid_in_block); } #endif // #ifndef _EXAMPLE_2_KERNEL_H_ // 主函数 int main( int argc, char** argv) { runTest( argc, argv); } void runTest( int argc, char** argv){ unsigned int mem_size = sizeof( float) * 4 * 4;//数据大小,这里我们用每一个线程计算一个单精度浮点数 // 在host端分配内存 float* h_idata; cudaMallocHost( (void**) &h_idata, mem_size); // 初始化线程中的值 for( unsigned int i = 0; i < 4; i++){ for( unsigned int j = 0; j < 4; j++){ h_idata[i * 4 + j] = 1.0f; } } // 在device端分配显存 float* d_idata; cudaMalloc( (void**) &d_idata, mem_size); // 将内存中的值读入显存 cudaMemcpy( d_idata, h_idata, mem_size,cudaMemcpyHostToDevice ); // 在device端分配显存,用于存储结果 float* d_odata; cudaMalloc( (void**) &d_odata, mem_size); // 设置运行参数,即网格的形状和线程块的形状 dim3 grid( 2, 2, 1); dim3 threads( 2, 2, 1); // 运行核函数,调用GPU进行运算 testKernel<<< grid, threads >>>( d_idata, d_odata, 4, 4); // 在host端分配内存,用于存储结果 float* h_odata; cudaMallocHost( (void**) &h_odata, mem_size); // 将结果从显存写入内存 cudaMemcpy( h_odata, d_odata, mem_size,cudaMemcpyDeviceToHost ); // 打印结果 for( unsigned int i = 0; i < 4; i++){ for( unsigned int j = 0; j < 4; j++){ printf( "%5.0f", h_odata[ i * 4 + j]); } printf(" "); } // 释放存储器 cudaFreeHost(h_idata); cudaFreeHost(h_odata); cudaFree(d_idata); cudaFree(d_odata); }