开发者
下载

函数执行空间限定符

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, ...);

该函数有以下约束:

  • 入参仅支持Ascend C的内置数据类型int32_t、uint32_t、float、half等)及对应的指针类型。
  • 函数返回值只能是Ascend C的内置数据类型int32_t、uint32_t、float、half等)及对应的指针类型。
  • 函数内只能调用__simt_callee__函数或constexpr __aicore__函数。