11.8 编译器怎样支持C++ AMP的线程块划分

线程划分属于GPU优化中的一个技巧。根据抽象的等级，编程模型可以显式的或隐式的支持线程划分。隐式的方式可能会自动的减少在一个线程束中对内存的访问，并且通过透明或半透明的线程划分，来达到最佳内存形式，并且获取最佳的性能。与之相反，显式的方式需要用户显式定义不同的内存对象(可能有些内存在片上共享，有些内存是离散的)，相关的数据移动也由用户进行控制。C++ AMP、CUDA和OpenCL都属于显式编程模型。剩下的内容，我们将一个编译器作者的角度，来考虑如何显式的支持线程划分。

对于支持线程划分的显式编程模型，我们通常都可以看到如下的特点：

  • 通过某种方式将计算域划分成固定大小的小块
  • 通过显式的方式进行数据内存的指定，通常是片上、离散或线程私有这几种形式。其在OpenCL中的对应为local, global和__private
  • 为固定大小的计算块提供同步机制，以便于不同工作项之间的协同工作

对于GPU有些了解的读者，可能对于C++ AMP不是很了解。C++ AMP中需要使用extent类中尺寸描述符，对计算区域的大小进行确定。另外，tile_extent描述了和对计算区域进行分块。其划分方式与OpenCL类似。

11.8.1 划分计算域

C++ AMP中有extent类中一个成员函数名为“tile”，用来计算tile_extent。该模板函数的参数就说明了划分区域的大小。现在C++ AMP就与OpenCL不太相同了，C++ AMP中的划分方式静态参数化的。

为了让库和编译器了解划分区域的信息，我们使用了一个带有一些参数Lambda内核)，将一些参数一次转换为tiled_index。tiled_index与OpenCL中的get_gloab_id(), get_local_id()和get_group_id()所获取的值一致。

  1. void mxm_amp_tiled(
  2.   int M, int N, int W,
  3.   const std::vector<float> &va,
  4.   const std::vector<float> &vb,
  5.   std::vector<float> &result){
  7.   extent<2> e_a(M, N), e_b(N, W), e_c(M, W);
  9.   array_view<const float, 2> av_a(e_a, va);
  10.   array_view<const float, 2> av_b(e_b, vb);
  11.   array_view<float, 2> av_c(e_c, result);
  13.   extent<2> compute_domain(e_c);
  14.   parallel_for_each(compute_domain.tile<TILE_SIZE, TILE_SIZE>(),
  15.     [=]($tiled_index<TILE_SIZE, TILE_SIZE> tidx) restrict(amp) {
  16.       mxm_amp_kernel(tidx, av_a, av_b, av_c);
  17.     });
  18. }

11.8.2 指定地址空间和栅栏

C++ AMP中的内核函数中，使用tile_static限定符用来声明对应内存对象在片上内存上分配(OpenCL中对应的是本地内存)。为了强制C++ AMP划分块中的线程同步，C++ AMP为tile_static对象提供了barrier.wait函数。和OpenCL中的概念一样，当该函数调用时，所有线程都要在同步调用点处停止。

C++ AMP和OpenCL有个有趣的区别在于地址空间处。OpenCL中，其是指针类型的一种，其指针使用local进行声明(不能使用private对一段内存进行声明)。C++ AMP中地址空间时指针值的一部分。可以使用通用指针进行指定：

  1. float *foo;

指针foo可以指向一个使用tile_static创建的内存(与OpenCL中的__local等价)，因为一定的局限性[1]同一个指针，只能指向全局内存中的一个值。

我们可以尝试在定义一个C++ AMP中tilestatic的宏，用以宏扩展Clang/LLVM中的`_attribute((address_space()))`限定符。其作为嵌入式C的一种扩展，将类型作为指针和内存对象类型的一部分。不过，在下面的代码片段中，这种方式可能无法产生foo指针正确的地址空间信息：

  1. tile_static float bar;
  2. float *foo = &bar;

因为我们没有办法嵌入地址空间描述符，让其作为指针类型的一部分，不过我们需要解决这种形式的使用，需要将地址空间信息作为变量定义的一部分。编译器使用模板的方式，显然已经不能对这种使用方式进行区分。

一种替代的方法是，通过将地址空间以变量属性方式对这些特殊的变量进行标记(而不是其类型定义的一部分)。这种属性可以作为编译器的一种扩展，可以用来指定二进制文件中对应数据段的变量定义。这里的属性是与定义的变量相关，与这个变量的类型无关：比如有两个整型书在不同的数据段，并且有指针可以指向不会出现类型错误的那个整型数。在CLamp中，我们使用的方法是——通过数据流分析进行简单映射，从而减少对地址空间信息的依赖，不过这些代码在C++中依旧是合法的：

  • 定义C++ AMP中的tile_static作为变量属性。
  • 所有指针进行了初始化，不过没有进行地址空间的指定。
  • 基于静态单赋值分析的方式，减少指向变量的属性。

这里的分析只是解决了众多问题中比较简单的，有些问题使用这种方式解决会出现不可预料的结果。下一节我们将仔细说明一下，如何进行地址空间的推断。

[1] 在C++ AMP1.2中，这种限制是为了能让编译器能够更好的推断相应内存对象的地址空间信息