CUDA学习3

昨天试着在GPU上输出了"HELLO WORD！"以及cuda的存储问题，那么今天先开始看第一个：线程管理：

核函数在主机端启动的时候，他的执行会移动到设备上去，这个时候，设备上会产生大量的线程，而且每一个线程都在执行由核函数指定的语句，所以，怎么样去组织线程是CUDA的一个关键部分，CUDA明确了线程层次抽象的概念，以便于我们组织线程，这里我们展示书上的一个两层的线程层次结构：

图一

由一个内核启动所产生的线程统称为一个网格，同意网格中的所有线程共享相同的全局存储空间，一个网格由多个线程块构成，一个线程块内又包含了一组线程，同一线程块内的线程写作可以通过以下方式来实现：

1. 同步

2. 共享内存

但是，不同块内的线程不能够协作！！！

线程依靠两个坐标变量来找：

1. blockidx：线程块在线程网格内的索引

2. threadidx：线程块内的线程索引

这两个变量是核函数中需要预先初始化的，当执行一个核函数的时候，CUDA运行时就会给每个线程分配以上两个变量，基于这种思想，我们就可以很轻松的把数据给分配到不同的线程中去，

CUDA是可以组织三维的网格和块的：

blockIdx.x
blockIdx.y
blockIdx.z
threadIdx.x
threadIdx.y
threadIdx.z

他的结构是一个包含二维块的二维网格，网格和快的维度由下面两个内置变量指定：

1. blockDim：线程快的维度：每个线程块中的线程数

2. gridDim：线程格的维度，每个线程格中的线程数来表示

和上面一样也是三维的；

通常，一个线程格会被组装成线程块的二维数组形式，一个线程块会被组织成线程的三维数组形式。

CUDA程序中有两种不同的网格和块变量：

1. 手动定义的dim3数据类型（仅主机端可见）

2. 预定义的uint3数据类型（仅设备端可见）

主机端作为内核调用的一部分，可以用dim3来定义一个网格和块的维度，执行核函数的时候，CUDA运行时会自动生成相应的内置预初始化的网格，块和线程变量，在核函数内均可以被访问且都是uint3类型。网格大小是块的倍数，跟着书上的学着写一下：checkDimension.cu文件：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<stdlib.h>
#include<string>
#include<time.h>
__global__ void checkDimension() {
    printf("threadindex:(%d,%d,%d) blockindex:(%d,%d,%d)    blockdimindex:(%d,%d,%d)"
        "griddim:(%d,%d,%d)\n", threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y, blockIdx.z, blockDim.x, blockDim.y, blockDim.z, gridDim.x, gridDim.y, gridDim.z
    );

}
int main() {
    int nElem = 6;
    dim3 block(3);
    dim3 grid((nElem + block.x - 1) / block.x);
    printf("grid.x%d,grid.y%d,grid.z%d\n", grid.x, grid.y, grid.z);
    printf("block.x%d,block.y%d,block.z%d\n", block.x, block.y, block.z);
    checkDimension << <grid, block >> > ();
    cudaDeviceReset();
    return 0;
}

图二：运行结果

对于一个给定的数据大小，确认网格和块尺寸的一般步骤是：

1. 确定块的大小

2. 在已知数据大小和块大小的基础上计算网格维度

要确定块的尺寸，通常要考虑的是

1. 内核的性能特性

2. GPU资源的限制