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位浮点数的矢量从标量数组中加载与存储也做了详细说明,用法与上述的讲述类似。

赞(1)
未经允许不得转载:极客笔记 » 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映射