latency & bandwidth
首先从车道的角度理解数据传输: 在一个高速道路上, 限制了车辆的速度, 怎么能提高传输效率?
一种方式是增加车道的数量, 我们也有另一种方式: 不等待前车跑完整条道路, 而是保持一定距离后让后车进入道路, 这样也能提高道路的运输效率
两种方式都没有改变: 一辆车跑完整条道路的时间
从硬件的角度来理解这件事:
Memory bandwidth : 内存系统向处理器传输数据的速度, 例如 20 GB/s
那么增加线程数能一直提高效率吗?
而对于计算机的指令串, 通常各个指令的执行时间是不同的 我们考虑一个多线程处理器核心 对于每个线程, 执行算子的操作是很快的, 但是从内存获得数据需要很长时间。处理器可以让多个 memory requests 同时处于进行中, 多线程也可以把一个线程的等待与另一个线程的执行重叠起来；如果下一条指令需要的数据还没有准备好, 核心就会出现 stall
增加线程可以隐藏 latency 并提高吞吐量, 直到最慢的模块达到饱和；当内存带宽已经被完全利用时, 再增加线程也无法提高稳态吞吐量 进一步分析, 我们发现处理器的大部分时间其实都消耗在了 stall 上
对于稳态的情形, 核心的利用率应该是被内存的吞吐率限制, 而不是单个数据的传输延迟
This computation is bandwidth limited!
对于现代计算, bandwidth 是核心的资源, 对于一个出色的并行程序, 应当
- 多重复使用在同一线程中加载过的数据
- 在线程之间共享数据
- 更偏向执行多余的计算, 而不是储存再获取值 (执行算子的消耗时间短得多)
程序必须尽可能低频率地获取内存, 来高效地利用现代处理器
Abstraction vs. implementation
我们需要将抽象层面和具体操作层面分开: 程序的:
- 抽象(semantics): 现代编程语言的操作的含义
- 具体(implementation / scheduling): 在并行机器上, 这些操作被怎么执行?
我们的目标: 给定程序和我们已知的对于并行编程的具体实现, 我们可以去追踪并行计算机每一步的具体操作
一个例子是 ISPC (Intel SPMD Program Compiler)
SPMD: Single program , Multiple data
回忆我们前面用泰勒级数计算正弦值的例子, 在 cpp 的语言中, 我们写了 sinx 这个函数, 并在 main 中进行调用
在 ISPC 的代码中, 这是类似的
export void ispc_sinx(
uniform int N,
uniform int terms,
uniform float* x,
uniform float* result)
{
// assume N % programCount == 0
for (uniform int i = 0; i < N; i += programCount)
{
int idx = i + programIndex;
float value = x[idx];
float numer = x[idx] * x[idx] * x[idx];
uniform int denom = 6; // 3!
uniform int sign = -1;
for (uniform int j = 1; j <= terms; j++)
{
value += sign * numer / denom;
numer *= x[idx] * x[idx];
denom *= (2*j + 2) * (2*j + 3);
sign *= -1;
}
result[idx] = value;
}
}programCount 和 programIndex 是 ISPC 的内置变量: 前者表示一个 gang 中并行运行的程序实例数, 后者表示当前逻辑实例的编号；它们不是从 C++ 调用时传入的参数
一个 SPMD 程序的抽象:
- 调用 ISPC 程序生成一系列 ISPC 程序实例
- 所有实例同时运行同一个 ISPC 代码
- 每个实例具有他们各自的局域变量的副本
- 在返回函数值后, 所有实例执行完毕
而在 main 中调用 ISPC 程序是直接的
#include "sinx_ispc.h"
int main(int argc, char** argv) {
int N = 1024;
int terms = 5;
float* x = new float[N];
float* result = new float[N];
// initialize x here
// execute ISPC code
ispc_sinx(N, terms, x, result);
return 0;
}在这个例子中, 假设编译目标使用 8-wide SIMD, 因而 programCount = 8, 8 个逻辑实例的编号用 programIndex 标记；具体数值取决于编译目标
uniform 是 ISPC 的类型修饰, 表示一个 gang 中的所有程序实例观察到同一个值；对于这里的 N、terms 和指针, 它表达了这些值在实例之间不变, 也帮助编译器优化
具体流程是以下抽象
C/C++ main()
│
│ sequential
▼
ispc_sinx()
│
├── instance 0 ───▶
├── instance 1 ───▶
├── instance 2 ───▶
├── instance 3 ───▶
├── instance 4 ───▶
├── instance 5 ───▶
├── instance 6 ───▶
└── instance 7 ───▶
concurrently
│
│ all complete
▼
return to C/C++
│
▼
sequential execution当然, 上一个 ISPC 程序内部的实现只是一种手写的结果, 程序的运行的例子是程序实例 0 计算项 0, 8, 16 … , 按照 8 的顺序 当然也可以有更多 ISPC 内部实现, 比如程序实例 0 计算 0, 1, 2 这三项…
当然, ISPC 提供了一种更方便的语言抽象 foreach , 不必需要手写 ISPC 逻辑来分配每个程序实例的计算对象, 而让 ISPC 自己分配
export void ispc_sinx(
uniform int N,
uniform int terms,
uniform float* x,
uniform float* result)
{
foreach (i = 0 ... N)
{
float value = x[i];
float numer = x[i] * x[i] * x[i];
uniform int denom = 6; // 3!
uniform int sign = -1;
for (uniform int j = 1; j <= terms; j++)
{
value += sign * numer / denom;
numer *= x[i] * x[i];
denom *= (2*j + 2) * (2*j + 3);
sign *= -1;
}
result[i] = value;
}
}foreach 告诉编译器下面的循环需要用一系列的程序实例并行, 而 ISPC 的具体实现会自动把任务分配给不同的程序实例
在一些简单的问题, foreach 可以让我们把问题变成一个和串行问题几乎一致的问题
foreach 的语义要求各次迭代彼此独立, 这是程序员对编译器作出的保证。编译器通常无法证明这种独立性；如果不同迭代之间存在冲突的内存访问, 程序仍可能编译, 但输出可能是未定义的
先看下面的例子:
// ISPC code:
export void shift_negative(
uniform int N,
uniform float* x,
uniform float* y)
{
foreach (i = 0 ... N)
{
if (i >= 1 && x[i] < 0)
y[i-1] = x[i];
else
y[i] = x[i];
}
}这个程序的输出不是良定义的, 因为不同的 foreach 迭代可能写入同一个内存位置。例如某次迭代写 y[i-1], 另一次迭代可能同时写同一个位置；这与指针本身被声明为 uniform 无关
同样, 另一个例子试图让多个实例共同更新一个 uniform 变量:
export uniform float sum_incorrect_2(
uniform int N,
uniform float* x)
{
uniform float sum = 0.0f;
foreach (i = 0 ... N)
{
sum += x[i];
}
return sum;
}这里 x[i] 在不同程序实例中具有不同的值, 是 varying value；但 sum 是整个 gang 共享一个值的 uniform 变量。把多个不同的 varying values 同时赋给一个 uniform value 没有唯一含义, 因此 ISPC 会在编译期报告类型错误
那么如何正确处理? 一种处理方式是每个程序实例维护自己的 partial_sum , 再用 ISPC 自己的 reduce 函数把各自的 partial sum 加起来, 即
export uniform float sum_array(
uniform int N,
uniform float* x)
{
uniform float sum;
float partial = 0.0f;
foreach (i = 0 ... N)
{
partial += x[i];
}
sum = reduce_add(partial);
return sum;
}当然我们也有 AVX 向量化的实现
#include <immintrin.h>
float sum_summary_AVX(int N, float* x)
{
float tmp[8];
__m256 partial = _mm256_setzero_ps();
for (int i = 0; i < N; i += 8)
partial = _mm256_add_ps(partial, _mm256_loadu_ps(&x[i]));
_mm256_storeu_ps(tmp, partial);
float sum = 0.f;
for (int i = 0; i < 8; i++)
sum += tmp[i];
return sum;
}这里为了突出 reduction 与 SIMD lanes 的对应, 假设 N 是 8 的倍数；实际代码还需要处理不足 8 个元素的尾部
可以看到, ISPC 程序实例与 AVX 的 SIMD lane 之间其实在这个例子中是对应的
ISPC:
instance 0 ─┐
instance 1 ─┤
... ├─> varying partial
instance 7 ─┘
|
reduce_add
|
v
uniform sum当然 ISPC 提供了很多跨程序实例的操作
uniform int64 reduce_add(int32 x);
// 计算各程序实例中某变量的和
uniform int32 reduce_min(int32 a);
// 计算各程序实例中某变量的最小值
int32 broadcast(int32 value, uniform int index);
// 将某程序实例的某个变量值广播到所有程序实例上
int32 rotate(int32 value, uniform int offset);
// 对所有 i, 把实例 i 的值传递到 (i+offset) % programCount 上Summary
ISPC: abstraction vs. implementation
抽象上是 SPMD, 具体实现是 SIMD : ISPC 将 SPMD 程序映射到 SIMD 的向量 lane 上的操作, 对于条件控制类型的判断, 也会像前面一样设置 mask 来实现
ISPC 抽象: SPMD 抽象 + uniform values
Single thread of control
|
v
Call SPMD function
|
+------------------+
| ↓ ↓ ↓ ↓ ↓ ↓ ↓ |
| ↓ ↓ ↓ ↓ ↓ ↓ ↓ | SPMD execution
| | multiple logical instances
+------------------+
|
SPMD function returns
|
v
Resume single thread
of control不使用 task 时, 一个基本的 ISPC gang 由 CPU 上一个线程中的 SIMD 指令实现, 所以上面的代码实际上只会占用一颗核心
ISPC 也有针对多核的设计, 被称为 “task” , 在 ass1 中会看到应用
ISPC 语言虽然看起来是一个抽象的并行语言, 可以让我们把 foreach 内部的程序看作一个个独立运行的程序, 但是实际上也暴露了一些底层细节: programIndex 和 programCount 让程序员看到 gang 的组织方式, 标准库中的 reduction、broadcast 和 rotate 等函数则提供跨程序实例的通信
这也意味着我们可能写出一些危险的操作, 比如之前的一些错误例子
借助 ISPC , 我们也可以实现一些高级操作
当然, 我们在抽象层面可以一直抬高, 最终只看到类似在 numpy 或者 torch 中看到的纯向量化操作, 它们的抽象层级是这样的
低层
↓
手动 thread / SIMD intrinsic
↓
ISPC programIndex / programCount
↓
ISPC foreach
↓
map / array / tensor operations
↑
高层Summary
- 编程模型提供了一种组织并行程序的方式
- 它们提供了大量有效的并行计算的 implementation 的 abstraction
- 总是去思考 abstraction vs. implementation 是有益的