为了让程序利用GPU进行并行运算. 可以利用Nvidia 提供的CUDA.学习新的知识最权威的当然是官方提供的文档
简单的CUDA helloworld.cu 程序
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 |
// includes, system #include <stdio.h> // includes CUDA Runtime #include <cuda_runtime.h> __global__ void helloworld() { printf("%s \n", "helloworld"); } int main() { helloworld <<<1, 1024>>> (); return 0; } |
刚看到这个程序时,我们通常会对<<<1,1024>>> 感到奇怪.我们会想这样的写法我们熟知的编译器可没有这种语义结构,那这段程序用什么编译的呢?
– 从上往下阅读,有个cuda_runtime.h头文件,我们搜索官方文档,会发现这个头文件是host端(CPU)运行时,他提供了
- device (GPU)管理
- context管理
- 内存管理
- 代码模块管理
- 纹理引用管理
- 与OpenGL和Direct3D 互操作
其他的库如cuBLAS(贝叶斯) cuFFT(快速傅里叶变换)也依赖于cuda runtime
cuda_runtime.h(C++-style?) 还有个父亲叫cuda_runtime_api.h(C-style)
继续往下看__global__ 是什么?
在一个函数前面用 __global__限定符:表示将一个函数声明为内核
- 表明这个函数在device(GPU)上执行
- 可以从host(CPU)调用
- 可以从device(GPU)调用cc3.2以上的device(GPU)
这个函数必须是 void 返回类型
执行该函数必须指定调用该函数的配置,即下面的<<<1,1024>>>
对__global__ 函数的调用是异步的,这意味着device在执行完之前返回
<<<1,1024>>> 是什么?
官方称这个叫 执行配置,那配置里面的参数都是什么意思
通过上面的链接我们发现 格式是这个样子的<<< Dg, Db, Ns, S >>>
- Dg类型为dim3 指定了 grid的维度和大小 Dg.x*Dg.y*Dg.z 等于启动的 block
- Db类型为dim3 指定了 block的维度和大小 Db.x*Db.y*Db.z 等于每个block启动的thread数量
- Ns类型为size_t?可选参数,?程序中我们除了静态分配内存外,利用这个函数可以为每个block 动态分配共享内存 number of bytes in shared memory,默认为0
- S类型为cudaStream_t 可选参数,?指定关联的流默认为0
这里我们涉及了grid,block,thread, memory 和stream.这些知识以后会讲解.
上面的<<<1,1024>>>如何配置的参数超过了自己机器上的限制,或者共享内存超过了限制,函数调用失败.
不同的GPU 计算能力可能不一样,通过上面这个链接的表查看参数的限制.也可以通过其他工具查看.
最后我们试着在VS上跑一下这个程序,并人为制造一些bug看一下.
执行后输出了1024条helloworld
修改<<<1,1024>>>为<<<1,1025>>>,程序正常执行,没有输出任何东西
修改<<<1,1024>>>为<<<2147483648,1>>>程序正常执行,没有输出任何东西
修改<<<1,1024>>>为<<<2,2>>> 输出4条
总结:由于Dg和Db我们这样写相当于一个一维的数,即Dg.y?Dg.z Db.y?Db.z 我们并没有指定.<<<1,1025>>> 由于我的GPU Db.x 阈值为1024,写成1025超过了,当然不调用kernel函数 其他同理, 并行执行的总数当然是Dg*Db.
