概述

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"