Function Execution Space Qualifier

The function execution space qualifier indicates whether the function is executed on the host or device, and whether the function can be called from the host or device.

• __global__

  The execution space qualifier __global__ declares a kernel function. The kernel function is executed on the device, and called only by the function on the host. __global__ only indicates the entrypoint of the function on the device, rather than the specific device type. The specific device type is marked by [aicore]. The restrictions are as follows:

  • A __global__ function must return the void type, and cannot be a member function of class.
  • The host must call the __global__ function using the <<<>>> heterogeneous calling syntax.
  • The __global__ call is asynchronous. The return of the function does not mean that the kernel function has been completely executed on the device. For explicit synchronization, use the AscendCL Runtime synchronization API, such as aclrtSynchronizeStream.
• [aicore]

  The execution space qualifier aicore declares a function with the following attributes:

  • Executed on the device.
  • Called only by the __global__ function or other aicore functions.
  • Same as the meaning of __aicore__ in Ascend C. It is used for developing custom operators, but does not use Ascend C header files.

  // 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.
  }