异构编程

基于昇腾多核并行架构，CCE编程提供了显式Host+Device异构编程，Kernel函数是一种SPMD（Single Program Multiple Data）的并行编程模型，直观地映射了芯片的多核并行能力。针对昇腾DSA（Domain Specific Architecture）加速指令，提供底层intrinsic接口访问硬件加速能力。

CCE异构编程通过最小化扩展原则，基于C++语言的扩展，将达芬奇原生机器模型和编程模型映射到高级语言。CCE异构编程语法主要扩展包括异构函数执行空间，地址空间，异构调用语法三部分；并行编程模型扩展主要包括编译器内建变量，如block_idx, block_dim等。完整的CCE异构编程需配套使用ACL运行时库，ACL运行时API请参考《应用开发接口》文档中的运行时管理章节。

异构编程允许源码文件同时包含运行于主机侧和设备侧的执行代码，但是通常设备侧芯片主要用于特定领域的加速计算，因此通常和主机侧具备完全不同的微架构和指令集。异构编程通过显式的函数执行空间区分主机代码和设备代码，并基于异构函数语义隐藏复杂的Host和Device交互ABI。

// QuickStartDemoACL.cce
#include "acl/acl.h"
#include <stdio.h>
#include <stdlib.h>

#define BLOCKS 4
#define CACHELINE_SZ 64

// Define a kernel
__global__[aicore] void foo(__gm__ uint8_t *Out, int Stride) {
  Out[block_idx * Stride] = block_idx;
}

int main(int argc, char *argv[]) {
  aclInit(nullptr);
  aclrtSetDevice(0);
  aclrtStream stream;
  aclrtCreateStream(&stream);

  uint8_t ExpectedValue[] = {0, 1, 2, 3};
  uint8_t *OutputValue = nullptr;
  aclrtMalloc((void **)&OutputValue, BLOCKS, ACL_MEM_MALLOC_NORMAL_ONLY);

  uint8_t InitValue[BLOCKS] = {0};
  aclrtMemcpyAsync((void *)OutputValue, sizeof(InitValue), InitValue,
                   sizeof(InitValue), ACL_MEMCPY_HOST_TO_DEVICE, stream);
  aclrtSynchronizeStream(stream);

  // Invoke a kernel, with BLOCKS number of logical blocks
  foo<<<BLOCKS, nullptr, stream>>>(OutputValue, CACHELINE_SZ);

  uint8_t *OutHost = nullptr;
  aclrtMallocHost((void **)&OutHost, BLOCKS * CACHELINE_SZ);
  aclrtMemcpyAsync(OutHost, BLOCKS * CACHELINE_SZ, OutputValue,
                   BLOCKS * CACHELINE_SZ, ACL_MEMCPY_DEVICE_TO_HOST, stream);
  aclrtSynchronizeStream(stream);

  for (int I = 0; I < sizeof(ExpectedValue) / sizeof(uint8_t); I++) {
    printf("i%d\t Expect: 0x%04x\t\t\t\tResult: 0x%04x\n", I, ExpectedValue[I],
           OutHost[I * CACHELINE_SZ]);
  }

  aclrtFreeHost(OutHost);
  aclrtFree(OutputValue);

  aclrtDestroyStream(stream);
  aclrtResetDevice(0);
  aclFinalize();
  return 0;
}

如上用例，CCE程序通过 __global__关键字标识设备侧入口函数（Kernel函数），通过<<<>>>异构调用语法执行Kernel函数。<<<BlockNum, SmDesc, Stream>>>总共包含三个可配置参数，其中第一个参数指定Kernel函数运行的实例化份数，第二个参数用于配置片上L2缓存的使用，第三个参数用于具体绑定执行该Kernel的Stream队列。

CCE异构程序的具体运行模式如上图所示，主要步骤包括：

1. 主机侧将Kernel参数经由Runtime以特定方式传递到设备侧内存。
2. 主机侧将Kernel任务通过Runtime提交到运行队列stream中。
3. 主机侧API返回。
4. 设备侧获取待运行kernel任务信息，关键信息为BlockNum、Kernel代码基址、Kernel参数基址等。
5. AICore硬件调度器根据BlockNum实例化具体任务给每一个AICore, 分配实例Id: block_idx。
6. 调度器将Kernel具体运行所需的信息配置给待执行任务的空闲AICore核。
7. AICore核心从配置的PC开始执行程序。
8. 调度器等待所有BlockNum个实例执行完成。
9. 异步kernel调用执行完成。