CANNBot SIMT API总览

SIMT C API 总览【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skillsSIMT API基于AI Core硬件能力通过AscendC::Simt::VF_CALL启动SIMT VF子任务。每32个线程组成一个WarpWarp内每个线程称为Lane编号0~31。头文件包含#include simt_api/asc_simt.h // 通用非half/bfloat16/fp8类型 #include simt_api/asc_fp16.h // half/half2类型API #include simt_api/asc_bf16.h // bfloat16/bfloat16x2_t类型API #include simt_api/asc_fp8.h // hifloat8x2_t/float8_e4m3x2_t/float8_e5m2x2_t类型APIAPI分类类别功能头文件索引文件核函数定义启动SIMT VF子任务common_functions.hmisc/asc_vf_call.md同步函数线程同步与内存可见性device_sync_functions.h01_同步函数.md数学函数三角/指数/对数/幂等运算math_functions.h02_数学函数.md精度转换取整(rint/round/floor/ceil)math_functions.h03_精度转换.md比较函数判断有限数/NaN/Infmath_functions.h04_比较函数.md原子操作UB/GM上的原子读写device_atomic_functions.h05_原子操作.mdWarp函数Warp内数据交换/归约device_warp_functions.h06_Warp函数.md类型转换float/half/bf16/int间转换device_functions.h07_类型转换.md向量构造make_int2/float2等vector_functions.h08_向量构造函数.mdCache Hints带缓存提示的Load/Storedevice_functions.h09_Cache_Hints.md调测接口printf/assert/trapasc_simt.h10_调测接口.md内置宏特殊值常量和数学常数asc_simt.h/asc_fp16.h/asc_bf16.h11_内置宏.md核函数定义与启动asc_vf_call在SIMD与SIMT混合编程场景启动SIMT VFVector Function子任务通过参数配置启动指定数目的线程执行指定的SIMT核函数。注意asc_vf_call启动SIMT VF子任务时子任务函数不能是类的成员函数推荐使用普通函数或类静态函数且入口函数必须使用__simt_vf__修饰宏。传递的参数只支持裸指针和常见基本数据类型不支持传递结构体、数组等。函数原型:template auto funcPtr, typename... Args __aicore__ inline void asc_vf_call(dim3 threadNums, Args ...args)模板参数:参数名描述funcPtr用于指定SIMT入口核函数Args定义可变参数用于传递实参到SIMT入口核函数参数说明:参数名输入/输出描述threadNums输入dim3结构{dimx,dimy,dimz}指定SIMT线程块内线程数量。线程总数dimxdimydimz必须2048args输入可变参数传递实参到SIMT入口核函数返回值: 无需要包含的头文件:#include simt_api/common_functions.h调用示例:__simt_vf__ __launch_bounds__(2048) inline void SimtCompute( __gm__ float* dst, __gm__ float* src0, __gm__ float* src1, int count) const { for(int idx threadIdx.x blockIdx.x * blockDim.x; idx count; idx gridDim.x * blockDim.x) { dst[idx] src0[idx] src1[idx]; } } __global__ __aicore__ void SimtComputeShell(__gm__ float* x, __gm__ float* y, __gm__ float* z, const int size) { __gm__ float* dst x; __gm__ float* src0 y; __gm__ float* src1 z; asc_vf_callSimtCompute(dim3{1024, 1, 1}, dst, src0, src1, size); }SIMD与SIMT混合编程辅助函数以下函数用于SIMD与SIMT混合编程场景中的辅助操作。GetRuntimeUBSize获取运行时UB空间的大小单位为byte。开发者根据UB的大小来计算循环次数等参数值。函数原型:__aicore__ inline uint32_t GetRuntimeUBSize()返回值: 运行时UB空间的大小字节。Ascend 950PR/Ascend 950DT架构下SIMD与SIMT混合场景中UB大小上限为216KB非混合场景返回固定值248KB。调用示例:uint32_t totalLength 126976; uint32_t tileLength AscendC::GetRuntimeUBSize() / sizeof(half) / 2; uint32_t tileNum totalLength / tileLength;BlockReduceMax / BlockReduceMin / BlockReduceSum对每个datablock内所有元素分别求最大值、最小值、求和。这些是SIMD层面的归约指令在SIMD与SIMT混合编程场景中使用。BlockReduceMax: 对每个datablock内所有元素求最大值BlockReduceMin: 对每个datablock内所有元素求最小值BlockReduceSum: 对每个datablock内所有元素求和二叉树方式两两相加支持的数据类型: half, float函数原型(mask连续模式):template typename T, bool isSetMask true __aicore__ inline void BlockReduceMax(const LocalTensorT dst, const LocalTensorT src, const int32_t repeatTime, const int32_t mask, const int32_t dstRepStride, const int32_t srcBlkStride, const int32_t srcRepStride)更多详细的参数说明mask模式、stride语义等请参考 AscendC SIMD API 文档。【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考