Function Execution Space Qualifier

The function execution space qualifier of AI CPU indicates whether a function is an AI CPU kernel function.

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

• The function is executed on the device and can be called only by functions on the host. Therefore, it must be declared with __global__.
• An __aicpu__ __global__ function cannot return the void type, and the input parameter can only be a pointer.
• An __aicpu__ __global__ function cannot be a member function of a class and cannot exist in an anonymous space.
• An __aicpu__ __global__ function cannot be defined in the .cce/.asc file and can only be declared, and extern must be used.
• When the __global__ __aicpu__ function is called on the host, the <<<>>> heterogeneous call syntax must be used. The input function parameter must include the size of the data read from the pointer based on the input parameter pointer.
• 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 runtime synchronization API, such as 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__ 
__global__ __aicpu__  void foo() {} // Error, return type is void
__global__ __aicpu__  int foo(void *a) {} // OK 
__global__ __aicpu__  int foo(int a) {} // Error, input param is not pointer
__global__ __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
__global__ __aicpu__  int foo(void *a); // Error, declare AI CPU kernel function in host file without extern
extern __global__ __aicpu__  int foo(void * a) {} // Error, define AI CPU kernel function in cce file
extern __global__ __aicpu__  int foo(void *a);  // OK

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 behavior
}