CUDA C Best Practices Guide 阅读笔记(二) Heterogeneous Computing

CUDA 编程涉及到在不同的平台上同时运行代码:包含CPU的host 和包含GPU的device.

所以了解host和device的对性能优化是非常重要的。

2.1. Differences between Host and Device

Threading resources

host 上能同时运行的线程数目比较少(比如24个)

device上能同时运行的线程数目比较多(数量级通常为1E3,1E4等)

Threads

操作系统必须交换CPU执行通道上和下的线程以提供多线程功能。因此,上下文切换(当交换两个线程时)既慢又昂贵。

相比之下,GPU上的线程非常轻量级。在典型的系统中,成千上万的线程排队等待工作(每个线程有32个线程)。如果GPU必须等待 one warp of threads,它只需开始在另一个线程上执行工作。

简而言之,CPU内核被设计为每次最小化一个或两个线程的等待时间,而GPU被设计为处理大量并发的轻量级线程以最大化吞吐量。

RAM

host和device 各自具有各自不同的附接物理存储器。host和device内存由PCI Express ( PCIe )总线分隔,因此host内存中的项目必须偶尔通过总线传送到device内存,反之亦然

 

2.2. What Runs on a CUDA-Enabled Device?

 

下面谈谈应该把应用的哪些部分放在device 上运行

  • 大数据集上的算术运算
  • 为了获得最佳性能,设备上运行的相邻线程的内存访问应该具有一定的一致性。某些内存访问模式使硬件能够将多个数据项的读或写组合并到一个操作中。当在CUDA上的计算中使用时,无法布局以实现合并的数据,或者没有足够的局部性来有效地使用L1或纹理缓存的数据,将倾向于看到较小的加速比。
  • host和device之间的数据交换尽可能少
    • 换到device上执行的数据一定会被做足够多的运算,不然数据从Host传送到device的代价 可能与该运算在device上并行计算的优势向抵消,甚至得不偿失。
    • 数据应尽可能长时间保存在设备上。因为传输应该最小化,所以在同一数据上运行多个内核的程序应该倾向于在内核调用之间将数据保留在设备上,而不是将中间结果传输到主机,然后再将它们发送回设备进行后续计算。就是说,如果有一段连续的操作要处理某些数据,就算其中的部分操作在host上运行要比在device上快(比如不是算数运算而是逻辑处理),那么考虑到数据传输的巨大代价,将所有数据都放在device上处理可能会更好。这种处理原则即使相对较慢的内核也可能是有利的,如果它避免了一个或多个PCIe传输。

 

 

 

CMake Error at gpuxgboost_generated_updater_gpu_hist_experimental.cu.obj.Release.cmake:282 的解决办法

请教了同事…果然是身经百战见得多啊

直接告诉我cuda 8.0.27 对应版本的thrust有bug.

解决办法是从 thrust 的github搞一份最新的下来,并覆盖掉/usr/local/cuda/include/thrust/

 

CUDA C Best Practices Guide 阅读笔记(1) 并行计算方法论(APOD)

APOD指的是Assess, Parallelize, Optimize, Deploy

Assess, Parallelize, Optimize, Deploy.

如图所示,APOD过程是一个循环的过程,每次只进行一部分,从A到P到O到D,然后再进行下一轮的APOD

Assess

对于一个项目,第一步要做的是接触(Assess)项目,得到项目代码中每部分的执行时间。有了这部分内容,开发者就可以找到并行优化的瓶颈所在,并开始尝试用GPU加速。

根据Amdahl’s and Gustafson’s laws,开发者可以确定并行优化的性能上界。

Parallelize

找到瓶颈所在并确定了优化的目标和期望,开发者就可以优化代码了。调用一些如cuBLAScuFFT, or Thrust 的 GPU-optimized library  可能会很有效。

另一方面,有些应用需要开发者重构代码,以此让可以被并行优化得部分被暴露出来。

Optimize

确定了需要被并行优化的部分之后,就要考虑具体得实现方式了。

具体得实现方式通常不过只有一种,所以充分理解应用的需求是很有必要的。

要记得,APOD是一个反复迭代的过程(找到可以优化的点,实现并测试优化,验证优化效果,然后再重复)

因为对于开发者来说,没有必要最初就找到解决所有性能瓶颈的策略。

优化可以在不同的level上进行,配合性能分析工具是很有帮助的。

Deploy

大的原则是当优化完一处之后,立刻将这一部分部署到生产环境,而不是再去寻找其他可以优化的地方。

这样做有很多重要的原因,比如这会使得用户尽早从这个优化中收益(尽管提速只是部分的,但是仍然有价值)

此外,每次有改动就重新部署,也使得变化是平稳而不是激进的,这有助于减少风险。

Recommendations and Best Practice

优化根据对性能影响程度的不同有不同的优先级。

在先要处理低优先级的优化之前,一定要确保其他所有的高优先级优化都做完了。

这种方法将倾向于为所投入的时间提供最佳结果,并且将避免过早优化的陷阱。

 

需要说明的一点是,教程中的代码为了方便简洁没有关于任何 check error的部分。

在实际上这是不可取的(这不同于编写C++代码!

CUDA 7.5: 用指令级性能分析精确找到性能问题

原文:

CUDA 7.5: Pinpoint Performance Problems with Instruction-Level Profiling

 

主要是介绍了CUDA 7.5 以上的版本的 NVIDIA Visual Profiler 加入的新特性

可以细粒度到指令级,分析出性能的瓶颈(在这之前,只能分析到kernel级别)

顺便了解了一下nvidia-visual-profiler

原理大概是用PC(程序计数器)采样目前活跃的warp,然后得到这些warp的PC和状态,以此来得到更细粒度的性能分析。

Instruction-level profiling in CUDA 7.5 uses program counter (PC) sampling with a per-streaming-multiprocessor (SM) fixed-frequency sampler. At each sample period, the collector picks an active warp from the SM and captures the warp’s PC and warp state. Warp selection is performed round-robin across all active warps on the SM.

文中还给出了分析完性能瓶颈之后做出的相应优化。

下面是对一段存在内存依赖延迟(?Memory Dependency Stalls )问题的代码给出优化的过程。

我们注意到代码第一行定义了一个size为7的float类型的数组,性能分析表明这一行涉及了 LDL(“load local”)  指令,是性能的瓶颈之一。

为什么?

因为NVIDIA 的GPU 没有下标 寄存器文件,所以如果一个数组被以动态的下标访问,编译器不得不为这个数组分配local memory.

在Maxwell 架构中,local memory 没有被加入在L1 缓存中,因此访问内存会有很大的延迟。

所以为了优化这部分,就把数组拆分成相同个数的单独的变量。

然后由于没有定义成数组,因此也需要把循环展开。

优化后的代码如下:

 

 

 

 

 

 

 

cuda 学习笔记 术语篇

最近在学习cuda,遇到的新词汇有点多,所以开一篇记录一下.

所记录的不一定和cuda有关,只是在学习cuda中遇到的陌生概念.

  • SIMD(Single Instruction Multiple Data)  单指令流多数据流  是一种采用一个控制器来控制多个处理器,同时对一组数据(又称“数据向量”)中的每一个分别执行相同的操作从而实现空间上的并行性的技术。图形处理器(GPU)拥有强大的并发处理能力和可编程流水线,面对单指令流多数据流时,运算能力远超传统CPU。OpenCLCUDA分别是目前最广泛使用的开源和专利通用图形处理器(GPGPU)运算语言。
  • SAXPY (Scalar Alpha X Plus Y)operation, 其实就是{\displaystyle \mathbf {y}=\alpha \mathbf {x} +\mathbf {y} ,\,}  其中标量矢量

cuda error checking 学习笔记

由于发现cuda c++ 的 debug方式和c++ 差别很大,因此打算再开一篇,专门记录一些和error checking 以及debug有关的内容.

Error checks in CUDA code can help catch CUDA errors at their source. There are 2 sources of errors in CUDA source code:

  1. Errors from CUDA API calls. For example, a call to  cudaMalloc() might fail.
  2. Errors from CUDA kernel calls. For example, there might be invalid memory access inside a kernel.

All CUDA API calls return a  cudaError value, so these calls are easy to check:

CUDA kernel invocations do not return any value. Error from a CUDA kernel call can be checked after its execution by calling  cudaGetLastError():

 

以及下面是一个check error的宏…

 

cuda 学习笔记

uodate:有毒吧。kernel中出问题原来是不会报错的。。。。

请教了组里的hust学长orz..、

学到了cuda-memcheck命令和cudaGetLastError来查看问题。。可以参考What is the canonical way to check for errors using the CUDA runtime API?

 

先放一波资料。

 

 

cuda 提出的目的是能够让程序员透明地使用GPU来高效地进行并行运算。

kernel和c语言中的函数相似,函数名字前通常用global来标识。

 

下面考虑一个2个大小的1M的数组相加的例子。

总的思路是通过并行,来观察到计算速度的加快。

如果不考虑并行,2个数组相加的代码,如下:

 

如果用cuda的方式来搞,代码如下:

除了代码注释,还有几个地方要说明:

  • 函数名字前面的global 是 cuda kernel  标识符
  • cuda kernel的调用方式是 <<<,>>>  更具体地说,是add<<<numBlocks,blockSize>>>(N,x,y)
  • .cu是 cuda C++ 文件的后缀,类似.cpp
  • nvcc是cuda C++ 的编译器,其将source code分成host codedevice code两部分。前者通过c++编译器编译,后者通过nvidia编译器编译。

 

关于devide code 和 host code,参考下图。

 

现在我们单线程地跑了一个cuda kernel, 接下来是如何使它并行.关键在于<<<1,1>>>这部分。

这行代码告诉了cuda runtime 有多少个并行的线程要被执行。
这里有2个参数,不过我们可以先改变第二个,也就是一个线程block中线程的个数。
cuda GPu的kernel 使用的blocks中线程的个数应该是32的倍数(后面会解释32代表什么),所以256看起来很合理。

 

 

不过如果只是修改了<<<1,1>>> 到  <<<1,256>>>
那实际上是对于每个线程都算了整个array的相加,而没有将整个计算任务分给多个并行的线程。
为了解决这个问题,我们需要修改kernel的代码。
cuda C++ 提供了关键字,允许kernel得到正要执行的thread是哪一个
threadIdx.x  表示当前运行的thread是block中的哪一个
blockDim.x 表示block中的线程个数

关于threadIdx.x等 下标问题,参考下图。

 

我们需要观察到使用cuda的方法之后时间的变化。

可以使用nvprof命令

 

 

可以看出时间的变化,从167.48ms到3.5144ms,再到1.8084ms

 

 

我们注意到,对线程的管理实际上是三维的 :grid,block,thread.

为什么要这样设计?

一个这样做的目的是,在一个block中,thread可以通过share_memory 来共享数据。

通过在声明的变量前面添加 shared  来表示,这个变量是声明在share memory部分了。

share memory类似与缓存,容量小,但是速度快。不过这个cache是可以编程控制的。

在一个block中共享的data,对于其他block是不可见的。

我们不妨考虑一个例子,有2个数组a,b

进行如下运算:

 

为了加快运行速度,我们还是考虑多线程的办法。

让每一个线程处理一个输出。

然而我们发现,in中除了边界元素,每一个元素都被读了7次。

这显然是没有必要的。

问题的关键在于,不同的线程之间,不知道某个元素已经被读入了。

更进一步,不同线程之间可以共享数据吗?

答案是可以的。也就是上面提到的shared_memory

然而这样就可以了吗。。

由于线程的访问顺序是不固定的(?

会发生如下的问题:

 

解决办法很无脑。。。因为cuda并没有想象中得那么底层。。。

就是使用__syncthreads(); 来同步一个block中的所有进程。。。

完整代码如下:

 

 

需要特别强调的是,cuda代码的debug问题,很多错误,不用特定的工具查看是不会显示的。
以及,虽然cuda代码是在c++上添加了一些东西,但是device code部分用的是nvidia的编译器。
所以c/cpp中,对于访问非法内存不敏感的特点,在cuda代码中不存在(我猜是因为编译器…)
在访问之前,一定要check访问地址合法性。。。。

在访问之前,一定要check访问地址合法性。。。。

在访问之前,一定要check访问地址合法性。。。。