基本编程指导

预定义宏

毕昇编译器同其他语言的编译器一样，会有一些预定义宏方便开发者编写程序。

函数执行空间限定符

函数执行空间限定符（Function Execution Space Qualifier）指示函数是在host侧执行还是在device侧执行，以及它是否可从host侧或device侧调用。

• __global__

  __global__执行空间限定符声明一个kernel函数。kernel函数有如下性质：在device上执行；只能被host侧函数调用；__global__只是表示这是device侧函数的入口，并不表示具体的设备类型，具体的设备类型由[aicore]标记。具有如下使用约束：

  • 一个__global__函数必须返回void类型，并且不能是class的成员函数。
  • 主机侧调用__global__函数必须使用<<<>>>异构调用语法。
  • __global__的调用是异步的，意味着函数返回，并不表示kernel函数在device侧已经执行完成，如果需要同步，需要使用AscendCL Runtime同步接口显式同步，如aclrtSynchronizeStream接口。
• [aicore]

  aicore执行空间限定符声明一个函数，它具有如下属性：

  • 在device侧执行
  • 只能被__global__函数，或者其他aicore函数调用
  • 该限定符与Ascend C中限定符 __aicore__ 含义一致。适用于开发自定义算子，但未使用Ascend C头文件场景。

  // Only callable from device functions with same kind
  // of execution space
  [aicore] void bar() {}
  
  // Define a kernel function execute on AICore device
  __global__ [aicore] void foo() {
    bar(); // OK.
  }

内建变量

以下变量为毕昇编译器内建变量，开发者程序中定义的变量名不能和内建变量名相同；同时内建变量通常提供给框架使用，开发者不直接感知，如需获取核id或者核数需要使用对应的Ascend C API获取。

设备侧支持C++高级特性

下文对设备侧支持的C++高级特性（template模板和constexpr表达式）进行介绍。

constexpr表达式是指值不会改变并且在编译过程就能得到计算结果的表达式。声明为constexpr的变量一定是一个const变量，而且必须用常量表达式初始化。

constexprint mf = 20;  //20是常量表达式
constexprint limit = mf + 1; // mf + 1是常量表达式
constexprint sz = size(); //之后当size是一个constexpr函数时才是一条正确的声明语句

在if条件语句中，条件的值必须是按语境可以转换到bool类型的常量表达式，有如下表达和约束：

• 修饰表达式

  template<typename T>
  auto inline __aicore__ get_value(T t)
  {
      if constexpr (std::is_pointer_v<T>)
          return *t; // 对 T = int* 推导返回类型为 int
      else
          return t;  // 对 T = int 推导返回类型为 int
  }
  
  void inline __aicore__  f(__gm__ int* Out)
  {
      if constexpr (sizeof(int) == 3)
          *Out = 1;      // 该代码编译期被舍去
      else
          *Out = 2;}

• 修饰bool值

  void inline __aicore__  f(__gm__ int* Out)
  {
      if constexpr (true)
          *Out = 1;
      else
          *Out = 2;      // 该代码编译期被舍去
  }

• 类中实现

  class ConstClass
  {
  public:
    int a = 2;
    int b = 3;
    int c = a * b;
  };
  int inline __aicore__ f()
  {
      int res = 0;
      constexpr ConstClass t1;
      if constexpr (t1.c > 0) // t1 是编译期常量，可以正常编译
             res = 100;
      else res = 200;
  
      ConstClass t2;
      if constexpr (t2.c > 0) // t2 不是编译期常量，编译失败
             res = 100;
      else res = 200;
  
      return res;
  }

• 舍入分支语法检查

  被舍去的语句同样会经过完整的代码检查，语法错误的代码并不会因为被舍弃而被忽略。

  void inline __aicore__  f()
  {
      if constexpr (false) // 该段代码不会被编译
      {
         int i = 0;
         int *p = i;   // 语法错误，导致编译失败
      }
  }

自动同步

自动同步编译选项能够正确插入同步强依赖于LocalTensor间依赖关系，特别做如下说明。

临时LocalTensor直接逻辑依赖（赋值、引用、指针操作）支持自动同步。非直接逻辑依赖（利用容器传递，手动tensor成员赋值，函数传值等）不支持自动同步，需要开发者手动改为支持场景。

下面以vector间PIPE_V的插入举例，其中Vector_OP指任意的vector指令：

• 直接逻辑依赖场景
  • LocalTensor初始化通过赋值方式

    // 定义dstLocal1
    LocalTensor<T> dstLocal1;
    Vector_OP(dstLocal1, params); 
    // 赋值定义dstLocal2
    LocalTensor<T> dstLocal2 = dstLocal1; 
    // 自动插入PIPE_V
    Vector_OP(dstLocal2, params);

  • LocalTensor初始化通过切片赋值方式

    // 定义dstLocal1
    LocalTensor<T> dstLocal1;
    Vector_OP(dstLocal1, params); 
    // 切片赋值定义dstLocal2
    LocalTensor<T> dstLocal2 = dstLocal1[512]; 
    // 自动插入PIPE_V
    Vector_OP(dstLocal2, params); 

  • LocalTensor初始化通过引用方式

    // 定义dstLocal1
    LocalTensor<T> dstLocal1;
    Vector_OP(dstLocal1, params); 
    // 引用定义dstLocal2
    LocalTensor<T> &dstLocal2 = dstLocal1; 
    // 自动插入PIPE_V
    Vector_OP(dstLocal2, params); 

  • LocalTensor初始化通过指针方式

    // 定义dstLocal1
    LocalTensor<T> dstLocal1;
    Vector_OP(dstLocal1, params); 
    // 指针定义dstLocal2
    LocalTensor<T> *dstLocal2 = &dstLocal1; 
    // 自动插入PIPE_V
    Vector_OP(*dstLocal2, params); 

• 非直接逻辑依赖场景
  • LocalTensor初始化通过容器Que方式

    // 定义dstLocal1
    LocalTensor<T> dstLocal1;
    Vector_OP(dstLocal1, params);
    inQueueSrc.EnQue(dstLocal1); 
    // 容器Que定义dstLocal2
    LocalTensor<T> dstLocal2 = inQueueSrc.DeQue<T>(); 
    // 无法获取依赖，需手动改为支持场景
    Vector_OP(dstLocal2, params);

  • LocalTensor初始化通过其内存成员的接口方式

    // 定义dstLocal1
    LocalTensor<T> dstLocal1;
    Vector_OP(dstLocal1, params);
    // LocalTensor操作内存成员接口定义dstLocal2内存地址
    dstLocal2.SetAddrWithOffset(dstLocal1, offset);
    // 无法获取依赖，需手动改为支持场景
    Vector_OP(dstLocal2, params);

  • 函数接口场景：

    LocalTensor<T> dstLocal1;
    Vector_OP(dstLocal1, params);
    
    __aicore__ void Foo (dstLocal1) { 
    // 无法获取依赖，需手动改为支持场景
       Vector_OP(dstLocal1, params);
    }

  毕昇编译器提供“--cce-auto-sync-log=<file>”编译选项可以输出同步插入信息到<file>文件中，帮助开发者显式地识别编译器在算子文件中插入的同步指令信息。需debug模式（添加-g编译选项）编译算子，用于获取算子代码文件行号。

  • 直接使用毕昇编译器的场景，可以直接在编译命令中添加该编译选项
  • 使用Ascend C kernel直调算子工程，可以通过ascendc_compile_options添加该编译选项
  • 使用Ascend C自定义算子开发工程，可以通过add_ops_compile_options添加该编译选项

  如下的代码文件sync_log_test.h：

  LocalTensor<T> dstLocal;
  T ave_tmp = 0;
  Vector_OP1(dstLocal, params); 
  ave_tmp = dstLocal.GetValue(0);
  Vector_OP2(dstLocal, params); 
  for (int i = 0; i < ave_tmp; ++i) {
      dstLocal.SetValue(i,0);
  }

  开启自动同步后，同步指令的插入位置如下：

  LocalTensor<T> dstLocal;
  T ave_tmp = 0;
  Vector_OP1(dstLocal, params); 
  SetFlag<HardEvent::V_S>(EVENT_ID0);
  WaitFlag<HardEvent::V_S>(EVENT_ID0);
  ave_tmp = dstLocal.GetValue(0);
  PipeBarrier<PIPE_V>();
  SetFlag<HardEvent::S_V>(EVENT_ID0);
  WaitFlag<HardEvent::S_V>(EVENT_ID0);
  Vector_OP2(dstLocal, params); 
  SetFlag<HardEvent::V_S>(EVENT_ID0);
  WaitFlag<HardEvent::V_S>(EVENT_ID0);
  for (int i = 0; i < ave_tmp; ++i) {
      dstLocal.SetValue(i,0);
  }

  开启自动同步debug日志功能后，输出日志如下：

  The BiSheng Auto Sync log of sync_log_test :  
  Position: absolute-path/sync_log_test.h:4 : line before insert sync : SetFlag<HardEvent::V_S>(EVENT_ID0);
  Position: absolute-path/sync_log_test.h:4 : line before insert sync : WaitFlag<HardEvent::V_S>(EVENT_ID0);
  Position: absolute-path/sync_log_test.h:5 : line before insert sync : PipeBarrier<PIPE_V>();
  Position: absolute-path/sync_log_test.h:5 : line before insert sync : SetFlag<HardEvent::S_V>(EVENT_ID0);
  Position: absolute-path/sync_log_test.h:5 : line before insert sync : WaitFlag<HardEvent::S_V>(EVENT_ID0);
  Position: absolute-path/sync_log_test.h:6 : line before insert sync : SetFlag<HardEvent::V_S>(EVENT_ID0);
  Position: absolute-path/sync_log_test.h:6 : line before insert sync : WaitFlag<HardEvent::V_S>(EVENT_ID0);

  其中，line before表示紧接着当前行前面插入的同步指令。