CUDA的结构指南简述
2016-12-16 16:41
253 查看
http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
1 核函数(kernel)
指的是在gpu上运行的函数,通常和C函数一样。使用__global__定义。通过<<<...>>>来实现执行。(参见C语言扩展)
每个执行核函数的线程都有一个独一无二的线程id,可以通过内置的threadIdx变量来访问。
下面这个函数实现了向量的加法。
这里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
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执行。其实一样的。
例如
表示有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操作恰好相反的数据。看你怎么喜欢了。
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 , float B , float C ) { 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操作恰好相反的数据。看你怎么喜欢了。
相关文章推荐
- 第16周 oj Problem B: 结构体---职工信息结构体
- 外观模式
- the code of MH backpack en/decipher algorithm
- Linux安装ftp组件
- 横竖屏切换的时候,Activity生命周期的变化
- 人家的js笔记
- 小记4
- 怎样在onCreate()里面得到View的宽和高
- Solr6.2.1 学习笔记(二)从数据库导入数据
- 在使用DrawerLayout中遇到的问题
- emmt html生成
- Vbox中Ubuntu的安装和共享文件夹设置
- 用java实现文件下载,提示java.lang.IllegalStateException: getOutputStream() has already been called for this response
- 三角形三心和特点
- [故障引起的故事]URL中带加号的处理
- string hash 函数
- linux shell--tee
- 你知道 Ctrl + Alt + Delete 是怎么来的吗?
- Nginx开发一个简单的HTTP过滤模块
- 黑客上榜时代周刊2016年度风云人物,排名第一的是特朗普