程序博客网 > 程序员必须掌握的算法

Programming with OpenCL C

来源：互联网 发布：程序员必须掌握的算法 编辑：程序博客网 时间：2024/05/21 18:14

原文地址：http://blog.csdn.net/breakawayroad/article/details/8227450

Programming with OpenCL C

Qualifiers - 修饰符

OpenCL C支持4中修饰符：函数修饰符(function qualifiers)、地址空间修饰符(address space qualifiers)、访问修饰符(access qualifiers)和类型修饰符(type qualifiers)。

Function Qualifiers - 函数修饰符

OpenCL C 添加了函数修饰符kernel(或__kernel)。这个修饰符用来说明程序源代码中的函数为kernel函数。下面的例子说明了这个函数修饰符的用法：
[cpp] view plaincopyprint?
  1. kernel void  
  2. parallel_add(global float *a, global float *b, global float *result)  
  3. {  
  4.     ...  
  5. // The following example is an example of an illegal kernel  
  6. // declaration and will result in a compile-time error.  
  7. // The kernel function has a return type of int instead of void.  
  8. kernel int  
  9. parallel_add(global float *a, global float *b, global float *result)  
  10. {  
  11.     ...  
  12. }  
kernel函数必须满足以下的规则：
  • 返回值必须是void，否则会产生编译错误。
  • host可以通过将在队列中添加一个执行该kernel函数的指令，从而在一个设备上执行这个kernel。
  • 如果该函数被一个kernel函数调用，那么它的表现和一般的函数是一样的。唯一的限制是：内部用local修饰符来申明的变量的kernel函数不能被其他kernel函数调用。
在下面的例子中，一个kernel函数调用了另外一个kernel函数，而这个被调用函数中有变量是被local修饰符申明的。这种调用产生的结果是依赖于具体的实现的(implementation-defined)，所以为了让程序能够在所有实现上都能正常运行，应该避免这样的调用。
[cpp] view plaincopyprint?
  1. kernel void  
  2. my_func_a(global float *src, global float *dst)  
  3. {  
  4.     local float l_var[32];  
  5.     ...  
  6. }  
  7. kernel voidmy_func_b(global float * src, global float *dst){  
  8.     my_func_a(src, dst); // implementation-defined behavior  
  9. }  
为了适用于所有实现，代码进行了修改，将local变量作为参数传给kernel函数：
[cpp] view plaincopyprint?
  1. kernel void  
  2. my_func_a(global float *src, global float *dst, local float *l_var)  
  3. {  
  4.     ...  
  5. }  
  6. kernel voidmy_func_b(global float * src, global float *dst, local float *l_var){  
  7.     my_func_a(src, dst, l_var);  
  8. }  

Kernel Attribute Qualifiers

修饰符kernel可以和关键字__attribute__一起使用，用来说明关于kernel的一些附加信息：
  • __attribute__((work_group_size_hint(X, Y, Z)))：提示编译器，说明最有可能的work-group的大小，也就是在local_work_size作为参数给clEnqueueNDRangeKernel的值。
  • __attribute__((reqd_work_group_size(XYZ)))：说明要使用的work-group的大小，也就是在local_work_size作为参数给clEnqueueNDRangeKernel的值。这样编译器就可以根据已知的work-group大小进行优化。
  • __atrribute__((vec_type_hint(<type>)))：告诉编译器kernel的宽度，也就是kernel所操作的数据类型的大小。这个是针对自动向量化的编译器的提示。默认的<type>是int，指明kernel本来是标量的，自动向量化程序就能通过向量单元的SIMD通道，为多个work-item把代码向量化。

Address Space Qualifiers - 地址空间修饰符

执行kernel的work-item可以访问四个不同的存储区域。这些存储区域是通过一个修饰符来指明的，它们是：global(或__global)，local(或__local)，constant(或__constant)和private(或__private)。
使用一个地址空间名(address space name)来修饰一个对象时，这个对象将被分配到特定的地址空间中。如果没有指定地址空间名，这个对象将被分配在一般的地址空间。针对函数参数和函数的局部变量，一般的（默认的）地址空间名是private。
下面的例子中可以看到如何指定地址空间名。
[cpp] view plaincopyprint?
  1. // declares a pointer p in the private address space that points to  
  2. // a float object in address space global  
  3. global float *p;  
  4. // declares an array of integers in the private address space  
  5. int f[4];  
  6. // for my_func_a function we have the following arguments:  
  7. //  
  8. // src - declares a pointer in the private address space that  
  9. // points to a float object in address space constant  
  10. //  
  11. // v - allocate in the private address space  
  12. //  
  13. int  
  14. my_func_a(constant float *src, int4 v)  
  15. {  
  16.     float temp; // temp is allocated in the private address space.  
  17. }  

如果kernel函数的参数是一个指针，那么这个指针必须指向以下的地址空间：global，local，或者constant。如果它没有指明，会发生编译错误。这个限制不适用于非kernel函数。

下面是一些例子。
[cpp] view plaincopyprint?
  1. kernel void my_func(int *p) // illegal because generic address space  
  2.                 // name for p is private.  
  3. kernel void  
  4. my_func(private int *p) // illegal because memory pointed to by  
  5.                         // p is allocated in private.  
  6. void  
  7. my_func(int *p)        // generic address space name for p is private.  
  8.                        // legal as my_func is not a kernel function  
  9. void  
  10. my_func(private int *p) // legal as my_func is not a kernel function  

Globle Address Space - 全局地址空间

这个地址空间名用来说明分配在全局存储区域(global memory region)的存储对象(memory object，包括缓冲区buffer和图像image)。kernel的所有work-group的所有work-item可以在这个存储区域进行读写。这个地址空间通过修饰符global来指明。
一个缓冲区对象可以申明为：一个指向标量的指针，一个向量，或者一个用户定义的结构体。如下面的例子：
[cpp] view plaincopyprint?
  1. global float4 *color; // an array of float4 elements  
  2. typedef struct {  
  3.     float3 a;  
  4.     int2 b[2];  
  5. } foo_t;  
  6. global foo_t *my_info; // an array of foo_t elements  
global修饰符不能用于图像对象。
指向全局地址空间的指针可以作为函数包括kernel函数)的参数，也可以在函数内部申明变量。在函数内部的变量不能分配在全局的地址空间中。
下面是一些例子：
[cpp] view plaincopyprint?
  1. void  
  2. my_func(global float4 *vA, global float4 *vB)  
  3. {  
  4.     global float4 *p; // legal  
  5.     global float4 a; // illegal  
  6. }  

Constant Address Space - 常量地址空间

这个地址空间名用来描述在全局存储中的、可以在kernel中访问的、只读变量。这个存储区在所有执行的kernel中的所有work-group中的所有work-item的只读访问。

图像类型不能分配在常量地址空间中。下面的例子将图像定义在常量地址空间中，这是非法的，将会导致编译错误。

[cpp] view plaincopyprint?
  1. kernel void  
  2. my_func(constant image2d_t imgA)  
  3. {  
  4.     ...  
  5. }  

指向常量地址空间的指针可以作为函数的参数，也可以在函数内部定义变量。

在kernel函数内部(函数作用域的最外层)的变量可以分配到常量地址空间。在program范围的变量只能被分配到常量地址空间。所有这些变量都需要被初始化，所有用于初始化的值都必须是编译是确定的常量。对这些变量进行写操作将会产生编译错误。

在program中申明的字符串也会存储在常量地址空间。

下面是一些例子。

[cpp] view plaincopyprint?
  1. // legal - program scope variables can be allocated only  
  2. // in the constant address space  
  3. constant float wtsA[] = { 0, 1, 2, . . . }; // program scope  
  4. // illegal - program scope variables can be allocated only  
  5. // in the constant address space  
  6. global float wtsB[] = { 0, 1, 2, . . . };  
  7. kernel void  
  8. my_func(constant float4 *vA, constant float4 *vB)  
  9. {  
  10.     constant float4 *p = vA;    // legal  
  11.     constant float a;       // illegal – not initialized  
  12.     constant float b = 2.0f;    // legal – initialized with a compile-time constant  
  13.     p[0] = (float4)(1.0f);  
  14.     // illegal – p cannot be modified  
  15.     // the string "opencl version" is allocated in the  
  16.     // constant address space  
  17.     char *c = "opencl version";  
  18. }  

注意：一个kernel使用的、常量地址空间中申明的变量的数量是受设备的CL_DEVICE_MAX_CONSTANT_ARGS参数限制的。OpenCL 1.1 要求所有实现必须支持的最小值为8。也就是说，一个kernel中使用在常量地址空间中的变量不大于8个时，程序在所有实现上都会正常的运行。这8个参数的大小由CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE给出，并被设置为64KB。这样就可以把多个常量声明混合到一个常量buffer中，只要所有的大小不超过CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE即可。这中将多个变量聚集到一起声明到常量地址空间的做法不是必须的实现，并不是所有的OpenCL实现都支持它。为了程序的通用性，开发者应该假定变量不能聚集到一个单独的常量缓冲区中。

Local Address Space - 局部地址空间

这个地址空间的名字描述的是分配在local memory中的变量，它再同一个work-group中的所有work-item共享，但不能一个kernel中的多个work-group之间共享。一个work-group的所有work-item可以对其进行读写操作。

local memory非常类似于用户管理缓冲区。如果一个work-item或者一个work-group中的多个work-item从global memory中的同一个地方读取数据，使用local memory会极大地改善性能。比如对一个图像进行高斯滤波，多个work-item读取图像的重叠区域。重叠区域的大小是有滤波宽度决定的。比起多次从global memory读取多次的方法，一次从global memory中读入到local memory，然后再由多个work-item从local memory分别读取，这样会极大的改善性能。

指向局部地址空间的指针可以作为函数的参数，可以作为函数内部的变量。

在一个kernel中的申明的变量能够分配在局部地址空间中，但是有一些限制：

  • 这些变量申明必须在kernel函数的范围
  • 这些变量不能被初始化。

需要注意的是：在局部地址空间中的变量作为指针变量传给kernel函数时，或者是在kernel函数内部时，这些变量的生命周期和work-group一样。

一些例子

[cpp] view plaincopyprint?
  1. kernel void  
  2. my_func(global float4 *vA, local float4 *l)  
  3. {  
  4.     local float4 *p;// legal  
  5.     local float4 a; // legal  
  6.     a = 1;  
  7.     local float4 b = (float4)(0);   // illegal – b cannot be initialized  
  8.     if (...)  
  9.     {  
  10.         local float c; // illegal – must be allocated at kernel function scope  
  11.         ...  
  12.     }  
  13. }  

Private Address Space - 私有地址空间

这个地址空间中的变量只在单个的work-item中使用，不能再work-item之间共享。在一个kernel函数中、没有被地址空间修饰词说明的变量，所有在非kernel函数中申明的变量，和所有函数参数均是在私有的地址空间中。

Casting between Address Space - 地址空间之间的转换

一个地址空间的指针只能被同一个地址空间的指针赋值。将一个地址空间中的指针转换为另一个地址空间的指针是不合法的。比如：

[cpp] view plaincopyprint?
  1. kernel void  
  2. my_func(global float4 *particles)  
  3. {  
  4.     // legal – particle_ptr & particles are in the  
  5.     // same address space  
  6.     global float *particle_ptr = (global float *)particles;  
  7.     // illegal – private_ptr and particle_ptr are in different  
  8.     // address spaces  
  9.     float *private_ptr = (float *)particle_ptr;  
  10. }  

Access Qualifiers - 访问修饰符

访问修饰符可以用来指明图像参数，是只读的(read_only)或者是只写(write_only)的。这是因为当前GPUs的一个限制：不允许一个kernel中对同一图像的读和写。因为读取的图像是被缓冲到一个纹理缓冲区中，而对图像的写操作并不更新这个纹理缓冲区。

下面的例子中imageA是只读的2D图像对象，而imageB是只写的图像对象。

[cpp] view plaincopyprint?
  1. kernel void  
  2. my_func(read_only image2d_t imageA, write_only image2d_t imageB)  
  3. {  
  4.     ...  
  5. }  

用read_only修饰符修饰的图像可以被内建的读取图像的函数使用，但是不能被内建的写图像的函数使用。类似的，用write_only修饰的图像只能写而不能读。可以参见下面的例子：

[cpp] view plaincopyprint?
  1. kernel void  
  2. my_func(read_only image2d_t imageA, write_only image2d_t imageB, sampler_t sampler)  
  3. {  
  4.     float4 clr;  
  5.     float2 coords;  
  6.     clr = read_imagef(imageA, sampler, coords); // legal  
  7.     clr = read_imagef(imageB, sampler, coords); // illegal  
  8.     write_imagef(imageA, coords, &clr); // illegal    
  9.     write_imagef(imageB, coords, &clr); // legal  
  10. }  

imageA是只读的独享，它不能传给write_imagef。类似的，imageB被申明为只写的图像，不能作为参数穿给read_imagef。

读写修饰符read_write是被保留的。使用这个操作符会导致编译时错误。

Type Qualifiers - 类型修饰符

OpenCL C支持由C99规范定义的类型修饰符const, restrict和volatile。这些修饰符不能用来修饰image2d_t和image3d_t。只有指针类型可以使用修饰符restrict。
程序员必须掌握的算法 程序员必须掌握的算法
原创粉丝点击
热门IT博客
接口授权源码接口授权源码
mac软件开发mac软件开发
海豚号码生成软件
软件工业化标准
网上淘宝客服招聘
python wasp eureka
mac和pc是什么意思
网络管理与维护技术
mac更换ssd 重装系统
网上电子申报软件
淘宝捉猫猫技巧
php如何读取json数据
国际象棋有哪些软件
局域网ip扫描软件
如何找到网络打印机
mac sierra pdf 签名
双色球软件中奖怎么办
知乎和豆瓣哪个比较好
微信群控软件破解版
网络歌手格子兮照片
热门问题 老师的惩罚 人脸识别 我在镇武司摸鱼那些年 重生之率土为王 我在大康的咸鱼生活 盘龙之生命进化 天生仙种 凡人之先天五行 春回大明朝 姑娘不必设防，我是瞎子 梦见黑伞 雨伞黑色 黑色雨伞好吗 思源黑体 黑体 黑体字 黑体转换器 思源黑体tff 黑体字图片 黑体字体 黑体字大全 黑体炉 大黑体字 黑体辐射源 黑体字字帖 黑体字转换器 思源黑体可以商用吗 黑体字体转换器 思源黑体下载 美术黑体字大全 思源黑体字体下载 黑体繁体字体下载 黑体字体在线生成器 思源黑体 下载 黑侠 网络黑侠 黑侠2 武神黑侠 变形黑侠 黑侠 国语 网络黑侠作品 hk4941黑侠影院 女黑侠木兰花 免费观看黑侠1 蛮荒仙界 网络黑侠 绝世丹神秦风网络黑侠 网游之从头再来 网络黑侠 真黑便与假黑便图片 黑便图片 黑便原因 拉黑便

程序博客网，程序员的互联网技术博客家园。csdn论坛精品 msdn技术资料都在这里