Skip to content

Latest commit

 

History

History
49 lines (36 loc) · 1.55 KB

File metadata and controls

49 lines (36 loc) · 1.55 KB

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

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

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

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

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

{%ace edit=false, lang='c_cpp'%} _kernel void ZZ6vecAddPfS_S_iEN3__019__cxxamp__trampolineEiiS_N11Concurrency11access_typeEiiS_S2_iiS_S2( __global float *llvm_cbe_tmp__1, unsigned int llvm_cbe_tmp__2, __global float *llvm_cbe_tmp__3, unsigned int llvm_cbe_tmp__4, __global float *llvm_cbe_tmp__5, unsigned int llvm_cbe_tmp__6){

unsigned int llvm_cbe_tmp__7; float llvm_cbe_tmp__8; float llvm_cbe_tmp__9; llvm_cbe_tmp__7 = /tail/get_global_id(0u); llvm_cbe_tmp__10 = *((&llvm_cbe_tmp__1[((signed int)llvm_cbe_tmp__7)])); llvm_cbe_tmp__11 = *((&llvm_cbe_tmp__3[((signed int)llvm_cbe_tmp__7)])); *((&llvm_cbe_tmp__5)[((signed int)llvm_cbe_tmp__7)])) = (((float)(llvm_cbe_tmp__10 + llvm_cbe_tmp__11))); return; } {%endace%}

图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]