从苹果提案到共同标准
OpenCL的全称是Open Computing Language,是把多核中央处理器、图形处理器和其他计算设备纳入共同编程接口的标准。它由苹果提出初始草案,随后在Khronos组织中集体制定。2008年12月的原始发布公告列出了苹果、AMD、英特尔、英伟达、IBM等参与企业,工作组由尼尔·特雷维特主持。
2008年12月9日,Khronos在新加坡宣布1.0规范获准并公开发布;苹果约在此前六个月提交了草案。2009年,OpenCL进入面向用户的软件产品,规范也继续维护。现存1.0规范第48修订版署2009年10月6日。
异构处理器带来的编程问题
当提高时钟频率受到功耗等条件限制时,增加处理器核心和利用GPU的并行能力成为提高计算吞吐量的重要方式。但CPU与GPU在指令、内存层次和适合的工作负载上有明显差别。过去的程序工具往往面向特定厂商或某类硬件,开发者难以在统一源代码框架中调度多种设备。
OpenCL把协调计算的API与描述计算的语言结合起来。OpenCL C以C99的一个子集为基础,增加并行执行所需的限定符和内建函数。宿主程序负责选择设备、建立运行环境和安排操作,设备上的内核函数则处理分配给它们的数据。
主机、内核与工作项
1.0规范的执行模型分成主机程序与内核两部分。主机建立包含设备、程序和内存对象的上下文,再通过命令队列安排数据传输、内核运行和同步。命令队列既可以按顺序执行,也可以在实现支持的模式下由程序显式约束依赖关系。
一次内核启动会定义一维、二维或三维的索引空间。每个索引对应一个工作项,它们执行同一份内核代码,但使用各自的编号访问数据。工作项再组成工作组。内存则区分全局、常量、组内共享的局部区域和各工作项私有区域,使程序能够表达数据共享范围。
参考:[1]
把数组相加拆成独立计算
以两个长度为n的数组相加为例,主机把输入放入内存对象,并启动覆盖这些元素的一维工作项。内核用get_global_id(0)取得编号i,在i小于n时执行c[i]=a[i]+b[i]。每个工作项只负责一个输出位置,许多元素便可以由设备并行处理,而无需把整个串行循环交给单一核心。
这个例子按规范的索引和内存模型构造。实际应用还需要建立缓冲区、设置内核参数、提交命令并等待结果可用,再把输出交还主机。若分配的工作项数超过数组长度,边界判断可以避免访问越界;不同操作之间的数据依赖则由队列顺序或同步事件表达。
参考:[1]
2009年的软件落地
苹果在2009年8月24日发布公告,宣布Mac OS X Snow Leopard将于8月28日上市,并把OpenCL列为其中的技术组成。开发者可以通过它把GPU计算能力用于图形绘制之外的任务。这个公告呈现了标准从公开规范走向面向用户的软件环境的一条早期路径。
2010年,斯通、戈哈拉和施国春的论文已比较多核CPU、GPU及Cell处理器上的OpenCL模型与实现。论文讨论的运行时编译方式允许应用把内核源代码交给目标设备的编译工具,由相应实现生成可执行代码,而不是要求所有设备执行同一种机器指令。
可移植接口与硬件差异并存
共同接口降低了为每种设备重新组织整套应用的负担,但同一内核在不同硬件上获得最佳速度,仍需要具体分析。2010年的论文指出,某种GPU适合的访存组织,未必适合Cell或多核CPU;程序可以查询设备能力,再选择合适的数据布局、工作组和内核版本。
OpenCL早期的意义因此在于建立跨厂商的低层并行计算约定:它统一了设备发现、程序编译、内存管理和任务提交的基本表达,同时让实现者保留面向硬件优化的空间。2008年的标准发布、2009年的产品采用与随后针对实际架构的研究,共同构成其早期形成过程。
参考:[4]