Kernel函数SPMD并行

Kernel函数block并行是一种经典的SPMD并行编程范式。简而言之，Single Program表明不同Kernel实例执行的是同一份二进制代码，Multiple Data表明每个Kernel实例处理不同的数据块，从而这些Kernel实例可以被分发到不同的物理核上去执行，并且通常情况下不感知Kernel实例调度执行的顺序。

__global__ [aicore] void foo(__gm__ int *Buffer, int SliceSize) {
  // 示意程序，往每个数据分片的第一个DWORD写一个整型值
  Buffer[block_idx * SliceSize] = block_idx;
}

每个Kernel实例拥有一个唯一的block_idx内建变量，由Runtime在加载Kernel函数的时候根据配置的BlockNum产生，其值为 {0，1，2，..., BlockNum - 1}。因此，每个Kernel实例拥有了自己的执行实例，除了用于索引每个block独立的数据分片，也可以控制kernel实例的执行路径。

__global__ [aicore] void foo() {
  if (block_idx % 2 == 0) {
    // even block_idx execute this path
  } else {
    // odd block_idx execute this path
  }
}

如图2所示，并行执行的AICORE核拥有完整独立的上下文，含独立PC、Register File等。