CUDA

CUDA-几点要记住-更新ID

Bare in mind:

  • 不管你的数据是一维的二维的还是更高维度的，在GPU端，高维被扁平化，都将被看成一维的，所以么有必要在Device上开辟，比如，一个二维数组。
  • CUDA code 需要你并行地思考：Think parallel.
  • 当你在写CUDA code， 实际上你是在为一个thread 写串行code，而每一个thread都执行这个段相同的串行code。看下图体会。
  • 可以这样理解，对于简单问题，把CPU code的for 循环去掉，其实就得到了GPU code。每个thread 有自己唯一的ID，其他都一样。

配置kernel

  • 当GPU可用的thread非常多，而当前所需解决的任务规模并不大时，可以一次invoke 足量的threads，这样便不用更新threads ID。如下：

    1
    2
    3
    4
    5
    6
    #define N 1<<7
    ...
    	int block(1024);
    	int grid ((N + block.x - 1) / block.x );

    	kernal_func <<<grid, block >>>(d_c, d_a, d_b);

    或者二维的configuration：

    1
    2
    3
    4
    5
    6
    #define N 1<<7
    ...
    	dim3 block (1024, 1024);
    	dim3 grid (( N + block.x -1) / block.x, (N + block.y -1) / block.y );

    	kernel_func <<<grid, block >>>(dev_m, dev_mt);

    当然所有的configuration都应该在你的GPU硬件极限内。

  • 当数据量很大时，所有CUDA cores 就需要工作不止一波，第一波后，就需要更新threads ID 继续工作下一波：

    1
    2
    3
    4
    5
    6
    7
    8
    __global__ void add(int *a, int *b, int *c) {
    	int tid = threadIdx.x + blockIdx.x * blockDim.x;
    	while (tid < N) {
    		c[tid] = a[tid] + b[tid];
    		// OPERATIONS
    		tid += blockDim.x * gridDim.x;     // update id when #threads are less than #elements 
    	}
    }

    此处的while循环 表示，只要threads ID 还小于元素个数N，就更新ID。