---
title: 函数
description: "__simt_vf__标记的SIMT入口函数需要遵循以下约束和限制：- 标量参数不允许使用指针和引用。"
url: https://www.hiascend.com/document/detail/zh/canncommercial/latest/programug/Ascendcopdevg/atlas_ascendc_10_00118.html
sourcePath: /source/zh/canncommercial/900/programug/Ascendcopdevg/atlas_ascendc_10_00118.html
indexId: 68e5fb144b64adaeee3bd8cb769cea4becc2bd22cc2e9da77fce8459e595651581
---
# 函数

#### SIMT入口函数

__simt_vf__标记的SIMT入口函数需要遵循以下约束和限制：- 标量参数不允许使用指针和引用。
- 不允许使用数组作为参数，以下为错误示例。
  1 2 3 4 5 6 7 8 9 __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); }

- 如果传参中出现多级指针，不允许使用内层栈地址指针访问，以下为错误示例。
  1 2 3 4 5 6 7 8 9 10 11 12 __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标识符。
- 不允许使用结构体作为参数，以下为错误示例。
  1 2 3 __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标识符。