11.6 编译之后的C++ AMP代码

我们再回顾一下C++ AMP以Lambda函数实现的向量相加，如图11.6所示。

  1. [=](index<1> idx) restrict(amp){cv[idx]=av[idx]+bv[idx];}

图11.6 C++ AMPLambda函数的向量相加

使用CLamp将其翻译成OpenCL内核代码时，就如图11.7所示。

  1. __kernel void
  2. ZZ6vecAddPfS_S_iEN3__019__cxxamp__trampolineEiiS_N11Concurrency11access_typeEiiS_S2_iiS_S2_(
  3.   __global float *llvm_cbe_tmp__1,
  4.   unsigned int llvm_cbe_tmp__2,
  5.   __global float *llvm_cbe_tmp__3,
  6.   unsigned int llvm_cbe_tmp__4,
  7.   __global float *llvm_cbe_tmp__5,
  8.   unsigned int llvm_cbe_tmp__6){
  10.   unsigned int llvm_cbe_tmp__7;
  11.   float llvm_cbe_tmp__8;
  12.   float llvm_cbe_tmp__9;
  13.   llvm_cbe_tmp__7 = /*tail*/get_global_id(0u);
  14.   llvm_cbe_tmp__10 = *((&llvm_cbe_tmp__1[((signed int)llvm_cbe_tmp__7)]));
  15.   llvm_cbe_tmp__11 = *((&llvm_cbe_tmp__3[((signed int)llvm_cbe_tmp__7)]));
  16.   *((&llvm_cbe_tmp__5)[((signed int)llvm_cbe_tmp__7)])) = (((float)(llvm_cbe_tmp__10 + llvm_cbe_tmp__11)));
  17.   return;
  18. }

图11.7 OpenCL SPIR版的向量相加

看起来编译器处理过之后的代码可读性差很多，不过这里我们依旧不难找到下面一些对关系：

  • 第1行：生成对应的内核名
  • 第2-3行：序列化array_view va
  • 第4-5行：序列化array_view vb
  • 第6-7行：序列化array_view vc
  • 第11行：获取全局工作项索引，通过C++ AMPLambda函数中的idx获取
  • 第12行：加载va[idx]
  • 第13行：加载vb[idx]
  • 第14行：计算va[idx]+vb[idx]，并将结果保存在vc[idx]