OpenCL 矢量数据拷贝

在之前教程中，我们讲述了OpenCL设备分为多种不同存储器，能否高效地利用这些大小不同，读写速度也不同的存储器对我们程序性能有很大影响。在本教程中，我们分为矢量数据拷贝以及异步拷贝和预取来讲解。

矢量数据拷贝

使用矢量数据拷贝相比使用标量数据拷贝，在单条指令内可以获得更多的数据，也就是说处理相同数据量的情况下，减少了指令数据。这对于支持矢量操作的设备来说，提高了程序性能。

对于同类型的矢量数据，如从int4数据中加载一个int4数据，我们可以使用最简单的“=”符号，除了实现对等号左边变量赋值之外，隐式地实现了数据从一种存储器向另一种存储器的拷贝。例如，在内核中需要把数据从全局存储器中拷贝到局部存储器中，处理这些数据之后再把数据从局部存储器拷贝到全局存储器中，如下代码：

kernel void test(global int4 ＊In, global int4 ＊Out)
{
    local int4 local_data;
    int id = get_global_id(0);
local_data = In[id];
……
Out[id] = local_data;
}

例子中，使用了两次“=”操作，第一次把数据从全局存储器拷贝到局部存储器，第二次把处理之后的数据从局部存储器拷贝到全局存储器中。

直接通过“=”操作把矢量数据拷贝到另一矢量变量的用法，毕竟不能适用我们实际开发中的情况。很多时候，我们的初始数据是存在于标量数组中，在内核中我们可以使用如下函数实现加载多个标量数据到矢量数据中：

gentypen vloadn(size_t offset, const gentype ＊p);

gentype表示内建的标量数据类型，可以为char、uchar、short、ushort、int、uint、long、ulong、f loat或double。n表示矢量数据分量的个数，可以为2、3、4、8或16。矢量和标量可以在不同的地址空间。矢量gentypen的数据起始位置为（p+（offset*n））读取的sizeof（gentypen）字节的数据。
例如，int_vec是一个int4类型的矢量数据，int_array为int类型的一个标量数据，如下代码展示了从int_array中加载前四个整型数据到int_vec中：

int4 int_vec=vload4(0,int_array)

参数offset决定了标量数组中哪些元素会被拷贝到矢量中。需要注意的是offset偏移量，指的是矢量数据偏移量，而不是标量元素的偏移量。如上代码如果把偏移量从0改成1，int_vec中保存的是标量数组中第5～8个元素。下图形象地说明了这个。

从标量数组中加载矢量，必须按照sizeof（gentype）*4位对齐，也就是说char或uchar类型，则必须按照8位对齐；short或ushort类型，则必须按照16位对齐；int、uint或f loat类型，则必须按照32位对齐；long、ulong或double类型，则必须按照64位对齐。而图4-1中下方图，把标量数组地址+2来加载非对齐数据。对于GPU而言，在一条指令周期内不能完成所需的非对齐数据加载，这会影响程序性能。
计算以后的矢量数据，我们可以通过如下函数实现把矢量数据传输到标量数组中：

gentypen vstoren(gentypen data, size_offset, gentype ＊p)

data为矢量数据，其他参数与函数vloadn（）一样。
下面的例子展示了把一个int4类型的int_vec矢量数据保存到标量数组int_array中：

vstore4(int_vec, 0, int_array);

通过改变offset的值，我们可以把矢量存储到标量数组不同矢量宽度偏移量的位置。改变指针p所指向的地址，我们也可以把矢量存储到标量数组的非对齐的位置。

在OpenCL 2.0的文档中对16位浮点数的矢量从标量数组中加载与存储也做了详细说明，用法与上述的讲述类似。