GPGPU OpenCL/CUDA 高性能编程的10大注意事项

转载自:http://hc.csdn.net/contents/content_details?type=1&id=341

1.展开循环

如果提前知道了循环的次数,可以进行循环展开,这样省去了循环条件的比较次数。但是同时也不能使得kernel代码太大。


 1 #include
2 using namespace std;
3
4 int main(){
5 int sum=0;
6 for(int i=1;i<=100;i++){
7 sum+=i;
8 }
9
10 sum=0;
11 for(int i=1;i<=100;i=i+5){
12 sum+=i;
13 sum+=i+1;
14 sum+=i+2;
15 sum+=i+3;
16 sum+=i+4;
17 }
18 return 0;
19 }

2.避免处理非标准化数字

OpenCL中非标准化数字,是指数值小于最小能表示的正常值。由于计算机的位数有限,表示数据的范围和精度都不可能是无限的。(具体可以查看IEEE
754标准,http://zh.wikipedia.org/zh-cn/IEEE_754)

在OpenCL中使用非标准化数字,可能会出现“除0操作”,处理很耗时间。

如果在kernel中“除0”操作影响不大的话,可以在编译选项中加入-cl-denorms-are-zero,如:

clBuildProgram(program, 0, NULL, "-cl-denorms-are-zero", NULL,
NULL);

3.通过编译器选项传输常量基本类型数据到kernel,而不是使用private
memory

如果程序中需要给kernel
传输常量基本类型数据,最好是使用编译器选项,比如宏定义。而不是,每个work-item都定义一个private
memory变量。这样编译器在编译时,会直接进行变量替换,不会定义新的变量,节省空间。

如下面代码所示(Dmacro.cpp):

?




1 #include

2 int
main()

3 {

4     int
a=SIZE;

5     printf("a=%d, SIZE=%d\n",a,SIZE);

6     return
0;

7 }

编译:

g++ -DSIZE=128 -o A Dmacro.cpp

4.如果共享不重要的话,保存一部分变量在private memory而不是local
memory

work-item访问private memory速度快于local memory,因此可以把一部分变量数据保存在private
memory中。当然,当private memory容量满时,GPU硬件会自动将数据转存到local memory中。

5.访问local memory应避免bank conflicts

local
memory被组织为一个一个的只能被单独访问的bank,bank之间交叉存储数据,以便连续的32bit被保存在连续的bank中。如下图所示:

(1)如果多个work-item访问连续的local
memory数据,他们就能最大限度的实现并行读写。

(2)如果多个work-item访问同一个bank中的数据,他们就必须顺序执行,严重降低数据读取的并行性。因此,要合理安排数据在local
memory中的布局。

(3)特殊情况,如果一个wave/warp中的线程同时读取一个local memory中的一个地址,这时将进行广播,不属于bank
冲突。

6.避免使用”%“操作

"%"操作在GPU或者其他OpenCL设备上需要大量的处理时间,如果可能的话尽量避免使用模操作。

7.kernel中重用(Reuse) private
memory,为同一变量定义不同的宏

如果kernel中有两个或者以上的private
variable在代码中使用(比如一个在代码段A,一个在代码段B中),但是他们可以被数值相同。

也就是当一个变量用作不同的目的时,为了避免代码中的命名困惑,可以使用宏。在一个变量上定义不同的宏。

?





1

2

3

4

5

6

7

8

9

10

11

1 #include

 2 int
main(){

 3     int
i=4;

 4     #define EXP i

 5             printf("EXP=%d\n",EXP);

 6    

 7     #define COUNT i

 8             printf("COUNT=%d\n",COUNT);

 9     getchar();

10     return
0;

11 }

8.对于(a*b+c)操作,尽量使用 fma
function

如果定义了“FP_FAST_FMAF”宏,就可以使用函数fma(a,b,c)精确的计算a*b+c。函数fma(a,b,c)的执行时间小于或等于计算a*b+c。

9.在program file
文件中对非kernel的函数使用inline

inline修饰符告诉编译器在调用inline函数的地方,使用函数体替换函数调用。虽然会使得编译后的代码占用memory增加,但是省去了函数调用时上下、函数调用栈的切换操作,节省时间。

10.避免分支预测惩罚,应该尽量使得条件判断为真的可能性大

现代处理器一般都会进行“分支预测”,以便更好的提前“预取”下一条要执行的指令,使得“取指令、译码分析、执行、保存”尽可能的并行。

在“分支预测”出错时,提前取到的指令,不是要执行的指令,就需要根据跳转指令,进行重新取指令,就是“分支预测惩罚”。

看如下的代码:


1 #include
2 int main()
3 {
4 int i=1;
5 int b=0;
6 if(i == 1)
7 b=1;
8 else
9 b=0;
10 return 1;
11 }

对应的汇编代码:

(movl 赋值,cmpl 比较,jne 不等于跳转,jmp 无条件跳转)

从上面的汇编指令代码看出,如果比较(<main+24>)结果相等,则执行<main+26>也就是比较指令的下一条指令,对应b=1顺序执行;如果比较(<main+24>)结果不相等,则执行跳转到<main+35>,不是顺序执行。

当然,有的处理器可能会根据以往“顺序执行”与“跳转执行”的比例来进行分支预测,但是这也是需要积累的过程。况且并不是,每个处理器多能这样只能。

最后,上面的10个tips,能过提升kernel函数的性能,但是你应该进行具体的性能分析知道程序中最耗时的地方在哪里。当然了,只有通过实验才能真正学会OpenCL高性能编程。

GPGPU OpenCL/CUDA 高性能编程的10大注意事项,布布扣,bubuko.com

时间: 2024-10-14 01:00:19

GPGPU OpenCL/CUDA 高性能编程的10大注意事项的相关文章

CUDA高性能编程中文实战11章例子中多设备的例子编译提示问题

提示的问题如下: error : argument of type "void *(*)(void *)" is incompatible with parameter of type "CUT_THREADROUTINE" 解决方法是将下面的代码 void* routine(void *pvoidData) 修改为下面的代码 unsigned WINAPI routine(void *pvoidData) 原因是为什么,我也不知道,同求解.

BS界面设计最重要的10大注意事项

如果您想订阅本博客内容,每天自动发到您的邮箱中,请点这里 BS界面是基于浏览器的界面,随着人们对于用户体验要求的不断提高,BS界面的设计要求也越来越高,那么在BS界面设计中有哪些注意事项呢?总结如下: 界面色彩要求:计算机屏幕的发光成像和普通视觉成像有很大的不同,应该注意这种差别作出恰当的色彩搭配.对于用户长时间使用的系统,应当使用户在较长时间使用后不至于过于感到视觉疲劳为宜.例如轻松的淡彩为主配色,灰色系为主配色等等.切忌色彩过多,花哨艳丽,严重妨碍用户视觉交互. 界面平面版式要求:系统样式排

Java字符串的10大热点问题盘点

下面我为大家总结了10条Java开发者经常会提的关于Java字符串的问题,如果你也是Java初学者,仔细看看吧: 1.如何比较字符串,应该用"=="还是equals()? 总的来说,"=="是用来比较字符串的引用地址,而equals()才是比较字符串的值.两个值相同的字符串用"=="比较结果有可能是false,而用equals()则一定为true.除非两个字符串是同一个new出来的对象,不然比较字符串值是否相同应该要用equals(). 2.对那

CUDA __shared__ thread、block、grid之间的一维关系 (例子chapter5 dot点积(GPU高性能编程))

chapter5里重要的例子是dot,来解释一个block内多个thread的共享内存和同步. __shared__共享内存:“对于在GPU上启动的每个线程块,cuda c编译器都将创建该变量的一个副本.线程块中的每个线程都共享这块内存,并和其他线程块无关,这使一个线程块中多个线程能够在计算上进行通信和协作” __syncthreads():确保线程块中的每个线程都执行完__syncthreads()前面的语句后,在往下执行. 例子是Grid->一维Block->一维Thread: 通过实例代

《GPU高性能编程CUDA实战》中代码整理

CUDA架构专门为GPU计算设计了一种全新的模块,目的是减轻早期GPU计算中存在的一些限制,而正是这些限制使得之前的GPU在通用计算中没有得到广泛的应用. 使用CUDA C来编写代码的前提条件包括:(1).支持CUDA的图形处理器,即由NVIDIA推出的GPU显卡,要求显存超过256MB:(2).NVIDIA设备驱动程序,用于实现应用程序与支持CUDA的硬件之间的通信,确保安装最新的驱动程序,注意选择与开发环境相符的图形卡和操作系统:(3).CUDA开发工具箱即CUDA Toolkit,此工具箱

2017 编程趋势预测:10大技术大热,10大技术遇冷

在编程世界,流行的是科学,严谨和精确的准则.这不是说编程是一种缺乏趋势的行业.不同的是,由于编程重视的是高效,用户化和使用方便,因此编程技术需要不断更新换代. 那么,在即将到来的2017年和未来的几年内,会有哪些新技术出现,成为新的潮流,又有哪些技术的热潮会渐渐消退呢?下面所列出的10种技术,摘自infoworld.com的一位编辑Peter Wayner撰写的文章21 hot programming trends -- and 21 going cold.本文旨在与大家一起讨论这个话题. 1.

漂亮的10大编程字体

10. Courier 基本上所有的系统都自带了此种字体,有时候它又以Courier New的名字出现.不幸的是很多终端和编辑器都默认使用此种字体,虽然这不会影响使用,但这会影响心情,它太枯燥了.以前看到这样一句评论:久不见牡丹会以仙人掌为美.这句评论形Courier字体非常合适.所以如果你还有其他选择的话,请勿使用此字体.更不幸的事情是最后你还是会继续使用它,那3秒只能强烈建议你调整一下字体大小并消除锯齿. 9. Andale Mono 稍微比Courier好一点的字体,有些时候它也被用作默认

Master the 10 Most Common Python Programming Problems - 10大最常见的Python编程错误

http://blog.csdn.net/pipisorry/article/details/45175457 Introduction 本文介绍python编程中很难捕捉10大错误 (Note: This article is intended for a more advanced audience than Common Mistakes of Python Programmers, which is geared(适合) more toward those who are newer t

每个程序员都该知道的10大编程格言

每个程序员都该知道的10大编程格言 编程格言1:无风不起浪 (There is no smoke without fire) 编程格言2:预防为主,治疗为辅(An ounce of prevention is worth a pound of cure:) 编程格言3:不要把鸡蛋都放在一个篮子(Don't put all your eggs in one basket) 编程格言4:种瓜得瓜,种豆得豆(As you sow,so shoul you reap) 编程格言5:欲速则不达(Great