函数空间执行限定符

AI CPU函数执行空间限定符（Function Execution Space Qualifier）指示函数是否为AI CPU Kernel函数。

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

• 在Device侧执行且只能被Host侧函数调用，因此必须与__global__同时声明。
• 一个__aicpu__ __global__函数不能是void返回类型，并且入参只能是一个指针。
• 一个__aicpu__ __global__函数不能是class的成员函数，不能存在于匿名空间下。
• 一个__aicpu__ __global__函数不能在.cce/.asc文件中进行定义，只能声明，且需要使用extern。
• Host侧调用__global__ __aicpu__函数时必须使用<<<>>>异构调用语法，输入的函数入参在入参指针的基础上需要输入从指针中读取的数据大小。
• __global__的调用是异步的，意味着函数返回，并不表示kernel函数在Device侧已经执行完成，如果需要同步，需要使用Runtime同步接口显式同步，如aclrtSynchronizeStream接口

foo.aicpu
// Define a AI CPU kernel function in AI CPU device file
__aicpu__ void foo() {} // Error, single __aicpu__ identifier without _global__
__global__ void foo() {} // Error, single __global__ identifier without __aicpu__ 
__aicore__ __aicpu__  void foo() {} // Error, return type is void
__aicore__ __aicpu__  int foo(void *a) {} // OK 
__aicore__ __aicpu__  int foo(int a) {} // Error, input param is not pointer
__aicore__ __aicpu__  int foo(void *a, void *b) {} // Error, input param num is not one

int bar() {}  // OK, normal funciton on aicpu device 

__global__ __aicpu__ int test() {
  bar(); // OK.
  foo(); // Error
}

foo.cce
// Declare a AI CPU kernel function in AI CPU host file
// foo has already been defined in foo.aicpu
__aicore__ __aicpu__  int foo(void *a); // Error, declare AI CPU kernel function in host file without extern
extern __aicore__ __aicpu__  int foo(int a) {} // Error, define AI CPU kernel function in cce file
extern __aicore__ __aicpu__  int foo(void *a);  // OK

__global__ __aicpu__ int test() {
  aclrtStream stream;
  aclrtCreateStream(&stream);
  int a[100];
  foo<<<1, nullptr, stream>>>(a, sizeof(a)); // OK
  foo<<<1, nullptr, stream>>>(a); // Error, lack argument for paramsize
  foo<<<5, nullptr, stream>>>(a, sizeof(a)); // OK with compiler and run, but only simplely run foo for 5 times, there may be undefined behavious
}