CUDA10.0 官方手册 阅读笔记 章二 CUDA编程模型

以下为本人读CUDA 10.0的programming guide的读书笔记，部分为直接翻译重点内容。

2.1 kernel

       Kernel函数是CUDA C拓展了C，使用户可以通过定义C函数来编写CUDA函数。

       与只执行一次的C函数不同，如果kernel执行N次，则它会被分配到N个不同的CUDA thread中。

        kernel通过”__global__”来定义，”<<<…>>>”来配置并调用kernel。每一个执行 kernel的thread都被赋予了一个独一无二的 内建的 thread id（threadIdx）。

    __global__ void vectorAdd(float *a, float *b, float *c) {//定义kernel函数

        Int it = threadId.x; c[i] = a[i] + b[i];

    }

    Int main() {

        //…

        vectorAdd << <1, N >> > (A, B, C);//调用N个线程的kernel函数，N个线程，各自执行一边kernel函数

        //…

    }

2.2 thread 层次

    方便起见，threadIdx是一个有3个分量的向量，所以均可用一维、二维和三维方式来构建thread形成一个线程block，这种方式可以方便的解决 矢量、矩阵和体积。

    Thread block的构建方式是：一维时略；二维时 设block大小为(D­­­x,Dy) 则(x,y)的线程索引号是(x + y*Dy);三维时 设 block size = (D­­­x,Dy,D­­­z), 则(x,y,z)索引号为 (x + y*Dy + z*Dx*Dy)。

    下面的例子说明了矩阵加法的方式

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {

    int i = threadIdx.x;

    int j = threadIdx.y;

    if (i < N&&j < N)

        C[i][j] = B[i][j] + C[i][j];

}

int main() {

    //...

    int numBlocks = 1;

    dim3 threadsPerBlock(N, N);//共三个参数，默认为1，故每个块的线程数是N*N*1

    MatAdd <<<numBlocks, threadsPerBlock >> > (A, B, C);//1个threadBlock，每个threadBlock中有N*N*1个线程

    //...

}

注：每个线程块中线程数 在现行架构中 要小于1024个线程。且每一个线程块贮存在同一个处理器中，公用有限的内存。所以要在 线程块中线程数 和 线程块数量之间衡量。Grid中线程块的数量常由处理数据量 或 设备所有计算单元数 所决定的。

线程块结构图示

每个线程块中的 线程数 和 每个线程网格 在<<<>>>语法中由int或dim3的参数说明。

线程块在线程网格中可由内建变量blockIdx通过一维or二维或三维方式索引到。

线程块大小可由blockDim获得。

下面介绍grid构建block的例子

#define N 32

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {

    int i = blockIdx.x * blockDim.x + threadIdx.x;//行

    int j = blockIdx.y * blockDim.y + threadIdx.y;//列

    if (i < N && j < N)

        C[i][j] = A[i][j] + B[i][j];

}

int main() {

    //...

    dim3 threadsPerBlock(16, 16);//每一个线程块由16*16*1构成==每个线程块有256个线程

    dim3 numBlock(N / threadsPerBlock.x, N / threadsPerBlock.y);//线程网格中的线程块组成一个二维的样子

    MatAdd <<<numBlock, threadsPerBlock >>> (A, B, C);

    //...

}

注意：线程块必须能独立执行，也就是说线程块之间执行是没有顺序的。线程块中的线程能使用共享内存和执行”__syncthread()”来同步。共享内存的访问速度近似layer 1 cache。

 

2.3 存储层次

 

线程执行时，能从不同地方访问内存。他们所专有的内存由下图所示。每个线程有私有本地内存。每个线程块有块内所有线程可见的共享内存，并且生命周期与线程块生命周期相同。

所有线程还能访问两种只读内存——常数内存、纹理内存。

2.4 混合编程

CUDA编程模型如下，通常是主机上运行一段代码准备设备，然后在设备上运行并行代码。并且假设设备和主机有自己独立的DRMA空间。所以有内存在主机和设备之间搬运的过程。

统一的内存提供了管理主机和设备内存的桥梁，可以统一管理CPUs和GPUs的内存，将相同内存映射到同一个地址去。这种统一内存管理方式，使内存操作变得简单，消除了显式内存操作的复杂性。

2.5 计算能力

SM版本代表了计算能力。SM版本由版本号(X)和修订号(Y)构成X.Y。相同版本号有着相同的架构，7是Volta架构，6是Pascal架构。5是Maxwell架构，3是Kepler架构，2是Fermi架构，1是Tesla架构。Turing架构≥7.5，是Volta架构的升级版。

不要将CUDA版本号与硬件版本号搞混，CUDA版本是软件平台，CUDA平台的更新包含独立于硬件的软件更新。