第11章 高级语言映射到OpenCL2.0 —— 从编译器作者的角度 - 11.9 地址空间的推断

上节中声明的每个OpenCL变量都具有自己地址空间限定符，用来区分这个变量是在哪端内存区域上分配的。地址空间对于OpenCL来说是十分中要的特性。将数据放入不同的内存区域，OpenCL程序在获得高性能的同时，保证了数据一致性。这个特性通常不会出现一些比较高级的语言中，比如C++ AMP。高级语言将数据放入通用地址空间内，从而就不用显式的说明这些内存是在哪里开辟的。OpenCL中声明的变量如果没有限定符，那么默认在私有内存上进行分配，这就违反了C++ AMP中的既定行为。举个例子，如果将tile_static的声明限定于私有，那么这个对象上的数据将不会与其他工作项共享，并且计算得到的结果是错误的。为了解决这个矛盾，就需要为每个声明和内存访问添加正确的地址空间信息。

CLamp中，生成OpenCL位码之后，需要在通过一次LLVM的转换，为相应的变量声明添加上正确的地址空间信息。理论上每个声明进行地址空间的推断是不可行的，因为分析器看不到整个程序，所以无法判断哪些内核要和哪些变量进行交互。不过，实际使用的程序中，推断地址空间是可行的。

array和array_view的实现都为推断地址空间提供着线索。C++ AMP中，只有通过array和array_view才能将大量的数据传入内核。C++ AMP运行时为内核的参数列表预留了指针。内核在使用这些数据时，只需要访问相关的指针即可。这些指针都会描述成全局的，因为这些数据时要每个工作项都可见的。推断过程的依据就是内核函数的参数列表，相关指针限定为全局，并且通过这些指针对更新所有内存操作。

tile_static数据的声明不能通过模式分析进行判别，所以CLamp的前端编译器要保存这些声明。当前的CLamp实现中，限定符声明tile_static的部分，使用一段特殊的位码进行表示。推断过程会将tile_static属性传递给任意一个指针，这些指针能获取这些变量的地址，然后将其添加到对应的OpenCL声明中。

我们看一个简单的C++ AMP实例，通过这个实例我们来了解转换是如何进行的：

void mm_kernel(int *p, int n){
  tile_static int tmp[30];
  int id = get_global_id(0);
  tmp[id] = 5566;
  barrier(0);
  p[id] = tmp[id];
}

通过CLamp初始化之后，代码将完全转化成LLVM IR。这个阶段中，地址空间是缺失的，并且这段代码会产生一个不正确的结果。注意变量tmp会放在一个特殊的ELF字段中(“clamp_opencl_local”):

@mm_kernel.tmp = internal unnamed_addr global[30xi32] zeroinitializer, align 16, section "clamp_opencl_local"
define void @mm_kernel(i32 *nocapture %p, i32 %n){
  %1 = tail call i32 bitcast (i32(...)* @get_global_id to i32 (i32)*)(i32 0)
  %2 = sext i32 % to i64
  %3 = getelementptr inbounds[30 x i32]* @mm_kernel.tmp, i64 0, i64 %2
  %4 = tail call i32 bitcast (i32(...)* @barrier to i32(i32)*)(i32 0) #2
  %5 = load i32 *%3, align 4, !tbaa!1
  %6 = getelementptr inbounds i32* %p, i64 %2
  store i32 %5, i32 * %6, align 4, !tbaa !1
  ret void
}

CLamp分析完成后，正确的地址空间信息就添加到对应的声明中去(mm_kernel.tmp中的一些内存操作)。正确的LLVM IR如下所示：

@mm_kernel.tmp = internal addrspace(3) unnamed_addr global[30xi32] zeroinitializer, align 4
define void @mm_kernel(i32 addrspace(1)*nocapture %p, i32 %n){
  %1 = tail call i32 bitcast(i32 (...)* @get_global_id to i32(i32)*)(i32 0)
  %2 = getelementptr inbounds[30 x i32] addrspace(3)* @mm_kenrel.tmp, i32 0, i32 %1 store i32 5566, i32 addrspace(3) %2, align4, !tbaa!2
  %3 = tail call i32 bitcast (i32(...) * @barrier to i32(i32)*)(i32 0)
  %4 = load i32 addrspace(3)* %2, align4, !tbaa !2
  %5 = getelementptr inbounds i32 addrspace(1) * %p, i32 %1
  store i32 %4, i32 addrspace(1)* %5, align 4, !tbaa !2
  ret void
}