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扩展有效，就可以在程序中定义双精度浮点数变量。