OpenCL编程指南-4.2矢量数据类型

矢量数据类型

OpenCL C还增加了对矢量数据类型的支持。矢量数据类型如下定义，首先是类型名，具体包括char、uchar、short、ushort、int、uint、float、long或ulong，后面是一个字面值n来定义矢量中的元素个数。对于所有矢量数据类型，支持的n值包括2、3、4、8和16。另外，还可以为double和half定义矢量数据类型，不过只有当设备支持双精度和半精度扩展时才可用。支持的矢量数据类型见表4-2。
在这里插入图片描述
声明为一个标量或矢量数据类型的变量要按所用数据类型的大小（字节数）对齐。内置数据类型大小必须按2的幂（字节数）对齐。如果一个内置数据类型的大小（字节数）不是2的幂，则要按紧邻的下一个2的幂值对齐（即大于它的最小的2的幂值)。这个规则不适用于结构（struct）或联合（union)。

例如，一个float4变量要按16字节边界对齐，char2变量要按2字节边界对齐。对于一个包含3个分量的矢量数据类型，这个数据类型的大小为4×sizeof(分量)。这说明包含3个分量的矢量数据类型要按4×sizeof(分量)边界来对齐。

OpenCL编译器负责根据数据类型适当地将数据项对齐。只有声明为一个数据类型指针的内核函数参数例外。对于这些函数，编译器假设指针所指向的对象总会按数据类型适当地对齐。

为了便于应用并确保数据存储能正确地对齐，表4-3列出了应用程序可用的一些数据类型。
在这里插入图片描述

矢量字面量

可以用矢量字面量由一组标量、矢量或标量和矢量的组合创建矢量。矢量字面量可以作为矢量初始化语句，也可以作为基本表达式，但不能作为左值。

矢量字面量的写法如下：首先是一个用括号括起的矢量类型，后面是一个用括号括起的参数列表，各参数之间用逗号分隔。矢量字面量会作为一个重载函数处理。可用的函数形式为可能的参数列表集合，其中所有参数与结果矢量有相同的元素类型，而且元素总数等于结果矢量中元素的个数。另外，还可以采用一种单个标量的形式，该标量的类型要与矢量的元素类型相同。例如，对于float4可以采用以下形式：

(float4) (float, float, float, float)
(float4) (float2, float, float)
(float4) (float, float2, float)
(float4) (float, float, float2)
(float4) (float2, float2)
(float4) (float3, float)
(float4) (float, float3)
(float4) (float)

操作数按函数计算的标准规则来计算，除非没有适用的隐式标量宽化规则。操作数会按其在内存中出现的顺序指定到结果矢量中相应的位置上。也就是说，第一个操作数的第一个元素指定到result.x，第一个操作数的第二个元素指定到result.y(或者如果第一个操作数是一个标量，则将第二个操作数的第一个元素指定到result.y)，以此类推。如果操作数是一个标量，那么这个操作数会复制到结果矢量的所有分量。

下面的例子显示了由一个标量列表创建的矢量float4：

float4 f = (float4)(1.0f, 2.0f, 3.0f, 4.0f);

下面的例子显示了由一个标量创建的矢量uint4，这个标量会复制到矢量的所有分量中：

uint4 u = (uint4)(1); //u will be (1,1,1,1)

下面几个例子展示了使用一个标量和一些较小矢量类型创建更为复杂的矢量组合：

float4 f = (float4)((float2)(1.0f, 2.0f), (float2)(3.0f, 4.0f));
float4 f = (float4)(1.0f, (float2)(2.0f, 3.0f), 4.0f);

下面展示了非法创建矢量字面量的几个例子。所有这些例子都会导致一个编译错误。

float4 f = (float4)(1.0f, 2.0f);
float4 f = (float2)(1.0f, 2.0f);
float4 f = (float4)(1.0f, (float2)(2.0f, 3.0f));

矢量分量

在有1~4个分量（即元素）的矢量数据类型中，其分量可以表示为<vector>.xyzw。表4-4列出了不同矢量类型可访问的分量。

char2 uchar2 short2 ushort2 int2 uint2 long2 ulong2 float2        .xy
char3 uchar3 short3 ushort3 int3 uint3 long3 ulong3 float3        .xyz
char4 uchar4 short4 ushort4 int4 uint4 long4 ulong4 float4        .xyzw
double2 half2                                                     .xy
double3 half3                                                     .xyz
double4 half4                                                     .xyzw

如果要访问那些超出矢量类型声明的分量，会得到一个错误。下面给出了合法和非法访问矢量分量的一些例子：

float2 pos;
pos.x = 1.0f; //is legal
pos.z = 1.0f; //is illegal

float3 pos;
pos.z = 1.0f; //is legal
pos.w = 1.0f; //is illegal

利用分量选择语法，可以在点号(.)后面追加分量名来选择多个分量。下面给出几个例子来了解如何使用分量选择语法：

float4 c;
c.xyzw = (float4)(1.0f, 2.0f, 3.0f, 4.0f)
c.z = 1.0f;
c.xy = (float2)(3.0f, 4.0f);
c.xyz = (float3)(3.0f, 4.0f, 5.0f);

利用分量选择语法还允许交换或复制分量，如下面的例子所示：

float4 pos = (float4)(1.0f, 2.0f, 3.0f, 4.0f);
float4 swiz = pos.wzyx; //swiz = (4.0f, 3.0f, 2.0f, 1.0f)
float4 dup = pox.xxyy; //dup = (1.0f, 1.0f, 2.0f, 2.0f)

还可以使用数值索引（指向矢量中相应的元素）访问矢量分量。数值索引的用法如表4-5所示
在这里插入图片描述
所有数值索引前面必须加上字母s或s。对于下面的例子，f.s0指示float8变量f的第1个元素。f.s7指示float8变量f的第8个元素。

float8 f

在下面的例子中，x.sa（或x.sA)指示float16变量x的第11个元素，x.sf(或x.sF)指示float16变量x的第16个元素：

float16 f

数值索引不能与.xyzw记法混合使用。例如：

float4 f;
float4 v_A = f.xs123; //is illegal
float4 v_B = f.s012w; //is illegal

矢量数据类型可以使用.lo (或.odd)和.hi(或.even)后缀来得到较小的矢量类型，或者把较小的矢量类型结合为一个较大的矢量类型。还可以使用多层.lo(或.odd）和.hi(或.even)后缀，直至指向一个标量类型。

.lo后缀指示一个给定矢量的低半部分。.hi后缀指示给定矢量的高半部分。.odd后缀指示给定矢量的奇数元素。.even后缀指示给定矢量的偶数元素。下面给出一些例子来说明这个概念：

float4 vf;

float2 low = vf.lo; //returns vf.xy
float2 high = vf.hi; //returns vf.zw
float x = low.low; //returns low.x
float y = low.hi; //returns low.y
float2 odd = vf.odd; //returns vf.yw
float2 even = vf.even; //returns vf.xz

对于一个包含3个分量的矢量，处理后缀.lo(或.odd)和.hi(或.even)时会把这个3分量的矢量当做一个4分量的矢量来处理，其w分量中的值未定义。

其他数据类型

OpenCL C支持的其他数据类型

image2d_t      2D图像类型
image3d_t      3D图像类型
sampler_t      图像采样器类型
event_t        事件类型。
               完成从全局内存到局部内存（以及从局部内存到全局内存）异步复制的内置函数会使用这些事件类型。
               每个异步复制操作会返回一个事件，并接收一个所等待的事件(该事件标识之前的一个异步复制操作)

对于图像和采样器类型的使用，存在一些限制：
1）只有当设备支持图像时才能定义图像和采样器类型。
2）图像和采样器类型不能声明为数组。下面的几个例子展示了一些不合法的用法。

kernel void foo(image2d_t imgA[10]) //error images cannot be declared as arrays
{
     image2d_t imgB[4]; //error images cannot be declared as arrays
     ...
}

kernel void foo(sampler_t smpA[10]) //error samplers cannot be declared as arrays
{
     sampler_t smpB[4]; //error samplers cannot be declared as arrays
     ...
}

1）不能在结构(struct）中声明image2d_t、image3d_t 和 sampler_t数据类型
2）变量不能声明为image2d_t、image3d_t和sampler_t数据类型的指针。

衍生类型

对于这些衍生类型的使用有一些限制：
1）如果结构或结构指针用做一个内核函数的参数类型，则该结构类型不能包含任何指针。例如，下面的用法是不合法的。

typedef struct {
    int x;
    global float *f;
} mystruct_t;
kernel void foo(global mystruct_t *p) //error mystruct_t contains a pointer
{
    ...
}

2）只有当结构或结构指针用做一个非内核函数的参数类型，或者声明为一个内核或非内核函数中的变量时，该结构类型才能包含指针。例如，以下用法是合法的。

void my_func(mystruct_t *p)
{
   ....
}

kernel void foo(global int *p1, global float *p2)
{
     mystruct_t s;
     s.x = p1[get_global_id(0)];
     s.f = p2;
     my_func(&s);
}