函数执行空间限定符

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

• __global__

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

  • 一个__global__函数必须返回void类型，并且不能是class的成员函数。
  • 主机侧调用__global__函数必须使用<<<>>>异构调用语法。
  • __global__的调用是异步的，意味着函数返回，并不表示kernel函数在device侧已经执行完成，如果需要同步，需要使用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.
  }