GPU编程自学4 —— CUDA核函数运行参数

深度学习的兴起，使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题。这里主要记录自己的GPU自学历程。

《GPU编程自学1 —— 引言》
《GPU编程自学2 —— CUDA环境配置》
《GPU编程自学3 —— CUDA程序初探》
《GPU编程自学4 —— CUDA核函数运行参数》
《GPU编程自学5 —— 线程协作》
《GPU编程自学6 —— 函数与变量类型限定符》
《GPU编程自学7 —— 常量内存与事件》
《GPU编程自学8 —— 纹理内存》
《GPU编程自学9 —— 原子操作》
《GPU编程自学10 —— 流并行》

四、 CUDA核函数运行参数

在前面的章节中，我们不止一次看到了在调用定义的核函数时采用了类似下面的形式：

kernel<<<1,1>>>(param1,param2,...)

“<<< >>>”中参数的作用是告诉我们该如何启动核函数(比如如何设置线程)。下面我们先直接介绍参数概念，然后详细说明其意义。

4.1 核函数运行参数

当我们使用 gloabl 声明核函数后

__global__ void kernel(param list){  }

在主机端(Host)调用时采用如下的形式：

kernel<<<Dg,Db, Ns, S>>>(param list);

Dg： int型或者dim3类型(x,y,z)。用于定义一个grid中的block是如何组织的。 int型则直接表示为1维组织结构。
Db： int型或者dim3类型(x,y,z)。用于定义一个block中的thread是如何组织的。 int型则直接表示为1维组织结构。
Ns： size_t类型，可缺省，默认为0。用于设置每个block除了静态分配的共享内存外，最多能动态分配的共享内存大小，单位为byte。 0表示不需要动态分配。
S： cudaStream_t类型，可缺省，默认为0。表示该核函数位于哪个流。

4.2 线程结构

关于CUDA的线程结构，有着三个重要的概念： Grid, Block, Thread

GPU工作时的最小单位是 thread。
多个 thread 可以组成一个 block，但每一个 block 所能包含的 thread 数目是有限的。因为一个block的所有线程最好应当位于同一个处理器核心上，同时共享同一块内存。于是一个 block中的所有thread可以快速进行同步的动作而不用担心数据通信壁垒。
执行相同程序的多个 block，可以组成 grid。不同 block 中的 thread 无法存取同一块共享的内存，无法直接互通或进行同步。因此，不同 block 中的 thread 能合作的程度是比较低的。不过，利用这个模式，可以让程序不用担心显示芯片实际上能同时执行的 thread 数目限制。例如，一个具有很少量执行单元的显示芯片，可能会把各个 block 中的 thread 顺序执行，而非同时执行。不同的 grid 则可以执行不同的程序(即 kernel)。

下图是一个结构关系图：

此外，Block, Thread的组织结构可以是可以是一维，二维或者三维。以上图为例，Block, Thread的结构分别为二维和三维。

CUDA中每一个线程都有一个唯一标识ThreadIdx，这个ID随着组织结构形式的变化而变化。 (注意：ID的计算，同计算行优先排列的矩阵元素ID思路一样。)

回顾之前我们的矢量加法：

// Block是一维的，Thread也是一维的
__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = blockIdx.x *blockDim.x + threadIdx.x;  
    c[i] = a[i] + b[i];
}

// Block是一维的，Thread是二维的

__global__ void addKernel(int *c, int *a, int *b)
{
    int i = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

// Block是二维的，Thread是三维的
__global__ void addKernel(int *c, int *a, int *b)
{
    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int i = blockId * (blockDim.x * blockDim.y * blockDim.z)  
        + (threadIdx.z * (blockDim.x * blockDim.y))  
        + (threadIdx.y * blockDim.x) + threadIdx.x; 
    c[i] = a[i] + b[i];
}

下表是不同计算能力的GPU的技术指标(更多可参见 CUDA Toolkit Documentation)

当然也可以通过下面的代码来直接查询自己GPU的具体指标：

#include "cuda_runtime.h"
#include <iostream>

int main()
{
    cudaError_t cudaStatus;

    // 初获取设备数量
    int num = 0;
    cudaStatus = cudaGetDeviceCount(&num);
    std::cout << "Number of GPU: " << num << std::endl;

    // 获取GPU设备属性
    cudaDeviceProp prop;
    if (num > 0)
    {
        cudaGetDeviceProperties(&prop, 0);
        // 打印设备名称
        std::cout << "Device: " <<prop.name << std::endl;
    }

    system("pause");
    return 0;
}

其中 cudaDeviceProp是一个定义在driver_types.h中的结构体，大家可以自行查看其定义。

4.3 内存结构

如下图所示,每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个 thread 则有共享的一份 share memory。此外，所有的 thread(包括不同 block 的 thread)都共享一份 global memory、constant memory、和 texture memory。不同的 grid 则有各自的 global memory、constant memory 和 texture memory。

这种特殊的内存结构直接影响着我们的线程分配策略，因为需要通盘考虑资源限制及利用率。这些后续再进行讨论。

4.4 异构编程

如下图所示，是常见的GPU程序的处理流程，其实是一种异构程序，即CPU和GPU的协同。

主机上执行串行代码，设备上则执行并行代码。

参考资料：

《CUDA by Example: An Introduction to General-Purpose GPU Programming》中文名《GPU高性能编程CUDA实战》
详解CUDA核函数及运行时参数 http://blog.csdn/a925907195/article/details/39500915
CUDA之——深入理解threadidx http://blog.csdn/canhui_wang/article/details/51730264
CUDA Toolkit Documentation http://docs.nvidia/cuda/cuda-c-programming-guide/index.html#axzz4oh8uLanv

更多推荐

GPU编程自学4 —— CUDA核函数运行参数