kernel和OpenCL执行模型

《OpenCL异构计算》第2章OpenCL简介，本章介绍OpenCL，其编程要素允许我们快捷地编写并发运行的应用程序。熟悉C和C++的程序员应该不难理解OpenCL的语法。本节为大家介绍kernel和OpenCL执行模型。

AD：51CTO移动APP安全沙龙！马上要爆满，手慢没座位！

 

kernel和OpenCL执行模型

kernel是OpenCL程序在设备上实际运行的那部分代码。OpenCL API允许应用程序新建上下文来管理OpenCL命令的执行，包括描述主机端和OpenCL内存结构之间的数据传输以及kernel代码的执行(即对这些数据进行处理以完成特定的任务)。

像许多CPU并发模型一样，OpenCL kernel的语法类似于标准的C函数，关键区别在于它额外有一个关键字集和OpenCLkernel实现的执行模型。使用OS线程API或OpenMP在CPU上开发并发程序时，编程人员需要考虑可用的计算资源(例如CPU核数)以及线程数超过资源限制时所带来的新建和切换开销。使用OpenCL，编程人员的目标是尽可能细粒度地表示程序中的并行性。OpenCL接口的通用性和底层的kernel语言使应用程序能够高效地映射到大多数硬件。以下将讨论向量相加函数的三个实现版本：串行C版本、多线程C版本和OpenCL版本。

向量相加的串行C实现代码执行的是一个循环，其迭代次数与需要计算的数组元素一样多。每次循环迭代都将输入数组中循环下标相应的元素相加，并将结果保存到输出数组中：

 

对于一个简单的多核设备，我们可以使用低层次的粗粒度线程API(如Win32或POSIX)或使用数据并行模型(如OpenMP)。编写一个具有相同功能的粗粒度多线程版本需要在线程之间划分任务(即循环迭代)。因为可能有大量的循环迭代，而每次迭代的工作量很小，所以我们需要将循环迭代按块划分成更大的粒度(称为"循环切分"(strip mining)，Cooper and Torczon [2011])。多线程版本的代码可能像下面这样：

相比Win32和POSIX线程API， OpenCL更接近于OpenMP，支持数据并行执行，但不足之处在于不太容易控制。OpenCL C并发执行单位是一个work-item 。如同前两个例子，每个work-item执行kernel函数体。不用手动切分循环，我们往往将单层循环迭代映射到一个work-item，使OpenCL运行时产生的work-item数目和输入输出元素数目保持一致，并允许在运行时将这些work-item映射到底层硬件，因此CPU或GPU核，以它们认为适当的任何方式执行程序。从概念上讲，这是非常类似 map操作函数 固有的并行性或针对循环的数据并行模型，如OpenMP。当OpenCL设备开始执行kernel时，会提供内置(intrinsic)函数让work-item表明它自身。在以下代码中，编程人员通过调用get_global_id(0)来获取当前work-item所在的位置，在这个简单的例子中用来重新获取循环计数器。

 

 

由于OpenCL能够以细粒度work-item的方式来描述执行，并能够将大量的work-item分发到硬件上支持细粒度线程的架构，因此很容易引起大家对OpenCL程序可扩展能力的担忧。OpenCL的层次并发模型确保了在支持大量work-item执行的同时也能取得良好的可扩展性。在kernel执行期间，编程人员通过指定一个n维索引空间(NDRange)来声明需要新建多少个work-item。一个NDRange是work-item的1维、2维或3维索引空间，往往被映射到输入或输出数据的索引空间上。NDRange的维度被定义为一个数据类型为size_t、长度为N的数组，其中N代表用来描述所新建的work-item的维度。

在向量相加实例中，我们的数据是一维的，并假设有1024个元素，大小可以被定义为一个1维，2维或者3维的向量。用于为1024个元素指定一个NDRange的主机端代码如下：

 

 

可扩展性的获得来自于将NDRange的work-item划分成更小的部分，等同于改变workgroup的大小(图2.1)。N维度的索引空间需要指定同样具有N维度的workgroup。因此，一个3维的索引空间需要有3维的workgroup。

一个workgroup中各个work-item之间具有特定的关联：它们可以执行barrier操作进行同步，可以访问同一段共享的内存地址空间。由于workgroup的大小是固定的，并且通信量不大，所以并不影响大规模并发调度的可扩展性。

对于向量相加实例，workgroup的大小可能设置如下：

 

 

如果每个数组的work-item总数为1024，结果便是新建16个workgroup(1024 个work-item/(每个workgroup中64个work-item)= 16个workgroup)。注意，OpenCL要求

 

图2.1 work-item以NDRange的形式新建并被划分为多个workgroup

索引空间中每个维度的长度能被workgroup的大小整除。考虑到硬件效率，workgroup的规模通常固定为一个合适的大小，我们通过对每个维度索引空间的长度取整来满足这种不可分要求。在kernel代码中，可以在每维中指定额外的work-item使函数立即返回而不输出任何数据。

该向量相加程序中的work-item是独立运行的(甚至在同一个workgroup内)，OpenCL允许编程人员将传递的参数设置为NULL来忽略本地workgroup的大小而靠实现来自动生成。

转载于:https://www.cnblogs.com/lifan3a/articles/4615816.html