《cuda c编程权威指南》04 - 使用块和线程索引映射矩阵索引

目录

1. 解决的问题

2. 分析

3. 方法

4. 代码示例

---

1. 解决的问题

利用块和线程索引，从全局内存中访问指定的数据。

2. 分析

通常情况下，矩阵是用行优先的方法在全局内存中线性存储的。如下。

8列6行矩阵（nx,ny）=（8，6）。

3. 方法

这里建立二维网格(2,3)+二维块(4,2)为例，使用其块和线程索引映射矩阵索引。

（1）第一步，可以用以下公式把线程和块索引映射到矩阵坐标上；

（2）第二步，可以用以下公式把矩阵坐标映射到全局内存中的索引/存储单元上；

比如要获取矩阵元素(col,row) = (2,4) ,其全局索引是34，映射到矩阵坐标上，

ix = 2 + 0*3=2； iy = 0 + 2*2=4. 然后再映射到全局内存idx = 4*8 + 2 = 34.

4. 代码示例

#include "cuda_runtime.h"
#include "device_launch_parameters.h"  // threadIdx#include <stdio.h>    // io
#include <time.h>     // time_t
#include <stdlib.h>  // rand
#include <memory.h>  //memset#define CHECK(call)                                   \
{                                                     \const cudaError_t error_code = call;              \if (error_code != cudaSuccess)                    \{                                                 \printf("CUDA Error:\n");                      \printf("    File:       %s\n", __FILE__);     \printf("    Line:       %d\n", __LINE__);     \printf("    Error code: %d\n", error_code);   \printf("    Error text: %s\n",                \cudaGetErrorString(error_code));          \exit(1);                                      \}                                                 \
}void initiaInt(int* p, const int N)
{for (int i = 0; i < N; i++){p[i] = i;}
}/// <summary>
/// 
/// </summary>
/// <param name="c">全局内存中线性存储的二维矩阵</param>
/// <param name="nx">列</param>
/// <param name="ny"></param>
void printMatrix(int* c, const int nx, const int ny)
{int* ic = c;printf("\n matrix: [%d, %d] \n", nx, ny);for (int i = 0; i < ny; i++){for (int j = 0; j < nx; j++){int cur_ele = ic[i * nx + j];printf("%d ", cur_ele);}printf("\n");}printf("\n");
}/// <summary>
/// 
/// </summary>
/// <param name="a">全局内存中是线性存储的</param>
/// <param name="nx">col</param>
/// <param name="ny"></param>
/// <returns></returns>
__global__ void printThreadIdx(int* a, const int nx, const int ny)
{// 矩阵行列int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;  // 全局索引unsigned int idx = iy * nx + ix;  // 前面有iy行，每行有nx个数。printf("thread_idx (%d, %d) block_idx (%d, %d) coordinate (%d, %d) global index %d val %d\n",threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, ix, iy, idx, a[idx]);
}int main(void)
{// get device infoint device = 0;cudaDeviceProp deviceProp;CHECK(cudaGetDeviceProperties(&deviceProp, device));printf("Using device: %d %s", device, deviceProp.name);  // 卡号0的显卡名称。CHECK(cudaSetDevice(device));  // 设置显卡号// set matrix dimensionint nx = 8, ny =6, nxy = nx * ny;int nBytes = nxy * sizeof(int);// malloc host memoryint* h_a;h_a = (int*)malloc(nBytes);// init datainitiaInt(h_a, nxy);printMatrix(h_a, nx, ny);// malloc device memoryint* d_Mat_a;cudaMalloc((void**)&d_Mat_a, nBytes);// transfer data from host to devicecudaMemcpy(d_Mat_a, h_a, nBytes, cudaMemcpyHostToDevice);// configdim3 block(4, 2);  // 二维线程块(x,y)=(4,2)dim3 grid((nx+block.x-1) / block.x, (ny+block.y-1)/block.y); // 二维网格(2,3)// 直接nx/block.x = 8/4=2. (8+4-1)/4=2.// invoke kernelprintThreadIdx << <grid, block >> > (d_Mat_a, nx, ny);cudaDeviceSynchronize();// free memorycudaFree(d_Mat_a);free(h_a);// reset devicecudaDeviceReset();return 0;
}

可以看到，全局索引值就是矩阵中存储的值。