技术开发 频道

CUDA 技巧与经验 关于block、thread

  【IT168 文档在使用CUDA进行GPGPU计算时,global + shared的黄金组合在速度上远远超过了texture,只有在以下两种情况下使用texture:

  1.需要图像输出时

  2.需要反复随机访问的,但内容不变的大块内容,如索引表,查找表

  其他时候尽量使用global + shared

  block、thread划分的简单原则:

  用不同的block处理完全不相关的数据可以获得最好的性能。此时只需要在block内进行数据交换。例如可以用每个block处理一个独立的矩阵,或者大型矩阵中的一部分(如一行,或者一个子矩阵)。需要处理一列数据时尽量将其转化为按行存储后再处理,以避免bank conflict。

  CUDA的thread是一种轻量级线程,比较适合处理小粒度问题,在block一级也仍然只适合处理中小粒度的问题。而使用MPI构造的系统,一般是在比较大的粒度上进行并行操作,以减少数据交换。因此即便是已经经过检验的传统并行算法,在CUDA上移植时也要注意是否有可能进行进一步的划分。

  block与thread的划分原则与可以并行处理的最小子问题占用的存储器大小和子问题的数量都有关系。不同尺寸的同一问题,block和thread划分方法往往不同。例如子问题占用存储器资源有以下几种情况:

  在同一block的shared memory可以储存多个并行子问题的数据,此时可以尽量用一个block处理相邻的并行子问题,或者若干个子问题合并成一个稍大的子问题。

  每个block的shared memory只能储存一个并行子问题的数据,此时就是简单根据输入数据算输出

  每个block的shared memory不足以处理一个并行子问题,此时可以使用texture和global作为存储器,或者将子问题进一步分拆为数个串行过程,用若干个kernel完成。

  每个block的thread数量以192-256为宜,处理的矩阵或者图像尺寸最好长宽都是16的整数倍,如果可能尽量将其补成16的整数倍。

  我们这里尽量只介绍使用CUDA进行计算的方法,而少涉及具体的算法。

  核内的几种操作:

  缩减问题(reduction),推荐使用二叉树的方式来做。

  例如:缩减加操作为可以用以下代码描述:

#define ADD     if (tid < 128) { temp[tid] += temp[tid + 128];} __syncthreads();\
    
if (tid <  64) { temp[tid] += temp[tid +  64];} __syncthreads();\
    
if (tid <  32) { temp[tid] += temp[tid +  32];} __syncthreads();\
    
if (tid <  16) { temp[tid] += temp[tid +  16];} __syncthreads();\
    
if (tid <   8) { temp[tid] += temp[tid +   8];} __syncthreads();\
    
if (tid <   4) { temp[tid] += temp[tid +   4];} __syncthreads();\
    
if (tid <   2) { temp[tid] += temp[tid +   2];} __syncthreads();\
    
if (tid <   1) { temp[tid] += temp[tid +   1];}\

  基于同样的树状结构进行缩减操作,还可以进行排序,累乘,寻找最大最小值,求方差等等

  邻域问题:

  在图像处理中,经常可以遇到邻域问题,如4-邻域,8-邻域等。基本思想就是使用图像中的几个元素按照一定的算子计算输出矩阵中的一个元素

  常见的8-邻域算子如:

  -1 -1 -1 1 1 1

  -1 8 -1 1 1 1

  -1 -1 -1 高通滤波器 1 1 1 低通滤波器

  等等,可以完成锐化、平滑、去噪,查找边缘等功能

  cuda的imagedenoising例子使用的是texture,其实也可以使用shared memory来做。

0
相关文章