已合并
feat: 再生成 85 个索引导航摘要完成两步索引迁移 #58
feat: 再生成 85 个索引导航摘要完成两步索引迁移 #58
已合并
Yifan Wang创建于 9月22日
共 88 个文件变更+576-438
@@ -17,7 +17,7 @@
17- [Reg矢量计算接口精度标准汇总](reg_vector_compute_interface_precision_standard_summary.md):本节汇总介绍了Reg矢量计算接口精度标准的情况,本节内容仅针对如下型号生效:17- [Reg矢量计算接口精度标准汇总](reg_vector_compute_interface_precision_standard_summary.md):本节汇总介绍了Reg矢量计算接口精度标准的情况,本节内容仅针对如下型号生效:
18- [Reg矢量计算接口软仿情况汇总](reg_vector_compute_interface_soft_simulation_summary.md):本节汇总介绍了Reg矢量计算接口软仿数据类型的情况,本节内容仅针对如下型号生效:18- [Reg矢量计算接口软仿情况汇总](reg_vector_compute_interface_soft_simulation_summary.md):本节汇总介绍了Reg矢量计算接口软仿数据类型的情况,本节内容仅针对如下型号生效:
19- [Reg矢量计算接口理论性能汇总](reg_vector_compute_interface_theoretical_perf_summary.md):本节汇总介绍了Reg矢量计算接口的理论性能,本节内容仅针对如下型号生效:19- [Reg矢量计算接口理论性能汇总](reg_vector_compute_interface_theoretical_perf_summary.md):本节汇总介绍了Reg矢量计算接口的理论性能,本节内容仅针对如下型号生效:
20-- [Release Notes](release_notes/index.md):CANN 9.0.0首次面向Ascend 950PR/950DT提供试验性C语言扩展层SIMD API,覆盖数据搬运、标量操作、矩阵计算、同步控制等能力。20+- [Release Notes](release_notes/index.md):共 2 张知识卡(如 Cann9ReleaseNotes、C API)。
21- [SIMT接口精度标准汇总](simt_interface_precision_standard_summary.md):本节按照数据类型和接口类型对SIMT浮点计算接口进行精度标准的分类汇总。对于非软仿接口,给出对应硬件指令的精度标准;对于软仿接口,由于软仿程度较低且使用到的硬件指令较少,给出其所用硬件指令的综合性精度标准。21- [SIMT接口精度标准汇总](simt_interface_precision_standard_summary.md):本节按照数据类型和接口类型对SIMT浮点计算接口进行精度标准的分类汇总。对于非软仿接口,给出对应硬件指令的精度标准;对于软仿接口,由于软仿程度较低且使用到的硬件指令较少,给出其所用硬件指令的综合性精度标准。
22- [SIMT接口软仿情况汇总](simt_interface_soft_simulation_summary.md):本节按照SIMT API文档所在目录分级汇总各接口的软仿情况。22- [SIMT接口软仿情况汇总](simt_interface_soft_simulation_summary.md):本节按照SIMT API文档所在目录分级汇总各接口的软仿情况。
23- [Vector指令理论性能汇总](vector_instruction_theoretical_perf_summary.md):本节汇总介绍了一些主要的Vector指令的理论性能,以下内容针对如下型号生效:23- [Vector指令理论性能汇总](vector_instruction_theoretical_perf_summary.md):本节汇总介绍了一些主要的Vector指令的理论性能,以下内容针对如下型号生效:
@@ -1,7 +1,7 @@
1# Ascend C API1# Ascend C API
2 2 
3-- [AI CPU API](ai_cpu_api/index.md):该接口实现AI CPU算子Kernel调试场景下的assert断言功能。 算子执行中,如果assert内部条件判断不为真,则输出assert条件、触发文件名、行号等信息。3+- [AI CPU API](ai_cpu_api/index.md):共 3 张知识卡(如 assert(AI CPU API)、DataStoreBarrier、printf(AI CPU API))。
4-- [附录](appendix/index.md):本节汇总介绍了一些主要的API的流水类型,以下内容针对如下型号生效:4+- [附录](appendix/index.md):共 23 张知识卡(如 API流水类型汇总(附录)、原子操作边界值汇总(附录)、C-API-Reg矢量计算接口理论性能汇总 等)。
5-- [SIMD API](simd_api/index.md):说明 SIMD 基础与高阶 API 的头文件、链接方式、存储单元对齐要求及 TPosition 内存映射。5+- [SIMD API](simd_api/index.md):共 1890 张知识卡(如 通用说明和约束(SIMD API) 等)。
6-- [SIMT API](simt_api/index.md):SIMT API是面向AI处理器的并行计算编程接口,可以实现高效的数据并行计算。SIMT API支持两种编程模型:SIMT编程、SIMD与SIMT混合编程,用户可先阅读SIMT编程简介和SIMD与SIMT混合编程简介,以了解编程基础,后续章节将详细介绍API接口。6+- [SIMT API](simt_api/index.md):SIMT API是面向AI处理器的并行计算编程接口,可以实现高效的数据并行计算。SIMT API支持两种编程模型:SIMT编程、SIMD与SIMT混合编程,用户可先阅读SIMT编程简介和SIMD与SIMT混合编程简介,以了解编程基础,后续章节将详细介绍API接口。(共 766 张知识卡,分布在 10 个子目录)
7-- [Utils API](utils_api/index.md):比较相同数据类型的两个数,返回最大值。7+- [Utils API](utils_api/index.md):共 183 张知识卡,分布在 11 个子目录(cpp_stdlib、log、platform_info 等)。
@@ -1,11 +1,11 @@
1# 激活函数1# 激活函数
2 2 
3-- [GeGLU接口](ge_glu_interface/index.md):采用GELU作为激活函数的GLU变体,按元素对两源Tensor计算GeGLU结果3+- [GeGLU接口](ge_glu_interface/index.md):共 3 张知识卡(如 GeGLU、GetGeGLUMaxMinTmpSize、GetGeGLUTmpBufferFactorSize)。
4-- [Gelu接口](gelu_interface/index.md):GELU激活函数的FastGelu化简版本,通过公式化简大幅提升计算性能,支持高精度和高性能模式。4+- [Gelu接口](gelu_interface/index.md):共 4 张知识卡(如 FasterGelu、FasterGeluV2、Gelu 等)。
5-- [LogSoftMax接口](log_soft_max_interface/index.md):对输入tensor做LogSoftmax计算。5+- [LogSoftMax接口](log_soft_max_interface/index.md):共 2 张知识卡(如 LogSoftMax、LogSoftMaxTiling)。
6-- [ReGlu接口](re_glu_interface/index.md):用于在host侧获取ReGlu接口所需的最大最小临时空间大小,作为Tiling参数传递到kernel侧使用。6+- [ReGlu接口](re_glu_interface/index.md):共 2 张知识卡(如 GetReGluMaxMinTmpSize、ReGlu)。
7-- [Sigmoid接口](sigmoid_interface/index.md):用于在host侧获取Sigmoid接口所需的最大最小临时空间大小,作为Tiling参数传递到kernel侧使用。7+- [Sigmoid接口](sigmoid_interface/index.md):共 2 张知识卡(如 GetSigmoidMaxMinTmpSize、Sigmoid)。
8-- [Silu接口](silu_interface/index.md):用于获取Silu Tiling参数:Silu接口能完成计算所需最大临时空间大小和最小临时空间大小。8+- [Silu接口](silu_interface/index.md):共 2 张知识卡(如 GetSiluTmpSize、Silu)。
9-- [SoftMax接口](soft_max_interface/index.md):调整SoftMax计算结果为指定值,当maxTensor中存在指定值时,按行将对应softMaxRes结果替换为自定义值,用于SoftMax后处理。9+- [SoftMax接口](soft_max_interface/index.md):共 15 张知识卡(如 AdjustSoftMaxRes、IsBasicBlockInSoftMax、SimpleSoftMax 等)。
10-- [SwiGLU接口](swi_glu_interface/index.md):在host侧获取SwiGLU接口预留临时空间的最大和最小值,用于Tiling参数选择10+- [SwiGLU接口](swi_glu_interface/index.md):共 3 张知识卡(如 GetSwiGLUMaxMinTmpSize、GetSwiGLUTmpBufferFactorSize、SwiGLU)。
11-- [Swish接口](swish_interface/index.md):用于获取Swish Tiling参数:Swish接口能完成计算所需最大临时空间大小和最小临时空间大小。11+- [Swish接口](swish_interface/index.md):共 2 张知识卡(如 GetSwishTmpSize、Swish)。
@@ -1,8 +1,8 @@
1# 卷积计算1# 卷积计算
2 2 
3-- [Conv3DBackpropFilter核函数(Kernel)侧接口](conv3_d_backprop_filter_kernel/index.md):Ascend C提供一组Conv3DBackpropFilter高阶API,便于用户快速实现卷积的反向运算,求解反向传播的误差。3+- [Conv3DBackpropFilter核函数(Kernel)侧接口](conv3_d_backprop_filter_kernel/index.md):共 9 张知识卡(如 Conv3DBackpropFilter使用说明、卷积计算End、GetTensorC 等)。
4-- [Conv3DBackpropFilter Tiling侧接口](conv3_d_backprop_filter_tiling/index.md):用于创建一个Conv3DBackpropFilter单核Tiling对象。4+- [Conv3DBackpropFilter Tiling侧接口](conv3_d_backprop_filter_tiling/index.md):共 13 张知识卡(如 Conv3DBackpropFilter Tiling构造函数、Conv3DBackpropFilterTiling、GetTiling 等)。
5-- [Conv3DBackpropInput核函数(Kernel)侧接口](conv3_d_backprop_input_kernel/index.md):提供实现3D卷积反向运算的高阶API,求解反向传播误差,也可用于实现转置卷积Conv3DTranspose算子。5+- [Conv3DBackpropInput核函数(Kernel)侧接口](conv3_d_backprop_input_kernel/index.md):共 9 张知识卡(如 Conv3DBackpropInput、卷积计算End、GetTensorC 等)。
6-- [Conv3DBackpropInput Tiling侧接口](conv3_d_backprop_input_tiling/index.md):用于创建一个Conv3DBackpropInput单核Tiling对象。6+- [Conv3DBackpropInput Tiling侧接口](conv3_d_backprop_input_tiling/index.md):共 14 张知识卡(如 Conv3DBackpropInput Tiling构造函数、Conv3DBackpropInputTiling、GetTiling 等)。
7-- [conv3 d kernel](conv3_d_kernel/index.md):创建Conv3D对象时需要传入: - Input、Weight、Output和Bias(可选)的参数类型信息,类型信息通过ConvType来定义,包括:内存逻辑位置、数据格式、数据类型。 - Conv3dParam信息(可选),用于开启不同场景的性能优化模板。当前暂不支持使用。7+- [conv3 d kernel](conv3_d_kernel/index.md):共 10 张知识卡(如 Conv3D模板参数、Conv3DUsage、卷积计算End 等)。
8-- [conv3 d tiling](conv3_d_tiling/index.md):Ascend C提供一组Conv3D Tiling API,方便用户获取Conv3D正向算子Kernel计算时所需的Tiling参数。用户只需要传入Input/Weight/Bias/Output的Position位置、Format格式和DType数据类型及相关参数等信息,调用API接口,即可获取Init中TConv3DApiTiling结构体中的相关参数。8+- [conv3 d tiling](conv3_d_tiling/index.md):共 16 张知识卡(如 Conv3D Tiling使用说明、Conv3D Tiling构造函数、GetTiling 等)。
@@ -1,5 +1,5 @@
1# 矩阵计算1# 矩阵计算
2 2 
3-- [获取Matmul计算所需空间](get_matmul_compute_space/index.md):BatchMatmul Tiling调用GetTiling接口获取Tiling参数后,根据Tiling结构体信息获取L1 Buffer/Unified Buffer/L0C Buffer的使用大小。3+- [获取Matmul计算所需空间](get_matmul_compute_space/index.md):共 6 张知识卡(如 BatchMatmulGetTmpBufSize、BatchMatmulGetTmpBufSizeV2、MatmulGetTmpBufSize 等)。
4-- [Matmul核函数(Kernel)侧接口](matmul_kernel/index.md):获取Iterate接口异步计算的结果矩阵。该接口功能已被GetTensorC覆盖,建议直接使用GetTensorC异步接口。4+- [Matmul核函数(Kernel)侧接口](matmul_kernel/index.md):共 50 张知识卡(如 AsyncGetTensorC、ClearBias、DisableBias 等)。
5-- [Matmul Tiling类](matmul_tiling/index.md):设置Bias是否参与运算,设置的信息必须与Kernel侧保持一致。5+- [Matmul Tiling类](matmul_tiling/index.md):共 40 张知识卡(如 EnableBias、EnableL1BankConflictOptimise、EnableMultiCoreSplitK 等)。
@@ -1,4 +1,4 @@
1# experimental1# experimental
2 2 
3-- [bessel i0 interface](bessel_i0_interface/index.md):本接口为试验接口,在后续版本中可能会调整或改进,不保证后续兼容性。请开发者在使用过程中关注后续版本更新。3+- [bessel i0 interface](bessel_i0_interface/index.md):共 3 张知识卡(如 BesselI0、GetBesselI0MaxMinTmpSize、GetBesselI0TmpBufferFactorSize)。
4- [ndtri interface](ndtri_interface/index.md):Ndtri当前属于实验性接口。使用前需要定义CANN_ASC_USE_EXPERIMENTAL,并显式包含头文件:4- [ndtri interface](ndtri_interface/index.md):Ndtri当前属于实验性接口。使用前需要定义CANN_ASC_USE_EXPERIMENTAL,并显式包含头文件:
@@ -1,5 +1,5 @@
1# HCCL通信类1# HCCL通信类
2 2 
3-- [HCCL Context](hccl_context/index.md):获取指定Index通信域的context(消息区)地址。3+- [HCCL Context](hccl_context/index.md):共 3 张知识卡(如 GetHcclContext、HCCL Context简介、SetHcclContext)。
4-- [HCCL核函数(Kernel)侧接口](hccl_kernel/index.md):集合通信算子AllGather的任务下发接口,返回该任务的标识handleId给用户。AllGather的功能为:将通信域内所有节点的输入按照rank id重新排序,然后拼接起来,再将结果发送到所有节点的输出。 !4+- [HCCL核函数(Kernel)侧接口](hccl_kernel/index.md):共 25 张知识卡(如 AllGather、AllReduce、AlltoAll 等)。
5-- [HCCL Tiling侧接口](hccl_tiling/index.md):根据输入的Tiling信息,对指定的通信算法所需资源进行预校验,返回检查结果。5+- [HCCL Tiling侧接口](hccl_tiling/index.md):共 18 张知识卡(如 CheckOpResSufficient、GetTiling、HCCL Tiling构造函数 等)。
@@ -1,18 +1,18 @@
1# 高阶API1# 高阶API
2 2 
3-- [激活函数](activation_functions/index.md):采用GELU作为激活函数的GLU变体,按元素对两源Tensor计算GeGLU结果3+- [激活函数](activation_functions/index.md):共 35 张知识卡,分布在 9 个子目录(ge_glu_interface、gelu_interface、log_soft_max_interface 等)。
4-- [卷积计算](convolution_compute/index.md):Ascend C提供一组Conv3DBackpropFilter高阶API,便于用户快速实现卷积的反向运算,求解反向传播的误差。4+- [卷积计算](convolution_compute/index.md):共 71 张知识卡,分布在 6 个子目录(conv3_d_backprop_filter_kernel、conv3_d_backprop_filter_tiling、conv3_d_backprop_input_kernel 等)。
5-- [矩阵计算](cube_compute/index.md):BatchMatmul Tiling调用GetTiling接口获取Tiling参数后,根据Tiling结构体信息获取L1 Buffer/Unified Buffer/L0C Buffer的使用大小。5+- [矩阵计算](cube_compute/index.md):共 96 张知识卡,分布在 3 个子目录(get_matmul_compute_space、matmul_kernel、matmul_tiling)。
6-- [数据过滤](data_filter/index.md):根据MaskTensor对SrcTensor进行过滤,支持字节模式和比特模式两种过滤方式6+- [数据过滤](data_filter/index.md):共 4 张知识卡(如 DropOut、GetDropOutMaxMinTmpSize、GetSelectMaxMinTmpSize 等)。
7-- [数据结构](data_structures/index.md):汇总高阶API在Host侧Tiling接口中使用的Ascend C公共类型,包含TensorShape、TensorDataType等数据结构定义。7+- [数据结构](data_structures/index.md):汇总高阶API在Host侧Tiling接口中使用的Ascend C公共类型,包含TensorShape、TensorDataType等数据结构定义。(共 3 张知识卡:TensorDataType、TensorShape)
8-- [HCCL通信类](hccl_communication/index.md):获取指定Index通信域的context(消息区)地址。8+- [HCCL通信类](hccl_communication/index.md):共 46 张知识卡,分布在 3 个子目录(hccl_context、hccl_kernel、hccl_tiling)。
9-- [索引计算](index_compute/index.md):给定起始值,等差值和长度,返回一个等差数列。9+- [索引计算](index_compute/index.md):共 2 张知识卡(如 Arange(索引计算)、GetArangeMaxMinTmpSize)。
10-- [数学计算](math_compute/index.md):下面的样例展示了数学库kernel侧API和Tiling API GetXxxMaxMinTmpSize的配套使用方法,具体流程如下:10+- [数学计算](math_compute/index.md):共 144 张知识卡(如 更多样例(数学计算) 等)。
11-- [归一化操作](normalization/index.md):对每一层输入做规范化处理,在batch维度上计算均值和方差进行归一化,再通过缩放和平移参数调整分布,加速训练并提升泛化能力。11+- [归一化操作](normalization/index.md):共 20 张知识卡(如 BatchNorm、BatchNorm Tiling、DeepNorm 等)。
12-- [量化操作](quantization/index.md):按元素做伪量化计算,将int8_t等低精度数据类型伪量化为half等高精度数据类型,支持PER_TENSOR、PER_CHANNEL、PER_TOKEN、PER_GROUP四种量化策略,最多支持二维输入。12+- [量化操作](quantization/index.md):共 18 张知识卡(如 AntiQuantize、AscendAntiQuant、AscendDequant 等)。
13- [随机函数](random_functions/index.md):基于Philox随机数生成算法,给定随机数种子,生成若干的随机数。 Philox随机数生成的核心算法是一个基于记数的伪随机数生成算法,输入为一个128bit的记数器C,两个32bit的key(k0和k1),输出为4个32bit的整数。13- [随机函数](random_functions/index.md):基于Philox随机数生成算法,给定随机数种子,生成若干的随机数。 Philox随机数生成的核心算法是一个基于记数的伪随机数生成算法,输入为一个128bit的记数器C,两个32bit的key(k0和k1),输出为4个32bit的整数。
14-- [归约操作](reduction_operations/index.md):在host侧获取Mean接口计算所需最大和最小临时空间大小,供Tiling参数选择。14+- [归约操作](reduction_operations/index.md):共 21 张知识卡,分布在 10 个子目录(mean_interface、reduce_all_interface、reduce_any_interface 等)。
15-- [排序操作](sort_operations/index.md):对数据进行预处理,将要排序的源操作数src一一对应的合入目标数据concat中,数据预处理完成后,可以进行Sort。15+- [排序操作](sort_operations/index.md):共 11 张知识卡(如 Concat(排序操作)、Extract(排序操作)、GetConcatTmpSize 等)。
16-- [张量变换](tensor_transform/index.md):将输入按照输出shape进行广播。 比如A的shape为\(2,1\),广播的目标shape为\(2,16\),则会将原来的一列扩展为相同的16列。16+- [张量变换](tensor_transform/index.md):共 11 张知识卡(如 Broadcast(张量变换)、Fill(张量变换)、GetBroadCastMaxMinTmpSize 等)。
17 17 
18-- [experimental](experimental/index.md):本接口为试验接口,在后续版本中可能会调整或改进,不保证后续兼容性。请开发者在使用过程中关注后续版本更新。18+- [experimental](experimental/index.md):共 4 张知识卡,分布在 2 个子目录(bessel_i0_interface、ndtri_interface)。
@@ -1,51 +1,51 @@
1# 数学计算1# 数学计算
2 2 
3-- [Acos接口](acos_interface/index.md):按元素做反余弦函数计算,输入值域需在-1到1范围内3+- [Acos接口](acos_interface/index.md):共 3 张知识卡(如 Acos、GetAcosMaxMinTmpSize、GetAcosTmpBufferFactorSize)。
4-- [Acosh接口](acosh_interface/index.md):按元素做双曲反余弦函数计算,输入值域需在1到65504范围内4+- [Acosh接口](acosh_interface/index.md):共 3 张知识卡(如 Acosh、GetAcoshMaxMinTmpSize、GetAcoshTmpBufferFactorSize)。
5-- [Asin接口](asin_interface/index.md):按元素做反正弦函数计算,输入值域需在-1到1范围内5+- [Asin接口](asin_interface/index.md):共 3 张知识卡(如 Asin、GetAsinMaxMinTmpSize、GetAsinTmpBufferFactorSize)。
6-- [Asinh接口](asinh_interface/index.md):按元素做反双曲正弦函数计算,输入值域需在-65504到65504范围内6+- [Asinh接口](asinh_interface/index.md):共 3 张知识卡(如 Asinh、GetAsinhMaxMinTmpSize、GetAsinhTmpBufferFactorSize)。
7-- [Atan接口](atan_interface/index.md):按元素做反正切运算,支持泰勒展开和多项式逼近两种算法配置7+- [Atan接口](atan_interface/index.md):共 3 张知识卡(如 Atan、GetAtanMaxMinTmpSize、GetAtanTmpBufferFactorSize)。
8-- [Atanh接口](atanh_interface/index.md):按元素做反双曲正切函数计算,输入值域需在指定区间内8+- [Atanh接口](atanh_interface/index.md):共 3 张知识卡(如 Atanh、GetAtanhMaxMinTmpSize、GetAtanhTmpBufferFactorSize)。
9-- [Axpy接口](axpy_interface/index.md):源操作数各元素与标量求积后加到目的操作数对应元素,通过Muls和Add组合提升精度9+- [Axpy接口](axpy_interface/index.md):共 3 张知识卡(如 Axpy、GetAxpyMaxMinTmpSize、GetAxpyTmpBufferFactorSize)。
10-- [BitwiseAnd接口](bitwise_and_interface/index.md):逐比特对两个输入进行与操作。 !10+- [BitwiseAnd接口](bitwise_and_interface/index.md):共 3 张知识卡(如 BitwiseAnd、GetBitwiseAndMaxMinTmpSize、GetBitwiseAndTmpBufferFactorSize)。
11-- [BitwiseNot接口](bitwise_not_interface/index.md):逐比特对输入进行取反。 !11+- [BitwiseNot接口](bitwise_not_interface/index.md):共 3 张知识卡(如 BitwiseNot、GetBitwiseNotMaxMinTmpSize、GetBitwiseNotTmpBufferFactorSize)。
12-- [BitwiseOr接口](bitwise_or_interface/index.md):逐比特对两个输入进行或操作。 !12+- [BitwiseOr接口](bitwise_or_interface/index.md):共 3 张知识卡(如 BitwiseOr、GetBitwiseOrMaxMinTmpSize、GetBitwiseOrTmpBufferFactorSize)。
13-- [BitwiseXor接口](bitwise_xor_interface/index.md):逐比特对两个输入进行异或操作。 !13+- [BitwiseXor接口](bitwise_xor_interface/index.md):共 3 张知识卡(如 BitwiseXor、GetBitwiseXorMaxMinTmpSize、GetBitwiseXorTmpBufferFactorSize)。
14-- [Ceil接口](ceil_interface/index.md):获取大于或等于输入值的最小整数,即向正无穷方向取整14+- [Ceil接口](ceil_interface/index.md):共 3 张知识卡(如 Ceil、GetCeilMaxMinTmpSize、GetCeilTmpBufferFactorSize)。
15-- [Clamp接口](clamp_interface/index.md):将输入中除nan值以外大于max的数替换为max,小于min的数替换为min,小于等于max和大于等于min的数保持不变,作为输出。当min大于max时,将除nan值外所有值替换为max。min和max可以为标量或LocalTensor。 ! !15+- [Clamp接口](clamp_interface/index.md):共 5 张知识卡(如 Clamp(数学计算)、ClampMax、ClampMin 等)。
16-- [Cos接口](cos_interface/index.md):按元素做余弦运算,支持多项式逼近和周期约减两种算法配置16+- [Cos接口](cos_interface/index.md):共 3 张知识卡(如 Cos、GetCosMaxMinTmpSize、GetCosTmpBufferFactorSize)。
17-- [Cosh接口](cosh_interface/index.md):按元素做双曲余弦函数计算17+- [Cosh接口](cosh_interface/index.md):共 3 张知识卡(如 Cosh、GetCoshMaxMinTmpSize、GetCoshTmpBufferFactorSize)。
18-- [CumSum接口](cum_sum_interface/index.md):对输入二维张量按行或列进行累加和操作,输出每个位置及之前所有行或列元素的累加结果,支持逐行和Sklansky二分两种算法。18+- [CumSum接口](cum_sum_interface/index.md):共 2 张知识卡(如 CumSum、GetCumSumMaxMinTmpSize)。
19-- [Digamma接口](digamma_interface/index.md):按元素计算x的gamma函数的对数导数19+- [Digamma接口](digamma_interface/index.md):共 3 张知识卡(如 Digamma、GetDigammaMaxMinTmpSize、GetDigammaTmpBufferFactorSize)。
20-- [Erf接口](erf_interface/index.md):按元素做误差函数计算,也称为高斯误差函数20+- [Erf接口](erf_interface/index.md):共 3 张知识卡(如 Erf、GetErfMaxMinTmpSize、GetErfTmpBufferFactorSize)。
21-- [Erfc接口](erfc_interface/index.md):返回输入x的互补误差函数结果,积分区间为x到无穷大21+- [Erfc接口](erfc_interface/index.md):共 3 张知识卡(如 Erfc、GetErfcMaxMinTmpSize、GetErfcTmpBufferFactorSize)。
22-- [Exp接口](exp_interface/index.md):按元素取自然指数,可选择是否使用泰勒展开公式进行计算22+- [Exp接口](exp_interface/index.md):共 3 张知识卡(如 Exp、GetExpMaxMinTmpSize、GetExpTmpBufferFactorSize)。
23-- [Floor接口](floor_interface/index.md):获取小于或等于x的最小整数值,即向负无穷取整23+- [Floor接口](floor_interface/index.md):共 3 张知识卡(如 Floor、GetFloorMaxMinTmpSize、GetFloorTmpBufferFactorSize)。
24-- [Fma接口](fma_interface/index.md):按元素计算两个输入相乘后与第三个输入相加的结果24+- [Fma接口](fma_interface/index.md):共 3 张知识卡(如 Fma、GetFmaMaxMinTmpSize、GetFmaTmpBufferFactorSize)。
25-- [Fmod接口](fmod_interface/index.md):按元素计算两个浮点数相除后的余数,Trunc为向零取整25+- [Fmod接口](fmod_interface/index.md):共 3 张知识卡(如 Fmod、GetFmodMaxMinTmpSize、GetFmodTmpBufferFactorSize)。
26-- [Frac接口](frac_interface/index.md):按元素做取小数计算,返回输入数据的小数部分,支持half和float数据类型。26+- [Frac接口](frac_interface/index.md):共 3 张知识卡(如 Frac、GetFracMaxMinTmpSize、GetFracTmpBufferFactorSize)。
27-- [Hypot接口](hypot_interface/index.md):在host侧获取Hypot接口预留或申请临时空间的最大和最小值,用于Tiling参数传递27+- [Hypot接口](hypot_interface/index.md):共 3 张知识卡(如 GetHypotMaxMinTmpSize、GetHypotTmpBufferFactorSize、Hypot)。
28- [IsFinite接口](is_finite_interface/index.md):按元素判断输入浮点数是否非NAN、非±INF,输出结果为浮点数或布尔值,支持half、bfloat16_t、float数据类型。28- [IsFinite接口](is_finite_interface/index.md):按元素判断输入浮点数是否非NAN、非±INF,输出结果为浮点数或布尔值,支持half、bfloat16_t、float数据类型。
29-- [IsInf接口](is_inf_interface/index.md):在host侧获取IsInf接口预留或申请临时空间的最大和最小值,用于Tiling参数传递29+- [IsInf接口](is_inf_interface/index.md):共 3 张知识卡(如 GetIsInfMaxMinTmpSize、GetIsInfTmpBufferFactorSize、IsInf)。
30-- [IsNan接口](is_nan_interface/index.md):在host侧获取IsNan接口预留或申请临时空间的最大和最小值,用于Tiling参数传递30+- [IsNan接口](is_nan_interface/index.md):共 3 张知识卡(如 GetIsNanMaxMinTmpSize、GetIsNanTmpBufferFactorSize、IsNan)。
31-- [Lgamma接口](lgamma_interface/index.md):在host侧获取Lgamma接口预留或申请临时空间的最大和最小值,用于Tiling参数传递31+- [Lgamma接口](lgamma_interface/index.md):共 3 张知识卡(如 GetLgammaMaxMinTmpSize、GetLgammaTmpBufferFactorSize、Lgamma)。
32-- [Log接口](log_interface/index.md):Host侧接口,用于获取Log接口能完成计算所需最小的临时空间大小,此空间为预留空间,即需要保证预留有足够的物理空间,用于执行计算。32+- [Log接口](log_interface/index.md):共 3 张知识卡(如 GetLogMaxMinTmpSize、GetLogTmpBufferFactorSize、Log)。
33-- [LogicalAnd接口](logical_and_interface/index.md):在host侧获取LogicalAnd接口所需的最大和最小临时空间大小,供开发者选择合适的Tiling参数传入kernel侧。33+- [LogicalAnd接口](logical_and_interface/index.md):共 3 张知识卡(如 GetLogicalAndMaxMinTmpSize、GetLogicalAndTmpBufferFactorSize、LogicalAnd)。
34-- [LogicalAnds接口](logical_ands_interface/index.md):在host侧获取LogicalAnds接口预留或申请临时空间的最大和最小值,用于Tiling参数传递34+- [LogicalAnds接口](logical_ands_interface/index.md):共 3 张知识卡(如 GetLogicalAndsMaxMinTmpSize、GetLogicalAndsTmpBufferFactorSize、LogicalAnds)。
35-- [LogicalNot接口](logical_not_interface/index.md):在host侧获取LogicalNot接口所需的最大和最小临时空间大小,供开发者选择合适的Tiling参数传入kernel侧。35+- [LogicalNot接口](logical_not_interface/index.md):共 3 张知识卡(如 GetLogicalNotMaxMinTmpSize、GetLogicalNotTmpBufferFactorSize、LogicalNot)。
36-- [LogicalOr接口](logical_or_interface/index.md):在host侧获取LogicalOr接口所需的最大和最小临时空间大小,供开发者选择合适的Tiling参数传入kernel侧。36+- [LogicalOr接口](logical_or_interface/index.md):共 3 张知识卡(如 GetLogicalOrMaxMinTmpSize、GetLogicalOrTmpBufferFactorSize、LogicalOr)。
37-- [LogicalOrs接口](logical_ors_interface/index.md):在host侧获取LogicalOrs接口预留或申请临时空间的最大和最小值,用于Tiling参数传递37+- [LogicalOrs接口](logical_ors_interface/index.md):共 3 张知识卡(如 GetLogicalOrsMaxMinTmpSize、GetLogicalOrsTmpBufferFactorSize、LogicalOrs)。
38-- [LogicalXor接口](logical_xor_interface/index.md):在host侧获取LogicalXor接口所需的最大和最小临时空间大小,供开发者选择合适的Tiling参数传入kernel侧。38+- [LogicalXor接口](logical_xor_interface/index.md):共 3 张知识卡(如 GetLogicalXorMaxMinTmpSize、GetLogicalXorTmpBufferFactorSize、LogicalXor)。
39- [更多样例(数学计算)](more_examples.md):下面的样例展示了数学库kernel侧API和Tiling API GetXxxMaxMinTmpSize的配套使用方法,具体流程如下:39- [更多样例(数学计算)](more_examples.md):下面的样例展示了数学库kernel侧API和Tiling API GetXxxMaxMinTmpSize的配套使用方法,具体流程如下:
40-- [Power接口](power_interface/index.md):在host侧获取Power接口所需的最大和最小临时空间大小,根据输入shape自动适配不同Power调用形式。40+- [Power接口](power_interface/index.md):共 3 张知识卡(如 GetPowerMaxMinTmpSize、GetPowerTmpBufferFactorSize、Power(数学计算))。
41-- [Rint接口](rint_interface/index.md):在host侧获取Rint接口预留或申请临时空间的最大和最小值,用于Tiling参数传递41+- [Rint接口](rint_interface/index.md):共 3 张知识卡(如 GetRintMaxMinTmpSize、GetRintTmpBufferFactorSize、Rint)。
42-- [Round接口](round_interface/index.md):在host侧获取Round接口预留或申请临时空间的最大和最小值,用于Tiling参数传递42+- [Round接口](round_interface/index.md):共 3 张知识卡(如 GetRoundMaxMinTmpSize、GetRoundTmpBufferFactorSize、Round)。
43-- [Sign接口](sign_interface/index.md):在host侧获取Sign接口计算所需最大和最小临时空间大小,供Tiling参数选择。43+- [Sign接口](sign_interface/index.md):共 3 张知识卡(如 GetSignMaxMinTmpSize、GetSignTmpBufferFactorSize、Sign)。
44-- [SinCos接口](sin_cos_interface/index.md):在host侧获取SinCos接口所需的最大和最小临时空间大小,供开发者选择合适的Tiling参数传入kernel侧。44+- [SinCos接口](sin_cos_interface/index.md):共 3 张知识卡(如 GetSinCosMaxMinTmpSize、GetSinCosTmpBufferFactorSize、SinCos)。
45-- [Sin接口](sin_interface/index.md):在host侧获取Sin接口计算所需最大和最小临时空间大小,供Tiling参数选择。45+- [Sin接口](sin_interface/index.md):共 3 张知识卡(如 GetSinMaxMinTmpSize、GetSinTmpBufferFactorSize、Sin)。
46-- [Sinh接口](sinh_interface/index.md):在host侧获取Sinh接口计算所需最大和最小临时空间大小,供Tiling参数选择。46+- [Sinh接口](sinh_interface/index.md):共 3 张知识卡(如 GetSinhMaxMinTmpSize、GetSinhTmpBufferFactorSize、Sinh)。
47-- [Tan接口](tan_interface/index.md):在host侧获取Tan接口计算所需最大和最小临时空间大小,供Tiling参数选择。47+- [Tan接口](tan_interface/index.md):共 3 张知识卡(如 GetTanMaxMinTmpSize、GetTanTmpBufferFactorSize、Tan)。
48-- [Tanh接口](tanh_interface/index.md):在host侧获取Tanh接口计算所需最大和最小临时空间大小,供Tiling参数选择。48+- [Tanh接口](tanh_interface/index.md):共 3 张知识卡(如 GetTanhMaxMinTmpSize、GetTanhTmpBufferFactorSize、Tanh)。
49-- [Trunc接口](trunc_interface/index.md):在host侧获取Trunc接口计算所需最大和最小临时空间大小,供Tiling参数选择。49+- [Trunc接口](trunc_interface/index.md):共 3 张知识卡(如 GetTruncMaxMinTmpSize、GetTruncTmpBufferFactorSize、Trunc)。
50-- [Where接口](where_interface/index.md):在host侧获取Where接口所需的最大和最小临时空间大小,供开发者选择合适的Tiling参数传入kernel侧。50+- [Where接口](where_interface/index.md):共 3 张知识卡(如 GetWhereMaxMinTmpSize、GetWhereTmpBufferFactorSize、Where(数学计算))。
51-- [Xor接口](xor_interface/index.md):在host侧获取Xor接口计算所需最大和最小临时空间大小,供Tiling参数选择。51+- [Xor接口](xor_interface/index.md):共 3 张知识卡(如 GetXorMaxMinTmpSize、GetXorTmpBufferFactorSize、Xor)。
@@ -1,12 +1,12 @@
1# 归约操作1# 归约操作
2 2 
3-- [Mean接口](mean_interface/index.md):在host侧获取Mean接口计算所需最大和最小临时空间大小,供Tiling参数选择。3+- [Mean接口](mean_interface/index.md):共 3 张知识卡(如 GetMeanMaxMinTmpSize、GetMeanTmpBufferFactorSize、Mean)。
4-- [ReduceAll接口](reduce_all_interface/index.md):在host侧获取ReduceAll接口计算所需的最大和最小临时空间大小,用于选择合适的空间作为Tiling参数传递到kernel侧。4+- [ReduceAll接口](reduce_all_interface/index.md):共 2 张知识卡(如 GetReduceAllMaxMinTmpSize、ReduceAll)。
5-- [ReduceAny接口](reduce_any_interface/index.md):在host侧获取ReduceAny接口计算所需的最大和最小临时空间大小,用于选择合适的空间作为Tiling参数传递到kernel侧。5+- [ReduceAny接口](reduce_any_interface/index.md):共 2 张知识卡(如 GetReduceAnyMaxMinTmpSize、ReduceAny)。
6-- [ReduceMax接口](reduce_max_interface/index.md):在host侧获取ReduceMax接口所需的最大和最小临时空间大小,当前最大值等于最小值,供开发者选择Tiling参数。6+- [ReduceMax接口](reduce_max_interface/index.md):共 2 张知识卡(如 GetReduceMaxMaxMinTmpSize、ReduceMax)。
7-- [ReduceMean接口](reduce_mean_interface/index.md):在host侧获取ReduceMean接口所需的最大和最小临时空间大小,当前最大值等于最小值,供开发者选择Tiling参数。7+- [ReduceMean接口](reduce_mean_interface/index.md):共 2 张知识卡(如 GetReduceMeanMaxMinTmpSize、ReduceMean)。
8-- [ReduceMin接口](reduce_min_interface/index.md):在host侧获取ReduceMin接口计算所需的最大和最小临时空间大小,用于选择合适的空间作为Tiling参数传递到kernel侧。8+- [ReduceMin接口](reduce_min_interface/index.md):共 2 张知识卡(如 GetReduceMinMaxMinTmpSize、ReduceMin)。
9-- [ReduceProd接口](reduce_prod_interface/index.md):在host侧获取ReduceProd接口计算所需的最大和最小临时空间大小,用于选择合适的空间作为Tiling参数传递到kernel侧。9+- [ReduceProd接口](reduce_prod_interface/index.md):共 2 张知识卡(如 GetReduceProdMaxMinTmpSize、ReduceProd)。
10-- [ReduceSum接口](reduce_sum_interface/index.md):在host侧获取ReduceSum接口所需的最大和最小临时空间大小,当前最大值等于最小值,供开发者选择Tiling参数。10+- [ReduceSum接口](reduce_sum_interface/index.md):共 2 张知识卡(如 GetReduceSumMaxMinTmpSize、ReduceSum)。
11-- [ReduceXorSum接口](reduce_xor_sum_interface/index.md):在host侧获取ReduceXorSum接口计算所需的最大和最小临时空间大小,用于选择合适的空间作为Tiling参数传递到kernel侧。11+- [ReduceXorSum接口](reduce_xor_sum_interface/index.md):共 2 张知识卡(如 GetReduceXorSumMaxMinTmpSize、ReduceXorSum)。
12-- [Sum接口](sum_interface/index.md):在host侧获取Sum接口所需的最大和最小临时空间大小,当前最大值等于最小值,供开发者选择Tiling参数。12+- [Sum接口](sum_interface/index.md):共 2 张知识卡(如 GetSumMaxMinTmpSize、Sum)。
@@ -1,11 +1,11 @@
1# 辅助数据结构1# 辅助数据结构
2 2 
3- [BinaryRepeatParams](binary_repeat_params.md):BinaryRepeatParams为用于控制操作数地址步长的数据结构。结构体内包含操作数相邻迭代间相同DataBlock的地址步长,操作数同一迭代内不同DataBlock的地址步长等参数。3- [BinaryRepeatParams](binary_repeat_params.md):BinaryRepeatParams为用于控制操作数地址步长的数据结构。结构体内包含操作数相邻迭代间相同DataBlock的地址步长,操作数同一迭代内不同DataBlock的地址步长等参数。
4-- [Coordinate](coordinate/index.md):Coordinate本质上是一个元组(tuple),用于表示张量在不同维度的位置信息,即坐标值。Coordinate(坐标)和Layout(布局)、Index(内存位置索引)之间存在紧密的关联:4+- [Coordinate](coordinate/index.md):共 3 张知识卡(如 Coordinate简介、Crd2Idx、MakeCoord)。
5-- [Layout](layout/index.md):返回描述张量形状的Shape对象。5+- [Layout](layout/index.md):共 10 张知识卡(如 GetShape、GetStride、is_layout 等)。
6- [ListTensorDesc](list_tensor_desc.md):ListTensorDesc用来解析符合以下内存排布格式的数据,并在kernel侧根据索引获取储存对应数据的地址及shape信息。 !6- [ListTensorDesc](list_tensor_desc.md):ListTensorDesc用来解析符合以下内存排布格式的数据,并在kernel侧根据索引获取储存对应数据的地址及shape信息。 !
7- [ShapeInfo(辅助数据结构)](shape_info.md):ShapeInfo用来存放LocalTensor或GlobalTensor的shape信息。7- [ShapeInfo(辅助数据结构)](shape_info.md):ShapeInfo用来存放LocalTensor或GlobalTensor的shape信息。
8- [TPosition](t_position.md):Ascend C管理不同层级物理内存的抽象逻辑位置,代替片上物理存储概念,包括矢量编程和矩阵编程所需的各类存储位置。8- [TPosition](t_position.md):Ascend C管理不同层级物理内存的抽象逻辑位置,代替片上物理存储概念,包括矢量编程和矩阵编程所需的各类存储位置。
9-- [TensorDesc](tensor_desc/index.md):创建与释放TensorDesc对象。9+- [TensorDesc](tensor_desc/index.md):共 8 张知识卡(如 构造和析构函数、GetDataObj、GetDataPtr 等)。
10-- [TensorTrait](tensor_trait/index.md):获取Layout实例化对象。10+- [TensorTrait](tensor_trait/index.md):共 7 张知识卡(如 GetLayout、is_tensorTrait、MakeTensorTrait 等)。
11- [UnaryRepeatParams](unary_repeat_params.md):UnaryRepeatParams为用于控制操作数地址步长的数据结构。结构体内包含操作数相邻迭代间相同DataBlock的地址步长,操作数同一迭代内不同DataBlock的地址步长等参数。11- [UnaryRepeatParams](unary_repeat_params.md):UnaryRepeatParams为用于控制操作数地址步长的数据结构。结构体内包含操作数相邻迭代间相同DataBlock的地址步长,操作数同一迭代内不同DataBlock的地址步长等参数。
@@ -1,12 +1,12 @@
1# 矩阵计算(ISASI)1# 矩阵计算(ISASI)
2 2 
3-- [矩阵计算分形介绍](cube_compute_fractal_intro/index.md):针对Ascend 950PR/Ascend 950DT产品:3+- [矩阵计算分形介绍](cube_compute_fractal_intro/index.md):共 3 张知识卡(如 辅助矩阵分形格式详解、背景与核心概念、关键分形格式详解)。
4-- [矩阵计算的搬入](cube_compute_load/index.md):将矢量数据广播到矩阵中,每个数据块中的每16个元素会被连续复制16次;当前支持的数据传输通路:Unified Buffer-\>L0C Buffer(VECIN/VECCALC/VECOUT-\>CO1)。 图1 功能示例 !4+- [矩阵计算的搬入](cube_compute_load/index.md):共 26 张知识卡(如 BroadCastVecToMMISASI、矩阵计算输入搬运约束、DataCopy 等)。
5-- [矩阵计算的搬出](cube_compute_store/index.md):将L0C Buffer中的矩阵计算结果搬运至Global Memory,支持随路量化、ReLU激活和NZ2ND格式转换。5+- [矩阵计算的搬出](cube_compute_store/index.md):共 11 张知识卡(如 DataCopyL0CToGM、DataCopyL0CToL1、DataCopy(L0C Buffer到UB数据搬运) 等)。
6-- [矩阵搬入辅助配置接口](cube_load_aux_config/index.md):将L1 Buffer/L0A Buffer/L0B Buffer上的LocalTensor初始化为指定数值,通过InitConstValueParams配置迭代次数、数据块个数和填充值。6+- [矩阵搬入辅助配置接口](cube_load_aux_config/index.md):共 11 张知识卡(如 Fill、LoadDataUnzip、LoadImageToLocal 等)。
7-- [矩阵搬出寄存器配置说明](cube_store_aux_config/index.md):本章节涉及寄存器在相同核函数内部会维持状态直到被人为改变,而在不同核函数内部不会继承之前的状态。7+- [矩阵搬出寄存器配置说明](cube_store_aux_config/index.md):本章节涉及寄存器在相同核函数内部会维持状态直到被人为改变,而在不同核函数内部不会继承之前的状态。(共 6 张知识卡:SetFixPipeAddr、SetFixPipeClipRelu、SetFixPipeConfig 等)
8-- [矩阵搬出关键特性说明](cube_store_key_features/index.md):矩阵计算搬出过程中支持随路cast和quant能力,包括Scalar和Tensor两种量化模式,支持多种量化类型。8+- [矩阵搬出关键特性说明](cube_store_key_features/index.md):共 10 张知识卡(如 AccompanyingQuantization、随路量化与随路ReLU场景组合、随路ReLU 等)。
9-- [矩阵计算](mmad_compute/index.md):该接口废弃,并将在后续版本移除,请不要使用该接口。 计算给定输入张量和权重张量的2-D卷积,输出结果张量。Conv2D卷积层多用于图像识别,使用过滤器提取图像中的特征。9+- [矩阵计算](mmad_compute/index.md):共 6 张知识卡(如 Conv2D(废弃)、Gemm(废弃)、Mmad 等)。
10-- [Mmad辅助配置说明](mmad_compute_aux_config/index.md):用于设置Mmad计算是否开启HF32模式,开启该模式后L0A Buffer/L0B Buffer中的FP32数据将在参与Mmad计算之前被舍入为HF32。10+- [Mmad辅助配置说明](mmad_compute_aux_config/index.md):共 3 张知识卡(如 SetHF32Mode、SetHF32TransMode、SetMMColumnMajor/SetMMRowMajor)。
11-- [Mmad计算关键特性说明](mmad_compute_key_features/index.md):Gemv的核心功能体现为:当矩阵A的维度M取值为1时,接口会自动启用Gemv功能,该操作退化为1×K维度的行向量与K×N维度矩阵之间的乘法运算。11+- [Mmad计算关键特性说明](mmad_compute_key_features/index.md):共 4 张知识卡(如 Cube计算(ISASI)GEMV、Cube计算(ISASI)HF32、K方向对齐约束 等)。
12-- [概述](overview/index.md):Ascend C针对矩阵计算编程模型提供了四类接口,分别承载Cube核中各个通路的搬运能力和计算能力,如下图所示:12+- [概述](overview/index.md):共 2 张知识卡(如 普通矩阵计算流程(概述)、矩阵计算单元(概述))。
@@ -1,5 +1,5 @@
1# Cube分组管理(ISASI)1# Cube分组管理(ISASI)
2 2 
3-- [CubeResGroupHandle](cube_res_group_handle/index.md):AIV从消息队列里申请消息空间,用于存放消息结构体,返回当前申请的消息空间的地址。消息队列的深度固定为4,申请消息空间的顺序为自上而下,然后循环。当消息队列指针指向的消息空间为FREE状态时,AllocMessage返回空间的地址,否则循环等待,直到当前空间的状态为FREE。 图1 AllocMessage示意图 !3+- [CubeResGroupHandle](cube_res_group_handle/index.md):共 11 张知识卡(如 AllocMessage、AssignQueue、CreateCubeResGroup 等)。
4-- [GroupBarrier](group_barrier/index.md):通知其他等待的AIV,本AIV已经完成其依赖的任务。4+- [GroupBarrier](group_barrier/index.md):共 5 张知识卡(如 Arrive(Cube核间管理)、GetWorkspaceLen、GroupBarrier构造函数 等)。
5-- [KfcWorkspace](kfc_workspace/index.md):用于创建KfcWorkspace对象,Kfc全称为kernel function call,表示核间通信调用。5+- [KfcWorkspace](kfc_workspace/index.md):共 3 张知识卡(如 构造函数与析构函数、GetKfcWorkspace、UpdateKfcWorkspace)。
@@ -6,5 +6,5 @@
6- [L1与Fixpipe Buffer数据搬运](l1_fixpipe_data_move.md):DataCopy能够实现L1 Buffer至Fixpipe Buffer的连续数据搬运,数据搬运时格式和内容保持不变。6- [L1与Fixpipe Buffer数据搬运](l1_fixpipe_data_move.md):DataCopy能够实现L1 Buffer至Fixpipe Buffer的连续数据搬运,数据搬运时格式和内容保持不变。
7- [L1与L0数据搬运](l1_l0_data_move.md):L1 Buffer与L0 Buffer之间的数据搬运提供了灵活的接口支持,能够充分适配多样化的计算场景(见表1)。其中,L1 Buffer-\>L0A Buffer、L1 Buffer-\>L0B Buffer的数据搬运用于矩阵计算(参考矩阵计算的搬入),L0C Buffer-\>L1 Buffer用于矩阵计算的搬出,完成矩阵计算的输出流程(参考矩阵计算的搬出)。7- [L1与L0数据搬运](l1_l0_data_move.md):L1 Buffer与L0 Buffer之间的数据搬运提供了灵活的接口支持,能够充分适配多样化的计算场景(见表1)。其中,L1 Buffer-\>L0A Buffer、L1 Buffer-\>L0B Buffer的数据搬运用于矩阵计算(参考矩阵计算的搬入),L0C Buffer-\>L1 Buffer用于矩阵计算的搬出,完成矩阵计算的输出流程(参考矩阵计算的搬出)。
8- [L1OrL0CUbDataMove](l1_or_l0_c_ub_data_move.md):提供L1 Buffer及L0C Buffer与Unified Buffer间数据搬运接口概述,涵盖连续搬运、高维切分、格式转换及随路量化激活等场景。8- [L1OrL0CUbDataMove](l1_or_l0_c_ub_data_move.md):提供L1 Buffer及L0C Buffer与Unified Buffer间数据搬运接口概述,涵盖连续搬运、高维切分、格式转换及随路量化激活等场景。
9-- [概览](overview/index.md):数据搬运类接口提供了全面的数据搬运功能,覆盖从最基础的连续搬运到复杂的随路转换和量化激活。按照功能复杂度递增,搬运能力可以分为以下几个类别:基础数据搬运、高维切分搬运、切片搬运、随路格式转换搬运、随路量化激活搬运、非对齐搬运、UB内部搬运(Copy)、矩阵分形搬运和多维数据搬运NDDMA。每一类在后续章节中都有对应的详细接口说明。9+- [概览](overview/index.md):数据搬运导览,介绍AI Core分层存储架构中的数据通路、TPosition映射、搬运功能分类及通用约束等核心概念。(共 4 张知识卡:数据搬运概念(概述)、数据通路(概述)、总体约束说明(概述))
10- [UB与UB数据搬运](ub_ub_data_move.md):Unified Buffer(UB)之间的数据搬运提供了灵活的接口支持,共有两组接口:DataCopy提供了基础的数据搬运能力,支持连续和非连续数据搬运;Copy除了支持连续和非连续数据搬运,还提供了mask操作和DataBlock间隔操作的能力,支持在接口内部和接口外部设置mask。具体接口功能如下表所示。10- [UB与UB数据搬运](ub_ub_data_move.md):Unified Buffer(UB)之间的数据搬运提供了灵活的接口支持,共有两组接口:DataCopy提供了基础的数据搬运能力,支持连续和非连续数据搬运;Copy除了支持连续和非连续数据搬运,还提供了mask操作和DataBlock间隔操作的能力,支持在接口内部和接口外部设置mask。具体接口功能如下表所示。
@@ -1,6 +1,6 @@
1# 数据结构1# 数据结构
2 2 
3- [内置数据类型(数据结构)](builtin_data_types.md):Ascend C提供b8\~b64(8bit\~64bit)四种不同位宽的数据类型,不同位宽对应的数据类型如下表所示。3- [内置数据类型(数据结构)](builtin_data_types.md):Ascend C提供b8\~b64(8bit\~64bit)四种不同位宽的数据类型,不同位宽对应的数据类型如下表所示。
4-- [GlobalTensor](global_tensor/index.md):GlobalTensor::GetPhyAddr 接口,用于获取全局数据的物理地址,支持全平台。4+- [GlobalTensor](global_tensor/index.md):共 13 张知识卡(如 GetPhyAddr、GetShapeInfo、GetSize 等)。
5-- [LocalTensor](local_tensor/index.md):获取LocalTensor数据长度。5+- [LocalTensor](local_tensor/index.md):共 20 张知识卡(如 GetLength、GetPhyAddr、GetPosition 等)。
6- [精度转换(数据结构)](precision_conversion.md):在了解不同类型之间的精度转换规则之前,需先了解二进制的精度舍入规则。6- [精度转换(数据结构)](precision_conversion.md):在了解不同类型之间的精度转换规则之前,需先了解二进制的精度舍入规则。
@@ -1,7 +1,7 @@
1# 调试接口1# 调试接口
2 2 
3-- [CPU孪生调试](cpu_twin_debug/index.md):CPU侧运行验证时创建共享内存,在/tmp目录下创建共享文件并返回映射指针。3+- [CPU孪生调试](cpu_twin_debug/index.md):共 5 张知识卡(如 GmAlloc、GmFree、IcpuRunKf 等)。
4- [关闭ASCENDC_DUMP说明](disable_ascendc_dump.md):调试接口(如本章的上板打印、异常检测接口)开启后会对算子性能产生一定影响,通常在调测阶段使用,生产环境建议关闭。ASCENDC_DUMP宏定义是控制调试接口打印的开关。4- [关闭ASCENDC_DUMP说明](disable_ascendc_dump.md):调试接口(如本章的上板打印、异常检测接口)开启后会对算子性能产生一定影响,通常在调测阶段使用,生产环境建议关闭。ASCENDC_DUMP宏定义是控制调试接口打印的开关。
5-- [异常检测](exception_detection/index.md):检查设定范围内UB读写行为的调试接口,有设定范围的读写行为会触发EXCEPTION报错。5+- [异常检测](exception_detection/index.md):共 2 张知识卡(如 CheckLocalMemoryIA、Trap)。
6-- [上板打印](onboard_print/index.md):打印指定Tensor内容并支持自定义标签,区别于DumpTensor,支持指定偏移位置开始打印Tensor数据。6+- [上板打印](onboard_print/index.md):共 3 张知识卡(如 DumpAccChkPoint、DumpTensor、PrintTimeStamp)。
7-- [性能统计](performance_stats/index.md):推荐使用asc_mark_stamp接口进行特定位置的标记,该接口同时适用于C语言和C++语言编程。 用户通过调用接口,用于在算子执行过程中标记特定位置,便于后期通过流水图分析代码执行路径与性能热点。7+- [性能统计](performance_stats/index.md):共 3 张知识卡(如 MarkStamp、MetricsProfStart、MetricsProfStop)。
@@ -1,20 +1,20 @@
1# 基础API1# 基础API
2 2 
3-- [原子操作](atomic_operations/index.md):对GM地址上的值执行原子加操作,将旧值与输入标量求和后写回,返回修改前的旧值。3+- [原子操作](atomic_operations/index.md):共 14 张知识卡(如 AtomicAdd、AtomicCas、AtomicExch 等)。
4-- [辅助数据结构](aux_data_structures/index.md):BinaryRepeatParams为用于控制操作数地址步长的数据结构。结构体内包含操作数相邻迭代间相同DataBlock的地址步长,操作数同一迭代内不同DataBlock的地址步长等参数。4+- [辅助数据结构](aux_data_structures/index.md):共 33 张知识卡(如 BinaryRepeatParams、ListTensorDesc、ShapeInfo(辅助数据结构) 等)。
5-- [缓存控制](cache_control/index.md):解决Data Cache与GM一致性问题,将脏数据写回GM并将指定Cache Line标记为无效,确保下次访问从GM重新加载。5+- [缓存控制](cache_control/index.md):共 6 张知识卡(如 DataCacheCleanAndInvalid、DataCachePreload、GetICachePreloadStatus 等)。
6-- [矩阵计算(ISASI)](cube_compute_isasi/index.md):针对Ascend 950PR/Ascend 950DT产品:6+- [矩阵计算(ISASI)](cube_compute_isasi/index.md):共 82 张知识卡,分布在 10 个子目录(cube_compute_fractal_intro、cube_compute_load、cube_compute_store 等)。
7-- [Cube分组管理(ISASI)](cube_group_mgmt_isasi/index.md):AIV从消息队列里申请消息空间,用于存放消息结构体,返回当前申请的消息空间的地址。消息队列的深度固定为4,申请消息空间的顺序为自上而下,然后循环。当消息队列指针指向的消息空间为FREE状态时,AllocMessage返回空间的地址,否则循环等待,直到当前空间的状态为FREE。 图1 AllocMessage示意图 !7+- [Cube分组管理(ISASI)](cube_group_mgmt_isasi/index.md):共 19 张知识卡,分布在 3 个子目录(cube_res_group_handle、group_barrier、kfc_workspace)。
8-- [数据搬运导览](data_move_guide/index.md):提供Global Memory与L1 Buffer及L0 Buffer间数据搬运接口概述,涵盖连续搬运、高维切分、格式转换及矩阵计算结果搬出等场景。8+- [数据搬运导览](data_move_guide/index.md):共 11 张知识卡(如 GmL1OrL0DataMove、GmUbDataMove、L1与BiasTable Buffer数据搬运 等)。
9-- [数据结构](data_structures/index.md):Ascend C提供b8\~b64(8bit\~64bit)四种不同位宽的数据类型,不同位宽对应的数据类型如下表所示。9+- [数据结构](data_structures/index.md):共 35 张知识卡(如 内置数据类型(数据结构)、精度转换(数据结构) 等)。
10-- [调试接口](debug_interface/index.md):调试接口(如本章的上板打印、异常检测接口)开启后会对算子性能产生一定影响,通常在调测阶段使用,生产环境建议关闭。ASCENDC_DUMP宏定义是控制调试接口打印的开关。10+- [调试接口](debug_interface/index.md):共 14 张知识卡(如 关闭ASCENDC_DUMP说明 等)。
11-- [Kernel Tiling](kernel_tiling/index.md):拷贝指定大小的数组内容到目标数组中,并返回指向拷贝后数组的指针。适用于拷贝一个结构体的数组成员变量的场景。该宏将指定数组拷贝至栈上,适用于频繁访问Tiling数据的场景,能够加快数据访问速度。11+- [Kernel Tiling](kernel_tiling/index.md):共 12 张知识卡(如 COPY_TILING_WITH_ARRAY、COPY_TILING_WITH_STRUCT、GET_TILING_DATA 等)。
12-- [Memory矢量计算](memory_vector_compute/index.md):单目矢量计算接口,对输入tensor按元素取绝对值。12+- [Memory矢量计算](memory_vector_compute/index.md):共 119 张知识卡,分布在 18 个子目录(basic_arithmetic、compare_and_select、composite_compute 等)。
13-- [Reg矢量计算](reg_vector_compute/index.md):Reg矢量计算API概述,直接操作Vector计算寄存器实现数据搬运与计算,较Memory矢量计算API提供更大开发自由度。13+- [Reg矢量计算](reg_vector_compute/index.md):Reg矢量计算API概述,直接操作Vector计算寄存器实现数据搬运与计算,较Memory矢量计算API提供更大开发自由度。(共 98 张知识卡,分布在 21 个子目录)
14-- [资源管理](resource_management/index.md):获取创建TPipe对象时设置的全局唯一TPipe指针,获取后可进行TPipe相关的操作如InitBuffer。14+- [资源管理](resource_management/index.md):共 68 张知识卡(如 GetTPipePtr、PopStackBuffer 等)。
15-- [标量计算](scalar_compute/index.md):将float类型标量数据转换为bfloat16_t类型标量数据。15+- [标量计算](scalar_compute/index.md):共 10 张知识卡(如 CastFloatToBfloat16、CastFloatToHalfInt32、CastMultiTypeToFloat 等)。
16-- [特殊寄存器访问](special_register_access/index.md):对指定的特殊寄存器进行清零操作。16+- [特殊寄存器访问](special_register_access/index.md):共 9 张知识卡(如 ClearSpr、GetCtrlSpr、GetSaturationFlag 等)。
17-- [同步控制](sync_control/index.md):在编程中,同步是指协调多个执行单元(如线程、进程等)对共享资源的访问顺序和时机来确保程序的正确性。如果没有同步来严格控制执行顺序,则会出现结果不一致、行为不可预测等多种问题。17+- [同步控制](sync_control/index.md):共 24 张知识卡(如 系统同步能力概述(同步控制) 等)。
18-- [工具接口](tool_interface/index.md):通过模板函数封装AIC与AIV分离模式的硬件条件分支,在指定执行单元下执行特定函数,避免代码中显式使用硬件条件分支。18+- [工具接口](tool_interface/index.md):共 24 张知识卡,分布在 4 个子目录(execution_mode、numeric_limits_utils、system_init 等)。
19 19 
20-- [deprecated interface](deprecated_interface/index.md):该接口已废弃,请使用ReadGmBypassDCache(ISASI)接口替代。20+- [deprecated interface](deprecated_interface/index.md):共 2 张知识卡(如 ReadGmByPassDCache(ISASI)(废弃)、WriteGmByPassDCache(ISASI)(废弃))。
@@ -1,20 +1,20 @@
1# Memory矢量计算1# Memory矢量计算
2 2 
3-- [基础算术](basic_arithmetic/index.md):单目矢量计算接口,对输入tensor按元素取绝对值。3+- [基础算术](basic_arithmetic/index.md):共 29 张知识卡(如 Abs、Add、Adds 等)。
4-- [比较与选择](compare_and_select/index.md):逐元素比较两个tensor大小,结果为真的对应比特位为1,否则为0,支持LT、GT、GE、EQ、NE、LE六种比较模式。4+- [比较与选择](compare_and_select/index.md):Compare、Compares称之为比较接口,用于逐元素比较Tensor或Scalar的大小,并根据比较结果生成掩码。Select、GatherMask称之为选择接口,用于根据掩码条件从源操作数中选取元素输出到目的操作数。推荐使用场景如下:(共 10 张知识卡:Compare、CompareStoreToRegister、Compares 等)
5-- [复合计算](composite_compute/index.md):将src0与src1相减再求绝对值,并将计算结果写入dst。5+- [复合计算](composite_compute/index.md):复合计算通常指的是在数学或编程中对复合函数或运算的处理。(共 15 张知识卡:AbsSub、AddDeqRelu、AddRelu 等)
6-- [数据排布转换](data_layout_conversion/index.md):将NCHW格式转换为NC1HWC0格式,也可用于二维矩阵数据块转置,单次repeat可处理512字节数据。6+- [数据排布转换](data_layout_conversion/index.md):共 2 张知识卡(如 TransDataTo5HD、Transpose)。
7-- [数据搬运](data_move/index.md):支持UB之间的连续数据搬运,数据搬运时格式和内容保持不变。7+- [数据搬运](data_move/index.md):共 12 张知识卡(如 Copy、Copy、DataCopy 等)。
8-- [数据搬运辅助配置接口](data_move_aux_config/index.md):刷新NDDMA DataCache,确保多核场景下DataCopy多维搬运前缓存数据为最新状态。8+- [数据搬运辅助配置接口](data_move_aux_config/index.md):共 4 张知识卡(如 NdDmaDci、ResetLoopModePara、SetLoopModePara 等)。
9-- [数据填充](data_padding/index.md):数据填充接口,每次取输入张量8个数,将每个数填充到结果张量一个DataBlock(32字节)中。9+- [数据填充](data_padding/index.md):共 4 张知识卡(如 Brcb、CreateVecIndex、Duplicate 等)。
10-- [数据重排(ISASI)](data_rearrange_isasi/index.md):将两个源操作数中的元素解交织存入两个结果操作数中,支持双输入和单输入两种模式。10+- [数据重排(ISASI)](data_rearrange_isasi/index.md):共 2 张知识卡(如 DeInterleave、Interleave)。
11-- [逻辑计算](logical_compute/index.md):对LocalTensor中的元素做按位与运算,支持整个tensor、前n个数据连续计算和高维切分计算三种调用方式。11+- [逻辑计算](logical_compute/index.md):共 9 张知识卡(如 And、Ands、Not 等)。
12-- [掩码操作](mask_operations/index.md):恢复mask值为默认值(全1),表示矢量计算中每次迭代内所有元素都参与运算。12+- [掩码操作](mask_operations/index.md):共 4 张知识卡(如 ResetMask、SetMaskCount、SetMaskNorm 等)。
13-- [归约计算](reduction_compute/index.md):概述 Tensor 归约计算及 ReduceSum、ReduceMax、ReduceMin、ReducePairElem 等接口能力。13+- [归约计算](reduction_compute/index.md):共 7 张知识卡(如 归约计算概述、ReduceDataBlock、ReduceMax 等)。
14-- [归约计算辅助配置接口](reduction_compute_aux_config/index.md):获取ReduceRepeat或ReduceMax/ReduceMin计算的最值及索引,结果以全局变量形式存储。14+- [归约计算辅助配置接口](reduction_compute_aux_config/index.md):共 2 张知识卡(如 GetReduceRepeatMaxMinSpr、GetReduceRepeatSumSpr)。
15-- [离散与聚合](scatter_gather/index.md):根据基地址和地址偏移张量确定索引,将输入张量中对应元素收集到结果张量中,支持前n个数据和高维切分计算。15+- [离散与聚合](scatter_gather/index.md):共 3 张知识卡(如 Gather、Gatherb、Scatter)。
16-- [SIMD计算说明](simd_compute/index.md):连续计算API,支持Tensor前n个数据计算。针对源操作数的连续n个数据进行计算并连续写入目的操作数,解决一维Tensor的连续计算问题。16+- [SIMD计算说明](simd_compute/index.md):矢量计算通过Vector计算单元完成,矢量计算的源操作数和目的操作数均通过Unified Buffer(UB)来进行存储。(共 4 张知识卡:连续计算(SIMD计算)、高维切分(SIMD计算)、SIMD计算掩码)
17-- [排序组合(ISASI)](sort_and_merge_isasi/index.md):与MrgSort接口配合使用,在队列耗尽停止后获取4个队列已处理的Region Proposal个数。17+- [排序组合(ISASI)](sort_and_merge_isasi/index.md):Region Proposal(简称RP)是目标检测算法中用于定位潜在目标区域的候选框。Region Proposal采用特定的数据结构进行存储和管理,以支持高效的排序和筛选操作。在业务场景下,为了从大量Region Proposals中筛选出最有可能包含目标的区域,需要根据每个RP的置信度score进行排序。(共 8 张知识卡:GetMrgSortResult、MrgSort、MrgSort4(排序与合并) 等)
18-- [类型转换](type_conversion/index.md):根据源操作数和目的操作数的数据类型进行精度转换,支持多种舍入模式。18+- [类型转换](type_conversion/index.md):共 2 张知识卡(如 Cast、Truncate)。
19- [类型转换辅助配置接口](type_conversion_aux_config/index.md):设置DEQSCALE寄存器的值,用于Vector计算单元上的量化计算,支持不同场景的量化参数配置。19- [类型转换辅助配置接口](type_conversion_aux_config/index.md):设置DEQSCALE寄存器的值,用于Vector计算单元上的量化计算,支持不同场景的量化参数配置。
20- [Vector逻辑架构](vector_logical_arch/index.md):Vector计算单元专用于执行向量计算。如下图所示,高亮部分展示了Vector计算单元及其关联的存储单元。20- [Vector逻辑架构](vector_logical_arch/index.md):Vector计算单元专用于执行向量计算。如下图所示,高亮部分展示了Vector计算单元及其关联的存储单元。
@@ -1,24 +1,24 @@
1# Reg矢量计算1# Reg矢量计算
2 2 
3-- [辅助数据类型](aux_data_types/index.md):枚举类型,用于指定寄存器数据写入模式,包括ZEROING置零和MERGING保留原值两种模式。3+- [辅助数据类型](aux_data_types/index.md):共 5 张知识卡(如 MaskMergeMode、PostLiteral、RegLayout 等)。
4-- [基础算术](basic_arithmetic/index.md):对srcReg中有效元素逐个取绝对值并写入dstReg,复数类型取模。4+- [基础算术](basic_arithmetic/index.md):寄存器矢量基本算术运算接口概述,涵盖Abs、Exp、Add、Mul等单目、双目及标量双目运算。(共 25 张知识卡:Abs、Add、AddC 等)
5-- [比较与选择](compare_and_select/index.md):逐元素比较两个RegTensor大小,根据比较模式将结果写入MaskReg对应比特位,支持EQ/GT/LT/GE/NE/LE模式。5+- [比较与选择](compare_and_select/index.md):比较与选择支持以下接口:(共 5 张知识卡:Compare、Compares、Select 等)
6-- [复合计算](composite_compute/index.md):将两个源操作数逐元素相减再求绝对值,结果写入目的操作数。6+- [复合计算](composite_compute/index.md):复合计算支持以下接口:(共 7 张知识卡:AbsSub、Axpy、ExpSub 等)
7-- [数据压缩](data_compression/index.md):将源操作数元素的低位部分写入目的寄存器的低半部分或高半部分,实现数据位宽压缩。7+- [数据压缩](data_compression/index.md):共 3 张知识卡(如 Pack、UnPack、Unsqueeze)。
8- [数据填充](data_padding/index.md):将标量或源寄存器最低位元素广播到目的寄存器,支持Scalar和Tensor两种模式,可配合mask使用。8- [数据填充](data_padding/index.md):将标量或源寄存器最低位元素广播到目的寄存器,支持Scalar和Tensor两种模式,可配合mask使用。
9-- [数据重排](data_reorder/index.md):对寄存器源操作数srcReg0和srcReg1中的元素进行解交织,存入结果寄存器dstReg0和dstReg1中。9+- [数据重排](data_reorder/index.md):共 2 张知识卡(如 DeInterleave、Interleave)。
10- [直方图计算](histogram_compute/index.md):对直方图数据进行统计,在目的寄存器原始数据上累加源寄存器数据的频率统计或累计统计结果。10- [直方图计算](histogram_compute/index.md):对直方图数据进行统计,在目的寄存器原始数据上累加源寄存器数据的频率统计或累计统计结果。
11- [索引操作](index_operations/index.md):以传入标量值为起始或终止值,生成递增或递减的索引序列并保存在目的寄存器中。11- [索引操作](index_operations/index.md):以传入标量值为起始或终止值,生成递增或递减的索引序列并保存在目的寄存器中。
12- [关键特性说明](key_features/index.md):Reg矢量计算关键特性说明,涵盖HardwareLoop硬件循环、指令双发优化、VF融合和数据搬运优化。12- [关键特性说明](key_features/index.md):Reg矢量计算关键特性说明,涵盖HardwareLoop硬件循环、指令双发优化、VF融合和数据搬运优化。
13-- [逻辑计算](logical_compute/index.md):根据mask对两个RegTensor或MaskReg执行按位与操作,将结果写入目的操作数,支持ZEROING模式。13+- [逻辑计算](logical_compute/index.md):逻辑计算支持以下接口:(共 9 张知识卡:And、Not、Or 等)
14-- [MaskReg计算](mask_reg_compute/index.md):将源操作数src0和src1中的元素解交织存入目的操作数dst0和dst1中。14+- [MaskReg计算](mask_reg_compute/index.md):共 7 张知识卡(如 DeInterleave、Interleave、Move 等)。
15- [Reg矢量计算概述](overview.md):Reg矢量计算API概述,直接操作Vector计算寄存器实现数据搬运与计算,较Memory矢量计算API提供更大开发自由度。15- [Reg矢量计算概述](overview.md):Reg矢量计算API概述,直接操作Vector计算寄存器实现数据搬运与计算,较Memory矢量计算API提供更大开发自由度。
16-- [归约计算](reduction_compute/index.md):将源寄存器中相邻两个数值相加,结果保存在目的寄存器的低位位置。16+- [归约计算](reduction_compute/index.md):共 3 张知识卡(如 PairReduceElem、Reduce、ReduceDataBlock)。
17-- [Reg数据搬入](reg_data_load/index.md):根据索引值index将源操作数按DataBlock(32B)收集到目的操作数dstReg中,index仅前8个元素有效,每个元素对应一个DataBlock。17+- [Reg数据搬入](reg_data_load/index.md):共 9 张知识卡(如 GatherB、Gather、Load 等)。
18-- [Reg数据搬出](reg_data_store/index.md):根据索引值index将源操作数srcReg中的元素分散到目的操作数UB中。18+- [Reg数据搬出](reg_data_store/index.md):共 7 张知识卡(如 ScatterDiscrete、Store、StoreAlignContinuous 等)。
19-- [寄存器数据类型](register_data_types/index.md):用于Reg矢量计算管理地址偏移量的专用寄存器,支持1至4层循环轴索引与步长绑定,逐层累加地址偏移。19+- [寄存器数据类型](register_data_types/index.md):Reg矢量计算API操作的基础寄存器数据类型介绍如下:(共 5 张知识卡:AddrReg、MaskReg、RegTensor 等)
20- [离散操作](scatter_operations/index.md):根据索引值将源操作数按元素收集到目的操作数中,支持b8、b16、b32数据类型。20- [离散操作](scatter_operations/index.md):根据索引值将源操作数按元素收集到目的操作数中,支持b8、b16、b32数据类型。
21- [同步控制](sync_control/index.md):Reg矢量计算宏函数内不同流水线之间的同步指令,目的流水线等待源流水线上所有指令完成才执行。21- [同步控制](sync_control/index.md):Reg矢量计算宏函数内不同流水线之间的同步指令,目的流水线等待源流水线上所有指令完成才执行。
22- [系统变量访问](system_variable_access/index.md):对指定的特殊目的寄存器进行清零操作,通过模板参数指定目标特殊寄存器,调用后该寄存器值被置为0。22- [系统变量访问](system_variable_access/index.md):对指定的特殊目的寄存器进行清零操作,通过模板参数指定目标特殊寄存器,调用后该寄存器值被置为0。
23-- [类型转换](type_conversion/index.md):Reg矢量计算的精度转换接口,支持浮点转整数、浮点转浮点、整数转浮点和整数转整数四种类型转换。23+- [类型转换](type_conversion/index.md):共 2 张知识卡(如 Cast、Truncate)。
24- [VF调用](vf_call/index.md):在SIMD编程场景下使用。用于启动SIMD VF(Vector Function)子任务。24- [VF调用](vf_call/index.md):在SIMD编程场景下使用。用于启动SIMD VF(Vector Function)子任务。
@@ -1,13 +1,13 @@
1# 资源管理1# 资源管理
2 2 
3-- [自定义TBufPool](custom_t_buf_pool/index.md):辅助开发者自定义TBufPool以管理Unified Buffer和L1 Buffer内存的宏,支持不连续内存块和跨TQue共享等自定义分配需求。3+- [自定义TBufPool](custom_t_buf_pool/index.md):共 8 张知识卡(如 ExternImplBufpool、GetBufHandle、GetCurAddr 等)。
4- [GetTPipePtr](get_t_pipe_ptr.md):获取创建TPipe对象时设置的全局唯一TPipe指针,获取后可进行TPipe相关的操作如InitBuffer。4- [GetTPipePtr](get_t_pipe_ptr.md):获取创建TPipe对象时设置的全局唯一TPipe指针,获取后可进行TPipe相关的操作如InitBuffer。
5-- [LocalMemAllocator](local_mem_allocator/index.md):根据用户指定的逻辑位置、数据类型、数据长度返回对应的LocalTensor对象。5+- [LocalMemAllocator](local_mem_allocator/index.md):共 4 张知识卡(如 Alloc、GetCurAddr、LocalMemAllocatorConstructor 等)。
6- [PopStackBuffer](pop_stack_buffer.md):在指定position(逻辑位置)申请临时空间,空间大小为指定position的全部剩余空间。6- [PopStackBuffer](pop_stack_buffer.md):在指定position(逻辑位置)申请临时空间,空间大小为指定position的全部剩余空间。
7-- [TBuf](t_buf/index.md):从TBuf上获取指定长度的Tensor,或者获取全部长度的Tensor。7+- [TBuf](t_buf/index.md):共 4 张知识卡(如 Get、GetWithOffset、TBufConstructor 等)。
8-- [TBufPool](t_buf_pool/index.md):从Tpipe或TBufPool划分资源池,支持非共享模式和共享模式(复用已有资源池地址)。8+- [TBufPool](t_buf_pool/index.md):共 5 张知识卡(如 InitBufPool、InitBuffer、Reset 等)。
9-- [TPipe](t_pipe/index.md):用于申请HardEvent(硬件类型同步事件)的TEventID,必须与ReleaseEventID搭配使用,调用后占用TEventID直至释放。9+- [TPipe](t_pipe/index.md):共 15 张知识卡(如 AllocEventID、Destroy、DestroyWithoutPipeAll 等)。
10-- [TQue](t_que/index.md):从Que中分配Tensor,Tensor大小为InitBuffer时设置的每块内存长度,支持non-inplace和inplace两种接口。10+- [TQue](t_que/index.md):共 9 张知识卡(如 AllocTensor、DeQue、EnQue 等)。
11-- [TQueBind](t_que_bind/index.md):从Que中分配Tensor,Tensor所占大小为InitBuffer时设置的每块内存长度。11+- [TQueBind](t_que_bind/index.md):共 13 张知识卡(如 AllocTensor、DeQue、EnQue 等)。
12-- [TSCM](tscm/index.md):从TSCM中分配tensor,大小为InitBuffer时设置的每块内存长度,分配内容可能为随机值。12+- [TSCM](tscm/index.md):共 5 张知识卡(如 AllocTensor、DeQue、EnQue 等)。
13-- [workspace](workspace/index.md):获取系统workspace指针。部分高阶API如Matmul需要使用系统workspace,相关接口需要传入系统workspace指针,此时可以通过该接口获取。使用系统workspace时,host侧开发者需要自行申请系统workspace的空间,其预留空间大小可以通过GetLibApiWorkSpaceSize接口获取。13+- [workspace](workspace/index.md):共 3 张知识卡(如 GetSysWorkSpacePtr、GetUserWorkspace、SetSysWorkSpace)。
@@ -1,6 +1,6 @@
1# 同步控制1# 同步控制
2 2 
3-- [核间同步](inter_core_sync/index.md):核间同步场景中向调度模块上报本核前置pipe流水任务已完成,与CrossCoreWaitFlag配对使用,通过flagId控制核间同步。3+- [核间同步](inter_core_sync/index.md):共 10 张知识卡(如 CrossCoreSetFlag、CrossCoreWaitFlag、IBSet 等)。
4-- [任务间同步](inter_task_sync/index.md):在SuperKernel子Kernel中调用,使其后指令与后续子Kernel并行执行,提升整体性能。4+- [任务间同步](inter_task_sync/index.md):共 2 张知识卡(如 SetNextTaskStart、WaitPreTaskEnd)。
5-- [核内同步](intra_core_sync/index.md):从框架获取并占用一个MutexID,与ReleaseMutexID配合使用,获取的MutexID可传入Mutex::Lock/Unlock接口用于流水同步。5+- [核内同步](intra_core_sync/index.md):共 11 张知识卡(如 AllocMutexID、DataSyncBarrier、核内同步能力概述 等)。
6- [系统同步能力概述(同步控制)](system_sync_overview.md):在编程中,同步是指协调多个执行单元(如线程、进程等)对共享资源的访问顺序和时机来确保程序的正确性。如果没有同步来严格控制执行顺序,则会出现结果不一致、行为不可预测等多种问题。6- [系统同步能力概述(同步控制)](system_sync_overview.md):在编程中,同步是指协调多个执行单元(如线程、进程等)对共享资源的访问顺序和时机来确保程序的正确性。如果没有同步来严格控制执行顺序,则会出现结果不一致、行为不可预测等多种问题。
@@ -1,6 +1,6 @@
1# 工具接口1# 工具接口
2 2 
3- [执行模式](execution_mode/index.md):通过模板函数封装AIC与AIV分离模式的硬件条件分支,在指定执行单元下执行特定函数,避免代码中显式使用硬件条件分支。3- [执行模式](execution_mode/index.md):通过模板函数封装AIC与AIV分离模式的硬件条件分支,在指定执行单元下执行特定函数,避免代码中显式使用硬件条件分支。
4-- [numeric limits utils](numeric_limits_utils/index.md):返回指定数据类型的最小正的非正规值,支持标量和矢量两种接口,矢量接口为dstLocal前count个元素赋该值。4+- [numeric limits utils](numeric_limits_utils/index.md):共 9 张知识卡(如 DeNormMin、Infinity、工具接口简介 等)。
5- [系统初始化](system_init/index.md):初始化AI Core全局状态,包括AIC与AIV的公共状态及各自特有状态;不同产品实际执行的初始化项存在差异。5- [系统初始化](system_init/index.md):初始化AI Core全局状态,包括AIC与AIV的公共状态及各自特有状态;不同产品实际执行的初始化项存在差异。
6-- [系统资源与变量](system_resources_and_variables/index.md):获取当前AI处理器架构版本号。6+- [系统资源与变量](system_resources_and_variables/index.md):共 13 张知识卡(如 GetArchVersion、GetBlockIdx、GetBlockNum 等)。
@@ -1,6 +1,6 @@
1# 原子操作1# 原子操作
2 2 
3- [原子操作概述](atomic_operations_overview.md):数据搬运随路原子操作接口用于对后续目的地址为GM的数据搬运指令开启原子操作,涉及的接口请参见表1。如下图1的左侧子图所示,未开启原子操作时,写入GM的数据搬运完成后,GM中原始数据将被新搬运数据完全覆盖。如图1右侧子图所示,当数据搬运随路原子操作接口被调用后,系统将为后续写入GM的数据搬运开启原子操作。此时,数据搬运完成后,GM中的最终数据由原始GM数据与新搬运数据共同决定。3- [原子操作概述](atomic_operations_overview.md):数据搬运随路原子操作接口用于对后续目的地址为GM的数据搬运指令开启原子操作,涉及的接口请参见表1。如下图1的左侧子图所示,未开启原子操作时,写入GM的数据搬运完成后,GM中原始数据将被新搬运数据完全覆盖。如图1右侧子图所示,当数据搬运随路原子操作接口被调用后,系统将为后续写入GM的数据搬运开启原子操作。此时,数据搬运完成后,GM中的最终数据由原始GM数据与新搬运数据共同决定。
4-- [搬运原子操作](datamove_atomic/index.md):头文件路径为:"c_api/atomic/atomic.h"。4+- [搬运原子操作](datamove_atomic/index.md):共 4 张知识卡(如 asc_disable_dma_atomic、asc_set_atomic_add、asc_set_atomic_max 等)。
5- [关键特性说明](key_features.md):确定性计算是指在相同输入条件下,无论执行次数或执行环境如何变化,总能产生完全一致输出结果的计算过程。确定性计算为系统稳定性和实验可验证性提供保障。5- [关键特性说明](key_features.md):确定性计算是指在相同输入条件下,无论执行次数或执行环境如何变化,总能产生完全一致输出结果的计算过程。确定性计算为系统稳定性和实验可验证性提供保障。
6-- [标量原子操作](scalar_atomic/index.md):对Global Memory中address指向的单个元素执行原子加操作。读取该元素的旧值old_value,将old_value + val写回该地址,并返回old_value。整个读取、计算和写回过程为原子操作。6+- [标量原子操作](scalar_atomic/index.md):共 13 张知识卡(如 asc_atomic_add、asc_atomic_and、asc_atomic_cas 等)。
@@ -2,23 +2,23 @@
2 2 
3- [asc_copy_gm2l0a](asc_copy_gm2l0a.md):将数据从GM搬运到L0A Buffer。3- [asc_copy_gm2l0a](asc_copy_gm2l0a.md):将数据从GM搬运到L0A Buffer。
4- [asc_copy_gm2l0b](asc_copy_gm2l0b.md):将矩阵数据从Global Memory搬运到L0B Buffer。4- [asc_copy_gm2l0b](asc_copy_gm2l0b.md):将矩阵数据从Global Memory搬运到L0B Buffer。
5-- [asccopygm2l1](asc_copy_gm2l1/index.md):将数据从Global Memory搬运到L1 Buffer。以512字节的分形为单位,从GM中按行列方向的起始位置、步长切分提取多个分形矩阵搬运到L1 Buffer,搬运过程中不支持分形格式转换,用于矩阵计算所需的2D格式数据加载。5+- [asccopygm2l1](asc_copy_gm2l1/index.md):共 3 张知识卡(如 asc_copy_gm2l1_arch_3510(2D矩阵搬运模式)、asc_copy_gm2l1、asc_copy_gm2l1_arch_3510(高维切分数据搬运))。
6- [asc_copy_gm2l1_align](asc_copy_gm2l1_align.md):对矩阵数据进行实时padding,完成padding后将数据从Global Memory搬运到L1 Buffer。需要与asc_set_gm2l1_loop_size、asc_set_gm2l1_loop1_stride、asc_set_gm2l1_loop2_stride和asc_set_gm2l1_pad配合使用。6- [asc_copy_gm2l1_align](asc_copy_gm2l1_align.md):对矩阵数据进行实时padding,完成padding后将数据从Global Memory搬运到L1 Buffer。需要与asc_set_gm2l1_loop_size、asc_set_gm2l1_loop1_stride、asc_set_gm2l1_loop2_stride和asc_set_gm2l1_pad配合使用。
7- [asc_copy_gm2l1_dn2nz](asc_copy_gm2l1_dn2nz.md):将数据从Global Memory搬运到L1 Buffer,并在搬运过程中进行DN到Nz格式的转换,使其满足后续矩阵计算对分形列连续性的要求。7- [asc_copy_gm2l1_dn2nz](asc_copy_gm2l1_dn2nz.md):将数据从Global Memory搬运到L1 Buffer,并在搬运过程中进行DN到Nz格式的转换,使其满足后续矩阵计算对分形列连续性的要求。
8-- [asccopygm2l1nd2nz](asc_copy_gm2l1_nd2nz/index.md):将数据从Global Memory (GM)搬运到Level 1 cache (L1),支持在数据搬运时进行ND格式到NZ格式的转换。8+- [asccopygm2l1nd2nz](asc_copy_gm2l1_nd2nz/index.md):共 2 张知识卡(如 asc_copy_gm2l1_nd2nz、asc_copy_gm2l1_nd2nz_arch_3510)。
9-- [asccopyl0c2gm](asc_copy_l0c2gm/index.md):矩阵计算完成后,对结果进行量化处理,之后将处理结果搬运到GM中。量化模式共分为9种,分别为: - NoQuant:不开启量化功能。 - F322F16:float量化成half。量化结果不支持INF_NAN模式。 - F322BF16:float量化成bfloat16_t。量化结果不支持INF_NAN模式。 - DEQF16:int32_t量化成half。量化结果不支持INF_NAN模式。 - VDEQF16:int32_t量化成half。量化结果不支持INF_NAN模式。 - QF322B8_PRE:float量化成uint8_t/int8_t。scalar量化。 - VQF322B8_PRE:float量化成uint8_t/int8_t。矢量量化。 - REQ8:int32_t量化成uint8_t/int8_t。scalar量化。 - VREQ8:int32_t量化成uint8_t/int8_t。矢量量化。9+- [asccopyl0c2gm](asc_copy_l0c2gm/index.md):共 2 张知识卡(如 asc_copy_l0c2gm、asc_copy_l0c2gm_arch_3510)。
10-- [asccopyl0c2l1](asc_copy_l0c2l1/index.md):矩阵计算完成后,对结果进行量化处理,之后将处理结果搬运到L1 Buffer中。量化模式共分为9种,分别为: - NoQuant:不开启量化功能。 - DEQF16:int32_t量化成half。量化结果不支持INF_NAN模式。 - F322BF16:float量化成bfloat16_t。量化结果不支持INF_NAN模式。 - F322F16:float量化成half。量化结果不支持INF_NAN模式。 - QF322B8_PRE:float量化成uint8_t/int8_t。scalar量化。 - REQ8:int32_t量化成uint8_t/int8_t。scalar量化。 - VDEQF16:int32_t量化成half。量化结果不支持INF_NAN模式。 - VQF322B8_PRE:float量化成uint8_t/int8_t。矢量量化。 - VREQ8:int32_t量化成uint8_t/int8_t。矢量量化。10+- [asccopyl0c2l1](asc_copy_l0c2l1/index.md):共 2 张知识卡(如 asc_copy_l0c2l1、asc_copy_l0c2l1)。
11- [asc_copy_l0c2ub](asc_copy_l0c2ub.md):将矩阵计算结果从L0C Buffer搬运至Unified Buffer,搬运过程中可同步支持随路量化、随路激活、随路格式转换(Nz2ND/Nz2DN)等能力组合。11- [asc_copy_l0c2ub](asc_copy_l0c2ub.md):将矩阵计算结果从L0C Buffer搬运至Unified Buffer,搬运过程中可同步支持随路量化、随路激活、随路格式转换(Nz2ND/Nz2DN)等能力组合。
12-- [asccopyl12bt](asc_copy_l12bt/index.md):将数据从L1 Buffer搬运到BiasTable Buffer中,BiasTable Buffer用于存放矩阵计算中的Bias。12+- [asccopyl12bt](asc_copy_l12bt/index.md):共 2 张知识卡(如 asc_copy_l12bt、asc_copy_l12bt)。
13-- [asccopyl12fb](asc_copy_l12fb/index.md):将数据从L1 Buffer搬运到Fixpipe Buffer中,Fixpipe Buffer用于存放量化参数。13+- [asccopyl12fb](asc_copy_l12fb/index.md):共 2 张知识卡(如 asc_copy_l12fb、asc_copy_l12fb)。
14- [asc_copy_l12gm](asc_copy_l12gm.md):将数据从L1 Buffer搬运到Global Memory。14- [asc_copy_l12gm](asc_copy_l12gm.md):将数据从L1 Buffer搬运到Global Memory。
15-- [asccopyl12l0a](asc_copy_l12l0a/index.md):负责完成矩阵计算所需的2D格式数据的搬运,搬运粒度为大小为512字节的数据分形,支持非转置与转置矩阵从L1 Buffer到L0A Buffer的数据搬运。15+- [asccopyl12l0a](asc_copy_l12l0a/index.md):共 3 张知识卡(如 asc_copy_l12l0a_arch_3510(2D矩阵搬运)、asc_copy_l12l0a_arch_3510(3D矩阵搬运)、asc_copy_l12l0a)。
16- [asc_copy_l12l0a_mx](asc_copy_l12l0a_mx.md):负责完成MX矩阵计算过程中所需的左矩阵对应的量化系数的搬运,数据通路为L1 Buffer->L0A_MX Buffer。其中左量化系数矩阵以32字节(固定数据类型为fp8_e8m0_t、分形大小为16×2的)的数据分形为单位进行搬运。16- [asc_copy_l12l0a_mx](asc_copy_l12l0a_mx.md):负责完成MX矩阵计算过程中所需的左矩阵对应的量化系数的搬运,数据通路为L1 Buffer->L0A_MX Buffer。其中左量化系数矩阵以32字节(固定数据类型为fp8_e8m0_t、分形大小为16×2的)的数据分形为单位进行搬运。
17- [asc_copy_l12l0a_trans](asc_copy_l12l0a_trans.md):该接口实现带转置的2D格式数据从L1 Buffer到L0A Buffer的加载。17- [asc_copy_l12l0a_trans](asc_copy_l12l0a_trans.md):该接口实现带转置的2D格式数据从L1 Buffer到L0A Buffer的加载。
18-- [asccopyl12l0b](asc_copy_l12l0b/index.md):负责完成矩阵计算所需的2D格式数据的搬运,搬运粒度为大小为512字节的数据分形,支持非转置与转置矩阵从L1 Buffer到L0B buffer的数据搬运。18+- [asccopyl12l0b](asc_copy_l12l0b/index.md):共 3 张知识卡(如 asc_copy_l12l0b_arch_3510(2D矩阵搬运)、asc_copy_l12l0b_arch_3510(3D矩阵搬运)、asc_copy_l12l0b)。
19- [asc_copy_l12l0b_mx](asc_copy_l12l0b_mx.md):负责完成MX矩阵计算过程中所需的右矩阵对应的量化系数的搬运,数据通路为L1 Buffer->L0B_MX Buffer。其中右量化系数矩阵以32字节(固定数据类型为fp8_e8m0_t、分形大小为2×16的)的数据分形为单位进行搬运。19- [asc_copy_l12l0b_mx](asc_copy_l12l0b_mx.md):负责完成MX矩阵计算过程中所需的右矩阵对应的量化系数的搬运,数据通路为L1 Buffer->L0B_MX Buffer。其中右量化系数矩阵以32字节(固定数据类型为fp8_e8m0_t、分形大小为2×16的)的数据分形为单位进行搬运。
20- [asc_copy_l12l0b_sparse](asc_copy_l12l0b_sparse.md):用于搬运存放在L1 Buffer里的512B大小的稠密权重矩阵到L0B Buffer里,同时读取128B大小的索引矩阵用于稠密矩阵的稀疏化。20- [asc_copy_l12l0b_sparse](asc_copy_l12l0b_sparse.md):用于搬运存放在L1 Buffer里的512B大小的稠密权重矩阵到L0B Buffer里,同时读取128B大小的索引矩阵用于稠密矩阵的稀疏化。
21-- [asccopyl12l0btrans](asc_copy_l12l0b_trans/index.md):该接口实现带转置的2D格式数据从L1 Buffer到L0B Buffer的加载。21+- [asccopyl12l0btrans](asc_copy_l12l0b_trans/index.md):共 2 张知识卡(如 asc_copy_l12l0b_trans、asc_copy_l12l0b_trans)。
22- [asc_copy_l12l0c](asc_copy_l12l0c.md):将矩阵由L1 Buffer搬运到L0C Buffer中。该搬运支持的数据类型如下: - bfloat16_t->bfloat16_t - half->half - float->half - float->bfloat16_t - float->float - int32_t->int32_t - uint32_t->uint32_t22- [asc_copy_l12l0c](asc_copy_l12l0c.md):将矩阵由L1 Buffer搬运到L0C Buffer中。该搬运支持的数据类型如下: - bfloat16_t->bfloat16_t - half->half - float->half - float->bfloat16_t - float->float - int32_t->int32_t - uint32_t->uint32_t
23- [asc_copy_l12ub](asc_copy_l12ub.md):将数据从L1 Buffer搬运到Unified Buffer(UB)中。23- [asc_copy_l12ub](asc_copy_l12ub.md):将数据从L1 Buffer搬运到Unified Buffer(UB)中。
24- [asc_fill_l0a](asc_fill_l0a.md):将value的二进制值赋值给dst,对L0A Buffer的Local Memory进行初始化。24- [asc_fill_l0a](asc_fill_l0a.md):将value的二进制值赋值给dst,对L0A Buffer的Local Memory进行初始化。
@@ -1,6 +1,6 @@
1# 基础定义1# 基础定义
2 2 
3-- [常量](constant/index.md):ASC_C_API_MRGSORT_ELEMENT_LEN表示MrgSort4操作支持的队列数量,值为4。该常量用于asc_mrgsort4接口的src数组长度,以及asc_get_vms4_sr接口的sorted_num数组长度。3+- [常量](constant/index.md):共 2 张知识卡(如 ASC_C_API_MRGSORT_ELEMENT_LEN、ASC_VDEQ_SIZE)。
4-- [枚举](enum/index.md):asc_hf32_round_mode用于Mmad计算开启HF32模式时由FP32舍入到HF32的舍入模式管理策略,在调用asc_set_hf32_round_mode时使用。4+- [枚举](enum/index.md):共 6 张知识卡(如 asc_hf32_round_mode、asc_l13d_repeat_direction、asc_load_l2_cache_mode 等)。
5- [寄存器类型](type/index.md):矢量数据寄存器用于存储矢量数据,其长度为VL(Vector Length)。在Ascend 950PR/Ascend 950DT产品中,VL为256字节,可存储VL / sizeof(dtype)个dtype类型的元素。例如,vector_float可存储256 / sizeof(float) = 64个float类型的元素。5- [寄存器类型](type/index.md):矢量数据寄存器用于存储矢量数据,其长度为VL(Vector Length)。在Ascend 950PR/Ascend 950DT产品中,VL为256字节,可存储VL / sizeof(dtype)个dtype类型的元素。例如,vector_float可存储256 / sizeof(float) = 64个float类型的元素。
6-- [联合体](union/index.md):asc_fill_value_config为fill_value的初始化参数结构体,包含asc_fill_l0a/asc_fill_l0b/asc_fill_l1接口需要配置的各种初始化参数。6+- [联合体](union/index.md):共 5 张知识卡(如 asc_fill_value_config、asc_l13d_fmatrix_config、asc_load3d_v2_config 等)。
@@ -1,20 +1,20 @@
1# API列表1# API列表
2 2 
3-- [原子操作](atomic/index.md):数据搬运随路原子操作接口用于对后续目的地址为GM的数据搬运指令开启原子操作,涉及的接口请参见表1。如下图1的左侧子图所示,未开启原子操作时,写入GM的数据搬运完成后,GM中原始数据将被新搬运数据完全覆盖。如图1右侧子图所示,当数据搬运随路原子操作接口被调用后,系统将为后续写入GM的数据搬运开启原子操作。此时,数据搬运完成后,GM中的最终数据由原始GM数据与新搬运数据共同决定。3+- [原子操作](atomic/index.md):共 19 张知识卡(如 原子操作概述、关键特性说明 等)。
4- [API列表](c_api.md):C API文档按最细粒度公共头文件归类。除试验接口外,整体使用时可以包含c_api/asc_simd.h。4- [API列表](c_api.md):C API文档按最细粒度公共头文件归类。除试验接口外,整体使用时可以包含c_api/asc_simd.h。
5-- [缓存控制](cache_ctrl/index.md):从GM地址预加载数据到Data Cache,每次调用预加载一个Cache Line大小的数据。5+- [缓存控制](cache_ctrl/index.md):共 9 张知识卡(如 asc_datacache_preload、AscDcci、asc_dci(缓存控制) 等)。
6-- [矩阵计算](cube_compute/index.md):本接口用于设置Mmad计算关闭HF32模式,其作用与asc_enable_hf32相反,两个接口不同时生效。关闭HF32模式后,L0A Buffer与L0B Buffer中的float数据在参与Mmad计算之前不做舍入处理。6+- [矩阵计算](cube_compute/index.md):共 13 张知识卡(如 asc_disable_hf32、asc_disable_hif8、AscEnableFp8 等)。
7-- [矩阵计算搬运](cube_datamove/index.md):将数据从GM搬运到L0A Buffer。7+- [矩阵计算搬运](cube_datamove/index.md):共 69 张知识卡(如 asc_copy_gm2l0a、asc_copy_gm2l0b、asc_copy_gm2l1_align 等)。
8-- [基础定义](defs/index.md):ASC_C_API_MRGSORT_ELEMENT_LEN表示MrgSort4操作支持的队列数量,值为4。该常量用于asc_mrgsort4接口的src数组长度,以及asc_get_vms4_sr接口的sorted_num数组长度。8+- [基础定义](defs/index.md):共 14 张知识卡,分布在 4 个子目录(constant、enum、type 等)。
9-- [废弃接口](deprecated_interface/index.md):已废弃的按元素取绝对值同步计算接口,请改用asc_abs和asc_sync实现精细化调优。9+- [废弃接口](deprecated_interface/index.md):废弃接口索引页,汇总矢量计算asc_abs_sync与原子操作asc_set_atomic_add_int等已废弃接口。(共 219 张知识卡:AscAbsSync、asc_add_relu_sync(废弃)、asc_add_scalar_sync(废弃) 等)
10- [其他操作](experimental/index.md):本接口为试验接口,在后续版本中可能会调整或改进,不保证后续兼容性。请开发者在使用过程中关注后续版本更新。10- [其他操作](experimental/index.md):本接口为试验接口,在后续版本中可能会调整或改进,不保证后续兼容性。请开发者在使用过程中关注后续版本更新。
11- [通用说明和约束(C API)](general_description_and_constraints.md):AI Core内外包含多级存储单元,各存储单元的主要用途如下:11- [通用说明和约束(C API)](general_description_and_constraints.md):AI Core内外包含多级存储单元,各存储单元的主要用途如下:
12-- [Reg矢量计算](reg_compute/index.md):从Unified Buffer(UB)中按dtype对齐的起始地址读取VL长度数据,并通过函数返回值返回或写入目的矢量数据寄存器,搬运过程中数据格式和内容保持不变。连续搬入时,需要在每次调用前手动更新源地址。该接口为易用性接口,对性能有要求时可使用asc_loadunalign或asc_loadunalign_postupdate。12+- [Reg矢量计算](reg_compute/index.md):共 187 张知识卡,分布在 19 个子目录(load、reg_addr_reg、reg_arith 等)。
13-- [标量计算](scalar_compute/index.md):对一个64位无符号整数执行位清除操作,将bits中第idx位清零,其余位保持不变。返回修改后的64位无符号整数。13+- [标量计算](scalar_compute/index.md):共 12 张知识卡,分布在 4 个子目录(scalar_bit、scalar_convert、scalar_load 等)。
14-- [特殊寄存器访问](spr/index.md):读取CTRL寄存器(控制寄存器)特定比特位上的值。14+- [特殊寄存器访问](spr/index.md):共 7 张知识卡(如 asc_get_ctrl、asc_get_saturation_flag、asc_get_saturation_strategy 等)。
15-- [同步控制](sync/index.md):用于AI Core内部异步流水线之间的同步,根据mutex_id获取Mutex,可按需阻塞指定流水线的执行。15+- [同步控制](sync/index.md):共 22 张知识卡(如 AscLock、asc_sync(同步)、AscSyncBlockArrive 等)。
16-- [工具接口](utils/index.md):本接口将当前AI Core上可能由前序算子遗留的全局状态恢复为接口约定的初始状态,避免这些状态影响当前算子的计算结果。本接口仅初始化控制状态,不会清零Global Memory、Local Memory或用户数据。16+- [工具接口](utils/index.md):共 17 张知识卡,分布在 3 个子目录(sys_init、sys_misc、sys_var)。
17-- [Memory矢量计算](vector_compute/index.md):按元素取绝对值,计算公式如下: $$ dst_i = |src_i| $$17+- [Memory矢量计算](vector_compute/index.md):共 100 张知识卡,分布在 11 个子目录(vector_arith、vector_broadcast、vector_compare 等)。
18-- [矢量数据搬运](vector_datamove/index.md):将数据从Unified Buffer(UB)搬运到UB。支持高维切分。18+- [矢量数据搬运](vector_datamove/index.md):共 23 张知识卡(如 asc_copy、asc_copy_ub2l1、asc_copy_ub2ub 等)。
19 19 
20- [C API](c_api_reference.md):C API开放芯片完备编程能力,支持以数组形式分配内存,一般基于指针编程。提供与业界一致的C语言编程体验。20- [C API](c_api_reference.md):C API开放芯片完备编程能力,支持以数组形式分配内存,一般基于指针编程。提供与业界一致的C语言编程体验。
@@ -1,21 +1,21 @@
1# Reg矢量计算1# Reg矢量计算
2 2 
3-- [Reg数据搬入](load/index.md):从Unified Buffer(UB)中按dtype对齐的起始地址读取VL长度数据,并通过函数返回值返回或写入目的矢量数据寄存器,搬运过程中数据格式和内容保持不变。连续搬入时,需要在每次调用前手动更新源地址。该接口为易用性接口,对性能有要求时可使用asc_loadunalign或asc_loadunalign_postupdate。3+- [Reg数据搬入](load/index.md):共 27 张知识卡(如 asc_load、asc_loadalign、asc_loadalign_brc_datablock 等)。
4- [地址寄存器操作](reg_addr_reg/index.md):更新地址寄存器(addr_reg),用于迭代寻址。支持1~4维偏移量设置,位宽支持b8、b16、b32三种模式:4- [地址寄存器操作](reg_addr_reg/index.md):更新地址寄存器(addr_reg),用于迭代寻址。支持1~4维偏移量设置,位宽支持b8、b16、b32三种模式:
5-- [基础算术](reg_arith/index.md):根据mask对源操作数src按元素取绝对值,得到计算结果。计算公式如下:5+- [基础算术](reg_arith/index.md):共 19 张知识卡(如 asc_abs、asc_add、asc_add_scalar 等)。
6-- [广播操作](reg_broadcast/index.md):将源操作数src的最低位元素(即下标为0的元素)广播到目的操作数dst中被mask筛选的位置。mask仅用于筛选dst中的元素位置,不影响从src中读取的元素;dst中未被mask筛选的位置置零。6+- [广播操作](reg_broadcast/index.md):共 6 张知识卡(如 asc_duplicate、asc_duplicate_highest、asc_duplicate_highest_merge 等)。
7-- [比较计算](reg_compare/index.md):eq表示相等(equal to),该接口用于逐元素比较两个源操作数是否相等,若相等则目的操作数对应结果位为1,否则为0。每个元素的比较结果占一个比特。计算公式如下:7+- [比较计算](reg_compare/index.md):共 12 张知识卡(如 asc_eq、asc_eq_scalar、asc_ge 等)。
8-- [类型转换](reg_convert/index.md):根据mask将src中参与计算的bfloat16_t类型元素转换为fp4x2_e1m2_t类型,结果写入dst。由于源操作数与目的操作数的类型位宽比为4:1,dst中的元素按索引分为四组,参数dst_quarter_pos从dst中离散选取每隔4个元素的索引位置写入转换结果,具体为索引0、4、8、…,索引1、5、9、…,索引2、6、10、…或索引3、7、11、…,其他位置置零。8+- [类型转换](reg_convert/index.md):共 55 张知识卡(如 asc_bfloat162e1m2x2、asc_bfloat162e2m1x2、asc_bfloat162float 等)。
9-- [Reg间搬运](reg_copy/index.md):完成数据在寄存器间的搬运,支持以下三种模式:9+- [Reg间搬运](reg_copy/index.md):共 3 张知识卡(如 asc_copy、asc_extract_mask、asc_mask)。
10-- [复合计算](reg_fused/index.md):根据mask对源操作数src0、src1按元素相减后取绝对值,得到计算结果。计算公式如下:10+- [复合计算](reg_fused/index.md):共 10 张知识卡(如 asc_abs_sub、asc_axpy、asc_exp_sub 等)。
11- [reg gather](reg_gather/index.md):根据索引位置index将源操作数src按元素收集到目的操作数dst中。11- [reg gather](reg_gather/index.md):根据索引位置index将源操作数src按元素收集到目的操作数dst中。
12-- [直方图](reg_histogram/index.md):对输入矢量数据寄存器中的元素进行累计频率统计,生成累计直方图。统计结果在目的矢量数据寄存器原有数据基础上累加。支持配置掩码用于指示参与统计的元素,掩码位为1时,对应元素参与统计;为0时不统计。掩码不影响目的矢量数据寄存器的写入行为。本接口在Vector Function(__simd_vf__标记的函数)内使用。12+- [直方图](reg_histogram/index.md):共 2 张知识卡(如 asc_cumulative_histogram、asc_frequency_histogram)。
13- [索引操作](reg_index/index.md):以传入的value为起始值,生成递增/递减的索引,并将生成的索引保存在dst中,Vector Length (VL)表示矢量数据寄存器的位宽,VL_T表示该寄存器可存储的元素数量。算法逻辑表示如下:13- [索引操作](reg_index/index.md):以传入的value为起始值,生成递增/递减的索引,并将生成的索引保存在dst中,Vector Length (VL)表示矢量数据寄存器的位宽,VL_T表示该寄存器可存储的元素数量。算法逻辑表示如下:
14-- [逻辑计算](reg_logic/index.md):根据mask对源操作数src0、src1执行按位与(&)操作,将结果作为返回值返回或写入目的操作数dst。14+- [逻辑计算](reg_logic/index.md):共 8 张知识卡(如 asc_and、asc_not、asc_or 等)。
15-- [掩码寄存器操作](reg_mask/index.md):根据传入的pat_mode生成对应的掩码寄存器,支持b8、b16、b32三种位宽模式。pat_mode参数指定mask的模式,即指定哪些位置的元素参与计算。15+- [掩码寄存器操作](reg_mask/index.md):共 3 张知识卡(如 asc_create_mask、asc_get_mask_spr、asc_update_mask)。
16-- [排布变换](reg_permute_sel/index.md):将src0和src1中的元素解交织存入dst0和dst1中,本接口提供矢量数据解交织和掩码解交织两种功能模式:16+- [排布变换](reg_permute_sel/index.md):共 11 张知识卡(如 asc_deintlv、asc_intlv、asc_pack_to_high 等)。
17-- [归约计算](reg_reduce/index.md):根据mask将源操作数src中每两个相邻元素相加,得到计算结果。结果连续保存在目的操作数dst的低半部分,dst的高半部分置0。以half数据类型为例,相加示意图如下:17+- [归约计算](reg_reduce/index.md):共 7 张知识卡(如 asc_pair_reduce_sum、asc_reduce_max、asc_reduce_max_datablock 等)。
18- [同步控制](reg_sync/index.md):Reg矢量计算内不同流水线之间的同步指令。该同步指令指定源流水线和目的流水线,如下图所示,目的流水线将等待源流水线上所有指令完成才进行执行。 读写场景下,当读指令使用的寄存器和写指令使用的寄存器相同时,可以触发寄存器保序,指令将会按照代码顺序执行,不需要插入同步指令,而当使用的寄存器不同时,如果要确保读写指令执行,则需要插入同步指令。写写场景同理。18- [同步控制](reg_sync/index.md):Reg矢量计算内不同流水线之间的同步指令。该同步指令指定源流水线和目的流水线,如下图所示,目的流水线将等待源流水线上所有指令完成才进行执行。 读写场景下,当读指令使用的寄存器和写指令使用的寄存器相同时,可以触发寄存器保序,指令将会按照代码顺序执行,不需要插入同步指令,而当使用的寄存器不同时,如果要确保读写指令执行,则需要插入同步指令。写写场景同理。
19- [Reg离散搬出](scatter/index.md):根据索引寄存器中的索引值,将源矢量数据寄存器中的元素分散搬出到Unified Buffer(UB)。mask用于指示参与搬出的元素,掩码为1的源元素写入dst中对应索引位置,掩码为0的源元素不写入目的地址,保留目的地址原数据。搬运过程中数据格式与内容保持不变。本接口在Vector Function(__simd_vf__标记的函数)内使用。19- [Reg离散搬出](scatter/index.md):根据索引寄存器中的索引值,将源矢量数据寄存器中的元素分散搬出到Unified Buffer(UB)。mask用于指示参与搬出的元素,掩码为1的源元素写入dst中对应索引位置,掩码为0的源元素不写入目的地址,保留目的地址原数据。搬运过程中数据格式与内容保持不变。本接口在Vector Function(__simd_vf__标记的函数)内使用。
20-- [Reg数据搬出](store/index.md):将矢量数据寄存器中的数据搬出到Unified Buffer(UB),目的地址需按照dtype对齐。底层经非对齐搬出指令完成单次连续写入,搬运过程中数据格式和内容保持不变。在追求极致性能时,对齐场景推荐使用asc_storealign接口。本接口在Vector Function(__simd_vf__标记的函数)内使用。本接口提供两种参数列表不同的功能模式:20+- [Reg数据搬出](store/index.md):共 17 张知识卡(如 asc_store、asc_storealign、asc_storealign_1st 等)。
21-- [ub gather](ub_gather/index.md):根据索引位置index将源操作数src按元素收集到目的操作数dst中。21+- [ub gather](ub_gather/index.md):共 2 张知识卡(如 asc_gather、asc_gather_datablock)。
@@ -1,6 +1,6 @@
1# 标量计算1# 标量计算
2 2 
3-- [位运算](scalar_bit/index.md):对一个64位无符号整数执行位清除操作,将bits中第idx位清零,其余位保持不变。返回修改后的64位无符号整数。3+- [位运算](scalar_bit/index.md):共 8 张知识卡(如 asc_clear_nthbit、asc_clz、asc_ffs 等)。
4-- [类型转换](scalar_convert/index.md):将float类型的标量转换为half类型并返回,舍入模式固定为ODD(最近邻奇数舍入)。4+- [类型转换](scalar_convert/index.md):共 2 张知识卡(如 asc_float2half、asc_float2int32)。
5- [数据搬入](scalar_load/index.md):不经过DCache直接从GM地址读取整型数据。 当多核操作GM地址且数据无法对齐到Cache Line时,经过DCache读写可能引入Cache Line粒度的数据覆盖。此时,可使用本接口绕过DCache读取GM数据。5- [数据搬入](scalar_load/index.md):不经过DCache直接从GM地址读取整型数据。 当多核操作GM地址且数据无法对齐到Cache Line时,经过DCache读写可能引入Cache Line粒度的数据覆盖。此时,可使用本接口绕过DCache读取GM数据。
6- [数据搬出](scalar_store/index.md):不经过DCache向GM地址上写数据。 当多核操作GM地址时,如果数据无法对齐到Cache Line,经过DCache的方式下,由于按照Cache Line大小进行读写,会导致多核数据随机覆盖的问题。此时,可以采用不经过DCache直接读写GM地址的方式,从而避免上述随机覆盖的问题。6- [数据搬出](scalar_store/index.md):不经过DCache向GM地址上写数据。 当多核操作GM地址时,如果数据无法对齐到Cache Line,经过DCache的方式下,由于按照Cache Line大小进行读写,会导致多核数据随机覆盖的问题。此时,可以采用不经过DCache直接读写GM地址的方式,从而避免上述随机覆盖的问题。
@@ -2,4 +2,4 @@
2 2 
3- [系统初始化](sys_init/index.md):本接口将当前AI Core上可能由前序算子遗留的全局状态恢复为接口约定的初始状态,避免这些状态影响当前算子的计算结果。本接口仅初始化控制状态,不会清零Global Memory、Local Memory或用户数据。3- [系统初始化](sys_init/index.md):本接口将当前AI Core上可能由前序算子遗留的全局状态恢复为接口约定的初始状态,避免这些状态影响当前算子的计算结果。本接口仅初始化控制状态,不会清零Global Memory、Local Memory或用户数据。
4- [系统其他操作](sys_misc/index.md):生成一条空操作指令,不执行实际计算。每调用一次阻塞一拍流水。4- [系统其他操作](sys_misc/index.md):生成一条空操作指令,不执行实际计算。每调用一次阻塞一拍流水。
5-- [系统资源和变量](sys_var/index.md):获取当前AI处理器架构版本号。5+- [系统资源和变量](sys_var/index.md):共 15 张知识卡(如 asc_get_arch_ver、asc_get_core_id、asc_get_ffts_base_addr 等)。
@@ -1,13 +1,13 @@
1# Memory矢量计算1# Memory矢量计算
2 2 
3-- [基础算术](vector_arith/index.md):按元素取绝对值,计算公式如下: $$ dst_i = |src_i| $$3+- [基础算术](vector_arith/index.md):共 17 张知识卡(如 asc_abs、asc_add、asc_add_scalar 等)。
4-- [广播操作](vector_broadcast/index.md):给定一个输入矢量,每一次取输入矢量中的8个数填充到结果矢量的8个datablock(32Bytes)中去,每个数对应一个datablock。4+- [广播操作](vector_broadcast/index.md):共 2 张知识卡(如 asc_brcb、asc_duplicate)。
5-- [比较计算](vector_compare/index.md):比较src0与src1在对应索引位置的元素大小。若比较结果为真,则输出结果的对应比特位设为1,否则设为0。5+- [比较计算](vector_compare/index.md):共 14 张知识卡(如 asc_eq、asc_eq_scalar、asc_ge 等)。
6-- [类型转换](vector_convert/index.md):将bfloat16_t类型数据转为float,无舍入模式。6+- [类型转换](vector_convert/index.md):共 24 张知识卡(如 asc_bfloat162float、asc_bfloat162int32、asc_float2bfloat16 等)。
7-- [复合计算](vector_fused/index.md):按元素求和,再进行Relu计算(结果和0对比取较大值),并提供转换最终结果的数据类型的功能(s162s8、f322f16、f162s8)。计算公式如下:7+- [复合计算](vector_fused/index.md):共 14 张知识卡(如 asc_add_relu、asc_axpy、asc_deq_int162b8 等)。
8-- [聚合操作](vector_gather/index.md):将src按照src_offset给定的偏移按元素收集到dst中。8+- [聚合操作](vector_gather/index.md):共 2 张知识卡(如 asc_gather、asc_gather_datablock)。
9-- [逻辑计算](vector_logic/index.md):执行矢量与运算。计算公式如下:9+- [逻辑计算](vector_logic/index.md):共 5 张知识卡(如 asc_and、asc_not、asc_or 等)。
10-- [掩码操作](vector_mask_config/index.md):设置Mask模式为Counter模式。该模式下,不需要开发者去感知迭代次数、处理非对齐的尾块等操作,可直接传入数据计算数量,实际迭代次数由Vector计算单元自动推断。在前n个数据计算类的接口内部会自动设置该模式,不需要自行设置。10+- [掩码操作](vector_mask_config/index.md):共 3 张知识卡(如 asc_set_mask_count、asc_set_mask_norm、asc_set_vector_mask)。
11-- [排布变换](vector_permute_sel/index.md):执行矢量选择操作,给定两个源操作数src0和src1,根据条件选择元素,得到目的操作数dst。此接口配合asc_set_cmp_mask(sel_mask)设置,规则为:当sel_mask的比特位是1时,从src0选取;当比特位是0时,从src1选择。11+- [排布变换](vector_permute_sel/index.md):共 5 张知识卡(如 asc_select、asc_set_va_reg、asc_squeeze 等)。
12-- [归约计算](vector_reduce/index.md):执行数据块内的求最大值规约(Reduce Max)操作。12+- [归约计算](vector_reduce/index.md):共 11 张知识卡(如 asc_datablock_reduce_max、asc_datablock_reduce_min、asc_datablock_reduce_sum 等)。
13-- [排序计算](vector_sort/index.md):位图排序函数,一次迭代可以完成32个元素的排序,数据按如下描述结构进行保存:13+- [排序计算](vector_sort/index.md):共 3 张知识卡(如 asc_bitsort、asc_get_vms4_sr、asc_mrgsort4)。
@@ -1,10 +1,10 @@
1# 矢量数据搬运1# 矢量数据搬运
2 2 
3- [asc_copy](asc_copy.md):将数据从Unified Buffer(UB)搬运到UB。支持高维切分。3- [asc_copy](asc_copy.md):将数据从Unified Buffer(UB)搬运到UB。支持高维切分。
4-- [asccopygm2ub](asc_copy_gm2ub/index.md):将数据从Global Memory (GM)搬运到Unified Buffer(UB)。4+- [asccopygm2ub](asc_copy_gm2ub/index.md):共 2 张知识卡(如 asc_copy_gm2ub、asc_copy_gm2ub)。
5-- [asccopygm2ubalign](asc_copy_gm2ub_align/index.md):提供数据非对齐搬运的功能,将数据从Global Memory (GM)搬运到Unified Buffer(UB),并支持8位/16位/32位数据类型搬运。5+- [asccopygm2ubalign](asc_copy_gm2ub_align/index.md):共 2 张知识卡(如 asc_copy_gm2ub_align、asc_copy_gm2ub_align)。
6-- [asccopyub2gm](asc_copy_ub2gm/index.md):将数据从Unified Buffer(UB)搬运到Global Memory (GM)。6+- [asccopyub2gm](asc_copy_ub2gm/index.md):共 2 张知识卡(如 asc_copy_ub2gm、asc_copy_ub2gm)。
7-- [asccopyub2gmalign](asc_copy_ub2gm_align/index.md):将数据从Unified Buffer(UB)搬运到Global Memory (GM),支持8位/16位/32位分块拷贝操作。7+- [asccopyub2gmalign](asc_copy_ub2gm_align/index.md):共 2 张知识卡(如 asc_copy_ub2gm_align、asc_copy_ub2gm_align)。
8- [asc_copy_ub2l1](asc_copy_ub2l1.md):提供数据搬运功能,将数据从Unified Buffer (UB)搬运到L1 Buffer。8- [asc_copy_ub2l1](asc_copy_ub2l1.md):提供数据搬运功能,将数据从Unified Buffer (UB)搬运到L1 Buffer。
9- [asc_copy_ub2ub](asc_copy_ub2ub.md):将数据从Unified Buffer(UB)搬运到UB。9- [asc_copy_ub2ub](asc_copy_ub2ub.md):将数据从Unified Buffer(UB)搬运到UB。
10- [asc_ndim_copy_gm2ub](asc_ndim_copy_gm2ub.md):多维数据搬运接口,将数据从Global Memory (GM)搬运到Unified Buffer(UB)。相比于基础数据搬移接口,可以更加自由配置搬入的维度信息。 本功能最多能处理5个维度,每个维度都可以单独配置Stride,且支持随路Padding功能。相关配置参数有很多,需要通过不同的接口传入。10- [asc_ndim_copy_gm2ub](asc_ndim_copy_gm2ub.md):多维数据搬运接口,将数据从Global Memory (GM)搬运到Unified Buffer(UB)。相比于基础数据搬移接口,可以更加自由配置搬入的维度信息。 本功能最多能处理5个维度,每个维度都可以单独配置Stride,且支持随路Padding功能。相关配置参数有很多,需要通过不同的接口传入。
@@ -1,7 +1,7 @@
1# SIMD API1# SIMD API
2 2 
3-- [高阶API](adv_api/index.md):采用GELU作为激活函数的GLU变体,按元素对两源Tensor计算GeGLU结果3+- [高阶API](adv_api/index.md):共 487 张知识卡,分布在 15 个子目录(activation_functions、convolution_compute、cube_compute 等)。
4-- [基础API](basic_api/index.md):对GM地址上的值执行原子加操作,将旧值与输入标量求和后写回,返回修改前的旧值。4+- [基础API](basic_api/index.md):共 580 张知识卡,分布在 17 个子目录(atomic_operations、aux_data_structures、cache_control 等)。
5-- [API列表](c_api/index.md):C API文档按最细粒度公共头文件归类。除试验接口外,整体使用时可以包含c_api/asc_simd.h。5+- [API列表](c_api/index.md):C API文档按最细粒度公共头文件归类。除试验接口外,整体使用时可以包含c_api/asc_simd.h。(共 715 张知识卡:C API、通用说明和约束(C API) 等)
6- [通用说明和约束(SIMD API)](general_description_and_constraints.md):说明 SIMD 基础与高阶 API 的头文件、链接方式、存储单元对齐要求及 TPosition 内存映射。6- [通用说明和约束(SIMD API)](general_description_and_constraints.md):说明 SIMD 基础与高阶 API 的头文件、链接方式、存储单元对齐要求及 TPosition 内存映射。
7-- [Tensor API](tensor_api/index.md):Tensor API的Algorithm层为基于Tensor的数据搬运和矩阵计算提供统一的调用入口。其中,copy和mmad接口分别用于执行数据搬运和矩阵计算,make_copy和make_mmad接口用于构造对应的Atom原子对象。7+- [Tensor API](tensor_api/index.md):Tensor API是一套面向Ascend C算子开发的基础编程接口,使用Tensor和Layout统一描述数据的存储位置、数据类型、尺寸大小和排列规则,并基于这些信息完成数据搬运、矩阵计算和矢量计算,减少开发者手工组织硬件指令参数和数据排布的工作量。(共 107 张知识卡:Algorithm 等)
@@ -2,10 +2,10 @@
2 2 
3- [Algorithm](algorithm.md):Tensor API的Algorithm层为基于Tensor的数据搬运和矩阵计算提供统一的调用入口。其中,copy和mmad接口分别用于执行数据搬运和矩阵计算,make_copy和make_mmad接口用于构造对应的Atom原子对象。3- [Algorithm](algorithm.md):Tensor API的Algorithm层为基于Tensor的数据搬运和矩阵计算提供统一的调用入口。其中,copy和mmad接口分别用于执行数据搬运和矩阵计算,make_copy和make_mmad接口用于构造对应的Atom原子对象。
4- [数据搬运](data_move/index.md):Tensor API通过copy统一执行不同存储位置之间的数据搬运。Tensor表达不同存储位置和类型的数据,Layout表达Tensor的尺寸大小和排列规则。4- [数据搬运](data_move/index.md):Tensor API通过copy统一执行不同存储位置之间的数据搬运。Tensor表达不同存储位置和类型的数据,Layout表达Tensor的尺寸大小和排列规则。
5-- [废弃接口](deprecated_interface/index.md):头文件路径为:tensor_api/tensor.h。5+- [废弃接口](deprecated_interface/index.md):共 16 张知识卡(如 layout::Capacity(废弃)、layout::Get(废弃)、layout::Rank(废弃) 等)。
6-- [layout](layout/index.md):头文件路径为:tensor_api/tensor.h。6+- [layout](layout/index.md):layout数据结构是描述多维张量内存布局的基础模板类,通过编译期的形状和步长信息,实现逻辑坐标空间到一维内存地址空间的映射。借助模板元编程技术,该类在编译期完成计算和代码生成,从而降低运行时开销。(共 23 张知识卡:capacity、coord、coshape 等)
7-- [矩阵计算(Tensor API)](matrix_compute/index.md):矩阵计算由数据搬入、mmad计算和结果搬出组成。各阶段通过Tensor的存储位置和Layout选择对应Arch通路,通过Atom配置Trait及运行时Params。7+- [矩阵计算(Tensor API)](matrix_compute/index.md):矩阵计算由数据搬入、mmad计算和结果搬出组成。各阶段通过Tensor的存储位置和Layout选择对应Arch通路,通过Atom配置Trait及运行时Params。(共 27 张知识卡:mmad计算辅助配置接口、概述 等)
8-- [reg矢量计算(Tensor API)](reg_vector_compute/index.md):头文件路径:"tensor_api/experimental/vector_compute.h"。8+- [reg矢量计算(Tensor API)](reg_vector_compute/index.md):共 26 张知识卡,分布在 12 个子目录(basic_arithmetic、compare_and_select、data_padding 等)。
9-- [tensor](tensor/index.md):头文件路径为:tensor_api/tensor.h。9+- [tensor](tensor/index.md):头文件路径为:tensor_api/tensor.h。(共 7 张知识卡:make_mem_ptr、make_tensor、pointer 等)
10- [Tensor API](tensor_api.md):Tensor API是一套面向Ascend C算子开发的基础编程接口,使用Tensor和Layout统一描述数据的存储位置、数据类型、尺寸大小和排列规则,并基于这些信息完成数据搬运、矩阵计算和矢量计算,减少开发者手工组织硬件指令参数和数据排布的工作量。10- [Tensor API](tensor_api.md):Tensor API是一套面向Ascend C算子开发的基础编程接口,使用Tensor和Layout统一描述数据的存储位置、数据类型、尺寸大小和排列规则,并基于这些信息完成数据搬运、矩阵计算和矢量计算,减少开发者手工组织硬件指令参数和数据排布的工作量。
11-- [utils](utils/index.md):头文件路径为:tensor_api/tensor.h。11+- [utils](utils/index.md):共 5 张知识卡(如 编译期整型别名、维度保留标记、get_layout_pattern 等)。
@@ -1,10 +1,10 @@
1# 矩阵计算(Tensor API)1# 矩阵计算(Tensor API)
2 2 
3-- [矩阵计算的搬入](cube_compute_load/index.md):头文件路径为:tensor_api/tensor.h。3+- [矩阵计算的搬入](cube_compute_load/index.md):共 9 张知识卡(如 copy(Global Memory到L1 Buffer数据搬运)、copy(L1 Buffer到BiasTable Buffer数据搬运)、copy(L1 Buffer到Fixpipe Buffer数据搬运) 等)。
4-- [矩阵计算的搬出](cube_compute_store/index.md):头文件路径为:tensor_api/tensor.h。4+- [矩阵计算的搬出](cube_compute_store/index.md):共 4 张知识卡(如 copy(L0C Buffer到Global Memory数据搬运)、copy(L0C Buffer到L1 Buffer数据搬运)、copy(L0C Buffer到Unified Buffer数据搬运) 等)。
5-- [矩阵搬出关键特性说明](cube_store_key_features/index.md):源张量为NZ格式,目的张量支持ND、DN或NZ格式。用户只需要在构造layout时传入Batch参数,copy调用形式与非batch模式保持一致。5+- [矩阵搬出关键特性说明](cube_store_key_features/index.md):共 6 张知识卡(如 batch搬运、F32 Channel Split、Int8 Channel Merge 等)。
6- [矩阵计算(Tensor API)](matrix_compute.md):矩阵计算由数据搬入、mmad计算和结果搬出组成。各阶段通过Tensor的存储位置和Layout选择对应Arch通路,通过Atom配置Trait及运行时Params。6- [矩阵计算(Tensor API)](matrix_compute.md):矩阵计算由数据搬入、mmad计算和结果搬出组成。各阶段通过Tensor的存储位置和Layout选择对应Arch通路,通过Atom配置Trait及运行时Params。
7- [矩阵计算](mmad_compute/index.md):头文件路径为:tensor_api/tensor.h。7- [矩阵计算](mmad_compute/index.md):头文件路径为:tensor_api/tensor.h。
8- [mmad计算辅助配置接口](mmad_compute_aux_config.md):mmad计算辅助配置接口列表如下:8- [mmad计算辅助配置接口](mmad_compute_aux_config.md):mmad计算辅助配置接口列表如下:
9-- [mmad计算关键特性说明](mmad_compute_key_features/index.md):GEMV的核心功能体现为:当矩阵A的维度M取值为1时,接口可以手动启用GEMV功能,该操作退化为1×K维度的行向量与K×N维度矩阵之间的乘法运算。9+- [mmad计算关键特性说明](mmad_compute_key_features/index.md):本节介绍mmad矩阵乘计算中的GEMV、HF32和unit_flag等关键特性,帮助理解不同计算模式、精度模式和细粒度同步机制。(共 4 张知识卡:GEMV、HF32、unit_flag)
10- [概述](overview.md):矩阵计算单元的详细描述请参考LINK。10- [概述](overview.md):矩阵计算单元的详细描述请参考LINK。
@@ -1,14 +1,14 @@
1# reg矢量计算(Tensor API)1# reg矢量计算(Tensor API)
2 2 
3-- [基础算术](basic_arithmetic/index.md):头文件路径:"tensor_api/experimental/vector_compute.h"。3+- [基础算术](basic_arithmetic/index.md):共 5 张知识卡(如 log、max、operator+ 等)。
4- [概述](compare_and_select/index.md):头文件路径:"include/tensor_api/experimental/arch/vector/compare_and_select.h"。4- [概述](compare_and_select/index.md):头文件路径:"include/tensor_api/experimental/arch/vector/compare_and_select.h"。
5- [数据填充](data_padding/index.md):头文件路径:"include/tensor_api/experimental/arch/vector/compare_and_select.h"。5- [数据填充](data_padding/index.md):头文件路径:"include/tensor_api/experimental/arch/vector/compare_and_select.h"。
6- [数据重排](data_reorder/index.md):头文件路径:"include/tensor_api/experimental/arch/vector/data_reorder.h"。6- [数据重排](data_reorder/index.md):头文件路径:"include/tensor_api/experimental/arch/vector/data_reorder.h"。
7- [逻辑计算](logical_compute/index.md):头文件路径:"tensor_api/experimental/vector_compute.h"。7- [逻辑计算](logical_compute/index.md):头文件路径:"tensor_api/experimental/vector_compute.h"。
8-- [mask寄存器计算](mask_reg_compute/index.md):头文件路径为:"tensor_api/arch/vector/experimental/mask_reg_compute.h",命名空间为asc::te::experimental。8+- [mask寄存器计算](mask_reg_compute/index.md):共 6 张知识卡(如 all_mask、deinterleave(掩码解交织)、make_mask 等)。
9-- [寄存器数据搬入](reg_data_load/index.md):头文件路径为:"tensor_api/experimental/arch/vector/reg_data_load.h",命名空间为asc::te::experimental。9+- [寄存器数据搬入](reg_data_load/index.md):共 2 张知识卡(如 load、load_broadcast)。
10- [寄存器数据搬出](reg_data_store/index.md):头文件路径为:"tensor_api/experimental/arch/vector/reg_data_store.h",命名空间为asc::te::experimental。10- [寄存器数据搬出](reg_data_store/index.md):头文件路径为:"tensor_api/experimental/arch/vector/reg_data_store.h",命名空间为asc::te::experimental。
11-- [reg_tensor](reg_tensor/index.md):头文件路径为:"tensor_api/arch/vector/experimental/reg_tensor.h",命名空间为asc::te::experimental。11+- [reg_tensor](reg_tensor/index.md):头文件路径为:"tensor_api/arch/vector/experimental/reg_tensor.h",命名空间为asc::te::experimental。(共 2 张知识卡:reg_pair)
12-- [类型转换](type_conversion/index.md):头文件路径:"tensor_api/arch/vector/experimental/type_conversion.h"。12+- [类型转换](type_conversion/index.md):共 2 张知识卡(如 cast、trunc)。
13-- [矢量计算的搬入](vector_compute_load/index.md):头文件路径为:tensor_api/tensor.h,命名空间为asc::te。13+- [矢量计算的搬入](vector_compute_load/index.md):共 3 张知识卡(如 copy(Global Memory到Unified Buffer的数据搬运)、copy(Unified Buffer到L1 Buffer的数据搬运)、copy(Unified Buffer到Unified Buffer的数据搬运))。
14- [矢量计算的搬出](vector_compute_store/index.md):头文件路径为:tensor_api/tensor.h,命名空间为asc::te。14- [矢量计算的搬出](vector_compute_store/index.md):头文件路径为:tensor_api/tensor.h,命名空间为asc::te。
@@ -1,12 +1,12 @@
1# 协作组1# 协作组
2 2 
3- [binary_partition](binary_partition.md):根据标签将父组划分为两个子组,标签相同的线程被分配到同一组中。3- [binary_partition](binary_partition.md):根据标签将父组划分为两个子组,标签相同的线程被分配到同一组中。
4-- [coalescedgroup](coalesced_group/index.md):判断coalesced_group组内所有线程的输入是否均不为0。4+- [coalescedgroup](coalesced_group/index.md):共 14 张知识卡(如 CoalescedGroup协作组all、CoalescedGroup协作组any、Ballot 等)。
5- [协作组简介](cooperative_groups_intro.md):协作组(cooperative_groups)是Ascend C SIMT编程模型的一种拓展,用于组织协作线程组。通过协作组能够更细粒度地管理线程,从而实现更高效的线程并行协作。5- [协作组简介](cooperative_groups_intro.md):协作组(cooperative_groups)是Ascend C SIMT编程模型的一种拓展,用于组织协作线程组。通过协作组能够更细粒度地管理线程,从而实现更高效的线程并行协作。
6-- [gridgroup](grid_group/index.md):获取当前线程所属线程块在Grid中的三维索引。6+- [gridgroup](grid_group/index.md):共 12 张知识卡(如 BlockIndex、BlockRank、DimBlocks 等)。
7- [sync(协作组)](sync.md):同步指定协作组内的所有线程,功能等价于调用该协作组对象的成员函数sync()。7- [sync(协作组)](sync.md):同步指定协作组内的所有线程,功能等价于调用该协作组对象的成员函数sync()。
8-- [threadblock](thread_block/index.md):获取当前线程所属线程块的线程配置,返回值与内置变量blockDim相同。8+- [threadblock](thread_block/index.md):共 10 张知识卡(如 DimThreads、GroupDim、GroupIndex 等)。
9-- [threadblocktile](thread_block_tile/index.md):判断thread_block_tile组内所有线程的输入是否均不为0。9+- [threadblocktile](thread_block_tile/index.md):共 15 张知识卡(如 ThreadBlockTile协作组all、ThreadBlockTile协作组any、Ballot 等)。
10- [tiled_partition](tiled_partition.md):将线程组划分为多个更小、固定大小的子组,提供模板和非模板版本,分别用于编译时和运行时确定划分大小,模板版本支持跨Warp。10- [tiled_partition](tiled_partition.md):将线程组划分为多个更小、固定大小的子组,提供模板和非模板版本,分别用于编译时和运行时确定划分大小,模板版本支持跨Warp。
11 11 
12-- [thread group](thread_group/index.md):头文件路径为:"simt_api/cooperative_groups.h"。12+- [thread group](thread_group/index.md):共 7 张知识卡(如 get_type、num_threads、size 等)。
@@ -1,13 +1,13 @@
1# SIMT API1# SIMT API
2 2 
3-- [地址空间转换函数](address_space_conversion_functions/index.md):将输入的指针转换为其指向的Global Memory内存空间的地址值并返回。3+- [地址空间转换函数](address_space_conversion_functions/index.md):共 6 张知识卡(如 __cvta_generic_to_global、__cvta_generic_to_local、__cvta_generic_to_ubuf 等)。
4-- [地址空间谓词函数](address_space_predicate_functions/index.md):判断输入的指针是否指向Global Memory内存空间的地址。4+- [地址空间谓词函数](address_space_predicate_functions/index.md):共 3 张知识卡(如 __isGlobal、__isLocal、__isUbuf)。
5-- [原子操作](atomic_operations/index.md):对Unified Buffer或Global Memory中的数据与指定数据执行原子加操作,即将指定数据累加到这些内存区域的数据中。5+- [原子操作](atomic_operations/index.md):共 12 张知识卡(如 asc_atomic_add、asc_atomic_and、asc_atomic_cas 等)。
6-- [协作组](cooperative_groups/index.md):根据标签将父组划分为两个子组,标签相同的线程被分配到同一组中。6+- [协作组](cooperative_groups/index.md):共 62 张知识卡(如 binary_partition、协作组简介、sync(协作组) 等)。
7-- [数学函数](math_functions/index.md):获取输入数据的绝对值。7+- [数学函数](math_functions/index.md):共 638 张知识卡,分布在 6 个子目录(bfloat16_type、data_type_conversion、float_math_functions 等)。
8-- [访存函数](memory_access_functions/index.md):该接口用于刷新核内的整个Data Cache,保证数据读取时Cache的一致性。 在SIMT编程中,写入数据时会立即写入Global Memory,使其他核可见,因此不存在核间一致性问题。当从Global Memory读取数据时,该数据可能已被其他核修改,此时应使用dcci接口直接访问Global Memory,以获取最新数据。 请注意,该接口耗时较大,性能敏感的场景慎用。8+- [访存函数](memory_access_functions/index.md):共 8 张知识卡(如 asc_dcci_entire、asc_dcci_single、AscLdca 等)。
9- [SIMT API概述](overview.md):SIMT API是面向AI处理器的并行计算编程接口,可以实现高效的数据并行计算。SIMT API支持两种编程模型:SIMT编程、SIMD与SIMT混合编程,用户可先阅读SIMT编程简介和SIMD与SIMT混合编程简介,以了解编程基础,后续章节将详细介绍API接口。9- [SIMT API概述](overview.md):SIMT API是面向AI处理器的并行计算编程接口,可以实现高效的数据并行计算。SIMT API支持两种编程模型:SIMT编程、SIMD与SIMT混合编程,用户可先阅读SIMT编程简介和SIMD与SIMT混合编程简介,以了解编程基础,后续章节将详细介绍API接口。
10-- [SIMD与SIMT混合编程简介](simd_simt_hybrid_programming_intro/index.md):SIMD与SIMT混合编程支持根据算子特征精细化映射,规则计算用SIMD获得高吞吐,不规则计算用SIMT缓解发散。10+- [SIMD与SIMT混合编程简介](simd_simt_hybrid_programming_intro/index.md):共 7 张知识卡(如 ProgrammingModel 等)。
11-- [SIMT编程简介](simt_programming_intro/index.md):SIMT编程模型采用Grid、Thread Block、Warp、Thread分层线程结构,适合控制流复杂和访存模式不规则的算子场景。11+- [SIMT编程简介](simt_programming_intro/index.md):共 7 张知识卡(如 编程模型 等)。
12-- [同步与内存栅栏](sync_and_memory_fence/index.md):SIMT程序中,不同Thread的执行进度和内存访问到达顺序可能不同。当线程之间存在数据依赖时,需要使用同步接口或内存栅栏接口约束执行顺序和内存可见性。12+- [同步与内存栅栏](sync_and_memory_fence/index.md):共 4 张知识卡(如 同步与内存栅栏简介(同步与内存屏障) 等)。
13-- [Warp函数](warp_functions/index.md):Warp是SIMT线程执行和调度的基本单位。一个线程块内的线程会按照线性线程号被划分为多个Warp,当前一个Warp包含32个线程。同一Warp内的线程执行相同的指令流,但每个线程拥有独立的寄存器和Lane ID,可处理不同的数据地址和分支路径。13+- [Warp函数](warp_functions/index.md):共 18 张知识卡(如 Warp函数 等)。
@@ -1,9 +1,9 @@
1# bfloat16类型1# bfloat16类型
2 2 
3-- [bfloat16类型算术函数](bfloat16_arithmetic_functions/index.md):获取输入数据的绝对值。3+- [bfloat16类型算术函数](bfloat16_arithmetic_functions/index.md):共 8 张知识卡(如 __habs、__hadd、__hdiv 等)。
4-- [bfloat16类型比较函数](bfloat16_compare_functions/index.md):比较两个bfloat16类型数据是否相等,相等时返回true。4+- [bfloat16类型比较函数](bfloat16_compare_functions/index.md):共 18 张知识卡(如 __heq、__hequ、__hge 等)。
5-- [bfloat16类型数学库函数](bfloat16_math_functions/index.md):获取大于或等于输入bfloat16_t数据的最小整数值。5+- [bfloat16类型数学库函数](bfloat16_math_functions/index.md):共 16 张知识卡(如 Hceil、BFloat16数学函数hcos、BFloat16数学函数hexp 等)。
6-- [bfloat16类型精度转换函数](bfloat16_precision_conversion_functions/index.md):将bfloat16x2的两个分量分别转换为float,并填充到float2后返回。6+- [bfloat16类型精度转换函数](bfloat16_precision_conversion_functions/index.md):共 108 张知识卡(如 __bfloat1622float2、__bfloat162bfloat162、__bfloat162bfloat16_rd 等)。
7-- [bfloat16x2类型算术函数](bfloat16x2_arithmetic_functions/index.md):计算输入bfloat16x2_t类型数据各分量的绝对值。7+- [bfloat16x2类型算术函数](bfloat16x2_arithmetic_functions/index.md):共 9 张知识卡(如 __habsx2、__haddx2、__hcmadd 等)。
8-- [bfloat16x2类型比较函数](bfloat16x2_compare_functions/index.md):比较两个bfloat16x2_t类型数据的两个分量是否相等,当两个分量均相等时返回true。若任一输入的分量为nan,该分量的比较结果为true。8+- [bfloat16x2类型比较函数](bfloat16x2_compare_functions/index.md):共 41 张知识卡(如 __hbequx2、__hbeqx2、__hbgeux2 等)。
9-- [bfloat16x2类型数学库函数](bfloat16x2_math_functions/index.md):获取大于或等于输入数据各元素的最小整数值。9+- [bfloat16x2类型数学库函数](bfloat16x2_math_functions/index.md):共 16 张知识卡(如 H2Ceil、H2Cos、H2Exp 等)。
@@ -1,4 +1,4 @@
1# 数据类型转换1# 数据类型转换
2 2 
3- [数据类型转换概述](overview.md):类型转换包括隐式转换和显式转换。隐式转换在类型不匹配时由系统自动执行。显式转换则通过调用类型转换API来主动完成,API会将输入数据从源类型转换为目标类型,并返回转换结果。整个转换过程遵循规定的处理规则。类型转换的结果受舍入模式和饱和模式的影响。关于浮点数的表示方式和舍入规则,请参见类型转换的相关说明。3- [数据类型转换概述](overview.md):类型转换包括隐式转换和显式转换。隐式转换在类型不匹配时由系统自动执行。显式转换则通过调用类型转换API来主动完成,API会将输入数据从源类型转换为目标类型,并返回转换结果。整个转换过程遵循规定的处理规则。类型转换的结果受舍入模式和饱和模式的影响。关于浮点数的表示方式和舍入规则,请参见类型转换的相关说明。
4-- [类型转换函数](type_conversion_functions/index.md):将浮点数四舍五入取整,并遵循CAST_FLOOR模式。4+- [类型转换函数](type_conversion_functions/index.md):共 49 张知识卡(如 __float2float_rd、__float2float_rn、__float2float_rna 等)。
@@ -1,4 +1,4 @@
1# fp8类型1# fp8类型
2 2 
3-- [fp8数据转换函数](fp8_conversion_functions/index.md):将float2的两个分量按CAST_RINT模式转换为FP8类型(e4m3或e5m2)的8位浮点数,支持饱和与非饱和溢出处理,结果以位级打包存储。3+- [fp8数据转换函数](fp8_conversion_functions/index.md):共 14 张知识卡(如 __asc_cvt_float2_to_fp8x2、__asc_cvt_float_to_fp8、__e4m3x22float2 等)。
4- [fp8数据类型简介(FP8类型)](fp8_data_type_intro.md):SIMT编程支持3种fp8数据类型,分别是float8_e4m3_t、float8_e5m2_t、hifloat8_t。其数据范围如下:4- [fp8数据类型简介(FP8类型)](fp8_data_type_intro.md):SIMT编程支持3种fp8数据类型,分别是float8_e4m3_t、float8_e5m2_t、hifloat8_t。其数据范围如下:
@@ -1,9 +1,9 @@
1# half类型1# half类型
2 2 
3-- [half2类型算术函数](half2_arithmetic_functions/index.md):计算输入half2类型数据各分量的绝对值。3+- [half2类型算术函数](half2_arithmetic_functions/index.md):共 9 张知识卡(如 __habsx2、__haddx2、__hcmadd 等)。
4-- [half2类型比较函数](half2_compare_functions/index.md):比较两个half2类型数据的两个分量是否相等,当两个分量均相等时返回true。若任一输入的分量为nan,该分量的比较结果为true。4+- [half2类型比较函数](half2_compare_functions/index.md):共 41 张知识卡(如 __hbequx2、__hbeqx2、__hbgeux2 等)。
5-- [half2类型数学库函数](half2_math_functions/index.md):获取大于或等于输入数据各元素的最小整数值。5+- [half2类型数学库函数](half2_math_functions/index.md):共 16 张知识卡(如 H2Ceil、H2Cos、H2Exp 等)。
6-- [half类型算术函数](half_arithmetic_functions/index.md):获取输入数据的绝对值。6+- [half类型算术函数](half_arithmetic_functions/index.md):共 8 张知识卡(如 __habs、__hadd、__hdiv 等)。
7-- [half类型比较函数](half_compare_functions/index.md):比较两个half类型数据是否相等,相等时返回true。7+- [half类型比较函数](half_compare_functions/index.md):共 18 张知识卡(如 __heq、__hequ、__hge 等)。
8-- [half类型数学库函数](half_math_functions/index.md):获取大于或等于输入half数据的最小整数值。8+- [half类型数学库函数](half_math_functions/index.md):共 16 张知识卡(如 Hceil、Half数学函数hcos、Half数学函数hexp 等)。
9-- [half类型精度转换函数](half_precision_conversion_functions/index.md):遵循CAST_FLOOR模式,将bfloat16类型数据转换为half类型数据,返回转换后的值。9+- [half类型精度转换函数](half_precision_conversion_functions/index.md):共 127 张知识卡(如 __bfloat162half_rd、__bfloat162half_rd_sat、__bfloat162half_rn 等)。
@@ -1,8 +1,8 @@
1# 数学函数1# 数学函数
2 2 
3-- [bfloat16类型](bfloat16_type/index.md):获取输入数据的绝对值。3+- [bfloat16类型](bfloat16_type/index.md):共 216 张知识卡,分布在 7 个子目录(bfloat16_arithmetic_functions、bfloat16_compare_functions、bfloat16_math_functions 等)。
4-- [数据类型转换](data_type_conversion/index.md):类型转换包括隐式转换和显式转换。隐式转换在类型不匹配时由系统自动执行。显式转换则通过调用类型转换API来主动完成,API会将输入数据从源类型转换为目标类型,并返回转换结果。整个转换过程遵循规定的处理规则。类型转换的结果受舍入模式和饱和模式的影响。关于浮点数的表示方式和舍入规则,请参见类型转换的相关说明。4+- [数据类型转换](data_type_conversion/index.md):类型转换包括隐式转换和显式转换。隐式转换在类型不匹配时由系统自动执行。显式转换则通过调用类型转换API来主动完成,API会将输入数据从源类型转换为目标类型,并返回转换结果。整个转换过程遵循规定的处理规则。类型转换的结果受舍入模式和饱和模式的影响。关于浮点数的表示方式和舍入规则,请参见类型转换的相关说明。(共 50 张知识卡,分布在 1 个子目录)
5-- [float类型数学库函数](float_math_functions/index.md):获取两个输入数据相除的结果。5+- [float类型数学库函数](float_math_functions/index.md):共 94 张知识卡(如 __fdividef、__saturatef、abs 等)。
6-- [fp8类型](fp8_type/index.md):SIMT编程支持3种fp8数据类型,分别是float8_e4m3_t、float8_e5m2_t、hifloat8_t。其数据范围如下:6+- [fp8类型](fp8_type/index.md):共 15 张知识卡(如 fp8数据类型简介(FP8类型) 等)。
7-- [half类型](half_type/index.md):计算输入half2类型数据各分量的绝对值。7+- [half类型](half_type/index.md):共 235 张知识卡,分布在 7 个子目录(half2_arithmetic_functions、half2_compare_functions、half2_math_functions 等)。
8-- [整型数学库函数](integer_math_functions/index.md):将输入数据的位序反转,返回反转后的值。8+- [整型数学库函数](integer_math_functions/index.md):共 28 张知识卡(如 __brev(整数数学函数)、__byte_perm、__clz(整数数学函数) 等)。
@@ -1,4 +1,4 @@
1# SIMD与SIMT混合编程简介1# SIMD与SIMT混合编程简介
2 2 
3-- [扩展语法](extended_syntax/index.md):SIMD与SIMT混合编程提供了一系列适用于Device侧的数据类型,包括标量、短向量和dim3结构体。3+- [扩展语法](extended_syntax/index.md):共 6 张知识卡(如 内置数据类型、扩展语法内置宏、扩展语法内置变量 等)。
4- [ProgrammingModel](programming_model.md):SIMD与SIMT混合编程支持根据算子特征精细化映射,规则计算用SIMD获得高吞吐,不规则计算用SIMT缓解发散。4- [ProgrammingModel](programming_model.md):SIMD与SIMT混合编程支持根据算子特征精细化映射,规则计算用SIMD获得高吞吐,不规则计算用SIMT缓解发散。
@@ -1,4 +1,4 @@
1# SIMT编程简介1# SIMT编程简介
2 2 
3-- [扩展语法](extended_syntax/index.md):SIMT编程提供了一系列适用于Device侧的数据类型,包括标量、短向量以及dim3结构体。3+- [扩展语法](extended_syntax/index.md):共 6 张知识卡(如 内置数据类型、扩展语法内置宏、扩展语法内置变量 等)。
4- [编程模型](programming_model.md):SIMT编程模型采用Grid、Thread Block、Warp、Thread分层线程结构,适合控制流复杂和访存模式不规则的算子场景。4- [编程模型](programming_model.md):SIMT编程模型采用Grid、Thread Block、Warp、Thread分层线程结构,适合控制流复杂和访存模式不规则的算子场景。
@@ -1,5 +1,5 @@
1# 同步与内存栅栏1# 同步与内存栅栏
2 2 
3-- [内存栅栏接口](memory_fence/index.md):内存栅栏接口,保证调用前的全局和共享内存写操作对其他线程可见,不阻塞线程,仅约束可见性顺序。3+- [内存栅栏接口](memory_fence/index.md):共 2 张知识卡(如 AscThreadfence、AscThreadfenceBlock)。
4- [同步与内存栅栏简介(同步与内存屏障)](sync_and_memory_fence_intro.md):SIMT程序中,不同Thread的执行进度和内存访问到达顺序可能不同。当线程之间存在数据依赖时,需要使用同步接口或内存栅栏接口约束执行顺序和内存可见性。4- [同步与内存栅栏简介(同步与内存屏障)](sync_and_memory_fence_intro.md):SIMT程序中,不同Thread的执行进度和内存访问到达顺序可能不同。当线程之间存在数据依赖时,需要使用同步接口或内存栅栏接口约束执行顺序和内存可见性。
5- [同步接口](sync_interface/index.md):阻塞当前线程块内所有线程直到所有线程到达同步点,确保同步点前的内存操作对块内所有线程可见,常用于数据共享和并行归约。5- [同步接口](sync_interface/index.md):阻塞当前线程块内所有线程直到所有线程到达同步点,确保同步点前的内存操作对块内所有线程可见,常用于数据共享和并行归约。
@@ -1,7 +1,7 @@
1# Warp函数1# Warp函数
2 2 
3-- [Lane ID类函数](lane_id_functions/index.md):获取线程在其所在Warp内的索引。3+- [Lane ID类函数](lane_id_functions/index.md):共 6 张知识卡(如 Laneid、LanemaskEq、LanemaskGe 等)。
4- [Warp函数](warp_functions_intro.md):Warp是SIMT线程执行和调度的基本单位。一个线程块内的线程会按照线性线程号被划分为多个Warp,当前一个Warp包含32个线程。同一Warp内的线程执行相同的指令流,但每个线程拥有独立的寄存器和Lane ID,可处理不同的数据地址和分支路径。4- [Warp函数](warp_functions_intro.md):Warp是SIMT线程执行和调度的基本单位。一个线程块内的线程会按照线性线程号被划分为多个Warp,当前一个Warp包含32个线程。同一Warp内的线程执行相同的指令流,但每个线程拥有独立的寄存器和Lane ID,可处理不同的数据地址和分支路径。
5-- [Warp Reduce类函数](warp_reduce_functions/index.md):对Warp内所有活跃线程输入的val求和。Warp内所有活跃线程返回相同的结果。5+- [Warp Reduce类函数](warp_reduce_functions/index.md):共 3 张知识卡(如 asc_reduce_add、asc_reduce_max、asc_reduce_min)。
6-- [Warp Shfl类函数](warp_shfl_functions/index.md):Warp级数据交换接口,按分组宽度获取指定线程持有的变量值,支持数据分发与动态交换场景6+- [Warp Shfl类函数](warp_shfl_functions/index.md):共 4 张知识卡(如 AscShfl、AscShflDown、AscShflUp 等)。
7-- [Warp Vote类函数](warp_vote_functions/index.md):查看Warp内所有线程是否为活跃状态。 返回一个32bit的无符号整数,若Warp内某个线程是活跃(已结束线程是非活跃状态)的,则返回值中与线程Lane ID对应的bit位为1,否则为0。Warp内所有活跃线程返回相同的结果。7+- [Warp Vote类函数](warp_vote_functions/index.md):共 4 张知识卡(如 asc_activemask、asc_all、asc_any 等)。
@@ -1,7 +1,7 @@
1# C++标准库1# C++标准库
2 2 
3-- [算法](algorithms/index.md):比较相同数据类型的两个数,返回最大值。3+- [算法](algorithms/index.md):共 2 张知识卡(如 Utils APImax、Utils APImin)。
4-- [容器函数](container_functions/index.md):get的作用是从tuple容器中提取指定位置的元素。4+- [容器函数](container_functions/index.md):共 3 张知识卡(如 Utils APIget、make_tuple、Tuple)。
5- [通用工具](general_utils/index.md):index_sequence是Ascend C提供的一个类模板,用于生成一个编译时的整数序列,适用于模板元编程。 make_index_sequence是Ascend C提供的一个模板,通常使用make_index_sequence创建一个index_sequence类型的对象,用于生成一个从0到N-1的整数序列。5- [通用工具](general_utils/index.md):index_sequence是Ascend C提供的一个类模板,用于生成一个编译时的整数序列,适用于模板元编程。 make_index_sequence是Ascend C提供的一个模板,通常使用make_index_sequence创建一个index_sequence类型的对象,用于生成一个从0到N-1的整数序列。
6-- [数学函数](math_functions/index.md):获取输入数据的绝对值。6+- [数学函数](math_functions/index.md):共 4 张知识卡(如 abs(数学函数)、CeilDiv、CeilDivision 等)。
7-- [Utils API类型特性](type_traits/index.md):在程序编译时,为指定类型添加const限定符,可以用于在编译时进行类型转换。7+- [Utils API类型特性](type_traits/index.md):type_traits是Ascend C标准库中的一个头文件(),它提供了一系列用于在编译时进行类型检查、类型转换和类型操作的模板元编程工具。这些工具可以帮助开发者在程序编译时进行类型相关的决策,从而提高代码的安全性和性能。(共 25 张知识卡:add_const、add_cv、add_lvalue_reference 等)
@@ -1,13 +1,13 @@
1# Utils API1# Utils API
2 2 
3-- [C++标准库](cpp_stdlib/index.md):比较相同数据类型的两个数,返回最大值。3+- [C++标准库](cpp_stdlib/index.md):共 35 张知识卡,分布在 5 个子目录(algorithms、container_functions、general_utils 等)。
4- [log](log/index.md):提供Host侧打印Log功能,支持在TilingFunc中输出日志;Tiling下沉场景下必须使用本接口,日志需Dump落盘解析。4- [log](log/index.md):提供Host侧打印Log功能,支持在TilingFunc中输出日志;Tiling下沉场景下必须使用本接口,日志需Dump落盘解析。
5-- [平台信息获取](platform_info/index.md):Kernel直调场景下获取硬件平台信息的单例管理类,通过GetInstance获取PlatformAscendC指针以访问核数等平台信息。5+- [平台信息获取](platform_info/index.md):共 17 张知识卡(如 PlatformAscendCManager 等)。
6-- [原型注册与管理](prototype_register_management/index.md):注册算子的原型定义,从而确保算子能够被框架正确识别、编译和执行。 算子原型主要描述了算子的输入输出、属性等信息以及算子在AI处理器上相关实现信息,并关联tiling实现等函数。算子原型通过自定义的算子类来承载,该算子类继承自OpDef类。完成算子的原型定义等操作后,需要调用OP_ADD接口,传入算子类型(自定义算子类的类名),进行算子原型注册。详细内容请参考算子原型定义。6+- [原型注册与管理](prototype_register_management/index.md):共 79 张知识卡(如 原型注册接口(OP_ADD)、OpAICoreConfig注册接口(REGISTER_OP_AICORE_CONFIG) 等)。
7-- [RTC](rtc/index.md):注册指定的名称表达式,该表达式代表了一个__global__函数。 主要用于__global__函数为模板函数的场景,调用本接口后,可以通过aclrtcGetLoweredName接口获取函数的mangling name,在launch阶段调用对应的模板__global__函数。7+- [RTC](rtc/index.md):共 11 张知识卡(如 aclrtcAddNameExpr、aclrtcCompileProg、aclrtcCreateProg 等)。
8- [SuperKernel](super_kernel/index.md):本接口为算子SuperKernel场景提供绑定原核函数和SK子函数的能力。核函数直调算子的完整适配方法请参见核函数直调算子额外适配说明。8- [SuperKernel](super_kernel/index.md):本接口为算子SuperKernel场景提供绑定原核函数和SK子函数的能力。核函数直调算子的完整适配方法请参见核函数直调算子额外适配说明。
9-- [Tiling数据结构注册](tiling_data_structure_register/index.md):定义TilingData类并添加成员变量以保存Tiling参数,继承TilingDef基类后提供set/get字段、序列化保存和长度获取等接口。9+- [Tiling数据结构注册](tiling_data_structure_register/index.md):共 2 张知识卡(如 TilingData结构定义、TilingData结构注册)。
10- [Tiling下沉](tiling_sink/index.md):在Tiling下沉场景中,该宏定义用于生成Tiling下沉的注册类,再通过调用注册类的成员函数来注册需要下沉的Tiling函数。10- [Tiling下沉](tiling_sink/index.md):在Tiling下沉场景中,该宏定义用于生成Tiling下沉的注册类,再通过调用注册类的成员函数来注册需要下沉的Tiling函数。
11-- [Tiling模板编程](tiling_template_programming/index.md):Tiling模板编程时,开发者通过调用此接口自动生成并配置TilingKey。 使用该接口需要包含定义模板参数和模板参数组合的头文件。详细内容请参考Tiling模板编程。11+- [Tiling模板编程](tiling_template_programming/index.md):共 3 张知识卡(如 ASCENDC_TPL_SEL_PARAM、GET_TPL_TILING_KEY、模板参数定义)。
12-- [Tiling调测](tiling_tuning/index.md):设置算子的属性以及对应值12+- [Tiling调测](tiling_tuning/index.md):共 22 张知识卡,分布在 2 个子目录(context_builder、op_tiling_registry)。
13-- [调测接口](tuning_interface/index.md):在SIMT代码中调用此接口会中断算子的运行,适用于Kernel侧异常场景的调试。13+- [调测接口](tuning_interface/index.md):共 11 张知识卡(如 __trap、asc_dump、asc_mark_stamp 等)。
@@ -1,4 +1,4 @@
1# 平台信息获取1# 平台信息获取
2 2 
3-- [PlatformAscendC](platform_ascend_c/index.md):针对Cube与Vector分离模式,计算融合算子启动的numBlocks值,即AIV与AIC组合数量,结果通过SetSimdNumBlocks设置。3+- [PlatformAscendC](platform_ascend_c/index.md):共 16 张知识卡(如 CalcTschNumBlocks、构造及析构函数(PlatformAscendC)、GetCoreMemBw 等)。
4- [PlatformAscendCManager](platform_ascend_c_manager.md):Kernel直调场景下获取硬件平台信息的单例管理类,通过GetInstance获取PlatformAscendC指针以访问核数等平台信息。4- [PlatformAscendCManager](platform_ascend_c_manager.md):Kernel直调场景下获取硬件平台信息的单例管理类,通过GetInstance获取PlatformAscendC指针以访问核数等平台信息。
@@ -1,12 +1,12 @@
1# 原型注册与管理1# 原型注册与管理
2 2 
3- [原型注册接口(OP_ADD)](op_add.md):注册算子的原型定义,从而确保算子能够被框架正确识别、编译和执行。 算子原型主要描述了算子的输入输出、属性等信息以及算子在AI处理器上相关实现信息,并关联tiling实现等函数。算子原型通过自定义的算子类来承载,该算子类继承自OpDef类。完成算子的原型定义等操作后,需要调用OP_ADD接口,传入算子类型(自定义算子类的类名),进行算子原型注册。详细内容请参考算子原型定义。3- [原型注册接口(OP_ADD)](op_add.md):注册算子的原型定义,从而确保算子能够被框架正确识别、编译和执行。 算子原型主要描述了算子的输入输出、属性等信息以及算子在AI处理器上相关实现信息,并关联tiling实现等函数。算子原型通过自定义的算子类来承载,该算子类继承自OpDef类。完成算子的原型定义等操作后,需要调用OP_ADD接口,传入算子类型(自定义算子类的类名),进行算子原型注册。详细内容请参考算子原型定义。
4-- [OpAICoreConfig](op_ai_core_config/index.md):用于标识该算子实现是否支持入图时的静态Shape编译。4+- [OpAICoreConfig](op_ai_core_config/index.md):共 11 张知识卡(如 DynamicCompileStaticFlag、DynamicFormatFlag、DynamicRankSupportFlag 等)。
5-- [OpAICoreDef](op_ai_core_def/index.md):注册算子支持的AI处理器型号以及OpAICoreConfig信息。5+- [OpAICoreDef](op_ai_core_def/index.md):共 8 张知识卡(如 AddConfig(OpAICoreDef)、LaunchWithZeroEleOutputTensors、SetCheckSupport 等)。
6-- [OpAICPUDef](op_aicpu_def/index.md):此参数为预留参数,请保持固定值“100”。6+- [OpAICPUDef](op_aicpu_def/index.md):共 9 张知识卡(如 ComputeCost、Engine(原型注册管理)、ExtendCfgInfo 等)。
7-- [OpAttrDef](op_attr_def/index.md):设置算子属性的注释。用于在自动生成算子原型头文件时,同步生成算子原型注释。 基于OpDef算子原型定义,自定义算子工程可以实现如下自动化能力:自动生成图模式场景使用的算子原型头文件,开发者可以使用生成的算子原型进行构图、图编译、图执行等操作。 生成的注释有助于辅助理解算子原型,并可以基于这些注释自动生成算子原型的文档说明。通常情况下,内置CANN算子使用较多。开发者可以按需使用。7+- [OpAttrDef](op_attr_def/index.md):共 2 张知识卡(如 Comment(原型注册管理)、OpAttrDef)。
8-- [OpDef](op_def/index.md):设置AICore的配置信息:包括Tiling处理回调函数、能力检查回调函数、基础的配置信息等。8+- [OpDef](op_def/index.md):共 13 张知识卡(如 AICore(OpDef)、AICPU(OpDef)、Attr(OpDef) 等)。
9-- [OpHostCPUDef](op_host_cpu_def/index.md):此参数为预留参数,请保持固定值“100”。9+- [OpHostCPUDef](op_host_cpu_def/index.md):共 9 张知识卡(如 ComputeCost、Engine(原型注册管理)、ExtendCfgInfo 等)。
10-- [OpMC2Def](op_mc2_def/index.md):配置通信域名称,每个名称对应一个通信域。配置后在Kernel侧调用GetHcclContext接口可获取通信域对应的context(消息区)地址。10+- [OpMC2Def](op_mc2_def/index.md):共 5 张知识卡(如 HcclGroup(OpMC2Def)、HcclServerType、OpMC2Def构造函数 等)。
11-- [OpParamDef](op_param_def/index.md):配置该参数,当单算子API(aclnnxxx)接口中的输入(aclTensor类型)是非连续tensor时,框架会自动将其转换为连续tensor。 该接口仅在如下场景支持: - 通过单算子API执行的方式开发单算子调用应用。 - 间接调用单算子API\(aclnnxxx\)接口:Pytorch框架单算子直调的场景。11+- [OpParamDef](op_param_def/index.md):共 20 张知识卡(如 AutoContiguous、Comment、构造&拷贝构造函数(OpParamDef) 等)。
12- [OpAICoreConfig注册接口(REGISTER_OP_AICORE_CONFIG)](register_op_aicore_config.md):不同的硬件形态算子原型定义不同的情况,可以通过新增OpAICoreConfig的方式,针对不同的AI处理器型号注册差异化的算子原型。REGISTER_OP_AICORE_CONFIG宏在不改变原有注册的基础上,允许单独新增文件来注册算子在不同硬件形态上的差异化信息。 使用该注册宏需要包含以下头文件:12- [OpAICoreConfig注册接口(REGISTER_OP_AICORE_CONFIG)](register_op_aicore_config.md):不同的硬件形态算子原型定义不同的情况,可以通过新增OpAICoreConfig的方式,针对不同的AI处理器型号注册差异化的算子原型。REGISTER_OP_AICORE_CONFIG宏在不改变原有注册的基础上,允许单独新增文件来注册算子在不同硬件形态上的差异化信息。 使用该注册宏需要包含以下头文件:
@@ -1,4 +1,4 @@
1# Tiling调测1# Tiling调测
2 2 
3-- [ContextBuilder](context_builder/index.md):设置算子的属性以及对应值3+- [ContextBuilder](context_builder/index.md):共 18 张知识卡(如 AddAttr、AddInputTd、AddOutputTd 等)。
4-- [OpTilingRegistry](op_tiling_registry/index.md):创建与释放OpTilingRegistry对象。4+- [OpTilingRegistry](op_tiling_registry/index.md):共 4 张知识卡(如 构造和析构函数、GetTilingFunc、Utils API简介 等)。
@@ -1,4 +1,4 @@
1# 3510 arch migration1# 3510 arch migration
2 2 
3- [2201到3510架构变更](2201_to_3510_arch_changes.md):NPU架构版本3510的架构图如图1所示,总体来看,3510架构新增了以下特性:3- [2201到3510架构变更](2201_to_3510_arch_changes.md):NPU架构版本3510的架构图如图1所示,总体来看,3510架构新增了以下特性:
4-- [2201 to 3510 guide](2201_to_3510_guide/index.md):Ascend C高阶API基本兼容NPU架构版本3510与NPU架构版本2201,部分API进行了扩展。当前3510架构不支持卷积计算类高阶API。4+- [2201 to 3510 guide](2201_to_3510_guide/index.md):共 3 张知识卡(如 高阶API迁移指导、基础API迁移指导、算子编译迁移指导)。
@@ -1,6 +1,6 @@
1# cross gen migration guide1# cross gen migration guide
2 2 
3-- [3510 arch migration](3510_arch_migration/index.md):NPU架构版本3510的架构图如图1所示,总体来看,3510架构新增了以下特性:3+- [3510 arch migration](3510_arch_migration/index.md):共 4 张知识卡(如 2201到3510架构变更 等)。
4- [Ascend C API兼容性策略](ascend_c_api_compatibility.md):本兼容性说明仅适用于Ascend C算子开发的兼容性迁移指导。总体兼容性策略见表1,兼容性范围不包含编译器BuiltIn API、Ascend C内部实现接口等。文档中涉及的兼容性分为两类:一是功能兼容,包括数据类型兼容、接口原型兼容和常量兼容;二是性能兼容,指对于同等数据量,新架构上执行API耗时不高于旧架构。4- [Ascend C API兼容性策略](ascend_c_api_compatibility.md):本兼容性说明仅适用于Ascend C算子开发的兼容性迁移指导。总体兼容性策略见表1,兼容性范围不包含编译器BuiltIn API、Ascend C内部实现接口等。文档中涉及的兼容性分为两类:一是功能兼容,包括数据类型兼容、接口原型兼容和常量兼容;二是性能兼容,指对于同等数据量,新架构上执行API耗时不高于旧架构。
5- [instructions for new features](instructions_for_new_features/index.md):本文介绍NPU架构版本3510新增特性,为开发者将算子程序从其他架构迁移至3510架构提供参考。有关3510架构的整体变更,请参见2201到3510架构变更。5- [instructions for new features](instructions_for_new_features/index.md):本文介绍NPU架构版本3510新增特性,为开发者将算子程序从其他架构迁移至3510架构提供参考。有关3510架构的整体变更,请参见2201到3510架构变更。
6- [跨代迁移指南概述](overview.md):Ascend C的核心设计目标之一是实现跨代兼容。然而,每一代芯片在微架构上都存在独特的硬件特性,因此无法保证所有接口仅通过重编译就能达成完全的跨代兼容。基于这一前提,本手册旨在为Ascend C算子开发者提供系统化指导,帮助将算子Kernel代码从当前芯片平滑移植到目标代际芯片上运行,实现高效、可控的迁移。6- [跨代迁移指南概述](overview.md):Ascend C的核心设计目标之一是实现跨代兼容。然而,每一代芯片在微架构上都存在独特的硬件特性,因此无法保证所有接口仅通过重编译就能达成完全的跨代兼容。基于这一前提,本手册旨在为Ascend C算子开发者提供系统化指导,帮助将算子Kernel代码从当前芯片平滑移植到目标代际芯片上运行,实现高效、可控的迁移。
@@ -2,4 +2,4 @@
2 2 
3- [Ascend C概述](ascend_c_overview_and_learning_path.md):介绍CANN和Ascend C的定位与多层级API体系,提供从快速入门到高级编程的学习路径导航。3- [Ascend C概述](ascend_c_overview_and_learning_path.md):介绍CANN和Ascend C的定位与多层级API体系,提供从快速入门到高级编程的学习路径导航。
4- [环境准备](environment_setup.md):介绍Ascend C算子开发的环境准备,包括驱动固件、CANN软件包与CMake安装。4- [环境准备](environment_setup.md):介绍Ascend C算子开发的环境准备,包括驱动固件、CANN软件包与CMake安装。
5-- [quick start](quick_start/index.md):> 本文作为算子编程快速入门章节,将从昇腾异构系统基础入手,逐步拆解编程模型核心要素,帮助您快速建立昇腾算子编程的整体认知,为后续实操开发奠定基础。5+- [quick start](quick_start/index.md):共 5 张知识卡(如 异构系统与编程模型 等)。
@@ -1,5 +1,5 @@
1# quick start1# quick start
2 2 
3- [异构系统与编程模型](heterogeneous_system_and_programming_model.md):> 本文作为算子编程快速入门章节,将从昇腾异构系统基础入手,逐步拆解编程模型核心要素,帮助您快速建立昇腾算子编程的整体认知,为后续实操开发奠定基础。3- [异构系统与编程模型](heterogeneous_system_and_programming_model.md):> 本文作为算子编程快速入门章节,将从昇腾异构系统基础入手,逐步拆解编程模型核心要素,帮助您快速建立昇腾算子编程的整体认知,为后续实操开发奠定基础。
4-- [基于SIMD编程](simd_programming/index.md):本示例是一个入门实践,基于Ascend C SIMD实现Add算子,帮助您快速上手。它完整呈现了Device端核函数实现、Host端调用及编译运行的全流程,助您建立整体认知。开始前,请先参考环境准备安装所需的CANN软件包。4+- [基于SIMD编程](simd_programming/index.md):共 2 张知识卡(如 Add算子快速入门、Hello World)。
5-- [基于SIMT编程](simt_programming/index.md):本入门示例基于Ascend C SIMT实现一维Gather算子,带你快速上手实践,涵盖Device端核函数实现、Host端调用以及编译运行的完整流程,帮助开发者建立整体认知。5+- [基于SIMT编程](simt_programming/index.md):共 2 张知识卡(如 Gather算子快速入门、Hello World)。
@@ -1,50 +1,50 @@
1# Ascend C 开发指南1# Ascend C 开发指南
2 2 
3- [Ascend C多层级编程接口选择参考](asc_how_to_choose_api.md):欢迎使用 Ascend C 进行昇腾AI处理器算子开发。Ascend C不仅致力于开放芯片完备编程能力支撑实现极致性能,同时通过多层级编程API设计,让您能够根据项目需求、团队技能与性能目标,灵活选择最合适的API,在开发效率与运行性能之间取得最佳平衡。3- [Ascend C多层级编程接口选择参考](asc_how_to_choose_api.md):欢迎使用 Ascend C 进行昇腾AI处理器算子开发。Ascend C不仅致力于开放芯片完备编程能力支撑实现极致性能,同时通过多层级编程API设计,让您能够根据项目需求、团队技能与性能目标,灵活选择最合适的API,在开发效率与运行性能之间取得最佳平衡。
4-- [cross gen migration guide](cross_gen_migration_guide/index.md):arch22→arch35 迁移与反向生成的参考总纲:给出 L1~L5 分层迁移模型的适用条件与范围(L1 基础适配必做、L2 RegBase 重写、L3 SIMT、L4 需 IsRegbaseSocVersion 或 UB 不足时升级、L5 仅供参考),以及目录地图、被 brief 构建层消费的方式与来源可复现性说明。4+- [cross gen migration guide](cross_gen_migration_guide/index.md):Ascend C的核心设计目标之一是实现跨代兼容。然而,每一代芯片在微架构上都存在独特的硬件特性,因此无法保证所有接口仅通过重编译就能达成完全的跨代兼容。基于这一前提,本手册旨在为Ascend C算子开发者提供系统化指导,帮助将算子Kernel代码从当前芯片平滑移植到目标代际芯片上运行,实现高效、可控的迁移。(共 16 张知识卡:arch22 → arch35 AscendC Migration and Backward-Generation Reference、Ascend C API兼容性策略、L1 迁移实施详细指南 等)
5-- [getting started](getting_started/index.md):介绍CANN和Ascend C的定位与多层级API体系,提供从快速入门到高级编程的学习路径导航。5+- [getting started](getting_started/index.md):共 7 张知识卡(如 Ascend C概述、环境准备 等)。
6- [msprof Agent 使用指南](msprof_agent_guide.md):让 Agent 用 msprof 做性能分析而不撑爆 context 的指南:先用系统级 msprof 定位瓶颈,落到单算子再用算子级 msopprof(等同 msprof op)展开;核心是分层读取——优先只读 op_statistic 聚合 csv,再按需 grep op_summary 取关键指标,避免读二进制 trace。6- [msprof Agent 使用指南](msprof_agent_guide.md):让 Agent 用 msprof 做性能分析而不撑爆 context 的指南:先用系统级 msprof 定位瓶颈,落到单算子再用算子级 msopprof(等同 msprof op)展开;核心是分层读取——优先只读 op_statistic 聚合 csv,再按需 grep op_summary 取关键指标,避免读二进制 trace。
7-- [operator practice](operator_practice/index.md):本文档主要用于指导开发者基于Ascend C编程语言在昇腾AI处理器上开发高性能的算子。在阅读此文档时,需要开发者具备如下能力:7+- [operator practice](operator_practice/index.md):共 129 张知识卡(如 本文档组织结构(算子实践)、异构计算(算子实践)、Host Tiling 单元测试开发 等)。
8-- [精度验证](precision/index.md):aclnn 接口调用指南:精度验证必须经 torch_aclnn_helper.h 的 EXEC_NPU_CMD 宏构建 PyTorch C++ Extension,Python 侧用 torch.ops.npu.<op> 调用;不存在 torch.ops.aclnn 命名空间。8+- [精度验证](precision/index.md):共 44 张知识卡(如 aclnn 接口调用指南、激活参考实现的权重状态、边界与稳定公式、APACE 编码 golden 与容差边界 等)。
9-- [programming guide](programming_guide/index.md):本指南为AscendC算子开发综合技术手册,系统涵盖SIMD/SIMT编程模型、语言扩展特性、C++类库API、编译运行机制、调试调优及高级编程等核心技术。文档兼顾入门学习与深度实践,适配不同层级研发人员。入门开发者可依托本指南完成技术体系搭建,资深算子工程师亦可借此深化技术能力,充分挖掘昇腾AI处理器硬件算力,高质量完成自定义算子的开发与性能优化。9+- [programming guide](programming_guide/index.md):共 121 张知识卡(如 本文档组织结构(编程指南) 等)。
10- [快速开始](quick_start.md):根据本地是否有NPU设备和使用目标选择对应的环境准备方式:10- [快速开始](quick_start.md):根据本地是否有NPU设备和使用目标选择对应的环境准备方式:
11- [AscendC SIMD Development Reference](simd_development_reference.md):AscendC SIMD 开发参考,主张 arch35/A5 上多算子计算链默认走 Regbase MicroAPI 而非 Membase LocalTensor(后者每步 UB 往返加 PipeBarrier,实测拖累可达数十倍);另含量化 kernel 常用 SIMD API 与 MXFP4 范式。11- [AscendC SIMD Development Reference](simd_development_reference.md):AscendC SIMD 开发参考,主张 arch35/A5 上多算子计算链默认走 Regbase MicroAPI 而非 Membase LocalTensor(后者每步 UB 往返加 PipeBarrier,实测拖累可达数十倍);另含量化 kernel 常用 SIMD API 与 MXFP4 范式。
12-- [technical appendix](technical_appendix/index.md):AscendC::TPosition::A1代表设备上用于矩阵计算的逻辑内存,用于存放左矩阵,物理存储对应AI Core的L1 Buffer。12+- [technical appendix](technical_appendix/index.md):共 8 张知识卡,分布在 2 个子目录(concepts_and_terms、cpp_standard_support)。
13 13 
14-- [regbase](regbase/index.md):补全 API 汇总中的模板、有效 lane、分组归约和局部存储约束,区分候选目录与实现证据。14+- [regbase](regbase/index.md):共 34 张知识卡(如 RegBase API 的操作数、掩码与存储合同、AscArange:固定 reference 调用约束、asc_bfloat162e1m2x2 (废弃):固定 reference 调用约束 等)。
15 15 
16-- [blaze](blaze/index.md):将 Block/Copy 的合法 base shape 与 Scheduler 的当前逻辑范围分开,按实际消费者校验 tiling、workspace 和融合后处理。16+- [blaze](blaze/index.md):共 21 张知识卡(如 Blaze 硬件基形与运行期尾块范围的区别、权重反量化的 Cast 顺序与 V→C 释放关系、Blaze 逐元素 Epilogue 的广播映射与加载边界 等)。
17-- [catlass](catlass/index.md):保存长度、布局、mask 和 MLA 语义,限制固定尺寸与通路结论。17+- [catlass](catlass/index.md):共 17 张知识卡(如 CATLASS FA 接口映射及参考配方冲突、CATLASS 编译架构与启动资源合同、CATLASS 的 Tile、Block 与 Kernel 组件契约 等)。
18-- [diagnostics](diagnostics/index.md):固定 reference 的调用约束与版本边界。该接口实现AI CPU算子Kernel调试场景下的assert断言功能。 算子执行中,如果assert内部条件判断不为真,则输出assert条件、触发文件名、行号等信息。18+- [diagnostics](diagnostics/index.md):共 24 张知识卡(如 assert(AI CPU API):固定 reference 调用约束、printf(AI CPU API):固定 reference 调用约束、Printf:固定 reference 调用约束 等)。
19-- [memory](memory/index.md):修正以元素数冒充字节数的多处说明,保存不同类型的物理布局。19+- [memory](memory/index.md):共 25 张知识卡(如 对齐单位、布局行距与填充值、BF16 别名缓冲与平台采样路径、Ascend C 缓冲所有权、流水交接与槽位布局 等)。
20-- [moe](moe/index.md):将特征、来源三元组、收发计数和量化 scale 按同一有效条目顺序关联,区分最大容量与本轮实际 token 数。20+- [moe](moe/index.md):共 19 张知识卡(如 MoE Dispatch 与 Combine 的元数据对应关系、MoE 容量公式与路由假设、MoE Combine 的可选路径与量化角色 等)。
21-- [reductions](reductions/index.md):保留 TwoPass、Welford 与分组归约的适用条件。21+- [reductions](reductions/index.md):共 34 张知识卡(如 归约算法选择与数值边界、归约轴映射与全载、分载设计、Blaze Pure Reduce 的 partial 状态与模板缺口 等)。
22-- [vector](vector/index.md):补充向量标量分流、scope 搬运和重载身份,限制机械映射。22+- [vector](vector/index.md):共 18 张知识卡(如 Ascend C API 家族与存储位置决定的搬运、字节 Mask、Bit Mask 与数值谓词、Ascend C 向量广播的地址语义与标量替换边界 等)。
23 23 
24-- [layouts](layouts/index.md):纠正BA/AB方向及Select极性,保留Softmax布局。24+- [layouts](layouts/index.md):共 34 张知识卡(如 广播 Mask 极性、步长与暂存复用、重排路线启发式与连续搬运证据、Gather 尾轴模式与有符号索引规范化 等)。
25-- [shmem](shmem/index.md):Blaze Heavy Kernel 覆盖与性能条件的来源实现、适用条件与审视结论。25+- [shmem](shmem/index.md):共 31 张知识卡(如 Blaze Heavy Kernel 覆盖与性能条件、Blaze-SHMEM Host Bootstrap 生命周期与失败传播、Blaze 本地 MX 模板的 L1 尺度窗口与事件 等)。
26-- [synchronization](synchronization/index.md):用读写集合区分 RAW、WAR 和 WAW,说明额外缓冲与局部屏障的实际作用。26+- [synchronization](synchronization/index.md):共 22 张知识卡(如 向量递推中的真实依赖与别名冲突、APACE ReduceScatter 时序与通知计数、原子模式切换与首次覆盖顺序 等)。
27-- [torch extensions](torch_extensions/index.md):对齐 ACLNN 参数、GE OpDef、PyTorch schema 与 Meta/图转换,避免将不同接口层的名称和顺序直接互换。27+- [torch extensions](torch_extensions/index.md):共 9 张知识卡(如 PyTorch 自定义 ACLNN 算子的四层接口契约、ACLNN 扩展辅助层的符号解析与异步资源生命周期、PyTorch Ascend C 扩展的编译单元与启动 ABI 等)。
28 28 
29-- [communications](communications/index.md):APACE AllGather 预触发与 scale 等待边界的来源实例与适用边界。29+- [communications](communications/index.md):共 28 张知识卡(如 APACE AllGather 预触发与 scale 等待边界、APACE AllGather 三区与输出归属、APACE AllToAll Matmul 的全局与局部坐标 等)。
30-- [matmul](matmul/index.md):补齐 Iterate 与 IterateAll 的不同调用、workspace 单位及消费前完成边界。30+- [matmul](matmul/index.md):共 20 张知识卡(如 异步 Matmul 的结果空间与完成等待、Matmul 基本块流量模型与输出对齐、Batch/Group Matmul 分区身份与分支顺序 等)。
31-- [sorting](sorting/index.md):保存小维度稀疏更新中按索引分组的机制,明确原始位置、跨块分组和累计顺序。31+- [sorting](sorting/index.md):共 9 张知识卡(如 排序聚合同索引更新的排列与所有权、硬件排序的行边界、索引与 Padding、多核排序的分层归并与工作区交接 等)。
32 32 
33-- [simt](simt/index.md):保留维度和数值语义映射,纠正 shuffle 参数及 VF 统一禁令。33+- [simt](simt/index.md):共 10 张知识卡(如 SIMT API 迁移的签名、lane 与版本边界、SIMT 调试上下文与缓存可见性、SIMT 调试输出的资源与执行影响 等)。
34 34 
35-- [tiling](tiling/index.md):保存多路径选择,纠正D为1直通和未执行Tiling的路径。35+- [tiling](tiling/index.md):共 39 张知识卡(如 自适应算法的退化域与 Fallback、自适应容量分块与多行缓冲预算、自适应分块搜索、可行性与安全容量 等)。
36 36 
37-- [environments](environments/index.md):纠正设备修饰符和命名空间的误判,保存Tensor、队列及Mask操作边界。37+- [environments](environments/index.md):共 17 张知识卡(如 API 示例清单与真实编译合同、Ascend C API 家族与文档版本的兼容边界、Blaze 设备地址空间与 VF 循环约束 等)。
38 38 
39-- [registration](registration/index.md):核对 L2、L0、OpDef、图推导与 Kernel 的实际支持域,补足广播和空输入检查。39+- [registration](registration/index.md):共 23 张知识卡(如 Add 注册模板的跨层 shape 与 dtype 合同、算子文档模板中的接口事实与占位值边界、DAG 输出与 Foreach Tensor 区间合同 等)。
40 40 
41-- [attention](attention/index.md):纠正设计中的二次除法,保存mask、分核和内存尺度条件。41+- [attention](attention/index.md):共 58 张知识卡(如 Attention 在线输出归一化与最终除法必须一致、反向 Cube 融合需保留全部梯度路径、因果窗口、GQA 与整块跳过的证明 等)。
42-- [performance](performance/index.md):APACE Host 参数接线与实际计时的来源实例与适用边界。42+- [performance](performance/index.md):共 40 张知识卡(如 APACE Host 参数接线与实际计时、APACE 调优假设与负收益的适用范围、asys 系统采样的参数范围与 msprof 下发身份 等)。
43 43 
44-- [pooling](pooling/index.md):区分物理行距、有效尾部和特定快速路径的支持范围。44+- [pooling](pooling/index.md):共 8 张知识卡(如 池化地址对齐守卫与算术反例、池化反向的梯度贡献与 gather/scatter 归属、深度池化 Direct Buffer 路径与同步修订 等)。
45 45 
46-- [data movement](data_movement/index.md):保存多维搬运和MX物理行距,补齐参数状态与反向映射。46+- [data movement](data_movement/index.md):共 7 张知识卡(如 批量搬运、LoopMode 与交错 Scale 偏移、DataCopyPad 多行搬运的有效长度、padding 与物理行距、DataCopyPad 的版本、重载与尾部真实范围 等)。
47 47 
48-- [testing](testing/index.md):APACE 多rank启动与 context 复用范围的来源实例与适用边界。48+- [testing](testing/index.md):共 30 张知识卡(如 APACE 多rank启动与 context 复用范围、APACE 对标路径与失败证据边界、算子测试的行为断点与属性交叉覆盖 等)。
49 49 
50-- [compilation](compilation/index.md):APACE 构建根目录与版本化符号的来源实例与适用边界。50+- [compilation](compilation/index.md):共 2 张知识卡(如 APACE 构建根目录与版本化符号、Softmax 模板ABI与独立头文件依赖)。
@@ -2,6 +2,6 @@
2 2 
3- [FlashAttention算子性能调优案例](flash_attention_tuning.md):本案例中的算子FlashAttentionScoreGrad,用于训练场景下计算注意力的反向输出,即FlashAttentionScore算子的反向计算。3- [FlashAttention算子性能调优案例](flash_attention_tuning.md):本案例中的算子FlashAttentionScoreGrad,用于训练场景下计算注意力的反向输出,即FlashAttentionScore算子的反向计算。
4- [GroupedMatmul算子性能调优案例](grouped_matmul_tuning.md):本案例对分组Matmul即GroupedMatmul算子的per-token量化场景进行性能分析和优化,GroupedMatmul算子计算过程(通过python代码表达)为:4- [GroupedMatmul算子性能调优案例](grouped_matmul_tuning.md):本案例对分组Matmul即GroupedMatmul算子的per-token量化场景进行性能分析和优化,GroupedMatmul算子计算过程(通过python代码表达)为:
5-- [matmul tuning cases](matmul_tuning_cases/index.md):矩阵乘算子中输入矩阵内轴非256字节对齐时,在AIV核上完成ND到NZ格式转换以提升随路搬运效率。5+- [matmul tuning cases](matmul_tuning_cases/index.md):共 14 张知识卡(如 AIV核上的ND2NZ格式转换、Matmul高阶API开启IBShare模板共享A和B矩阵数据、Matmul高阶API开启IBShare模板共享B矩阵数据 等)。
6- [MC²算子性能调优案例](mc_operator_tuning.md):MC2通算融合算子性能调优案例,通过数据切分实现计算与通信流水掩盖,介绍切分策略制定与优化方法。6- [MC²算子性能调优案例](mc_operator_tuning.md):MC2通算融合算子性能调优案例,通过数据切分实现计算与通信流水掩盖,介绍切分策略制定与优化方法。
7- [SIMT实现InsertHashTable算子案例](simt_insert_hash_table.md):本案例介绍InsertHashTable算子,展示基于SIMT实现的大规模线程并发访问Global Memory内存的算子实践。7- [SIMT实现InsertHashTable算子案例](simt_insert_hash_table.md):本案例介绍InsertHashTable算子,展示基于SIMT实现的大规模线程并发访问Global Memory内存的算子实践。
@@ -1,15 +1,15 @@
1# operator practice1# operator practice
2 2 
3-- [优秀实践](best_practices/index.md):本案例中的算子FlashAttentionScoreGrad,用于训练场景下计算注意力的反向输出,即FlashAttentionScore算子的反向计算。3+- [优秀实践](best_practices/index.md):共 18 张知识卡(如 FlashAttention算子性能调优案例、GroupedMatmul算子性能调优案例、MC²算子性能调优案例 等)。
4- [本文档组织结构(算子实践)](document_structure.md):本文档主要用于指导开发者基于Ascend C编程语言在昇腾AI处理器上开发高性能的算子。在阅读此文档时,需要开发者具备如下能力:4- [本文档组织结构(算子实践)](document_structure.md):本文档主要用于指导开发者基于Ascend C编程语言在昇腾AI处理器上开发高性能的算子。在阅读此文档时,需要开发者具备如下能力:
5-- [功能调试](functional_debug/index.md):性能优化的基础是算子运行得到正确的计算结果。评判计算结果正确性需要有一定的评判标准,即使用已知的正确的输出和实际结果进行比较。优化过程中,每次迭代修改后,都需要验证性能优化的新结果是否满足精度评判标准。5+- [功能调试](functional_debug/index.md):共 3 张知识卡(如 精度正常(功能调试)、算子调试(功能调试)、运行正常(功能调试))。
6- [Host Tiling 单元测试开发](host_tiling_ut.md):面向 ACLNN 等注册式算子包的 Host Tiling 单元测试,涵盖独立 oracle、TilingContext 构造、平台资源 mock、注册调用、隔离构建,以及测试先行和已有实现的回归验证。6- [Host Tiling 单元测试开发](host_tiling_ut.md):面向 ACLNN 等注册式算子包的 Host Tiling 单元测试,涵盖独立 oracle、TilingContext 构造、平台资源 mock、注册调用、隔离构建,以及测试先行和已有实现的回归验证。
7- [异构计算(算子实践)](heterogeneous_computing.md):使用Ascend C进行编程时,会涉及到在两个不同的平台(Host、Device)上开发代码。本章简单介绍Host、Device之间的差异,便于开发者宏观地了解这个异构系统;同时给出算子相关的数据流,结合异构架构的特点,开发者可以进一步了解应该如何合理安排算子代码的执行位置,便于得到更好的性能。7- [异构计算(算子实践)](heterogeneous_computing.md):使用Ascend C进行编程时,会涉及到在两个不同的平台(Host、Device)上开发代码。本章简单介绍Host、Device之间的差异,便于开发者宏观地了解这个异构系统;同时给出算子相关的数据流,结合异构架构的特点,开发者可以进一步了解应该如何合理安排算子代码的执行位置,便于得到更好的性能。
8-- [性能分析](performance_analysis/index.md):理论性能为算子实际性能的理想目标。不同的硬件平台的硬件规格各异,理论性能可以帮助我们了解硬件的潜能,从而设定性能优化的目标。8+- [性能分析](performance_analysis/index.md):共 2 张知识卡(如 分析性能数据(性能分析)、获取性能数据(性能分析))。
9-- [SIMD算子实现](simd_operator_impl/index.md):Ascend C的算子实现主要包含两个部分:9+- [SIMD算子实现](simd_operator_impl/index.md):Ascend C的算子实现主要包含两个部分:(共 36 张知识卡,分布在 4 个子目录)
10-- [SIMD算子性能优化](simd_operator_optimization/index.md):汇总 SIMD 算子的 Tiling、启动开销、流水编排、内存访问、矢量计算和矩阵计算优化建议。10+- [SIMD算子性能优化](simd_operator_optimization/index.md):共 36 张知识卡(如 优化建议总览表(SIMD算子优化) 等)。
11-- [SIMD与SIMT混合算子实现](simd_simt_hybrid_operator_impl/index.md):本节内容为使用Reg矢量计算API和SIMT API进行SIMD与SIMT混合编程的指导。11+- [SIMD与SIMT混合算子实现](simd_simt_hybrid_operator_impl/index.md):共 2 张知识卡(如 基础知识(SIMD与SIMT混合算子实现)、算子实现(SIMD与SIMT混合算子实现))。
12-- [SIMD与SIMT混合算子性能优化](simd_simt_hybrid_optimization/index.md):AI Core上SIMD与SIMT混合编程结合了SIMD多数据并行计算能力与SIMT离散访存、复杂分支处理能力,可以在同一算子中以Vector Function(VF)为基本调度单位灵活组织SIMD和SIMT执行逻辑。SIMD编程适合处理连续、规整的向量或矩阵计算,SIMT编程适合处理分支判断、索引映射、离散访存等不规则逻辑,在程序中结合使用两种编程方式,可以兼顾硬件高吞吐和编程灵活性。12+- [SIMD与SIMT混合算子性能优化](simd_simt_hybrid_optimization/index.md):AI Core上SIMD与SIMT混合编程结合了SIMD多数据并行计算能力与SIMT离散访存、复杂分支处理能力,可以在同一算子中以Vector Function(VF)为基本调度单位灵活组织SIMD和SIMT执行逻辑。SIMD编程适合处理连续、规整的向量或矩阵计算,SIMT编程适合处理分支判断、索引映射、离散访存等不规则逻辑,在程序中结合使用两种编程方式,可以兼顾硬件高吞吐和编程灵活性。(共 5 张知识卡,分布在 4 个子目录)
13-- [SIMT算子实现](simt_operator_impl/index.md):本节内容为使用SIMT API进行SIMT编程的指导。13+- [SIMT算子实现](simt_operator_impl/index.md):共 2 张知识卡(如 基础知识(SIMT算子实现)、算子实现(SIMT算子实现))。
14-- [SIMT算子性能优化](simt_operator_optimization/index.md):SIMT编程模式以线程(Thread)为基本编程单位,由开发者编写单个线程的执行逻辑,再通过>>配置将大量线程组织为线程块(线程块)与线程网格(Grid),调度到硬件的Vector Core(物理核,即AIV核)上以Warp为单位并行执行。14+- [SIMT算子性能优化](simt_operator_optimization/index.md):SIMT编程模式以线程(Thread)为基本编程单位,由开发者编写单个线程的执行逻辑,再通过>>配置将大量线程组织为线程块(线程块)与线程网格(Grid),调度到硬件的Vector Core(物理核,即AIV核)上以Warp为单位并行执行。(共 10 张知识卡,分布在 4 个子目录)
15-- [典型算子实践案例](typical_operator_cases/index.md):本节在中级性能优化的基础上,介绍Add算子的高阶极限打磨方法。15+- [典型算子实践案例](typical_operator_cases/index.md):共 12 张知识卡,分布在 1 个子目录(simd_operator_impl)。
@@ -1,4 +1,4 @@
1# 融合算子编程1# 融合算子编程
2 2 
3-- [cv fusion](cv_fusion/index.md):>学习融合算子编程之前,请确保已经掌握矩阵编程相关知识。3+- [cv fusion](cv_fusion/index.md):共 2 张知识卡(如 融合算子编程基础知识、融合算子编程算子实现)。
4-- [general fusion](general_fusion/index.md):>本节内容为通算融合算子的理论背景和开发指导,学习本节内容之前,请确保已经掌握矩阵编程和《HCCL集合通信库》中的相关知识。4+- [general fusion](general_fusion/index.md):共 3 张知识卡(如 通用融合基础知识、特性场景、通用融合算子实现)。
@@ -1,7 +1,7 @@
1# SIMD算子实现1# SIMD算子实现
2 2 
3-- [融合算子编程](fusion_operator_programming/index.md):>学习融合算子编程之前,请确保已经掌握矩阵编程相关知识。3+- [融合算子编程](fusion_operator_programming/index.md):共 5 张知识卡,分布在 2 个子目录(cv_fusion、general_fusion)。
4-- [矩阵编程(高阶API)](matrix_advanced_api/index.md):>本节内容为使用高阶API进行矩阵乘法的编程指导。使用高阶API进行实际的矩阵编程时,需要通过API参考确认接口支持的产品型号。4+- [矩阵编程(高阶API)](matrix_advanced_api/index.md):共 17 张知识卡(如 SIMD算子实现基础知识、算子实现 等)。
5-- [matrix basic api](matrix_basic_api/index.md):>本节内容为针对耦合模式,使用基础API进行矩阵乘法的编程指导。5+- [matrix basic api](matrix_basic_api/index.md):共 2 张知识卡(如 SIMD算子实现耦合模式、SIMD算子实现分离模式)。
6- [SIMD算子实现概述](overview.md):Ascend C的算子实现主要包含两个部分:6- [SIMD算子实现概述](overview.md):Ascend C的算子实现主要包含两个部分:
7-- [矢量编程](vector_programming/index.md):基于Ascend C方式实现基础矢量算子核函数的流程如下图所示。7+- [矢量编程](vector_programming/index.md):本节将以Add算子为例,带您快速构建Ascend C矢量算子程序,并学习矢量算子开发的典型场景以及处理方式。涉及的场景包括:(共 11 张知识卡:基础矢量算子、Broadcast场景、DoubleBuffer场景 等)
@@ -1,5 +1,5 @@
1# 矩阵编程(高阶API)1# 矩阵编程(高阶API)
2 2 
3- [SIMD算子实现基础知识](basics.md):>本节内容为使用高阶API进行矩阵乘法的编程指导。使用高阶API进行实际的矩阵编程时,需要通过API参考确认接口支持的产品型号。3- [SIMD算子实现基础知识](basics.md):>本节内容为使用高阶API进行矩阵乘法的编程指导。使用高阶API进行实际的矩阵编程时,需要通过API参考确认接口支持的产品型号。
4-- [feature scenarios](feature_scenarios/index.md):对稀疏矩阵B进行4:2稠密化后调用Matmul API完成矩阵乘,跳过零元素以减少内存占用和计算量。4+- [feature scenarios](feature_scenarios/index.md):共 15 张知识卡(如 4:2稀疏矩阵乘、AIC和AIV独立运行机制、异步场景处理 等)。
5- [算子实现](operator_impl.md):介绍使用Matmul高阶API实现矩阵乘算子的流程,包括Host侧Tiling参数获取与Kernel侧矩阵乘操作。5- [算子实现](operator_impl.md):介绍使用Matmul高阶API实现矩阵乘算子的流程,包括Host侧Tiling参数获取与Kernel侧矩阵乘操作。
@@ -3,7 +3,7 @@
3- [基础矢量算子](basic_vector_operator.md):基于Ascend C方式实现基础矢量算子核函数的流程如下图所示。3- [基础矢量算子](basic_vector_operator.md):基于Ascend C方式实现基础矢量算子核函数的流程如下图所示。
4- [Broadcast场景](broadcast_scenario.md):在某些场景下,可能会存在两个输入shape不相同的情况。由于Add接口只支持对shape相同的输入进行计算,因此需要先对输入进行shape变换,再进行Add计算。本节将对满足Broadcast条件的输入在算子实现中的Broadcast处理进行介绍,其他场景可以参考本章节中提供的思路。4- [Broadcast场景](broadcast_scenario.md):在某些场景下,可能会存在两个输入shape不相同的情况。由于Add接口只支持对shape相同的输入进行计算,因此需要先对输入进行shape变换,再进行Add计算。本节将对满足Broadcast条件的输入在算子实现中的Broadcast处理进行介绍,其他场景可以参考本章节中提供的思路。
5- [DoubleBuffer场景](double_buffer_scenario.md):因存在算子中多次搬入搬出数据的场景,为充分利用硬件资源,实现多流水并行,引入DoubleBuffer机制。DoubleBuffer是通过将输入数据分成大小相等的两块,充分利用AI Core的硬件资源,实现数据搬入、计算、数据搬出的并行执行方式。下面以“核间不均分,核内不均分”的样例为例,介绍算子中DoubleBuffer的实现。5- [DoubleBuffer场景](double_buffer_scenario.md):因存在算子中多次搬入搬出数据的场景,为充分利用硬件资源,实现多流水并行,引入DoubleBuffer机制。DoubleBuffer是通过将输入数据分成大小相等的两块,充分利用AI Core的硬件资源,实现数据搬入、计算、数据搬出的并行执行方式。下面以“核间不均分,核内不均分”的样例为例,介绍算子中DoubleBuffer的实现。
6-- [multi core tiling](multi_core_tiling/index.md):长度为totalLength的输入数据可以按参与计算的核数和主块长度mainTileLength均匀切分,主块均分后无剩余数据。各核处理的数据量相同,每个核处理tileNum个主块,核内只包含主块。6+- [multi core tiling](multi_core_tiling/index.md):Ascend C核函数是在Device侧执行的并行函数。为了提高算子的执行效率,通常采用多核并行计算,将输入数据切分后分配到不同的逻辑核上处理;同时,由于单个核的Local Memory容量有限,无法一次完整容纳算子的输入和输出数据,需要分批搬运数据进行计算。这种对输入数据进行切分、分块计算的过程称为Tiling。(共 5 张知识卡:SIMD算子实现主块均分、SIMD算子实现尾块均分、TailCoreSplit 等)
7- [SIMD算子实现概述](overview.md):本节将以Add算子为例,带您快速构建Ascend C矢量算子程序,并学习矢量算子开发的典型场景以及处理方式。涉及的场景包括:7- [SIMD算子实现概述](overview.md):本节将以Add算子为例,带您快速构建Ascend C矢量算子程序,并学习矢量算子开发的典型场景以及处理方式。涉及的场景包括:
8- [TBuf的使用](tbuf_usage.md):在大多数算子开发时,核函数计算过程需要使用临时内存来存储运算的中间结果,这些中间结果以临时变量表示,临时变量占用的内存可以使用TBuf数据结构来管理,具体介绍请参考TBuf。下文将以输入的数据类型为bfloat16_t、在单核上运行的Add算子为例,介绍TBuf的使用方式。本文中介绍的算子完整代码请参见使用临时内存的Add算子样例。8- [TBuf的使用](tbuf_usage.md):在大多数算子开发时,核函数计算过程需要使用临时内存来存储运算的中间结果,这些中间结果以临时变量表示,临时变量占用的内存可以使用TBuf数据结构来管理,具体介绍请参考TBuf。下文将以输入的数据类型为bfloat16_t、在单核上运行的Add算子为例,介绍TBuf的使用方式。本文中介绍的算子完整代码请参见使用临时内存的Add算子样例。
9- [非对齐场景](unaligned_scenario.md):本节介绍非32字节对齐数据的更多处理方式,包括数据搬入、计算和搬出的处理。用户在实际算子开发中,可以参考如下方案介绍和算子样例灵活地处理非对齐场景。9- [非对齐场景](unaligned_scenario.md):本节介绍非32字节对齐数据的更多处理方式,包括数据搬入、计算和搬出的处理。用户在实际算子开发中,可以参考如下方案介绍和算子样例灵活地处理非对齐场景。
@@ -1,9 +1,9 @@
1# SIMD算子性能优化1# SIMD算子性能优化
2 2 
3-- [矩阵计算](matrix_compute/index.md):【描述】算子中进行带bias的矩阵乘计算时,可将bias数据搬运至Bias Table Buffer(C2)上,调用一次Mmad接口实现矩阵乘加bias的计算,或者直接调用Matmul高阶API完成功能。相比于先将矩阵乘的结果从L0C Buffer(CO1)搬运到GM上,再搬运到UB上进行加bias的过程,减少了数据搬运的次数,可提升内存使用效率。数据流图对比如下:3+- [矩阵计算](matrix_compute/index.md):共 5 张知识卡(如 通过BT Buffer实现高效的bias计算、通过FP Buffer存放量化参数实现高效随路量化、通过L0C Buffer数据暂存实现高效的矩阵乘结果累加 等)。
4-- [内存访问](memory_access/index.md):>该性能优化指导适用于如下产品型号:4+- [内存访问](memory_access/index.md):共 14 张知识卡(如 避免同地址访问、EfficientTransferApi、GM地址尽量512B对齐 等)。
5- [优化建议总览表(SIMD算子优化)](optimization_summary.md):汇总 SIMD 算子的 Tiling、启动开销、流水编排、内存访问、矢量计算和矩阵计算优化建议。5- [优化建议总览表(SIMD算子优化)](optimization_summary.md):汇总 SIMD 算子的 Tiling、启动开销、流水编排、内存访问、矢量计算和矩阵计算优化建议。
6-- [头尾开销优化](overhead_optimization/index.md):【编译器背景知识】创建类对象时,会分配内存空间,用于存储类中的相关成员变量或函数。当类中变量需要参与计算时,变量值从内存被加载到寄存器,计算完成后,变量从寄存器存储回内存。Scalar常量折叠和常量传播是编译器编译时的优化方式,优化前编译器会判断变量是否只初始化过一次或只赋值过一次,若满足此编译优化的前提条件,变量值将会尽量驻留在寄存器中,从而在后续使用变量时,将减少读取内存的操作,提升运行性能。6+- [头尾开销优化](overhead_optimization/index.md):共 5 张知识卡(如 避免TPipe在对象内创建和初始化、设置合适的核数和算子Kernel类型、设置DCI编译选项来减少算子尾开销 等)。
7-- [流水编排](pipeline_scheduling/index.md):【描述】执行于AI Core上的指令队列主要包括如下几类,Vector指令队列(V)、Cube指令队列(M)、Scalar指令队列(S)和搬运指令队列(MTE1/MTE2/MTE3)。不同指令队列间的相互独立性和可并行执行特性,是DoubleBuffer优化机制的基石。7+- [流水编排](pipeline_scheduling/index.md):共 2 张知识卡(如 开启DoubleBuffer、启用Iterate或IterateAll异步接口避免AIC/AIV同步依赖)。
8- [Tiling策略](tiling_strategy/index.md):通过调整拖尾核位置实现核间负载均衡,避免部分核计算拖尾时其他核空闲,优化多核整体性能。8- [Tiling策略](tiling_strategy/index.md):通过调整拖尾核位置实现核间负载均衡,避免部分核计算拖尾时其他核空闲,优化多核整体性能。
9-- [矢量计算](vector_compute/index.md):指令执行延迟(Instruction Execution Latency)指的是一条指令从开始执行到完全完成(即所有操作结束,结果可用)所消耗的时间,它直接影响程序的响应速度和实时性。在延迟敏感的场景中,降低指令执行延迟是提升性能的关键。下文以归约操作为例,介绍了几种归约方案的性能对比,便于开发者在使用归约指令时,能够根据具体的数据规模和场景,选择性能更高的方案。9+- [矢量计算](vector_compute/index.md):共 8 张知识卡(如 选择低延迟指令,优化归约操作性能、基于全局掩码复用的计算性能优化、通过Unified Buffer融合实现连续vector计算 等)。
@@ -1,7 +1,7 @@
1# 内存访问1# 内存访问
2 2 
3- [避免同地址访问](avoid_same_address_access.md):>该性能优化指导适用于如下产品型号:3- [避免同地址访问](avoid_same_address_access.md):>该性能优化指导适用于如下产品型号:
4-- [避免UB的bank冲突](avoid_ub_bank_conflict/index.md):>该性能优化建议适用于如下产品型号:4+- [避免UB的bank冲突](avoid_ub_bank_conflict/index.md):为了提高数据访问的效率和吞吐量,Unified Buffer采用了大小相等的内存模块(bank)结构设计。当多条读写指令同时访问Unified Buffer时,由于硬件资源的限制,这些指令不能同时执行,从而引发bank冲突。在这种情况下,指令需要排队等待资源,无法在一个指令周期内完成。(共 3 张知识卡:避免bank冲突(NPU架构版本2201)、避免bank冲突(NPU架构版本3510))
5- [EfficientTransferApi](efficient_transfer_api.md):使用搬运API时应通过配置搬运控制参数实现连续或固定间隔搬运,避免使用for循环,提升搬运效率。5- [EfficientTransferApi](efficient_transfer_api.md):使用搬运API时应通过配置搬运控制参数实现连续或固定间隔搬运,避免使用for循环,提升搬运效率。
6- [GM地址尽量512B对齐](gm_address_512b_alignment.md):【描述】由于AI处理器内部设计约束,从GM向Local Memory搬运数据时,保证GM地址512B对齐可以最有效的发挥出带宽的效率。如下图示例,展示了在512B对齐以及32B对齐情况下单核的带宽效率:搬运同等数据量,带宽差距最大的情况,32B对齐场景只能达到512B对齐场景的70%。6- [GM地址尽量512B对齐](gm_address_512b_alignment.md):【描述】由于AI处理器内部设计约束,从GM向Local Memory搬运数据时,保证GM地址512B对齐可以最有效的发挥出带宽的效率。如下图示例,展示了在512B对齐以及32B对齐情况下单核的带宽效率:搬运同等数据量,带宽差距最大的情况,32B对齐场景只能达到512B对齐场景的70%。
7- [设置合理的L2 CacheMode](l2_cache_mode_config.md):>该性能优化指导适用于如下产品型号:7- [设置合理的L2 CacheMode](l2_cache_mode_config.md):>该性能优化指导适用于如下产品型号:
@@ -4,4 +4,4 @@
4- [基于全局掩码复用的计算性能优化](mask_reuse_optimization.md):>该性能优化建议适用于如下产品型号:4- [基于全局掩码复用的计算性能优化](mask_reuse_optimization.md):>该性能优化建议适用于如下产品型号:
5- [通过Unified Buffer融合实现连续vector计算](ub_fusion_vector_compute.md):【描述】算子实现中涉及多次vector计算,且前一次计算输出是后一次计算输入的情况下,可将前一次计算输出暂存在UB(Unified Buffer)上直接作为下一次计算的输入,不需要将前一次的计算输出从UB搬运到GM后再从GM搬运到UB。这种UB Buffer融合的方式可以减少搬入搬出次数,实现连续vector计算,提升内存使用效率。数据流图对比如下:5- [通过Unified Buffer融合实现连续vector计算](ub_fusion_vector_compute.md):【描述】算子实现中涉及多次vector计算,且前一次计算输出是后一次计算输入的情况下,可将前一次计算输出暂存在UB(Unified Buffer)上直接作为下一次计算的输入,不需要将前一次的计算输出从UB搬运到GM后再从GM搬运到UB。这种UB Buffer融合的方式可以减少搬入搬出次数,实现连续vector计算,提升内存使用效率。数据流图对比如下:
6- [Vector算子灵活运用Counter模式](vector_counter_mode.md):【描述】Normal模式下,通过迭代次数repeatTimes和掩码mask,控制Vector算子中矢量计算API的计算数据量;当用户想要指定API计算的总元素个数时,首先需要自行判断是否存在不同的主块和尾块,主块需要将mask设置为全部元素参与计算,并且计算主块所需迭代次数,然后根据尾块中剩余元素个数重置mask,再进行尾块的运算,在此过程中涉及大量Scalar计算。6- [Vector算子灵活运用Counter模式](vector_counter_mode.md):【描述】Normal模式下,通过迭代次数repeatTimes和掩码mask,控制Vector算子中矢量计算API的计算数据量;当用户想要指定API计算的总元素个数时,首先需要自行判断是否存在不同的主块和尾块,主块需要将mask设置为全部元素参与计算,并且计算主块所需迭代次数,然后根据尾块中剩余元素个数重置mask,再进行尾块的运算,在此过程中涉及大量Scalar计算。
7-- [vf optimization](vf_optimization/index.md):连续非对齐搬运场景的性能优化建议与接口介绍。7+- [vf optimization](vf_optimization/index.md):共 4 张知识卡(如 连续非对齐场景优化、指令双发优化、VfFusionOptimization 等)。
@@ -1,7 +1,7 @@
1# SIMT算子性能优化1# SIMT算子性能优化
2 2 
3- [control flow](control_flow/index.md):同一Warp内线程应避免分支跳转以防止分支发散,减少掩码屏蔽等待,提升SIMT算子执行效率。3- [control flow](control_flow/index.md):同一Warp内线程应避免分支跳转以防止分支发散,减少掩码屏蔽等待,提升SIMT算子执行效率。
4-- [execution config](execution_config/index.md):SIMT算子线程块数量应结合物理核数与数据规模设置,大数据量匹配物理核数,小数据量需实测最优值。4+- [execution config](execution_config/index.md):共 2 张知识卡(如 线程块数量配置优化、合理配置线程数避免寄存器溢出)。
5- [instruction optimization](instruction_optimization/index.md):【描述】SIMT编程模式下,多个线程更新同一个内存地址时需要使用原子操作接口保证更新结果正确。原子操作属于高开销指令,其耗时由两方面决定:一是同一地址上的原子操作只能排队依次执行,二是单条原子指令的处理开销。5- [instruction optimization](instruction_optimization/index.md):【描述】SIMT编程模式下,多个线程更新同一个内存地址时需要使用原子操作接口保证更新结果正确。原子操作属于高开销指令,其耗时由两方面决定:一是同一地址上的原子操作只能排队依次执行,二是单条原子指令的处理开销。
6-- [memory access](memory_access/index.md):【描述】SIMT编程模式下,一个Warp内的多个线程可能在同一条UB访问指令中同时访问UB。对于使用UB中转做转置、重排等操作的算子,即使GM读写已经连续,仍需要分析UB访问模式,避免UB内部访问成为新的性能瓶颈。6+- [memory access](memory_access/index.md):共 5 张知识卡(如 避免UB的Bank冲突、DCacheAccess、MemoryAccessMerge 等)。
7- [SIMT算子优化概述](overview.md):SIMT编程模式以线程(Thread)为基本编程单位,由开发者编写单个线程的执行逻辑,再通过>>配置将大量线程组织为线程块(线程块)与线程网格(Grid),调度到硬件的Vector Core(物理核,即AIV核)上以Warp为单位并行执行。7- [SIMT算子优化概述](overview.md):SIMT编程模式以线程(Thread)为基本编程单位,由开发者编写单个线程的执行逻辑,再通过>>配置将大量线程组织为线程块(线程块)与线程网格(Grid),调度到硬件的Vector Core(物理核,即AIV核)上以Warp为单位并行执行。
@@ -1,3 +1,3 @@
1# 典型算子实践案例1# 典型算子实践案例
2 2 
3-- [simd operator impl](simd_operator_impl/index.md):本节在中级性能优化的基础上,介绍Add算子的高阶极限打磨方法。3+- [simd operator impl](simd_operator_impl/index.md):共 12 张知识卡,分布在 4 个子目录(add_vector_operator、gelu_reg_vector_operator、matmul_leakyrelu_fusion_operator 等)。
@@ -1,6 +1,6 @@
1# simd operator impl1# simd operator impl
2 2 
3-- [add vector operator](add_vector_operator/index.md):本节在中级性能优化的基础上,介绍Add算子的高阶极限打磨方法。3+- [add vector operator](add_vector_operator/index.md):共 3 张知识卡(如 高阶性能打磨、入门功能落地、中级性能优化)。
4-- [gelu reg vector operator](gelu_reg_vector_operator/index.md):本节在中级性能优化的基础上,介绍GELU算子的高阶极限打磨方法。4+- [gelu reg vector operator](gelu_reg_vector_operator/index.md):共 3 张知识卡(如 高阶极限打磨、入门功能落地、中级性能优化)。
5-- [matmul leakyrelu fusion operator](matmul_leakyrelu_fusion_operator/index.md):本节在中级性能优化的基础上,介绍Matmul融合算子的高阶极限打磨方法。5+- [matmul leakyrelu fusion operator](matmul_leakyrelu_fusion_operator/index.md):共 3 张知识卡(如 高阶极限打磨、入门功能落地、中级性能优化)。
6-- [matmul matrix operator](matmul_matrix_operator/index.md):本节基于C++ Tensor编程范式,在上文中级性能优化的基础上,进一步介绍Matmul基础API算子的高级极限打磨方法,通过5个高级优化步骤(大包搬运、LoadData3D替代LoadData2D、常量Tiling、UnitFlag优化、L2Cache优化),深入挖掘硬件性能极限。6+- [matmul matrix operator](matmul_matrix_operator/index.md):共 3 张知识卡(如 高级极限打磨、入门功能落地、中级性能优化)。
@@ -1,8 +1,8 @@
1# aclnn operator development1# aclnn operator development
2 2 
3- [Aclnn算子工程化开发快速入门](aclnn_quick_start.md):本文是基础内容,通过一个完整的AddCustom算子示例,帮助您在5分钟内跑通从创建工程、补充算子实现、编译部署到单算子API调用验证的端到端流程。3- [Aclnn算子工程化开发快速入门](aclnn_quick_start.md):本文是基础内容,通过一个完整的AddCustom算子示例,帮助您在5分钟内跑通从创建工程、补充算子实现、编译部署到单算子API调用验证的端到端流程。
4-- [附录](appendix/index.md):本文是参考内容,整理算子工程编译过程涉及的CMake函数接口,供需要深入了解编译机制的高级开发者查阅。4+- [附录](appendix/index.md):共 3 张知识卡(如 CMake函数参考(附录)、外部样例链接汇总(附录)、命名转换规则对照表(附录))。
5-- [compilation and deployment](compilation_and_deployment/index.md):本文是基础内容,介绍算子Kernel侧和Host侧实现开发完成后,如何对算子工程进行编译、生成自定义算子安装包\.run。详细的编译操作包括:5+- [compilation and deployment](compilation_and_deployment/index.md):共 7 张知识卡(如 编译与部署基本流程、高级编程编译加速、编译过程调试 等)。
6-- [design and implementation](design_and_implementation/index.md):本文介绍如何编写并注册GenSimplifiedKey回调。该回调用于预编译多份核函数(Kernel)binary的场景:算子开发者根据本次调用的输入、输出和属性生成一个自定义simplified key片段,运行时框架使用该片段查找对应的binary。6+- [design and implementation](design_and_implementation/index.md):共 9 张知识卡(如 GenSimplifiedKey回调实现、高级编程基本流程、高级编程基本流程 等)。
7-- [invocation](invocation/index.md):本文是扩展内容,介绍算子开发完成后,编译产物如何被加载和执行,以及开发时的设计决策如何影响运行表现。本文从算子开发者视角梳理运行时调用的基本模型和代码写法对运行行为的影响。7+- [invocation](invocation/index.md):共 3 张知识卡(如 运行时加载机制、基于simplified key的运行时选择流程、单算子API调用)。
8- [高级编程概述](overview.md):本文是基础内容,说明Aclnn算子工程化开发这组文档的范围、阅读顺序和前置条件。建议先确认环境满足要求,再按快速入门跑通AddCustom示例,跑通后回到对应章节逐步补齐细节。8- [高级编程概述](overview.md):本文是基础内容,说明Aclnn算子工程化开发这组文档的范围、阅读顺序和前置条件。建议先确认环境满足要求,再按快速入门跑通AddCustom示例,跑通后回到对应章节逐步补齐细节。
@@ -1,3 +1,3 @@
1# 高级AI-Core编程模型1# 高级AI-Core编程模型
2 2 
3-- [simd simt hybrid programming](simd_simt_hybrid_programming/index.md):Ascend C支持SIMD与SIMT混合编程,实现向量级并行与线程级并行的高效协同。AI Core包含AIC和AIV,其中AIV核支持SIMT、SIMD两种工作模式,因此可以实现AIV内SIMD与SIMT的混合编程,也支持AIV的SIMT与AIC的SIMD融合编程,整体硬件架构如下图所示:3+- [simd simt hybrid programming](simd_simt_hybrid_programming/index.md):为提升向量计算在复杂控制流、离散访存等场景下的编程灵活性,在SIMD与SIMT深度融合的混合编程架构中:矩阵计算单元延续SIMD设计,向量计算单元则在SIMD基础上引入SIMT能力。这一设计确立了以SIMD为主、SIMT为辅的混合编程模型。(共 5 张知识卡:抽象硬件架构(SIMD与SIMT混合编程)、核函数与VF函数(SIMD与SIMT混合编程)、内存层级(SIMD与SIMT混合编程) 等)
@@ -1,6 +1,6 @@
1# ai framework adaptation1# ai framework adaptation
2 2 
3-- [onnx framework](onnx_framework/index.md):>针对Atlas A3 训练系列产品/Atlas A3 推理系列产品,暂不支持ONNX框架算子调用。3+- [onnx framework](onnx_framework/index.md):共 2 张知识卡(如 适配插件开发、高级编程调用样例)。
4- [高级编程概述](overview.md):本章节内容介绍AI框架调用自定义算子的方法。如下图所示,Pytorch支持单算子和图模式两种,TensorFlow仅支持图模式。4- [高级编程概述](overview.md):本章节内容介绍AI框架调用自定义算子的方法。如下图所示,Pytorch支持单算子和图模式两种,TensorFlow仅支持图模式。
5- [PyTorch框架](pytorch_framework.md):通过PyTorch框架进行模型的训练、推理时,会调用很多算子进行计算。开发者开发的自定义算子如果需要集成部署到PyTorch框架,有如下几种方式:5- [PyTorch框架](pytorch_framework.md):通过PyTorch框架进行模型的训练、推理时,会调用很多算子进行计算。开发者开发的自定义算子如果需要集成部署到PyTorch框架,有如下几种方式:
6- [TensorFlow框架](tensorflow_framework.md):>针对Ascend 950PR/Ascend 950DT,暂不支持TensorFlow框架算子调用。6- [TensorFlow框架](tensorflow_framework.md):>针对Ascend 950PR/Ascend 950DT,暂不支持TensorFlow框架算子调用。
@@ -1,5 +1,5 @@
1# 硬件实现1# 硬件实现
2 2 
3-- [架构规格](architecture_spec/index.md):本节介绍__NPU_ARCH__版本号为2002的硬件架构和其功能说明。对应的产品型号为:3+- [架构规格](architecture_spec/index.md):共 4 张知识卡(如 NPU架构版本2002、NPU架构版本2201、NPU架构版本3002 等)。
4- [基本架构(硬件实现)](basic_architecture.md):如下图所示,基于Ascend C开发的算子运行在AI Core上。Ascend C编程模型基于AI Core硬件架构的抽象进行介绍,了解硬件架构能够帮助开发者更好的理解编程模型;对于需要完成高性能编程的深度开发者,更需要了解硬件架构相关知识,算子实践参考中很多内容都以本章为基础进行介绍。4- [基本架构(硬件实现)](basic_architecture.md):如下图所示,基于Ascend C开发的算子运行在AI Core上。Ascend C编程模型基于AI Core硬件架构的抽象进行介绍,了解硬件架构能够帮助开发者更好的理解编程模型;对于需要完成高性能编程的深度开发者,更需要了解硬件架构相关知识,算子实践参考中很多内容都以本章为基础进行介绍。
5-- [hardware constraints](hardware_constraints/index.md):本节介绍硬件约束以及解决方案建议。对应的产品型号为:Atlas 推理系列产品。5+- [hardware constraints](hardware_constraints/index.md):共 2 张知识卡(如 NPU架构版本2002、NPU架构版本2201)。
@@ -1,11 +1,11 @@
1# 高级编程1# 高级编程
2 2 
3-- [aclnn operator development](aclnn_operator_development/index.md):本文是基础内容,通过一个完整的AddCustom算子示例,帮助您在5分钟内跑通从创建工程、补充算子实现、编译部署到单算子API调用验证的端到端流程。3+- [aclnn operator development](aclnn_operator_development/index.md):本文是基础内容,说明Aclnn算子工程化开发这组文档的范围、阅读顺序和前置条件。建议先确认环境满足要求,再按快速入门跑通AddCustom示例,跑通后回到对应章节逐步补齐细节。(共 24 张知识卡:Aclnn算子工程化开发快速入门 等)
4-- [高级AI-Core编程模型](advanced_ai_core_programming_model/index.md):Ascend C支持SIMD与SIMT混合编程,实现向量级并行与线程级并行的高效协同。AI Core包含AIC和AIV,其中AIV核支持SIMT、SIMD两种工作模式,因此可以实现AIV内SIMD与SIMT的混合编程,也支持AIV的SIMT与AIC的SIMD融合编程,整体硬件架构如下图所示:4+- [高级AI-Core编程模型](advanced_ai_core_programming_model/index.md):共 5 张知识卡,分布在 1 个子目录(simd_simt_hybrid_programming)。
5-- [ai framework adaptation](ai_framework_adaptation/index.md):本章节内容介绍AI框架调用自定义算子的方法。如下图所示,Pytorch支持单算子和图模式两种,TensorFlow仅支持图模式。5+- [ai framework adaptation](ai_framework_adaptation/index.md):本章节内容介绍AI框架调用自定义算子的方法。如下图所示,Pytorch支持单算子和图模式两种,TensorFlow仅支持图模式。(共 5 张知识卡:PyTorch框架、TensorFlow框架 等)
6- [AOT编译优化(高级编程)](aot_compilation_optimization.md):AOT(Ahead-of-Time,提前编译)编译优化是一种将运行时参数提前到编译阶段常量化的性能优化技术。通过在编译阶段对Kernel进行额外的常量化编译,在运行时根据Tiling参数自动匹配预编译的特化版本,让编译器能够进行更深度的优化(如循环展开、死代码消除、常量传播等),从而获得超越通用版本的执行性能。6- [AOT编译优化(高级编程)](aot_compilation_optimization.md):AOT(Ahead-of-Time,提前编译)编译优化是一种将运行时参数提前到编译阶段常量化的性能优化技术。通过在编译阶段对Kernel进行额外的常量化编译,在运行时根据Tiling参数自动匹配预编译的特化版本,让编译器能够进行更深度的优化(如循环展开、死代码消除、常量传播等),从而获得超越通用版本的执行性能。
7-- [硬件实现](hardware_implementation/index.md):如下图所示,基于Ascend C开发的算子运行在AI Core上。Ascend C编程模型基于AI Core硬件架构的抽象进行介绍,了解硬件架构能够帮助开发者更好的理解编程模型;对于需要完成高性能编程的深度开发者,更需要了解硬件架构相关知识,算子实践参考中很多内容都以本章为基础进行介绍。7+- [硬件实现](hardware_implementation/index.md):共 7 张知识卡(如 基本架构(硬件实现) 等)。
8-- [内存模型](memory_model/index.md):本文介绍NPU编程中常见的缓存一致性问题,说明DCache、ICache、SIMT DCache与GM的数据一致性。8+- [内存模型](memory_model/index.md):共 2 张知识卡(如 缓存一致性、内存一致性)。
9-- [operator graph development](operator_graph_development/index.md):该开发流程以工程化算子开发为基础,除了需要提供工程化算子开发中的算子实现文件外,还需要额外交付算子入图的代码文件。本节仅提供算子入图代码文件的开发指导。9+- [operator graph development](operator_graph_development/index.md):图模式是神经网络模型的一种运行模式,在图模式下用户首先将模型的计算过程构造成一张图,然后通过GE将图下发到昇腾硬件执行。相对于单个算子依次下发的方式,图模式下,GE可以通过计算图优化、多流并行、内存复用、模型下沉等技术手段,加速模型执行效率,减少模型内存占用。(共 4 张知识卡:基本开发流程、开启Tiling下沉、图编译和图执行)
10- [SIMT协作组简介(高级编程)](simt_cooperative_groups.md):协作组(cooperative_groups)是Ascend C SIMT编程模型的扩展,用于组织协作线程组。协作组使开发者能够控制线程协作的粒度,实现更丰富、更高效的并行处理逻辑。10- [SIMT协作组简介(高级编程)](simt_cooperative_groups.md):协作组(cooperative_groups)是Ascend C SIMT编程模型的扩展,用于组织协作线程组。协作组使开发者能够控制线程协作的粒度,实现更丰富、更高效的并行处理逻辑。
11-- [super kernel](super_kernel/index.md):对于使用>>方式开发的Ascend C算子,除了遵循算子适配说明中的通用约束外,还需要在算子Kernel入口侧增加SuperKernel入口函数。本文档介绍具体的适配方法。11+- [super kernel](super_kernel/index.md):共 4 张知识卡(如 核函数直调算子额外适配说明、算子适配说明、算子自验证说明 等)。
@@ -1,7 +1,7 @@
1# appendix1# appendix
2 2 
3-- [common operations](common_operations/index.md):动态输入算子是指算子的输入个数是动态的,例如AddN,将N个输入tensor累加到一起,输出一个tensor,输入tensor的个数是不固定的。动态输入算子的开发在构造和解析输入数据方面有差异:核函数的入参采用ListTensorDesc的结构存储输入数据信息,对应的,调用时需构造TensorList结构保存参数信息。下面基于kernel直调和工程化算子开发两种开发方式分别介绍具体开发流程。3+- [common operations](common_operations/index.md):共 5 张知识卡(如 如何开发动态输入算子、EnableVectorCore、如何使用workspace 等)。
4-- [faq](faq/index.md):通过NPU进行核函数的运行验证时,出现挂死现象;通过CPU进行核函数的运行验证时,出现AllocTensor/FreeTensor失败的报错,日志报错和调用栈打印如下:4+- [faq](faq/index.md):共 8 张知识卡(如 运行验证时AllocTensor/FreeTensor失败、调用算子时出现无法打开config.ini的报错、算子工程编译时出现文件名过长报错 等)。
5- [基于样例工程完成Kernel直调](kernel_direct_call_from_sample.md):>本章节介绍的基于样例工程完成Kernel直调的方式,后续不再演进。推荐开发者直接使用命令行或者编写Cmake文件进行编译,详细内容请参考AI Core SIMD编译。5- [基于样例工程完成Kernel直调](kernel_direct_call_from_sample.md):>本章节介绍的基于样例工程完成Kernel直调的方式,后续不再演进。推荐开发者直接使用命令行或者编写Cmake文件进行编译,详细内容请参考AI Core SIMD编译。
6- [msobjdump工具](msobjdump_tool.md):本工具主要针对工程编译生成的算子ELF文件(Executable and Linkable Format)提供解析和解压功能,并将结果信息以可读形式呈现,方便开发者直观获得kernel文件信息。当前支持解析融合编译工程、标准/简易自定义算子工程,以及Kernel直调工程生成的相关产物。6- [msobjdump工具](msobjdump_tool.md):本工具主要针对工程编译生成的算子ELF文件(Executable and Linkable Format)提供解析和解压功能,并将结果信息以可读形式呈现,方便开发者直观获得kernel文件信息。当前支持解析融合编译工程、标准/简易自定义算子工程,以及Kernel直调工程生成的相关产物。
7- [optype_collector工具](optype_collector_tool.md):optype_collector工具用于采集内置/自定义算子包中指定AI处理器的OpType信息,并检测自定义算子包与内置算子包、自定义算子包之间的OpType重名问题,帮助开发者在自定义算子安装或交付前提前发现命名冲突。7- [optype_collector工具](optype_collector_tool.md):optype_collector工具用于采集内置/自定义算子包中指定AI处理器的OpType信息,并检测自定义算子包与内置算子包、自定义算子包之间的OpType重名问题,帮助开发者在自定义算子安装或交付前提前发现命名冲突。
@@ -1,4 +1,4 @@
1# 编译与运行1# 编译与运行
2 2 
3- [异步执行(编译与执行)](async_execution.md):当开发者完成一个Kernel核函数的编写后,需要通过特定的下发/启动语法将其提交到指定的Stream(计算流)上执行。例如:3- [异步执行(编译与执行)](async_execution.md):当开发者完成一个Kernel核函数的编写后,需要通过特定的下发/启动语法将其提交到指定的Stream(计算流)上执行。例如:
4-- [operator compilation](operator_compilation/index.md):当开发者完成一个Kernel核函数的编写,并在Host侧通过\>\>方式完成核函数的启动调用后,就需要进行算子源码编译与算子运行,接下来我们介绍通过bisheng编译器完成算子编译的基本用法。4+- [operator compilation](operator_compilation/index.md):共 5 张知识卡(如 AI Core编译基本用法、AI CPU算子编译、毕昇编译器简介 等)。
@@ -1,5 +1,5 @@
1# 调试调优1# 调试调优
2 2 
3-- [功能调试](functional_debug/index.md):本节介绍CPU域调试的方法:CPU侧验证核函数,gdb调试、使用printf命令打印。当前SIMT编程场景不支持。3+- [功能调试](functional_debug/index.md):共 2 张知识卡(如 CPU域孪生调试(功能调试)、NPU域上板调试(功能调试))。
4- [调试与调优概述](overview.md):介绍Ascend C算子调试与调优整体方案,涵盖CPU域孪生调试、NPU域上板调试及性能调优工具。4- [调试与调优概述](overview.md):介绍Ascend C算子调试与调优整体方案,涵盖CPU域孪生调试、NPU域上板调试及性能调优工具。
5- [性能调优(调试与调优)](performance_tuning.md):通过毕昇编译器编译生成可执行程序后,使用msOpProf工具采集Ascend C算子在AI处理器上执行的性能数据,进行性能精细调优。5- [性能调优(调试与调优)](performance_tuning.md):通过毕昇编译器编译生成可执行程序后,使用msOpProf工具采集Ascend C算子在AI处理器上执行的性能数据,进行性能精细调优。
@@ -1,10 +1,10 @@
1# programming guide1# programming guide
2 2 
3-- [高级编程](advanced_programming/index.md):AOT(Ahead-of-Time,提前编译)编译优化是一种将运行时参数提前到编译阶段常量化的性能优化技术。通过在编译阶段对Kernel进行额外的常量化编译,在运行时根据Tiling参数自动匹配预编译的特化版本,让编译器能够进行更深度的优化(如循环展开、死代码消除、常量传播等),从而获得超越通用版本的执行性能。3+- [高级编程](advanced_programming/index.md):共 53 张知识卡(如 AOT编译优化(高级编程)、SIMT协作组简介(高级编程) 等)。
4-- [appendix](appendix/index.md):>本章节介绍的基于样例工程完成Kernel直调的方式,后续不再演进。推荐开发者直接使用命令行或者编写Cmake文件进行编译,详细内容请参考AI Core SIMD编译。4+- [appendix](appendix/index.md):共 17 张知识卡(如 基于样例工程完成Kernel直调、msobjdump工具、optype_collector工具 等)。
5-- [编译与运行](compilation_and_execution/index.md):当开发者完成一个Kernel核函数的编写后,需要通过特定的下发/启动语法将其提交到指定的Stream(计算流)上执行。例如:5+- [编译与运行](compilation_and_execution/index.md):共 6 张知识卡(如 异步执行(编译与执行) 等)。
6-- [调试调优](debug_and_tuning/index.md):介绍Ascend C算子调试与调优整体方案,涵盖CPU域孪生调试、NPU域上板调试及性能调优工具。6+- [调试调优](debug_and_tuning/index.md):介绍Ascend C算子调试与调优整体方案,涵盖CPU域孪生调试、NPU域上板调试及性能调优工具。(共 4 张知识卡:性能调优(调试与调优) 等)
7- [本文档组织结构(编程指南)](document_structure.md):本指南为AscendC算子开发综合技术手册,系统涵盖SIMD/SIMT编程模型、语言扩展特性、C++类库API、编译运行机制、调试调优及高级编程等核心技术。文档兼顾入门学习与深度实践,适配不同层级研发人员。入门开发者可依托本指南完成技术体系搭建,资深算子工程师亦可借此深化技术能力,充分挖掘昇腾AI处理器硬件算力,高质量完成自定义算子的开发与性能优化。7- [本文档组织结构(编程指南)](document_structure.md):本指南为AscendC算子开发综合技术手册,系统涵盖SIMD/SIMT编程模型、语言扩展特性、C++类库API、编译运行机制、调试调优及高级编程等核心技术。文档兼顾入门学习与深度实践,适配不同层级研发人员。入门开发者可依托本指南完成技术体系搭建,资深算子工程师亦可借此深化技术能力,充分挖掘昇腾AI处理器硬件算力,高质量完成自定义算子的开发与性能优化。
8-- [语言扩展层](language_extension/index.md):如其他语言一样,会提供一些内置的宏方便用户编写程序。预定义宏一节着重介绍一些用户在做异构编程时会经常用到的宏,以及宏的解释。8+- [语言扩展层](language_extension/index.md):共 5 张知识卡(如 SIMD BuiltIn关键字、SIMD语言扩展层C API、SIMD&SIMT BuiltIn关键字 等)。
9-- [类库API](library_api/index.md):>本章提供编程API的概述。具体API参考见Ascend C API。9+- [类库API](library_api/index.md):共 10 张知识卡(如 编程接口概述(库API)、Utils API(库API) 等)。
10-- [编程模型](programming_model/index.md):AI CPU是位于Device侧ARM64架构的处理器,其具备与AI Core相同的内存访问能力,可直接访问Device侧内存资源;也可以与Host侧的CPU一样,进行类似的数据计算,通常作为AI Core的补充,主要承担非矩阵类、逻辑比较复杂的分支密集型计算。AI CPU的运行环境为基础的Linux环境,编程时可使用libc库,C++标准库,STL模板库等。其硬件架构图如下所示:10+- [编程模型](programming_model/index.md):共 25 张知识卡(如 AI CPU编程(编程模型)、异构系统(编程模型)、编程模型概述 等)。
@@ -1,4 +1,4 @@
1# advanced api1# advanced api
2 2 
3- [库API概述](overview.md):高阶API基于单核对常见算法进行抽象和封装,实现了一些常用的计算算法,旨在提高编程开发效率。高阶API一般通过调用多种基础API实现。高阶API包括数学计算、矩阵计算、激活函数等API。3- [库API概述](overview.md):高阶API基于单核对常见算法进行抽象和封装,实现了一些常用的计算算法,旨在提高编程开发效率。高阶API一般通过调用多种基础API实现。高阶API包括数学计算、矩阵计算、激活函数等API。
4-- [quick reference](quick_reference/index.md):Kernel侧接口的内部实现一般涉及复杂的数学计算,需要额外的临时空间来存储计算过程中的中间变量。除矩阵计算、HCCL通信类、卷积计算等,对于多数高阶API中临时空间的处理,开发者可以通过Kernel侧接口的入参sharedTmpBuffer传入提前申请的临时空间、通过接口框架申请临时空间两种方式。4+- [quick reference](quick_reference/index.md):共 2 张知识卡(如 如何使用Kernel侧临时空间、如何使用Tiling依赖的头文件)。
@@ -1,5 +1,5 @@
1# basic api1# basic api
2 2 
3-- [interface classification](interface_classification/index.md):连续计算API,支持Tensor前n个数据计算。针对源操作数的连续n个数据进行计算并连续写入目的操作数,解决一维tensor的连续计算问题。3+- [interface classification](interface_classification/index.md):共 2 张知识卡(如 连续计算API、高维切分API)。
4- [基础API概述](overview.md):基础API实现对硬件能力的抽象,开放芯片的能力,保证完备性和兼容性。标注为ISASI(Instruction Set Architecture Special Interface,硬件体系结构相关的接口)类别的API,不能保证跨硬件版本兼容。4- [基础API概述](overview.md):基础API实现对硬件能力的抽象,开放芯片的能力,保证完备性和兼容性。标注为ISASI(Instruction Set Architecture Special Interface,硬件体系结构相关的接口)类别的API,不能保证跨硬件版本兼容。
5-- [quick reference](quick_reference/index.md):Mask用于控制矢量计算中参与计算的元素个数,支持以下工作模式及配置方式:5+- [quick reference](quick_reference/index.md):共 2 张知识卡(如 如何使用掩码操作API、如何使用归约计算API)。
@@ -1,6 +1,6 @@
1# 类库API1# 类库API
2 2 
3-- [advanced api](advanced_api/index.md):高阶API基于单核对常见算法进行抽象和封装,实现了一些常用的计算算法,旨在提高编程开发效率。高阶API一般通过调用多种基础API实现。高阶API包括数学计算、矩阵计算、激活函数等API。3+- [advanced api](advanced_api/index.md):高阶API基于单核对常见算法进行抽象和封装,实现了一些常用的计算算法,旨在提高编程开发效率。高阶API一般通过调用多种基础API实现。高阶API包括数学计算、矩阵计算、激活函数等API。(共 3 张知识卡,分布在 1 个子目录)
4-- [basic api](basic_api/index.md):基础API实现对硬件能力的抽象,开放芯片的能力,保证完备性和兼容性。标注为ISASI(Instruction Set Architecture Special Interface,硬件体系结构相关的接口)类别的API,不能保证跨硬件版本兼容。4+- [basic api](basic_api/index.md):基础API实现对硬件能力的抽象,开放芯片的能力,保证完备性和兼容性。标注为ISASI(Instruction Set Architecture Special Interface,硬件体系结构相关的接口)类别的API,不能保证跨硬件版本兼容。(共 5 张知识卡,分布在 2 个子目录)
5- [编程接口概述(库API)](programming_interface_overview.md):>本章提供编程API的概述。具体API参考见Ascend C API。5- [编程接口概述(库API)](programming_interface_overview.md):>本章提供编程API的概述。具体API参考见Ascend C API。
6- [Utils API(库API)](utils_api.md):Ascend C开发提供了丰富的通用工具类,涵盖标准库、平台信息获取、上下文构建、运行时编译及日志输出等功能,支持开发者高效实现算子开发与性能优化。6- [Utils API(库API)](utils_api.md):Ascend C开发提供了丰富的通用工具类,涵盖标准库、平台信息获取、上下文构建、运行时编译及日志输出等功能,支持开发者高效实现算子开发与性能优化。
@@ -1,8 +1,8 @@
1# AI-Core-SIMD编程1# AI-Core-SIMD编程
2 2 
3- [抽象硬件架构(AI Core SIMD编程)](abstract_hardware_architecture.md):AI Core是AI处理器的计算核心,AI处理器内部有多个AI Core。本章节将介绍AI Core的并行计算架构抽象,该抽象架构屏蔽了不同硬件之间的差异。使用Ascend C进行编程时,基于抽象硬件架构,可以简化硬件细节,显著降低开发门槛。如需了解更详细的硬件架构信息或者原理,请参考硬件实现。3- [抽象硬件架构(AI Core SIMD编程)](abstract_hardware_architecture.md):AI Core是AI处理器的计算核心,AI处理器内部有多个AI Core。本章节将介绍AI Core的并行计算架构抽象,该抽象架构屏蔽了不同硬件之间的差异。使用Ascend C进行编程时,基于抽象硬件架构,可以简化硬件细节,显著降低开发门槛。如需了解更详细的硬件架构信息或者原理,请参考硬件实现。
4-- [c pointer programming](c_pointer_programming/index.md):C语言编程是Ascend C SIMD编程的底层核心接口,直接映射NPU硬件指令,支持精细化控制向量、矩阵和存储等硬件资源。4+- [c pointer programming](c_pointer_programming/index.md):共 4 张知识卡(如 C语言编程概述、Cube矩阵计算编程、Memory矢量计算编程 等)。
5-- [cpp tensor programming](cpp_tensor_programming/index.md):Ascend C SIMD编程路径的基础编程接口,基于Tensor对象抽象封装内存地址,提供类型安全和编译期检查的高层次编程能力。5+- [cpp tensor programming](cpp_tensor_programming/index.md):共 5 张知识卡(如 CppTensorProgrammingOverview、Cube矩阵计算编程、MemoryVectorComputation 等)。
6- [核函数(AI Core SIMD编程)](kernel_function.md):核函数(Kernel Function)是在设备侧并行执行的C/C++函数。它由主机端代码调用,调用时被实例化为多份并行副本,分别调度到多个AI Core上运行(每个AI Core被抽象为一个Block)。这些Block利用AI Core的计算能力,协同完成某一算子的计算。6- [核函数(AI Core SIMD编程)](kernel_function.md):核函数(Kernel Function)是在设备侧并行执行的C/C++函数。它由主机端代码调用,调用时被实例化为多份并行副本,分别调度到多个AI Core上运行(每个AI Core被抽象为一个Block)。这些Block利用AI Core的计算能力,协同完成某一算子的计算。
7- [AI Core SIMD编程模型概述](overview.md):SIMD并行机制是AI Core的核心算力支撑,承担了整机90%以上的计算吞吐量。因此,本章将以SIMD编程模型为核心主线,遵循「宏观架构—硬件调度—代码落地」的递进逻辑,完整覆盖SIMD算子开发全链路:阐述多AI Core集群的并行任务分发机制,拆解单AI Core内部异构计算单元的任务调度逻辑,并逐层讲解如何基于差异化的C/C++编程接口,完成算子Kernel计算逻辑的开发。7- [AI Core SIMD编程模型概述](overview.md):SIMD并行机制是AI Core的核心算力支撑,承担了整机90%以上的计算吞吐量。因此,本章将以SIMD编程模型为核心主线,遵循「宏观架构—硬件调度—代码落地」的递进逻辑,完整覆盖SIMD算子开发全链路:阐述多AI Core集群的并行任务分发机制,拆解单AI Core内部异构计算单元的任务调度逻辑,并逐层讲解如何基于差异化的C/C++编程接口,完成算子Kernel计算逻辑的开发。
8-- [tpipe tque programming](tpipe_tque_programming/index.md):上一章阐述了TPipe‑TQue编程模型的底层设计原理:基于经典C/C++队列管道(Queue Pipeline)思想,将算子任务拆分为多级流水Stage,依靠TQue队列解决写后读数据依赖、依靠AllocTensor/FreeTensor解决读后写内存复用冲突,通过底层同步指令实现多单元异步并行的时序安全。8+- [tpipe tque programming](tpipe_tque_programming/index.md):共 2 张知识卡(如 TPipe和TQue编程范式、TPipe-TQue编程原理介绍)。
@@ -1,7 +1,7 @@
1# 编程模型1# 编程模型
2 2 
3-- [AI-Core-SIMD编程](ai_core_simd_programming/index.md):AI Core是AI处理器的计算核心,AI处理器内部有多个AI Core。本章节将介绍AI Core的并行计算架构抽象,该抽象架构屏蔽了不同硬件之间的差异。使用Ascend C进行编程时,基于抽象硬件架构,可以简化硬件细节,显著降低开发门槛。如需了解更详细的硬件架构信息或者原理,请参考硬件实现。3+- [AI-Core-SIMD编程](ai_core_simd_programming/index.md):SIMD并行机制是AI Core的核心算力支撑,承担了整机90%以上的计算吞吐量。因此,本章将以SIMD编程模型为核心主线,遵循「宏观架构—硬件调度—代码落地」的递进逻辑,完整覆盖SIMD算子开发全链路:阐述多AI Core集群的并行任务分发机制,拆解单AI Core内部异构计算单元的任务调度逻辑,并逐层讲解如何基于差异化的C/C++编程接口,完成算子Kernel计算逻辑的开发。(共 14 张知识卡:抽象硬件架构(AI Core SIMD编程)、核函数(AI Core SIMD编程) 等)
4-- [AI-Core-SIMT编程](ai_core_simt_programming/index.md):AI Core上的SIMT(Single Instruction Multiple Thread,单指令多线程)编程允许指令对数据进行独立寻址,从而支持线程级并行计算。这种方式特别适用于离散访问和复杂控制逻辑等场景。SIMT编程能够有效简化复杂算子和不规则控制流的开发,缓解包含分支的发散计算并控制程序复杂度,从而提高硬件利用率和能效。4+- [AI-Core-SIMT编程](ai_core_simt_programming/index.md):SIMT编程模型适用于离散数据访问和复杂控制流场景,基于Grid-Block-Thread三级线程层级在AIV核上实现线程级并行计算。(共 8 张知识卡:抽象硬件架构(AI Core SIMT编程)、原子操作(AI Core SIMT编程)、核函数(AI Core SIMT编程) 等)
5- [AI CPU编程(编程模型)](ai_cpu_programming.md):AI CPU是位于Device侧ARM64架构的处理器,其具备与AI Core相同的内存访问能力,可直接访问Device侧内存资源;也可以与Host侧的CPU一样,进行类似的数据计算,通常作为AI Core的补充,主要承担非矩阵类、逻辑比较复杂的分支密集型计算。AI CPU的运行环境为基础的Linux环境,编程时可使用libc库,C++标准库,STL模板库等。其硬件架构图如下所示:5- [AI CPU编程(编程模型)](ai_cpu_programming.md):AI CPU是位于Device侧ARM64架构的处理器,其具备与AI Core相同的内存访问能力,可直接访问Device侧内存资源;也可以与Host侧的CPU一样,进行类似的数据计算,通常作为AI Core的补充,主要承担非矩阵类、逻辑比较复杂的分支密集型计算。AI CPU的运行环境为基础的Linux环境,编程时可使用libc库,C++标准库,STL模板库等。其硬件架构图如下所示:
6- [异构系统(编程模型)](heterogeneous_system.md):一个基于昇腾处理器的异构系统通常包含CPU与昇腾NPU。其中,CPU及其内存称为Host与Host Memory;NPU及其内存称为Device与Device Memory。6- [异构系统(编程模型)](heterogeneous_system.md):一个基于昇腾处理器的异构系统通常包含CPU与昇腾NPU。其中,CPU及其内存称为Host与Host Memory;NPU及其内存称为Device与Device Memory。
7- [编程模型概述](programming_model_overview.md):介绍Ascend C编程模型,涵盖SIMD与SIMT并行执行模型、AI Core硬件基础及编程路径选择。7- [编程模型概述](programming_model_overview.md):介绍Ascend C编程模型,涵盖SIMD与SIMT并行执行模型、AI Core硬件基础及编程路径选择。
@@ -2,5 +2,5 @@
2 2 
3- [技术附录术语表](glossary.md):AscendC::TPosition::A1代表设备上用于矩阵计算的逻辑内存,用于存放左矩阵,物理存储对应AI Core的L1 Buffer。3- [技术附录术语表](glossary.md):AscendC::TPosition::A1代表设备上用于矩阵计算的逻辑内存,用于存放左矩阵,物理存储对应AI Core的L1 Buffer。
4- [memory access](memory_access/index.md):Scalar计算单元仅支持读写Global Memory和Unified Buffer,介绍其读写方式及同步机制。4- [memory access](memory_access/index.md):Scalar计算单元仅支持读写Global Memory和Unified Buffer,介绍其读写方式及同步机制。
5-- [neural networks and operators](neural_networks_and_operators/index.md):数据排布格式(Data Layout Format)是深度学习中对多维Tensor在内存中存储方式的描述。5+- [neural networks and operators](neural_networks_and_operators/index.md):共 2 张知识卡(如 数据排布格式、算子基本概念)。
6- [performance optimization](performance_optimization/index.md):执行于AI Core上的指令队列主要包括如下几类,即Vector指令队列、Cube指令队列和MTE指令队列。不同指令队列间的相互独立性和可并行执行性,是DoubleBuffer优化机制的基石。6- [performance optimization](performance_optimization/index.md):执行于AI Core上的指令队列主要包括如下几类,即Vector指令队列、Cube指令队列和MTE指令队列。不同指令队列间的相互独立性和可并行执行性,是DoubleBuffer优化机制的基石。
@@ -1,4 +1,4 @@
1# cpp standard support1# cpp standard support
2 2 
3- [技术附录概述](overview.md):Host侧与clang15一致,支持完整的C/C++标准。3- [技术附录概述](overview.md):Host侧与clang15一致,支持完整的C/C++标准。
4-- [语法限制](syntax_restrictions/index.md):C++标准支持特性说明,列举不支持的C++特性。4+- [语法限制](syntax_restrictions/index.md):共 2 张知识卡(如 特性、技术附录函数)。
@@ -1,4 +1,4 @@
1# technical appendix1# technical appendix
2 2 
3-- [concepts and terms](concepts_and_terms/index.md):AscendC::TPosition::A1代表设备上用于矩阵计算的逻辑内存,用于存放左矩阵,物理存储对应AI Core的L1 Buffer。3+- [concepts and terms](concepts_and_terms/index.md):共 5 张知识卡(如 技术附录术语表 等)。
4-- [cpp standard support](cpp_standard_support/index.md):Host侧与clang15一致,支持完整的C/C++标准。4+- [cpp standard support](cpp_standard_support/index.md):Host侧与clang15一致,支持完整的C/C++标准。(共 3 张知识卡,分布在 1 个子目录)
@@ -1,5 +1,54 @@
1# 2026-09-221# 2026-09-22
2 2 
3+## [21:13] governance | 最终采用方案 A:GitCode 强制合入前解决 CI 问题
4+ 
5+### Summary
6+ 
7+确认 GitCode 代码仓强制要求合入前解决 CI 问题(分支保护硬性阻断),方案 B 的 override 路径不可行,最终采用方案 A 两步合入。!58 恢复为第 1 步内容(仅生成器代码 + `IndexSummaryTests` + 日志,knowledge 与 master 零差异,CI 预期绿、无需 override);第 2 步(85 个 `knowledge/**/index.md` 再生,内容源 ccfd49a7)待第 1 步合入后基于新 master 提交,届时 master 已含新生成器,CI 同样预期绿。
8+ 
9+### Details
10+ 
11+- 本分支 tree 与 141a18e7(原第 1 步提交)在 knowledge/governance/evals 上完全一致(diff 为空),仅日志比其多 [19:54] 与本条决策记录。
12+- 验证:knowledge 与 master diff 为空;`IndexSummaryTests` 7/7;该 tree 的 lint 此前已实测(checked=7513,index.stale=85 预期瞬态、无其他发现,第 2 步合入后消除)。
13+- 窗口约束(重申):第 1 步合入至第 2 步合入之间,任何 PR 触发 FormatCheck 都会报 85 个 index.stale;!55 在第 1 步之前或第 2 步之后合入/触发 CI。
14+ 
15+## [19:54] knowledge | 恢复完整索引提交,改用方案 B 单 PR 合入
16+ 
17+### Summary
18+ 
19+用户决定采用方案 B:!58 恢复为完整单 PR(新生成器 + 85 个 `knowledge/**/index.md` 再生 + `IndexSummaryTests`),放弃两步拆分。CI FormatCheck 将以 master 旧生成器校验新格式索引,预期 85 个 `index.stale`(已 worktree 复现验证的既知集合),合入时一次性 override;合入后 master 直接从旧格式自洽迁移到新格式自洽,无方案 A 的窗口期协调成本。
20+ 
21+### Details
22+ 
23+- 在 141a18e7(第 1 步)基础上 `git restore --source=ccfd49a7 -- knowledge` 恢复 85 个索引;tree 与 ccfd49a7 在 knowledge/governance/evals 上完全一致,仅日志多两条记录([19:20] 重排、本条)。
24+- 验证:lint checked=7513 P0=P1=P2=0(新生成器 + 新索引自洽);IndexSummaryTests 7/7。
25+- 与 !55 无内容依赖,任意顺序合入;`logs/2026-09-22.md` 同日文件在第二个合入的 PR 上需人工合并条目。
26+- [19:20] 条目中的方案 A 窗口约束随本决定作废。
27+ 
28+## [19:20] governance | 索引摘要改进重排为两步合入(方案 A 第 1 步:生成器代码先行)
29+ 
30+### Summary
31+ 
32+采用两步合入方案 A:!58 重排为"生成器代码先行"(本分支,仅 `governance/contracts/indexes.py` 概述卡优先摘要 + `IndexSummaryTests` 7 测 + 日志,不含任何 knowledge 变更)与"索引再生随后"(第 2 步,85 个 `knowledge/**/index.md`,待本步合入后基于新 master 提交)。CI trusted snapshot 以 master 旧生成器校验本 PR 未变更的 knowledge,FormatCheck 预期绿;本分支自身 lint 将显示 85 个 index.stale(新生成器 + 旧格式索引的预期瞬态,第 2 步合入后消除)。
33+ 
34+### Details
35+ 
36+- 合入窗口约束:本步合入至第 2 步合入之间 master 处于自不一致状态(新生成器 + 旧格式索引),期间任何 PR 触发 FormatCheck 都会报 85 个 index.stale;!55 应在本步之前或第 2 步之后合入/触发 CI。
37+- 第 2 步索引内容源:ccfd49a7(原 !58 完整提交,本地保留),届时基于新 master `git restore --source=ccfd49a7 -- knowledge` 再生成并提交 MR。
38+- 验证:IndexSummaryTests 7/7;本分支 lint checked=7513,index.stale 恰为 85 个预期集合,无其他 P0/P1/P2;knowledge 与 master 零差异(diff 为空)。
39+ 
40+## [18:37] governance+knowledge | 索引导航摘要改进以独立 PR 提交(拆分自 !55)
41+ 
42+### Summary
43+ 
44+按维护者建议与 CI trusted snapshot 约束(FormatCheck 以 master 检查代码校验 PR knowledge,索引格式生成器与索引内容无法在同一 PR 演进),将 !55 拆出的索引改进以独立 PR 提交:`governance/contracts/indexes.py` 概述卡优先摘要(`{目录名}.md` > `overview.md`)+ 递归全量计数 + 85 个 `knowledge/**/index.md` 再生 + `IndexSummaryTests` 7 测。基于 origin/master(da6d4e6c)独立成枝,不依赖 !55 其他内容;`evals/test_governance.py` 基于主线版本仅追加 IndexSummaryTests(探针修复属 !55,不在本 PR 重复)。
45+ 
46+### Details
47+ 
48+- 预期 CI:FormatCheck 将以 master 旧生成器校验本 PR 新格式索引,出现 85 个 `index_sync/index.stale`(该配置已在 !55 流程中用 worktree 精确复现验证:master 生成器 + 新格式 knowledge = 85 P0,纯 master = 0);合入需两步路径(先合生成器代码、再合索引再生)或维护者一次性 override,方案对比见 MR 描述。
49+- 验证:本分支 lint checked=7513 P0=P1=P2=0(新生成器 + 新索引自洽);IndexSummaryTests 7/7;检索引擎不消费 index.md(所有索引不参与检索),检索行为与 master 一致。
50+- 与 !55 的关系:!55 平台状态仍为 open、master 仍为 da6d4e6c(用户已在别处手动合入 !55 内容并 CI 通过);本 PR 基于 master 独立构建,!55 后续合入不产生内容依赖,仅 logs/2026-09-22.md 同日文件可能需人工合并。
51+ 
3## [17:28] fix | 撤出索引语义摘要改进以适配 CI trusted snapshot 校验(MR !55 拆分)52## [17:28] fix | 撤出索引语义摘要改进以适配 CI trusted snapshot 校验(MR !55 拆分)
4 53 
5### Summary54### Summary
@@ -0,0 +1,55 @@
1+# 2026-09-24
2+ 
3+## [11:38] knowledge+governance | 按维护者意见改回原子单 PR:生成器 + 85 个索引再生成
4+ 
5+### Summary
6+ 
7+zhanghz1 在 !58 评审(09-24 11:08)明确反对两步合入:第 1 步单独合入会让
8+master 出现 85 个 stale index,`bash check.sh` 立即失败并阻断窗口期内所有
9+PR,且第 2 步不能修复第 1 步提交本身不可自检的问题;建议生成器、测试和
10+85 个再生索引**原子放在同一 PR**。按此意见重构 !58:squash 提交 0b3026dd
11+之上再生成 85 个 stale 索引并 amend 为单一原子提交。
12+ 
13+### Details
14+ 
15+- 规模核实:`index_updates` 对当前树(master ea29f07d + 深层计数修复)计算
16+ 出 **恰好 85 个 stale 索引**,全部位于 `ops/ascendc/{apis,guides}`;
17+ contrib 子树(57a8ccdd 新增 196 文件)零涉及(其索引非弱摘要格式,
18+ `refreshed_entry` 保留原样)。
19+- 再生成内容包含本日深层计数修复(一级子目录语义):如 `apis/simd_api`
20+ 的"高阶API…分布在 15 个子目录"(旧算法会计成 90);37 个目录计数文本
21+ 与 ccfd49a7 时代产物不同,故未复用旧产物、全部按当前生成器重新生成。
22+- 复验:再生成后 `index_updates` = 0 stale(662 目录);
23+ `IndexSummaryTests` 9/9;85 文件 +438/-438(均为条目摘要行重写,
24+ 索引不参与检索,检索行为不变)。
25+- 已知 CI 行为(09-22 worktree 复现):FormatCheck 将检查代码回退到
26+ master、仅采信 PR 的 knowledge/,故本原子 PR 的 /compile 预期在
27+ FormatCheck 报 85 个 index.stale——需维护者一次性 override 合入,
28+ 合入后 master 新生成器 + 新索引自洽、check.sh 恢复绿色。该说明将随
29+ 回复提交给 zhanghz1 确认。
30+ 
31+## [09:47] update | 索引导航摘要子目录计数修复为仅统计一级子目录
32+ 
33+### Summary
34+ 
35+修复 `governance/contracts/indexes.py` `summarize_subdir` 的深层嵌套计数缺陷:
36+原实现把每个后代文件的直接父目录当作"子目录"统计(`{path.parent ...}`),
37+两层及更深嵌套时会系统性多计(如 `ops/` 报"分布在 537 个子目录")。
38+改为按 `relative_to(directory).parts[0]` 统计一级子目录,与文案语义一致。
39+新增两个两层嵌套回归测试。
40+ 
41+### Details
42+ 
43+- 评审反馈指出的缺陷,实证复现:`hub/{a/y.md, a/b/x.md, c/z.md}` 三个文件、
44+ 两个一级子目录,原实现输出"分布在 3 个子目录(a、b、c)"(把二级目录 b
45+ 当作独立子目录);修复后输出"分布在 2 个子目录(a、c)"。
46+- 修复采用最小 diff:`child_dirs` 在分支前一次计算,summary 分支与无直接卡
47+ 分支共用;单卡目录、overview 优先、样例列表等既有行为不变。
48+- 全库影响面:37 个目录的"分布在 N 个子目录"计数文本改变
49+ (ops: 537→3、ops/ascendc: 514→7、ops/ascendc/apis: 344→5 等)。
50+ **第 2 步 85 个 index.md 的 MR(源 ccfd49a7)合入前必须用修复后生成器
51+ 重新生成**,否则 CI trusted snapshot 校验不匹配。
52+- 验证:`evals.test_governance.IndexSummaryTests` 9/9 通过
53+ (7 既有 + 2 新增:`test_deep_nesting_counts_first_level_subdirectories_only`、
54+ `test_deep_nesting_with_overview_counts_first_level_only`)。
55+ 
@@ -1,5 +1,39 @@
1# 2026-09-281# 2026-09-28
2 2 
3+## [17:00] update | !58 重构为索引专用 PR 并 rebase 至 0fe7f05d
4+ 
5+### Summary
6+ 
7+按维护者两步合入安排,索引摘要生成器与测试已经 PR #67(0fe7f05d)先行合入 master;!58 不再重复提交代码,仅保留 85 个再生成 index.md 与本工作流日志,rebase 到 0fe7f05d。
8+ 
9+### Details
10+ 
11+- `governance/contracts/indexes.py`、`evals/test_governance.py` 与 PR #67 合入版本逐字节一致(diff 为空),本 PR 不再重复提交;
12+- 85 个索引内容与原子版 098cb4a3 完全一致;master 侧 310P 迁移(7c53ff7b)未触及 apis/guides 子树;
13+- 评测机在 0fe7f05d 上复核 `sync-indexes --check`:662 索引 85 stale;再生成 Synchronized 85、unchanged 577,改动路径集与本 PR 85 个文件完全一致,apis/guides 内容与 098cb4a3 逐字节一致;
14+- `logs/2026-09-28.md` 按时间降序合并 PR #67 侧 [03:39]/[07:24] 条目、本分支 [09:26] 条目与本条目;`logs/2026-09-22.md` 沿用此前两侧合并版(master 侧在 1e015cdc 后无变化)。
15+ 
16+## [09:26] update | 按维护者意见清理日志并 rebase 最新 master(!58 终审准备)
17+ 
18+### Summary
19+ 
20+zhanghz1 09-25 在 !58 评审指出:日志写入内网 IP 等敏感信息应删除;!55 的
21+代码核验记录与本 PR 索引摘要变更无关应删除;并确认 trusted snapshot 机制下
22+原子 PR 的 FormatCheck 报 85 个 index.stale 属预期冲突——"最终失败严格只有
23+预期的 85 个 index.stale、其余门禁均通过,则核对后按一次性迁移处理"。
24+按此清理日志并 rebase 到最新 master。
25+ 
26+### Details
27+ 
28+- 删除 `logs/2026-09-24.md` 中 [09:30] !55 核验记录整段(与本 PR 无关);
29+- 删除 [11:38] 中经评测机转发推送的内部环境细节(含内网 IP);
30+- 顺带删除 [09:47] 中已过时的"两步合入第 1 步/RFC"顺序说明(已被原子
31+ 单 PR 结构取代);
32+- rebase 到 1e015cdc(含 !55 检索改动合入):`logs/2026-09-22.md` 与
33+ !55 合入带来的同名日志冲突,按时间降序合并两侧全部 9 条条目
34+ (21:13/19:54/19:20/18:37 + 17:28/15:04/14:13 + 12:00 + 09:24);
35+- 清理后重新 /compile,预期失败严格限于 85 个已知 index.stale。
36+ 
3## [07:24] governance | 先合入索引摘要生成器与回归测试37## [07:24] governance | 先合入索引摘要生成器与回归测试
4 38 
5### Summary39### Summary