第11章 高级语言映射到OpenCL2.0 —— 从编译器作者的角度 - 11.6 编译之后的C++ AMP代码

优质
小牛编辑
125浏览
2023-12-01

我们再回顾一下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所示。

{%ace edit=false, lang=’ccpp’%}
kernel void
ZZ6vecAddPfS_S_iEN3
019cxxamptrampolineEiiS_N11Concurrency11access_typeEiiS_S2_iiS_S2
(
global float *llvm_cbe_tmp1,
unsigned int llvm_cbe_tmp2, global float llvm_cbe_tmp3,
unsigned int llvm_cbe_tmp
4,
__global float
llvm_cbe_tmp5,
unsigned int llvm_cbe_tmp
6){

unsigned int llvm_cbe_tmp7;
float llvm_cbe_tmp
8;
float llvm_cbe_tmp9;
llvm_cbe_tmp
7 = /tail/get_global_id(0u);
llvm_cbe_tmp10 = *((&llvm_cbe_tmp1[((signed int)llvm_cbe_tmp7)]));
llvm_cbe_tmp
11 = ((&llvm_cbe_tmp3[((signed int)llvm_cbe_tmp7)])); ((&llvm_cbe_tmp5)[((signed int)llvm_cbe_tmp7)])) = (((float)(llvm_cbe_tmp10 + llvm_cbe_tmp11)));
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]