基本编程指导

预定义宏

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

函数执行空间限定符

函数执行空间限定符（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获取。

自动同步

上文描述的核内同步，自动同步编译选项能够正确插入同步强依赖于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);
    }