C++ AMP(C++ Accelerate Massive Parallelism )是微软提出异构计算开放标准。C++ AMP在CUDA已经形成气候,和OpenCL风生水起的背景下问世,无疑为异构计算和GPGPU领域增添了新的解决方案,也为CUDA和OCL带来新的挑战。作为操作系统以及其派生应用的提供者微软,以C++ AMP为其新的利器,以简洁的编程接口,快速的入门(有一定并行开发基础)以及强大的编译器与调试工具的(visual studio 2012)支持,相信可以吸引到不少在CUDA和OCL之间徘徊,甚至是已经在CUDA和OCL领域进行多年开发的程序员的广泛关注。
C++ AMP支持的操作系统是win7,win server 2008 R2。支持C++ AMP的显卡需要支持DirectX 11 /多核CPU应支持SIMD指令
实现两个一维数组相加的C++ AMP的代码看上去类似这样:
#include <amp.h>
... ...
float * arrayA = new float[N];
float * arrayB = new float[N];
float * arrayC = new float[N];
initialize(arrayA);
initialize(arrayB);
concurrency::array_view<float ,1> a(N,arrayA);
concurrency::array_view<float ,1> b(N,arrayB);
concurrency::array_view<float ,1> c(N,arrayC);
concurrency::parallel_for_each(c.extent,
[&](concurrency::index<1> i) restrict(amp)
{
c[i] = a[i] + b[i];
}
);
... ...
这里简单解释下代码:
amp.h包含了所有并行相关模板类的定义,需要包含进来。
concurrency是定义了并发相关类的命名空间。concurrency::array_view就是其中一个模板类,这个类可以简单地理解为内存对象,存储参与并发的数据。这里每个array_view就存储了host端的数组。
concurrency::parallel_for_each函数可以认为是执行kernel的函数,在ocl中相当于函数clEnqueueNDRangeKernel
上式中红色部分就是并行部分的核心代码,类似于OCL中的kernel.这是一个lamda表达式,简洁地表示一个函数对象。[&]表示以引用方式引用函数体内出现的a,b,c;
这个函数对象的参数concurrency::index<1> i 表示一个1维序号,实际上对应的是线程id. c.extent等价于N,表示global线程维度。
由这个代码片段不难看出,C++ AMP的最大特色就是以类似STL的编程方式来实现GPGPU编程,这是CUDA与OCL难以比拟的优势。尤其是与OCL相比,省去了建立环境的大量流程,让程序员更多地关注算法并行化本身。从C++ AMP提供的基本类,可以将其类比到CUDA与OCL的概念。所以有了OCL基础,入门C++ AMP是相当easy的,要做的功课主要是熟悉kernel的写法以及新的模板类的使用。
这里结合本例做一个类比:
(1). array_view ----> ocl memory object( buffer)
(2).parallel_for_each-----> clEnqueueNDRangeKernel
(3) function object by lamda expression ------> OCL kernel
(4). index<ND> -----> global id of thread
(5) tile -----> workgroup
(6) accelerator----> cl_device_id
关于5-6将会在以后的实例展开。
综上, C++ AMP的主要优势体现在:接口简洁;代码量小;对于C++程序员上手容易;VS直接编译... ...