8.1.C++ AMP简介
C++ AMP是专为设计支持C++的异构并行模型.
全程是: Accelerator Massive Parallelism
下面是一个Vector C++ AMP的代码,通过这段代码来解释C++ AMP的语法:
#include <amp.h> using namespace concurrency; void vecAdd(float* A, float* B, float* C, int n) { array_view<float,1> AV(n,A), BV(n,B); array_view<float,1> CV(n,C); CV.discard_data(); parallel_for_each(CV.get_extent(), [=](index<1> i) restrict(amp) { CV[i] = AV[i] + BV[i]; }); CV.synchronize(); }
1. 头文件amp.h 提供需要用到的函数等的声明
2. using namespace concurrency 作用域包含AMP的函数和类.
3. 模板array_view<T, D> AV<size, name>, T表示数据类型, D表示dimension维度, size表示变量大小,name表示变量名.这个模板的作用就是提供对C++数组的引用,提供一种新的方式访问现有的数组. 可以看出变量 A 和 B是一维的.
4. CV.discard_data(); 表示当前数组C中的数据值是无关紧要的,我理解为可以被覆盖.
AMP kernel函数
parallel_for_each(CV.get_extent(), [=](index<1> i)
restrict(amp) {
CV[i] = AV[i] + BV[i];
});
1. parallel_for_each指定一个计算应用到指定的数据集合中,
2. CV.get_extent(): 就是指定的数据集合,当前实例中,CV是一维的,数据集合的范围是[0,n-1].
3. 第二个参数: [=](index<1> i, 这是一个C++的lambda 表达式:
"=" 这个我理解的是: 把 i索引的数据拷贝到device空间中供kernel函数使用,这个有待进一步考证
关于index<1> i 表示一维的索引,i 表示的是threadIdx.x全局的线程id. 关于这个 index的用法,下面有几个示例有助于理解:
eg1: 这个表示一维索引第2个元素,索引从0开始算起
int aCPP[] = {1, 2, 3, 4, 5}; array_view<int, 1> a(5, aCPP); index<1> idx(2); std::cout << a[idx] << "\n"; // Output: 3
eg2: 这个表示二维索引,idx中的第一个元素表示行,第二个元素表示列,所以这个索引的是第二行第三列
int aCPP[] = {1, 2, 3, 4, 5, 6}; array_view<int, 2> a(2, 3, aCPP); index<2> idx(1, 2); std::cout << a[idx] << "\n"; // Output: 6
eg3:
int aCPP[] = { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12}; array_view<int, 3> a(2, 3, 4, aCPP); // Specifies the element at 3, 1, 0. index<3> idx(0, 1, 3); std::cout << a[idx] << "\n"; // Output: 8
4. restrict(amp) 类似cuda中的"__device__" 表示运行在device上的kernel函数.
5. CV.synchronize(); 同步,等待kernel函数执行完成.
关于C++ AMP的更多信息可以参考这个网址:
https://msdn.microsoft.com/zh-cn/library/hh265137.aspx