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 标量数据类型
分享到: 更多 (0)

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址