概 述

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文 件。

Text

#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"