OpenCL 标量数据类型

标量和矢量,对于物理学家、数学家和编程人员,不同的人会有不同的理解。所以笔者觉得很有必要解释下OpenCL中标量和矢量数据类型。

对于一个标量数据,每个数据包含了单个数据类型的值。而矢量数据类似于数组,包含了多个相同类型的值。一个矢量数据就是包含了多个标量数据的数据集合。一个矢量数据只能包含特定个数的标量数据。对一个矢量数据执行某个操作,如加法操作,矢量内的标量数据同时执行相同操作。

例如,我们需要把4个整数相加。对于标量数据操作,可以定义数组a和b,它们分别有4个整数,数组c为相加的结果,相应代码如下:

int a[4], b[4], c[4]
for(int i = 0; i 〈 4; i++)
{
    c[i] = a[i] + b[i];
}

对于矢量数据操作,a、b和c都为矢量数据,每个矢量包含4个标量数据,相应代码如下:

int4 a, b, c;
c = a + b;

矢量操作带来的不仅仅是代码更简洁,在具有SIMD处理单元或VLIW处理的单元的处理器上,处理速度也相应地提高了。

现在我们还是回到本节标量数据类型的内容中,在OpenCL中支持的标量数据类型如下表所示:

OpenCL 内建标量数据类型

OpenCL支持的标量数据类型是比较简单的,功能与C/C++中的数据类型是一样。大多数内置标量数据类型在OpenCL API(和头文件中)被声明为可供程序使用的适当类型。对于上表中n/a的类型,是无法确保计算设备端的数据类型所占字节数能与主机端取得一致。就拿size_t类型来说,如果计算设备当前为64位地址空间,而主机端为32位地址空间,那么设备端的size_t类型宽度为8字节(64位),而主机端则是4字节(32位),因此两者此时无法兼容。所以我们要把主机端的参数传递到内核函数上时,尽量避免使用上表中n/a的数据类型,而是使用可确定兼容的数据类型。

这里需要强调的是,双精度浮点数是一个可选数据类型。因为支持OpenCL的设备很多,但不是所有的设备都支持双精度浮点数,如高通Adreno GPU。可以查询OpenCL设备CL_DEVICE_DOUBLE_FP_CONFIG属性信息,如果结果为0,则说明设备不支持双精度。

对于支持双精度浮点运算的设备,为了在内核函数中启用这个功能,可以在内核代码最上方添加如下语句:

#pragma OPENCL EXTENSION cl_khr_fp64:enable

添加上述语句后,在内核函数中可以定义双精度浮点数变量,可以对这些变量进行相应的操作。例如:

#pragma OPENCL EXTENSION cl_khr_fp64:enable
kernel void DoubleProcess(global double *A, global double *B,
                              global double *C)
{
    int ID = get_global_id(0);
    double temp = 0.35;
    C[ID] = temp * A[ID] + B[ID];
    ……
}

出于计算精度要求,我们会声明一些变量为双精度浮点数。但是对于OpenCL设备而言,相比于处理单精度浮点数能力,处理双精度浮点数能力要更弱。例如,AMD FirePro W9100 GPU,理论双精度与单精度峰值计算能力比为2.62/5.24=1/2;NVIDIA Tesla K80 GPU,理论双精度与单精度峰值计算能力比为2.91/8.74=1/3。通俗的理解,对于上述设备,如果处理双精度浮点数,速度会比单精度慢2~3倍。所以笔者建议,在不需要高精度的情况下,使用单精度浮点数;在迫不得已需要更高精度时,中间临时变量声明为双精度浮点数,这么做的目的就是尽可能地减少设备对双精度浮点数的处理,提升整个程序性能。

下面的例子展示了如何判断设备是否支持双精度浮点数,然后利用宏定义选择内核部分代码。

主机端代码:

……
cl_device_fp_config DeviceDouble;
int Doubleflag;//1:设备支持双精度;0:设备不支持双精度
……
Doubleflag = 1;
//查询设备CL_DEVICE_DOUBLE_FP_CONFIG属性,如果返回值为0,则
//设备不支持双精度浮点数
err = clGetDeviceInfo(device, CL_DEVICE_DOUBLE_FP_CONFIG,
                          sizeof(cl_device_fp_config),
&DeviceDouble, NULL);
if(0 == DeviceDouble)
{
    Doubleflag = 0;
}
if(1 == Doubleflag)
{
    //设备支持双精度浮点数,编译时添加FP_64宏
    clBuildProgram(program, 1, &device, "-D FP_64", NULL, NULL)
}
else
{
    //设备不支持双精度浮点数,采用默认编译选项
    clBuildProgram(program, 1, &device, NULL, NULL, NULL)
}
……
内核代码:
#ifdef FP_64
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
kernel void DoubleTest(global float *a, global float *b,
                            global float *out)
{
#ifdef FP_64
    double c = (double)(*a/*b);
    *out=(float)c;
#else
    *out=(*a)*(*b)
#endif
}

上述代码,在主机端查询设备的CL_DEVICE_DOUBLE_FP_CONFIG属性,如果返回值为1,说明设备支持双精度浮点数;如果返回0,设备不支持双精度浮点数。如果设备支持双精度,在编译内核程序时添加编译选项“-D FP_64”,开启内核代码中FP_64宏定义,这样使得cl_khr_fp64扩展在内核代码中有效。一旦cl_khr_fp64扩展有效,就可以在程序中定义双精度浮点数变量。

赞(0)
未经允许不得转载:极客笔记 » OpenCL 标量数据类型

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址
OpenCL 基本概念
OpenCL 是什么OpenCL 平台模型OpenCL 执行模型OpenCL 上下文简介OpenCL 命令队列简介OpenCL 在设备上执行内核OpenCL 存储器区域OpenCL 存储器对象OpenCL 共享虚拟存储器OpenCL 与OpenGL
OpenCL 基础教程
OpenCL 在Windows上搭建开发环境OpenCL 在Linux上搭建开发环境OpenCL 在OS X上搭建开发环境OpenCL 第一个程序OpenCL 平台OpenCL 设备OpenCL 创建上下文OpenCL 创建命令队列OpenCL 创建Program对象OpenCL 编译Program对象OpenCL 查询和管理Program对象OpenCL 创建内核对象OpenCL 设置内核参数OpenCL 查询和管理内核对象OpenCL 执行内核OpenCL 编写内核代码OpenCL 错误处理
OpenCL C特性
OpenCL 地址空间修饰符OpenCL 函数修饰符OpenCL 对象访问修饰符OpenCL 标量数据类型OpenCL 为什么要有矢量数据类型OpenCL 矢量初始化OpenCL 读取和修改矢量分量OpenCL 运算符OpenCL 维度和工作项OpenCL 工作组OpenCL 矢量数据拷贝OpenCL 异步拷贝和预取OpenCL 数学函数OpenCL 公共函数OpenCL 几何函数OpenCL 整数函数OpenCL 关系函数OpenCL 杂项矢量函数OpenCL 同步函数OpenCL 原子函数OpenCL 内建图像读函数OpenCL 内建无采样器图像读函数OpenCL 内建图像写函数OpenCL 内建图像查询函数OpenCL 工作组函数OpenCL 内建管道读/写函数OpenCL 内建工作组管道读/写函数OpenCL 内建管道查询函数OpenCL 设备队列OpenCL Blocks语法OpenCL 设备队列相关函数OpenCL 子内核存储器可见性OpenCL 设备队列的使用示例
OpenCL 存储器对象
OpenCL 存储器对象OpenCL 分配缓冲区对象OpenCL 创建子缓冲区对象OpenCL 图像对象和采样器对象OpenCL 图像对象OpenCL 图像格式描述符OpenCL 图像描述符OpenCL 图像对象查询OpenCL 采样器对象OpenCL 主机端采样器对象OpenCL 设备端采样器对象OpenCL 图像旋转示例OpenCL 管道OpenCL 创建管道对象OpenCL 管道对象查询OpenCL 主机与设备间数据传输OpenCL 图像对象主机与设备间数据拷贝OpenCL 缓冲区对象数据填充OpenCL 图像对象数据填充OpenCL 缓冲区对象间数据传输OpenCL 图像对象和缓冲区对象间数据拷贝OpenCL 缓冲区对象映射OpenCL 图像对象映射OpenCL 解映射OpenCL 共享虚拟存储器OpenCL SVM缓冲创建与释放OpenCL SVM缓冲映射与解映射OpenCL SVM缓冲填充与拷贝OpenCL SVM类型OpenCL SVM特性OpenCL 共享虚拟存储器示例OpenCL 存储器一致性模型OpenCL 存储器次序规则OpenCL 原子操作的存储器次序规则OpenCL 栅栏操作的存储器次序规则OpenCL 工作组函数的存储器次序规则OpenCL 主机端与设备端命令的存储器次序规则OpenCL 关于存储器次序在实际OpenCL计算设备中的实现
OpenCL 同步及事件机制
OpenCL 同步及事件机制OpenCL 主机端的OpenCL同步OpenCL OpenCL事件机制OpenCL 对OpenCL事件的标记和栅栏OpenCL 内核程序中的同步OpenCL 工作组内同步OpenCL 原子操作OpenCL 1.2中的原子操作OpenCL 2.0中的原子操作OpenCL 局部存储器与全局存储器间的异步拷贝OpenCL 工作组间同步
OpenCL 与OpenGL互操作
OpenCL 与OpenGL互操作OpenCL 从一个OpenGL上下文来创建OpenCL上下文OpenCL 使用OpenGL共享的缓存对象OpenCL 使用OpenGL纹理数据OpenCL 共享OpenGL渲染缓存OpenCL 从一个OpenCL存储器对象查询OpenGL对象信息OpenCL 访问共享对象的OpenCL与OpenGL之间的同步OpenCL AMD Cayman架构GPUOpenCL AMD GCN架构的GPUOpenCL NVIDIA CUDA兼容的GPUOpenCL NVIDIA GPU架构的执行模型OpenCL NVIDIA GPU的全局存储器OpenCL NVIDIA GPU的局部存储器OpenCL ARM Mali GPU硬件架构OpenCL ARM Mali GPU存储器层次OpenCL ARM Mali GPU OpenCL映射