矩阵赋值实例二

最终输出实例为：

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("\n");
        }
        // 释放存储器
        cudaFreeHost(h_idata);
        cudaFreeHost(h_odata);
        cudaFree(d_idata);
        cudaFree(d_odata);
}