函数执行空间限定符
SIMD与SIMT混合编程涉及多种类型的函数,它们之间遵循严格的调用关系和层级约束。本章节将会介绍整体函数类型和调用关系,再详细展开每种函数的定义、调用语法及约束。
SIMD与SIMT混合编程中涉及的函数类型如下表所示:
函数执行空间限定符 |
函数功能 |
调用方式 |
|---|---|---|
__global__ __aicore__ |
算子入口,协调VF执行 |
Host侧通过<<<...>>>调用 |
__aicore__ |
Device侧辅助函数 |
核函数或同级函数调用 |
__simt_vf__ |
线程级并行计算任务 |
通过SIMT提供的asc_vf_call接口调用 |
__simd_vf__ |
向量级并行计算任务 |
通过SIMD提供的asc_vf_call接口调用 |
__simt_callee__ |
SIMT VF的子函数 |
SIMT VF内部调用 |
__simd_callee__ |
SIMD VF的子函数 |
SIMD VF内部调用 |
constexpr __aicore__ |
编译期执行的Device侧辅助函数 |
Device侧函数内部均可调用 |
函数调用关系如下图所示:

__simt_vf__
__simt_vf__用于标记SIMT VF入口函数,函数无返回值。使用asc_vf_call接口调用SIMT VF入口函数,启动VF子任务。函数内可使用SIMT内置变量:threadIdx、blockIdx、blockDim、gridDim。函数定义示例如下:
__simt_vf__ inline void function_name(
__gm__ uint32_t* gm_param,
__ubuf__ float* ubuf_param,
uint64_t scalar_param, ...);
SIMT VF函数定义中的关键修饰符说明如下:
修饰符 |
作用 |
|---|---|
__simt_vf__ |
函数标识符,标识SIMT VF函数。 |
inline |
建议内联,实际是否内联由编译器决定。 |
__gm__ |
内存空间修饰符,标识内存空间为GM。 |
__ubuf__ |
内存空间修饰符,标识内存空间为UB。 |
通过SIMT的asc_vf_call接口在核函数或__aicore__函数中调用,调用示例如下:
asc_vf_call<function_name>(dim3(blockDim), arg1, arg2, ...);
SIMT VF函数有以下约束:
- 入参仅支持Ascend C的内置数据类型(int32_t、uint32_t、float、half等)及其组成的指针、数组、结构体类型,且指针类型必须指向GM或者UB内存。
- 函数返回类型必须是void。
- SIMT VF内只能调用__simt_callee__函数或constexpr __aicore__函数。
__simt_callee__
__simt_callee__子函数是SIMT VF函数内部调用的辅助函数,函数内部可以使用SIMT内置变量。定义示例如下:
__simt_callee__ uint32_t simt_helper(__gm__ uint32_t* gm_param, __ubuf__ float* ubuf_param, uint64_t scalar_param, ...);
调用示例如下:
uint32_t result = simt_helper(arg1, arg2, ...);
该函数有以下约束: