SIMT Entrypoint Function

The SIMT entrypoint function marked by __simt_vf__ must comply with the following constraints:

Pointers and references cannot be used as scalar parameters.

Arrays cannot be used as parameters. The following is an incorrect example:

__simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, int* array) {
    int idx = blockIdx * blockDim.x + threadIdx.x;
    // error: Arrays are not allowed.
    a[idx] = b[idx] + c[idx] + array[0];
}
__global__ __aicore__ void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c) {
    int array[5] = {0,1,2,3,4};
    asc_vf_call<foo>(dim3{256}, a, b, c, array);
}

If a multi-level pointer is passed as a parameter, the inner stack address pointer cannot be used for access. The following is an incorrect example:

__simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, __ubuf__ uint64_t* s) {
    int idx = blockIdx * blockDim.x + threadIdx.x;
    int* stack = (int*)(s[0]);
    // error: *stack indicates that data is read from a multi-level pointer, which is not allowed.
    a[idx] = b[idx] + c[idx] + *stack;
}
__global__ __aicore__ void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c) {
    int stack = 0;
    __ubuf__ uint64_t* s = ...;
    s[0] = &stack;
    asc_vf_call<foo>(dim3{256}, a, b, c, s);
}

Indirect calls through function pointers are not supported. The __simt_vf__ function to be called must be determined during compilation.
The inline identifier must be added.

Structures cannot be used as parameters. The following is an incorrect example:

__simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, struct S s) {
    // error: s indicates a structure parameter, which is not allowed.
}

SIMT Subfunction

SIMT subfunctions marked by __simt_callee__ must comply with the following constraints:

The inline identifier must be added.