HALIDE

介绍

参考 https://halide-lang.org/ 和 http://stellar.mit.edu/S/course/6/sp15/6.815/courseMaterial/topics/topic2/lectureNotes/14_Halide_print/14_Halide_print.pdf。
Halide 的核心思想是把图像处理(可以理解为矩阵运算)的算法(需要计算什么内容)和调度(如何优化执行计算)分开。

WHY

大规模的矩阵运算性能优化空间很大，但是目前已有的人工优化和编译器优化都有一些问题。
人工优化，有两种范式，一种是针对具体的场景手工优化，一种是提供BLAS, IPP, MKL, OpenCV这类高度优化的库。前一种效率太低(场景众多，还要针对不同的后端硬件，优化工作量太大)，后者则只能提供局部最优的模块，无法在全局进行调度和融化优化。
编译器优化，可以看见完整的运算pipeline，但是优化的效果相对手动优化差了很多(就是一个最简单的矩阵乘法，编译器的输出都可能比手工优化要慢数倍)。同时，编译器中的很多核心优化决策都没有开放外部控制（比较典型的是，连循环展开的次数，GCC等编译器都是最近几年才通过#pragma unroll等方式提供了支持，更不要说直接控制cache block的大小等等），这导致在编译器的基础上人工再调优(纠正编译器的错误优化决策)很困难。
综合以上信息，高效和高质量的全局优化，还是要靠编译器。Halide的创造者也是沿着这个思路解决问题。

WHAT

Halide 是一种DSL(领域语言)，也是该DSL的编译器。
它的核心思想是把运算的逻辑和运算的过程分离。将运算过程剥离出后，再将各个典型优化决策的控制变量和控制逻辑暴露出来，以便人工或者黑盒优化算法(如遗传算法等)能持续调整优化决策，达到更好的性能。
它主页中的示例代码很好的说明了其思想，如下所示。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
Func blur_3x3(Func input) {
  Func blur_x, blur_y;
  Var x, y, xi, yi;

  // The algorithm - no storage or order
  blur_x(x, y) = (input(x-1, y) + input(x, y) + input(x+1, y))/3;
  blur_y(x, y) = (blur_x(x, y-1) + blur_x(x, y) + blur_x(x, y+1))/3;

  // The schedule - defines order, locality; implies storage
  blur_y.tile(x, y, xi, yi, 256, 32)
        .vectorize(xi, 8).parallel(y);
  blur_x.compute_at(blur_y, x).vectorize(x, 8);

  return blur_y;
}

上面代码将blur的运算逻辑，与具体运算的实施过程进行了分离。用很简单的几个方法就指定了tile/vectorize等重要优化，以及其对应的参数。直观看起来，代码很简洁，并且要修改优化的类型和参数，工作量也很小。
而要迫使编译器实现同样的优化，需要写下面一大段代码。
并且，这样的代码由于直接使用intel 的SIMD原语，可移植性大幅度下降。
更加严重的是，想要微调各项优化参数(为了针对不同的硬件做优化)，都需要对代码做大幅度的修改，工作量很大。
而上面的Halide代码，只需修改几个入参就可以完成优化决策的调整。哪怕Halide不提供内置的autotuner，使用一个简单的python脚本接入opentuner等黑盒优化框架也都会非常简单（毕竟只是改几个参数而已）。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
void box_filter_3x3(const Image &in, Image &blury) 
{
    __m128i one_third = _mm_set1_epi16(21846);
    #pragma omp parallel for
    for (int yTile = 0; yTile < in.height(); yTile += 32) {
        __m128i a, b, c, sum, avg;
        __m128iblurx[(256/8)*(32+2)]; // allocate tile blurx array
        for (int xTile = 0; xTile < in.width(); xTile += 256){
            __m128i *blurxPtr = blurx;
            for (int y = -1; y < 32+1; y++) {
                        const uint16_t *inPtr = &(in[yTile+y][xTile]);
                        for (int x = 0; x < 256; x += 8){
                             a = _mm_loadu_si128((__m128i*)(inPtr-1));
                             b = _mm_loadu_si128((__m128i*)(inPtr+1));
                             c = _mm_load_si128((__m128i*)(inPtr));
                             sum = _mm_add_epi16(_mm_add_epi16(a, b), c);
                             avg = _mm_mulhi_epi16(sum, one_third);
                             _mm_store_si128(blurxPtr++, avg);
                             inPtr += 8;
                             }}
                             blurxPtr = blurx;
                             for (int y = 0; y < 32; y++) {
                                 __m128i *outPtr = (__m128i *)(&(blury[yTile+y][xTile]));
                                 for (int x = 0; x < 256; x += 8) {
                                            a = _mm_load_si128(blurxPtr+(2*256)/8);
                                            b = _mm_load_si128(blurxPtr+256/8);
                                            c = _mm_load_si128(blurxPtr++);
                                            sum = _mm_add_epi16(_mm_add_epi16(a, b), c);
                                            avg = _mm_mulhi_epi16(sum, one_third);
                                            _mm_store_si128(outPtr++, avg);
                                            }}}}}

HOW

Halide介绍ppt中的一页示意图很好地展示了它的工作过程。

图中飘逸的寥寥几笔注释已经把核心的工作原理讲清楚了。
如果对编译技术比较熟悉，看完注释后最核心的几个疑问应该就豁然开朗了。
简单的说，Halide语言没有独立的语法定义，也就不需要独立的lexer和parser。
它利用了C++语言的元编程能力，直接构造出了Halide语言的中间表达IR。
具体情况，随后一节会有详细的分析。

语言实现分析

为了分析Halide编译器的具体实现，下载并编译了Halide的代码(使用Ubuntu18.04自带的clang/llvm8，按照官方命令编译，比较简单)。
然后编译、运行和调试其tutorial目录中的各个示例。
可以对其实现有一个大致的了解。

语言定义

从最简单的示例入手，理解整体概念更容易。
参考下面代码注释。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
//这个示例用Halide完成了一个矩阵
int main(int argc, char **argv) {
/*这里声明了三个核心概念
Func 是一系列运算(expr)的合集
Var 表达运算中涉及的变量
Expr 表达单个运算过程
*/
    Halide::Func gradient;
    Halide::Var x, y;
    Halide::Expr e = x + y;
//这里才完成了函数定义 f(x,y)  = x+ y
    gradient(x, y) = e;
/*
上面的声明式定义，其实已经体现出Halide是一个新的语言了。
有几个比较细节的点：
1 注意到我们没有对Var x和y进行赋值，就直接在expr中使用它们了。
可以这样做的原因是，它们只是对应二维数组的两个轴向而已，并不代表具体的值。
2 Halide::Expr e = x + y; 中的'='和'+'显然都不是常规语义。这一句实际上构建了一个Halide expr IR节点，op为+，LHS是x，RHS是y；
3 gradient(x, y) = e; 把expr关联到函数上，同样对应了IR上的操作。
*/

    //调用realize完成了编译和运行，并得到了结果
    Halide::Buffer<int32_t> output = gradient.realize(800, 600);
    //下面只是校验Halide和通常的运算结果一致
    for (int j = 0; j < output.height(); j++) {
        for (int i = 0; i < output.width(); i++) {
            if (output(i, j) != i + j) {
                printf("Something went wrong!\n"
                       "Pixel %d, %d was supposed to be %d, but instead it's %d\n",
                       i, j, i + j, output(i, j));
                return -1;
            }
        }
    }
    printf("Success!\n");
    return 0;
}

总结一下，Halide的核心概念就是Func、Var和Expr。它没有文本源代码的格式，直接是寄生在C++上。Func、Var和Expr都是C++的class。在完成声明和赋值的同时，利用对=、+和()的重载，完成了Halide的IR构建。

编译、调试和源码分析

编译准备

Halide的编译只需要依赖LLVM，在ubuntu18.04上安装llvm8就可以了。
clone下Halide代码后，进入目录后执行如下命令，可构建出带有调试信息的版本。

1
2
3
4
5
mkdir build
cd build
export CXXFLAGS="-O0 -g3"
export OPTIMIZE="-O0"
make -e -f ../Makefile -j 8

编译示例并调试

Halide 前端

在上一步构建完成的build目录中，继续执行如下命令。

1
2
3
cd distrib/tutorial/
g++ lesson_01*.cpp -g -I ../include -L ../bin -lHalide -lpthread -ldl -o lesson_01 -std=c++11 -g3
gdb ./lesson_01

这个lesson01就是前面语言定义一节中已经给出过的示例代码。
由于Function和Var的定义都没有传参，可以跳过，直接单步跟踪expr的赋值。这里给出的结果就非常典型了。
首先变量x和y被转换为了Expr。

1
2
3
4
15518	    /** A Var can be treated as an Expr of type Int(32) */
15519	    operator const Expr &() const {
15520	        return e;
15521	    }

然后Expr的operator+，就调用了Internal::Add::make构建了Expr这个IR。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
1139	Expr operator+(Expr a, Expr b) {
1140	    user_assert(a.defined() && b.defined()) << "operator+ of undefined Expr\n";
1141	    Internal::match_types(a, b);
1142	    return Internal::Add::make(std::move(a), std::move(b));
1143	}
--->
21	Expr Add::make(Expr a, Expr b) {
22	    internal_assert(a.defined()) << "Add of undefined\n";
23	    internal_assert(b.defined()) << "Add of undefined\n";
24	    internal_assert(a.type() == b.type()) << "Add of mismatched types\n";
25	
(gdb) 
26	    Add *node = new Add;
27	    node->type = a.type();
28	    node->a = std::move(a);
29	    node->b = std::move(b);
30	    return node;
31	}

调试到这里已经基本能确认Halide前端的工作原理了，通过operator重载，Halide直接构造了IR的，跳过了Lexer和Parser部分。
下一个问题是，Halide的IR如何，或者在什么时机进行Codegen。

Halide的代码生成流程

如前所述，Halide中通过realize方法完成了代码的编译和运行。接着调试上面程序的gradient.realize调用。
可以看到如下的调用链条。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
#0  Halide::Internal::lower (output_funcs=std::vector of length 1, capacity 1 = {...}, pipeline_name="f0", t=..., args=std::vector of length 1, capacity 1 = {...}, 
    linkage_type=Halide::LinkageType::ExternalPlusMetadata, requirements=std::vector of length 0, capacity 0, trace_pipeline=false, 
    custom_passes=std::vector of length 0, capacity 0) at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Lower.cpp:87
#1  0x00007ffff3cb7b6b in Halide::Pipeline::compile_to_module (this=0x7fffffffd920, args=std::vector of length 1, capacity 1 = {...}, fn_name="f0", target=..., 
    linkage_type=Halide::LinkageType::ExternalPlusMetadata) at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:506
#2  0x00007ffff3cb819b in Halide::Pipeline::compile_jit (this=0x7fffffffd920, target_arg=...)
    at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:573
#3  0x00007ffff3cbbda7 in Halide::Pipeline::realize (this=0x7fffffffd920, outputs=..., t=..., param_map=...)
    at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:1099
#4  0x00007ffff3cb98b0 in Halide::Pipeline::realize (this=0x7fffffffd920, sizes=std::vector of length 2, capacity 2 = {...}, target=..., param_map=...)
    at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:703
#5  0x00007ffff3ac078c in Halide::Func::realize (this=0x7fffffffdbf0, sizes=std::vector of length 0, capacity 0, target=..., param_map=...)
    at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Func.cpp:2922
#6  0x00007ffff3ac0a7d in Halide::Func::realize (this=0x7fffffffdbf0, x_size=800, y_size=600, target=..., param_map=...)
    at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Func.cpp:2937
#7  0x000055555555d56a in main (argc=1, argv=0x7fffffffdd98) at lesson_01_basics.cpp:78

看到lower，老司机应该已经心领神会找到门路了。一般lower意味着高层表达向硬件层级扩展，表达的内容将越发具体完整。

lower函数的过程，可以看到大致有两个主要的工作，第一个是补充最终程序需要的系列流程，如初始化环境，建立循环，已经插入一些等等；第二个是进行各项高层优化（优化越接近源码，执行起来越简单。）。但是lower部分看到结尾，仍然没有向另外一种IR或者机器指令转换。
从lower返回后，在compile_jit函数中继续向下调试，可以最终找到如下堆栈回溯中，Halide完成了IR到LLVM-IR的codegen过程(当然如果结合代码分析，查找LLVM的相关流程，找到这里会更快)。
#0 Halide::Internal::CodeGen_LLVM::compile (this=0x5555557b60e0, input=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/CodeGen_LLVM.cpp:637
#1 0x00007ffff399a42c in Halide::codegen_llvm (module=…, context=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/CodeGen_LLVM.cpp:46
#2 0x00007ffff3c326e1 in Halide::compile_module_to_llvm_module (module=…, context=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/LLVM_Output.cpp:381
#3 0x00007ffff3c13c91 in Halide::Internal::JITModule::JITModule (this=0x7fffffffbf90, m=…, fn=…, dependencies=std::vector of length 0, capacity 0)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/JITModule.cpp:251
#4 0x00007ffff3cb86dd in Halide::Pipeline::compile_jit (this=0x7fffffffd920, target_arg=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:607
#5 0x00007ffff3cbbda7 in Halide::Pipeline::realize (this=0x7fffffffd920, outputs=…, t=…, param_map=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:1099
#6 0x00007ffff3cb98b0 in Halide::Pipeline::realize (this=0x7fffffffd920, sizes=std::vector of length 2, capacity 2 = {…}, target=…, param_map=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Pipeline.cpp:703
#7 0x00007ffff3ac078c in Halide::Func::realize (this=0x7fffffffdbf0, sizes=std::vector of length 0, capacity 0, target=…, param_map=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Func.cpp:2922
#8 0x00007ffff3ac0a7d in Halide::Func::realize (this=0x7fffffffdbf0, x_size=800, y_size=600, target=…, param_map=…)
at /media/majiang/c6b38ac3-8b8a-4613-8259-dddbffe2f4cb/majiang/opensource/Halide/src/Func.cpp:2937
#9 0x000055555555d56a in main (argc=1, argv=0x7fffffffdd98) at lesson_01_basics.cpp:78
后面的流程更加直接一些，CodeGen_LLVM.cpp包含了主要的转换内容，compile_func中的 f.body.accept(this); 发起了LLVM-IR的发射动作。
后面就是CodeGen_LLVM.cpp中的一堆visit函数完成了针对不同类型Halide IR的LLVMIR代码生成。

遗留的学习

Halide自带的autotuner如何工作？

有意思的一些编程技巧

1 把可变参数的输入转成vector处理

1
2
3
4
5
6
template<typename... Args>
HALIDE_NO_USER_CODE_INLINE typename std::enable_if<Internal::all_are_convertible<Var, Args...>::value, FuncRef>::type
operator()(Args &&... args) const {
    std::vector<Var> collected_args{std::forward<Args>(args)...};
    return this->operator()(collected_args);
}

2 在父类中访问子类成员
https://en.wikipedia.org/wiki/Curiously_recurring_template_pattern

Typically, the base class template will take advantage of the fact that member function bodies (definitions) are not instantiated until long after their declarations, and will use members of the derived class within its own member functions, via the use of a cast; e.g.:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
template <class T> 
struct Base
{
    void interface()
    {
        // ...
        static_cast<T*>(this)->implementation();
        // ...
    }

    static void static_func()
    {
        // ...
        T::static_sub_func();
        // ...
    }
};

struct Derived : Base<Derived>
{
    void implementation();
    static void static_sub_func();
};