最终输出实例为:
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);
}