SIMT入口函数

__simt_vf__标记的SIMT入口函数需要遵循以下约束和限制：

标量参数不允许使用指针和引用。

不允许使用数组作为参数，以下为错误示例。

__simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, int* array) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // error: 不允许使用数组array
    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);
}

如果传参中出现多级指针，不允许使用内层栈地址指针访问，以下为错误示例。

__simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, __ubuf__ uint64_t* s) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int* stack = (int*)(s[0]);
    // error: *stack表示从多级指针中读取，不允许使用
    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);
}

不支持通过函数指针进行间接调用，被调用的__simt_vf__函数需要在编译期确定。
必须添加inline标识符。

不允许使用结构体作为参数，以下为错误示例。

__simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, struct S s) {
    // error: s表示结构体参数，不允许使用
}

SIMT子函数

__simt_callee__标记的SIMT子函数需要遵循以下约束和限制：

必须添加inline标识符。