函数执行空间限定符

__global__

__global__执行空间限定符声明一个kernel函数。该kernel函数有如下属性：

• 在设备上执行
• 只能被主机侧函数调用
• __global__只是表示这是设备侧函数的入口，并不表示具体的设备类型，具体的设备类型由[aicore]标记。

• __global__函数必须返回void类型，并且不能是class的成员函数。
• 主机侧可以使用<<<>>>直接进行__global__函数异构调用。
• __global__的调用是异步的，函数返回并不表示kernel函数在设备侧已经执行完成。如果需要同步，须使用Runtime提供的同步接口显式同步，如aclrtSynchronizeStream(...)。

AICore限定符

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

• 在AICore设备上执行。
• 只能被__global__函数，或者其他aicore函数调用。

代码示例：

// 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();
}

host限定符

host执行空间限定符声明一个主机侧函数，具有以下属性：

• 只能在主机侧执行
• 只能被主机侧函数调用
• __global__ 和[host]不能一起使用

注意事项：

• [host]限定符是可选项，无函数执行空间限定符定义的函数，默认是host函数。
• 非__global__函数可以同时被声明为host和aicore函数，编译器会为主机侧和设备侧都生成代码。

代码示例：

[aicore] void f() {} 
 
// defines a host side function 
void foo() {} 
 
// defines a host side function 
[host] void bar() { 
  f();     // Error. 
  foo();   // OK. 
} 
 
// define a function executing on both host and aicore 
// it should be able to compile on both host and aicore 
// otherwise it will report compiler error 
[host, aicore] void tee() { 
  f(); // Error. Function execution space does not match 
} 
 
// Error. 
__global__ [host] void kfunc() {} 
 
// Error. 
__global__ [host,aicore] void kfunc() {}