概述
SIMT API是面向AI处理器的并行计算编程接口,可以实现高效的数据并行计算。SIMT API支持两种编程模型:SIMT编程、SIMD与SIMT混合编程,用户可先阅读SIMT编程简介和SIMD与SIMT混合编程简介,以了解编程基础,后续章节将详细介绍API接口。
表 1 SIMT API分类列表
| 类别 | 功能 |
|---|---|
| 同步与内存栅栏 | 内存管理与同步接口,解决不同核内的线程间可能存在的数据竞争以及线程的同步问题。 |
| 原子操作 | 对Unified Buffer或Global Memory上的数据与指定数据执行原子操作的一系列接口。 |
| Warp函数 | 对单个Warp内32个线程的数据进行处理的相关操作的一系列API接口。 |
| 数学函数 | 用于处理数学运算的函数集合以及不同精度、数据类型的转换函数集合 |
| 地址空间谓词函数 | 判断输入指针是否为指定空间的地址。 |
| 地址空间转换函数 | 将指定地址空间的地址值转换为指针,或将输入的指针转换为对应内存空间的地址值的接口。 |
| 访存函数 | 数据加载和数据缓存相关接口。 |
| 协作组 | 提供一套标准且安全的机制,实现更高效的线程并行协作。 |
| 调测接口 | SIMT VF调试场景下使用的相关接口。 |
Ascend C SIMT API支持通过包含simt_api/asc_simt.h文件来调用输入数据为除half、half2、bfloat16_t、bfloat16x2_t、hifloat8x2_t、float8_e4m3x2_t、float8_e5m2x2_t以外类型的接口,调用输入数据为half和half2类型的SIMT API需要包含simt_api/asc_fp16.h文件,调用输入数据为bfloat16_t和bfloat16x2_t类型的SIMT API需要包含simt_api/asc_bf16.h文件,调用输入数据为hifloat8x2_t、float8_e4m3x2_t和float8_e5m2x2_t类型的SIMT API需要包含simt_api/asc_fp8.h文件。
#include "simt_api/asc_simt.h"
#include "simt_api/asc_fp16.h"
#include "simt_api/asc_bf16.h"
#include "simt_api/asc_fp8.h"
表 2 各类SIMT API需要包含的头文件
| 类别 | 除half、half2、bfloat16_t、bfloat16x2_t之外的类型需要包含的头文件 |
|---|---|
| 同步与内存栅栏 | #include "simt_api/device_sync_functions.h" |
| 原子操作 | #include "simt_api/device_atomic_functions.h" |
| Warp函数 | #include "simt_api/device_warp_functions.h" |
| 数学函数 | #include "simt_api/math_functions.h" #include "simt_api/device_functions.h" |
| 地址空间谓词函数 地址空间转换函数 访存函数 |
#include "simt_api/device_functions.h" |
| 协作组 | #include "simt_api/cooperative_groups.h" |
表 3 不同数据类型下使用接口需要包含的头文件
| 类别 | half、half2类型需要包含的头文件 | bfloat16_t、bfloat16x2_t类型需要包含的头文件 | hifloat8x2_t、float8_e4m3x2_t、float8_e5m2x2_t类型需要包含的头文件 |
|---|---|---|---|
| 任一类别API | #include "simt_api/asc_fp16.h" | #include "simt_api/asc_bf16.h" | #include "simt_api/asc_fp8.h" |