CUDA检查块和线程索引

1. 组织并行线程

    对于一个给定的线程， 首先可以通过把线程和块索引映射到矩阵坐标上来获取线程块和线程索引的全局内存偏移量， 然后将这些矩阵坐标映射到全局内存的存储单元中。
    第一步， 可以用以下公式把线程和块索引映射到矩阵坐标上：

    ix = threadIdx.x + blockDim.x*blockIdx.x;

    iy = threadIdx.y + blockDim.y*blockIdx.y;

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

    idx = iy * nx + ix;

    如下图说明了块和线程索引、 矩阵坐标以及线性全局内存索引之间的对应关系。

2. 举个栗子

    常情况下，一个矩阵用行优先的方法在全局内存中进行线性存储。如下图所示，这是一个8*6的矩阵。

    值得注意的是：CUDA中组织线程的方式与数组不同。

3. 代码

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <malloc.h>
#include <stdlib.h>

/* 宏定义的细节
	（1）对于有参数的宏定义，宏定义时，在宏名与带参数的括号之间不应加空格，否则将空格以后的字符都作为替代字符串的一部分。
	（2）带参数的宏定义只是进行简单的字符替换，宏展开则是在编译前进行的，在展开时并不分配内存单元，不进行值得传递处理，因此替换不会占用运行时间，只占用编译时间。
	（3）宏不存在类型问题，宏名无类型，宏的参数也无类型，只是一个符号代表，展开时代入指定的字符串即可，宏定义时，字符串可以是任何类型的数据。
	作者：Wonshington
	链接：https://www.jianshu.com/p/dd0c8938ea90
	来源：简书
	著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。*/

//如果不加这些“\”程序就会报错。其实这是宏定义的一个细节，理解为“继续符”的意思，表示本行与下一行连接起来。
#define CHECK(call)                                                      \
{                                                                        \
    const cudaError_t error = call;                                      \
    if(error!=cudaSuccess)                                               \
    {                                                                    \
        printf("Error: %s:%d, "__FILE__,__LINE__);                       \
        printf("code:%d, reason: %s\n",error,cudaGetErrorString(error)); \
        exit(1);                                                         \
    }                                                                    \
}

void initialInt(int *ip, int size)
{
	for(int i=0; i<size; i++)
	{
		ip[i]=i;
	}
}

void printMatrix(int *C, const int nx, const int ny)
{
	int *ic=C;
	printf("\nMatrix:(%d.%d\n)",nx,ny);
	for(int iy = 0; iy<ny; iy++)
	{
		for(int ix=0; ix<nx; ix++)
		{
			printf("%3d",ic[ix]);
		}
		ic+=nx;
		printf("\n");
	}
	printf("\n");
}

__global__ void printThreadIndex(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;

	printf("thread_id(%d,%d) block_id(%d,%d) coordinate(%d,%d) global index %2d ival %2d\n",
		    threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, ix, iy, idx, A[idx]);
}

int main(int agrc, char **argv)
{
	printf("%s, Starting...\n",argv[0]);

	// get device information
	int dev =0;
	cudaDeviceProp deviceProp;
	CHECK(cudaGetDeviceProperties(&deviceProp, dev));
	printf("Using Device %d:%s\n",dev,deviceProp.name);
	CHECK(cudaSetDevice(dev));

	// set matrix dimension
	int nx = 8;
	int ny = 6;
    int nxy = nx*ny;
	int nBytes = nxy*sizeof(int);

	// malloc host memory
	int *h_A;
	h_A = (int *)malloc(nBytes);
	
	// initialize host matrix with integer
	initialInt(h_A, nxy);

	// malloc device memory
	int *d_MatA;
	cudaMalloc((void**)&d_MatA, nBytes);

	// transfer data from host to device
	cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice);

	// set up execution configuration
	dim3 block(4,2);
	dim3 grid((nx+block.x-1)/block.x,(ny+block.y-1)/block.y);

	// invoke the kernel
	printThreadIndex<<<grid,block>>>(d_MatA,nx,ny);
	cudaDeviceSynchronize();

	// free host and devide memory
	cudaFree(d_MatA);
	free(h_A);

	// reset device
	cudaDeviceReset();

	system("pause");

    return 0;
}

4. 运行结果