CUDA的结构指南简述

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

1 核函数（kernel）

指的是在gpu上运行的函数，通常和C函数一样。使用__global__定义。通过<<<...>>>来实现执行。（参见C语言扩展）

每个执行核函数的线程都有一个独一无二的线程id，可以通过内置的threadIdx变量来访问。

下面这个函数实现了向量的加法。

// Kernel definition __global__ void VecAdd(float* A, float* B, float* C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main(){ ... // Kernel invocation with N threads VecAdd<<<1, N>>>(A, B, C); ...}

这里N个线程执行了vecAdd函数
上面是个简单的例子

为了方便起见。threadIdx是一个三元素的向量 x y z。因此一个线程(thread)可以使用一维 二维或者 三维的线程索引来表示一个线程。并形成线程块(block).因此有很自然的方式来计算一维向量。二维矩阵，和三维体积。线程的索引和该线程的id是很直接的关系：对于一维的线程块，都是相同的。对于二维的线程块（Dx,Dy）。线程的索引如果是(x,y)，那么线程的id是(x+yDx)。对于三维的线程块(Dx,Dy,Dz) ，如果线程的索引是(x,y,z)那么线程的id是（x+yDx+zDy）

如下一个A+B=C的矩阵加法，矩阵大小为N*N

// Kernel definition __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) { int i = threadIdx.x; int j = threadIdx.y;C[i][j] = A[i][j] + B[i][j]; }int main() { ... // Kernel invocation with one block of N * N * 1 threads int numBlocks = 1; dim3 threadsPerBlock(N, N); MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); ... }

MatAdd使用了一个线程块。每个线程块包含N*N个线程。

因此从线程ID中取出x y得到对应维度的坐标。进行计算。

每个线程块可以包含多达1024个线程。

由多个线程块可以组成线程格

就是由线程Thread组成线程块block。由线程块组成线程格grid。

<<< >>>中的参数可以是int类型或者dim3类型。

每个grid中的block可以通过blokidx来索引到。

每个线程块block的维度可以通过blockDim得到。blockDim就是定义了block的长度宽度和高度，是固定死的。

比如dim nblock(16,17)。那么blockDim.x=16。blockDim.y=16是固定死的。

而blockidx.x则从0-15这么变化。blockidx.y则从0-16这么变化

还有一个是gridDim定义了grid的宽度和高度。也是固定死的。

目前我看到的都只有一个grid。参数传递也只有block和thread的相关参数。如果以后还可以传递grid参数。那么一定会有grididx来定义特定的grid

这里说的有些模糊也很容易搞不清

下面结合两篇外文仔细分析一下 thread block grid以及之间的关系。

http://study.marearts.com/2015/03/meaning-of-threadidx-blockidx-blockdim.html

http://study.marearts.com/2015/03/meaning-of-threadidx-blockidx-blockdim_12.html

后续部分图片转自上面的博客

1首先了解执行GPU函数的调用方式两种。

一种是直接在代码中编写通过<<< Dg, Db, Ns, S >>>的形式调用

Dg是grid(xyz) Db是block(xyz)。（第一个Dg是block的数量grid的维度就是block的数量,Db是thread的数量，也就是block的维度）Ns S暂且不看默认为0

首先由多个thread构成 block。再由多个block构成grid

第二种就是先将cu执行编译。然后通过load的方式获取导出函数，然后通过cuLaunchKernel执行。其实一样的。

例如

#define N 15__global__ void increase(int *c){ int tid = threadIdx.x + blockIdx.x * blockDim.x;  if(tid < N)  c[tid] = tid;}

increase<<< 4, 3>>>(dev_c);

结果是

表示有4个block .每个block3个线程。一共12个线程

我们需要得到每个线程的唯一id。这个id怎么计算呢？就按照前面的方式计算。

blockId和threadid是存在于不同空间的。他们各自在自己的大一级范围内才独一无二。

比如threadid只在属于他的blockid里面独一无二。blockid只在属于他的grid里面独一无二。

输出线程id的时候可以看出 threadIdx.x;

线程id是 0 1 2 0 12 012 循环的。所以说线程id在当前的block的独一无二。

如果是

<<<1,12>>>

则是

如果是<<<2,6>>>则是

threadid告诉了我们当前线程的索引。

blockid告诉了我们当前block的索引。

blockDim告诉了我们一个block中thread的数量（维度）

gridDim告诉我们一个grid中block的数量（维度）

如下图。全局的唯一索引计算方式如下

如果是2D的话

dim3 blocks(2,3);
dim3 thread(3,2);
Kernel<<< blocks, threads >>>

2*3*2*3一共36个线程

如何访问第十五个线程呢？

threadidx.x和thread.idx.y只是在当前block中的索引。因此还要加上block的位置数据。

计算方式如上图。

所以给显卡相应的数据以后。我们都是通过这些内置的索引来操作对应的数据进行并行计算的。

至于哪个线程执行那一部分操作。完全由你自己决定。

比如<1,12>的12个线程。你可以直接操作 a[i]=xxxxx。也就是第i个线程操作第i个数据。也可以 a[11-i]=xxxx操作恰好相反的数据。看你怎么喜欢了。