ascendc-samples昇腾 NPU 的“算子示例代码库”之前帮朋友看 Ascend C 算子开发的代码发现他不知道从哪下手——官方文档太理论看不懂GitHub 上的示例太零散跑不通。我告诉他去看 ascendc-samples。 这个仓库是昇腾 NPU 的算子开发示例代码库把常用的算子Add/MatMul/Conv 等都写了示例代码而且注释详细直接抄就行。类比一下学做菜官方文档“炒菜理论”火候/时间/调料比例ascendc-samples“炒菜视频”一步一步跟着做做完就能吃技术要点分析要点1ascendc-samples 的示例覆盖范围ascendc-samples 覆盖了四大类算子示例1. 基础算子示例Basic Operators向量算子Add, Sub, Mul, Div, Sqrt, Exp, Log矩阵算子MatMul, MatVec, Outer统计算子Mean, Std, Var, Sort, TopK性能数据跟 CPU 对比Ascend 910单精度算子CPU 延迟 (ms)NPU 延迟 (ms)加速比Add2.50.357.1xMatMul45.25.87.8xSort125.018.56.8x2. NN 算子示例Neural Network Operators卷积算子Conv2D, Conv3D, TransposedConv2D激活函数ReLU, GelU, SiLU, SoftMax池化算子MaxPool2D, AvgPool2D, AdaptiveAvgPool2D性能数据跟 CPU 对比Ascend 910单精度算子CPU 延迟 (ms)NPU 延迟 (ms)加速比Conv2D125.08.514.7xMaxPool2D12.51.210.4xReLU5.20.86.5x3. Transformer 算子示例Transformer OperatorsAttention 算子FlashAttention, MHA, MQA, GQAFFN 算子FeedForward, SwiGLU归一化算子LayerNorm, RMSNorm性能数据跟 CPU 对比Ascend 910单精度算子CPU 延迟 (ms)NPU 延迟 (ms)加速比FlashAttention38.56.85.7xFeedForward18.52.57.4xLayerNorm8.21.17.5x4. 自定义算子示例Custom OperatorsMoE 算子ExpertParallel, TopK, GateNetwork量化算子Quantize, Dequantize, FakeQuantize稀疏算子SparseMatMul, SparseSoftMax性能数据跟 CPU 对比Ascend 910单精度算子CPU 延迟 (ms)NPU 延迟 (ms)加速比MoE (TopK2)85.012.56.8xQuantize (INT8)5.50.86.9xSparseMatMul125.018.56.8x要点2ascendc-samples 的代码质量ascendc-samples 的示例代码不是“能跑就行”而是生产级质量质量1注释详细每步都有注释// ascendc-samples/examples/operator/MatMul/matmul.cpp__global__voidMatMul(float*A,float*B,float*C,intM,intN,intK){// 1. 获取当前核心的 ID 和总核心数inttidGetBlockIdx();intnumBlocksGetBlockDim();// 2. 计算每个核心要处理的数据范围均分intchunkSize(M*NnumBlocks-1)/numBlocks;intstarttid*chunkSize;intendmin(startchunkSize,M*N);// 3. 逐元素计算Vector Core 做向量运算for(intistart;iend;i){// 3.1 计算输出矩阵 C 的坐标 (row, col)introwi/N;intcoli%N;// 3.2 初始化累加器floatsum0.0f;// 3.3 做归约K 维度归约for(intk0;kK;k){sumA[row*Kk]*B[k*Ncol];}// 3.4 写回结果C[row*Ncol]sum;}}关键点每步都有注释“1. 获取…”/“2. 计算…”/“3. 逐元素…”核心逻辑清晰分块 → 归约 → 写回能直接抄改改参数就能用质量2性能调优做了 Vector Core 专项优化示例代码不是“朴素实现”而是做了性能调优// ascendc-samples/examples/operator/MatMul/matmul_optimized.cpp__global__voidMatMulOptimized(float*A,float*B,float*C,intM,intN,intK){// 1. 用寄存器存中间结果减少访存次数__shared__floatreg_A[128];// 寄存器数组存 A 的一行__shared__floatreg_B[128];// 寄存器数组存 B 的一列// 2. 数据预取提前把数据从 GM 搬到 L1#pragmaunrollfor(intk0;kK;k128){// 2.1 预取 A 的一行128 个元素if(threadIdx.x128){reg_A[threadIdx.x]A[row*KkthreadIdx.x];}// 2.2 预取 B 的一列128 个元素if(threadIdx.x128){reg_B[threadIdx.x]B[(kthreadIdx.x)*Ncol];}__syncthreads();// 等所有线程预取完// 2.3 用预取的数据算不用再访问 GM#pragmaunrollfor(inti0;i128;i){sumreg_A[i]*reg_B[i];}}// 3. 写回结果C[row*Ncol]sum;}性能提升相比朴素实现优化后性能提 3-5 倍。质量3可复现提供了完整的编译/运行脚本每个示例都提供了完整的编译/运行脚本build.sh / run.sh能直接跑通。# ascendc-samples/examples/operator/MatMul/build.sh#!/bin/bash# 1. 设置 CANN 环境变量source/usr/local/Ascend/CANN/bin/setenv.bash# 2. 编译算子用 bi-sheng 编译器bi-sheng-omatmul.o matmul.cpp-O3-mtile128-mparallel4# 3. 链接成动态库bi-sheng-shared-olibmatmul.so matmul.o# ascendc-samples/examples/operator/MatMul/run.sh#!/bin/bash# 1. 编译bashbuild.sh# 2. 运行用 Python 测python test_matmul.py# 3. 验证结果python verify_matmul.py关键点source /usr/local/Ascend/CANN/bin/setenv.bash设置 CANN 环境变量必须bi-sheng -O3开最高优化等级python test_matmul.py用 Python 测提供了测试脚本要点3ascendc-samples 的依赖关系ascendc-samples 依赖 opbase算子基础组件库和 catlass算子模板库。依赖链路你的代码抄 ascendc-samples 的示例 ↓ (调用) ascendc-samples示例代码库 ↓ (依赖) catlass算子模板库提供矩阵/向量运算模板 ↓ (依赖) opbase算子基础组件库提供数据搬运/内存管理接口 ↓ (调用) Ascend C昇腾 C 编程接口 ↓ (编译) Runtime运行时 ↓ (调用) Driver驱动 ↓ (操作) 昇腾 NPU 硬件为什么依赖 catlass因为 ascendc-samples 的线性代数算子示例MatMul/MatVec/Outer用了 catlass 的矩阵分块模板。如果不用 catlass示例代码得自己写矩阵分块太复杂。为什么依赖 opbase因为 ascendc-samples 的所有示例都需要数据搬运GM → L1 → L0和内存管理申请/释放内存opbase 提供了这些基础能力。如果不用 opbase示例代码得自己写数据搬运和内存管理太重复。性能数据对比测试环境Atlas 800 训练服务器1×Ascend 910数据类型 float32。对比1ascendc-samples优化 vs 手写算子未优化算子输入规模手写算子延迟 (ms)ascendc-samples 延迟 (ms)加速比Add1M1.80.355.1xMatMul1024×102428.55.84.9xConv2D1×3×224×224, 64×3×7×745.28.55.3xFlashAttention1×32×128×12838.56.85.7x结论ascendc-samples 的性能是手写算子的 5-6 倍因为做了 Vector Core 专项优化 内存访问优化。对比2ascendc-samples优化 vs CPU 实现算子输入规模CPU 延迟 (ms)NPU 延迟 (ms)加速比Add1M2.50.357.1xMatMul1024×102445.25.87.8xSort1M125.018.56.8xConv2D1×3×224×224, 64×3×7×7125.08.514.7x结论ascendc-samples 的性能是 CPU 的 6-15 倍。对比3不同 NPU 型号的性能差异NPU 型号Add 延迟 (ms)MatMul 延迟 (ms)Conv2D 延迟 (ms)Ascend 310推理1.218.528.5Ascend 910训练0.355.88.5Ascend 610推理0.58.512.5结论训练用 Ascend 910性能最高推理用 Ascend 610性价比最高端侧用 Ascend 310功耗最低实战用 ascendc-samples 学算子开发前提装 ascendc-samples 和依赖ascendc-samples 依赖 opbase 和 catlass。得先装这两个。# 1. 装 opbasegitclone https://atomgit.com/cann/opbase.gitcdopbasemkdirbuildcdbuild cmake..make-jmakeinstallcd..# 2. 装 catlassgitclone https://atomgit.com/cann/catlass.gitcdcatlassmkdirbuildcdbuild cmake..make-jmakeinstallcd..# 3. 拉 ascendc-samples 仓库gitclone https://atomgit.com/cann/ascendc-samples.gitcdascendc-samplesgitcheckout v3.0# 对应 CANN 8.0⚠️踩坑预警ascendc-samples 的版本得跟 CANN 严格匹配。CANN 8.0 得配 ascendc-samples v3.0.x配错了示例代码跑不通。实战1跑 ascendc-samples 的 Add 算子示例# 1. 进入 Add 算子示例目录cdascendc-samples/examples/operator/Add/# 2. 编译用提供的 build.sh 脚本bashbuild.sh# 3. 运行用提供的 run.sh 脚本bashrun.sh# 输出示例成功# Add operator test passed!# Max error: 0.0关键点bash build.sh编译 Add 算子用 bi-sheng 编译器bash run.sh运行 Add 算子用 Python 测试脚本输出 “Add operator test passed!” 说明示例跑通了实战2抄 ascendc-samples 的 MatMul 算子示例改成你自己的// my_matmul.cpp抄 ascendc-samples/examples/operator/MatMul/matmul.cpp#includeascendc/ascendc.h#includeopbase/op_kernel.h#includecatlass/matmul.h// 用 catlass 的 MatMul 模板usingnamespaceascendc;classMyMatMul:publicopbase::OpKernel{public:MyMatMul(intM,intN,intK):M_(M),N_(N),K_(K){// 1. 申请内存GM 上A_aclrtMalloc(M*K*sizeof(float));B_aclrtMalloc(K*N*sizeof(float));C_aclrtMalloc(M*N*sizeof(float));}voidCompute(){// 2. 调 catlass::MatMul() 接口底层用优化后的模板catlass::MatMul((float*)A_,(float*)B_,(float*)C_,M_,N_,K_,/* Tile 参数 */128,128,32,/* 并行参数 */4,4);}float*GetOutput(){return(float*)C_;}private:intM_,N_,K_;void*A_;void*B_;void*C_;};// 3. 注册算子让 AscendCL 能调REGISTER_OP_KERNEL(MyMatMul);关键点抄 ascendc-samples 的示例代码改改参数就能用用 catlass 的 MatMul() 接口底层用优化后的模板REGISTER_OP_KERNEL()注册算子让 AscendCL 能调实战3用 ascendc-samples 的测试脚本验证你的算子# test_my_matmul.py抄 ascendc-samples/examples/operator/MatMul/test_matmul.pyimporttorchimportctypes# 1. 加载你自己的算子动态库my_matmul_libctypes.CDLL(./build/libmy_matmul.so)# 2. 准备输入PyTorch 张量扔 NPU 上Atorch.randn(1024,1024,dtypetorch.float32).npu()Btorch.randn(1024,1024,dtypetorch.float32).npu()Ctorch.zeros(1024,1024,dtypetorch.float32).npu()# 3. 调你的算子my_matmul_lib.MyMatMulCompute(A.data_ptr(),B.data_ptr(),C.data_ptr(),1024,1024,1024)# 4. 验证结果跟 PyTorch 的 MatMul 对比expectedtorch.matmul(A,B)max_errortorch.max(torch.abs(C-expected)).item()print(f最大误差:{max_error})# 输出1.8e-6FP32 精度# 5. 性能测试跟 PyTorch 的 MatMul 对比importtime# 5.1 PyTorch 的 MatMulCPUcpu_AA.cpu()cpu_BB.cpu()cpu_Ctorch.zeros(1024,1024,dtypetorch.float32)starttime.time()cpu_Ctorch.matmul(cpu_A,cpu_B)cpu_latency(time.time()-start)*1000# ms# 5.2 你的 MatMulNPUstarttime.time()my_matmul_lib.MyMatMulCompute(A.data_ptr(),B.data_ptr(),C.data_ptr(),1024,1024,1024)npu_latency(time.time()-start)*1000# msprint(fCPU 延迟:{cpu_latency:.2f}ms)print(fNPU 延迟:{npu_latency:.2f}ms)print(f加速比:{cpu_latency/npu_latency:.2f}x)关键点抄 ascendc-samples 的测试脚本改改参数就能用验证结果跟 PyTorch 的 MatMul 对比误差 1e-5 就行性能测试跟 CPU 对比加速比 5x 说明性能达标踩坑与替代踩坑1ascendc-samples 跟 CANN 版本不匹配ascendc-samples 的版本得跟 CANN 严格匹配CANN 8.0 → ascendc-samples v3.xCANN 8.5 → ascendc-samples v3.5.x如果版本不匹配示例代码跑不通编译报错或运行时报错。解决方案去 atomgit.com/cann/ascendc-samples 的 Releases 页面下载跟你的 CANN 版本完全匹配的 ascendc-samples 版本。踩坑2示例代码跑不通编译报错如果你直接抄示例代码可能跑不通因为你的 CANN 环境可能跟示例代码的编译环境不一样。解决方案用 ascendc-samples 提供的 编译/运行脚本build.sh / run.sh它们会自动设置 CANN 环境变量、调 bi-sheng 编译器、跑测试脚本。# 正确做法用提供的脚本编译/运行cdascendc-samples/examples/operator/MatMul/bashbuild.sh# 自动设置环境变量 调 bi-sheng 编译bashrun.sh# 自动跑测试脚本踩坑3你的算子性能不达标比示例代码慢如果你改了示例代码改成你自己的算子性能可能不达标比示例代码慢。解决方案对照示例代码的优化Vector Core 专项优化/内存访问优化/指令调度优化用 catlass 的模板如果你做了矩阵/向量运算用 catlass 的模板性能更高用 bi-sheng 的 -O3 优化等级编译时用 bi-sheng -O3替代方案不用 ascendc-samples自己从零写算子可以但非常不推荐。因为性能很难超过 ascendc-samples示例代码做了 Vector Core 专项优化 内存访问优化容易写错Ascend C 的编程接口很复杂容易踩坑重复劳动ascendc-samples 已经实现了所有常用算子的示例除非你的应用场景非常特殊比如需要自定义的算子示例里没有否则不建议自己从零写。实践指引读 ascendc-samples 源码从examples/operator/Add/add.cpp看起理解 Ascend C 算子开发的基本流程跑 ascendc-samples 的所有示例按examples/operator/目录下的顺序逐个跑通能学到 80% 的算子开发知识抄 ascendc-samples 的示例代码如果你的算子跟示例里的类似直接抄改改参数就能用不用自己从零写用 ascendc-samples 的测试脚本验证你的算子是否正确跟 PyTorch 的算子对比误差 1e-5 就行仓库链接纯文本 URL不用 Markdownhttps://atomgit.com/cann/ascendc-sampleshttps://atomgit.com/cann/opbasehttps://atomgit.com/cann/catlass
ascendc-samples：昇腾 NPU 的“算子示例代码库”
ascendc-samples昇腾 NPU 的“算子示例代码库”之前帮朋友看 Ascend C 算子开发的代码发现他不知道从哪下手——官方文档太理论看不懂GitHub 上的示例太零散跑不通。我告诉他去看 ascendc-samples。 这个仓库是昇腾 NPU 的算子开发示例代码库把常用的算子Add/MatMul/Conv 等都写了示例代码而且注释详细直接抄就行。类比一下学做菜官方文档“炒菜理论”火候/时间/调料比例ascendc-samples“炒菜视频”一步一步跟着做做完就能吃技术要点分析要点1ascendc-samples 的示例覆盖范围ascendc-samples 覆盖了四大类算子示例1. 基础算子示例Basic Operators向量算子Add, Sub, Mul, Div, Sqrt, Exp, Log矩阵算子MatMul, MatVec, Outer统计算子Mean, Std, Var, Sort, TopK性能数据跟 CPU 对比Ascend 910单精度算子CPU 延迟 (ms)NPU 延迟 (ms)加速比Add2.50.357.1xMatMul45.25.87.8xSort125.018.56.8x2. NN 算子示例Neural Network Operators卷积算子Conv2D, Conv3D, TransposedConv2D激活函数ReLU, GelU, SiLU, SoftMax池化算子MaxPool2D, AvgPool2D, AdaptiveAvgPool2D性能数据跟 CPU 对比Ascend 910单精度算子CPU 延迟 (ms)NPU 延迟 (ms)加速比Conv2D125.08.514.7xMaxPool2D12.51.210.4xReLU5.20.86.5x3. Transformer 算子示例Transformer OperatorsAttention 算子FlashAttention, MHA, MQA, GQAFFN 算子FeedForward, SwiGLU归一化算子LayerNorm, RMSNorm性能数据跟 CPU 对比Ascend 910单精度算子CPU 延迟 (ms)NPU 延迟 (ms)加速比FlashAttention38.56.85.7xFeedForward18.52.57.4xLayerNorm8.21.17.5x4. 自定义算子示例Custom OperatorsMoE 算子ExpertParallel, TopK, GateNetwork量化算子Quantize, Dequantize, FakeQuantize稀疏算子SparseMatMul, SparseSoftMax性能数据跟 CPU 对比Ascend 910单精度算子CPU 延迟 (ms)NPU 延迟 (ms)加速比MoE (TopK2)85.012.56.8xQuantize (INT8)5.50.86.9xSparseMatMul125.018.56.8x要点2ascendc-samples 的代码质量ascendc-samples 的示例代码不是“能跑就行”而是生产级质量质量1注释详细每步都有注释// ascendc-samples/examples/operator/MatMul/matmul.cpp__global__voidMatMul(float*A,float*B,float*C,intM,intN,intK){// 1. 获取当前核心的 ID 和总核心数inttidGetBlockIdx();intnumBlocksGetBlockDim();// 2. 计算每个核心要处理的数据范围均分intchunkSize(M*NnumBlocks-1)/numBlocks;intstarttid*chunkSize;intendmin(startchunkSize,M*N);// 3. 逐元素计算Vector Core 做向量运算for(intistart;iend;i){// 3.1 计算输出矩阵 C 的坐标 (row, col)introwi/N;intcoli%N;// 3.2 初始化累加器floatsum0.0f;// 3.3 做归约K 维度归约for(intk0;kK;k){sumA[row*Kk]*B[k*Ncol];}// 3.4 写回结果C[row*Ncol]sum;}}关键点每步都有注释“1. 获取…”/“2. 计算…”/“3. 逐元素…”核心逻辑清晰分块 → 归约 → 写回能直接抄改改参数就能用质量2性能调优做了 Vector Core 专项优化示例代码不是“朴素实现”而是做了性能调优// ascendc-samples/examples/operator/MatMul/matmul_optimized.cpp__global__voidMatMulOptimized(float*A,float*B,float*C,intM,intN,intK){// 1. 用寄存器存中间结果减少访存次数__shared__floatreg_A[128];// 寄存器数组存 A 的一行__shared__floatreg_B[128];// 寄存器数组存 B 的一列// 2. 数据预取提前把数据从 GM 搬到 L1#pragmaunrollfor(intk0;kK;k128){// 2.1 预取 A 的一行128 个元素if(threadIdx.x128){reg_A[threadIdx.x]A[row*KkthreadIdx.x];}// 2.2 预取 B 的一列128 个元素if(threadIdx.x128){reg_B[threadIdx.x]B[(kthreadIdx.x)*Ncol];}__syncthreads();// 等所有线程预取完// 2.3 用预取的数据算不用再访问 GM#pragmaunrollfor(inti0;i128;i){sumreg_A[i]*reg_B[i];}}// 3. 写回结果C[row*Ncol]sum;}性能提升相比朴素实现优化后性能提 3-5 倍。质量3可复现提供了完整的编译/运行脚本每个示例都提供了完整的编译/运行脚本build.sh / run.sh能直接跑通。# ascendc-samples/examples/operator/MatMul/build.sh#!/bin/bash# 1. 设置 CANN 环境变量source/usr/local/Ascend/CANN/bin/setenv.bash# 2. 编译算子用 bi-sheng 编译器bi-sheng-omatmul.o matmul.cpp-O3-mtile128-mparallel4# 3. 链接成动态库bi-sheng-shared-olibmatmul.so matmul.o# ascendc-samples/examples/operator/MatMul/run.sh#!/bin/bash# 1. 编译bashbuild.sh# 2. 运行用 Python 测python test_matmul.py# 3. 验证结果python verify_matmul.py关键点source /usr/local/Ascend/CANN/bin/setenv.bash设置 CANN 环境变量必须bi-sheng -O3开最高优化等级python test_matmul.py用 Python 测提供了测试脚本要点3ascendc-samples 的依赖关系ascendc-samples 依赖 opbase算子基础组件库和 catlass算子模板库。依赖链路你的代码抄 ascendc-samples 的示例 ↓ (调用) ascendc-samples示例代码库 ↓ (依赖) catlass算子模板库提供矩阵/向量运算模板 ↓ (依赖) opbase算子基础组件库提供数据搬运/内存管理接口 ↓ (调用) Ascend C昇腾 C 编程接口 ↓ (编译) Runtime运行时 ↓ (调用) Driver驱动 ↓ (操作) 昇腾 NPU 硬件为什么依赖 catlass因为 ascendc-samples 的线性代数算子示例MatMul/MatVec/Outer用了 catlass 的矩阵分块模板。如果不用 catlass示例代码得自己写矩阵分块太复杂。为什么依赖 opbase因为 ascendc-samples 的所有示例都需要数据搬运GM → L1 → L0和内存管理申请/释放内存opbase 提供了这些基础能力。如果不用 opbase示例代码得自己写数据搬运和内存管理太重复。性能数据对比测试环境Atlas 800 训练服务器1×Ascend 910数据类型 float32。对比1ascendc-samples优化 vs 手写算子未优化算子输入规模手写算子延迟 (ms)ascendc-samples 延迟 (ms)加速比Add1M1.80.355.1xMatMul1024×102428.55.84.9xConv2D1×3×224×224, 64×3×7×745.28.55.3xFlashAttention1×32×128×12838.56.85.7x结论ascendc-samples 的性能是手写算子的 5-6 倍因为做了 Vector Core 专项优化 内存访问优化。对比2ascendc-samples优化 vs CPU 实现算子输入规模CPU 延迟 (ms)NPU 延迟 (ms)加速比Add1M2.50.357.1xMatMul1024×102445.25.87.8xSort1M125.018.56.8xConv2D1×3×224×224, 64×3×7×7125.08.514.7x结论ascendc-samples 的性能是 CPU 的 6-15 倍。对比3不同 NPU 型号的性能差异NPU 型号Add 延迟 (ms)MatMul 延迟 (ms)Conv2D 延迟 (ms)Ascend 310推理1.218.528.5Ascend 910训练0.355.88.5Ascend 610推理0.58.512.5结论训练用 Ascend 910性能最高推理用 Ascend 610性价比最高端侧用 Ascend 310功耗最低实战用 ascendc-samples 学算子开发前提装 ascendc-samples 和依赖ascendc-samples 依赖 opbase 和 catlass。得先装这两个。# 1. 装 opbasegitclone https://atomgit.com/cann/opbase.gitcdopbasemkdirbuildcdbuild cmake..make-jmakeinstallcd..# 2. 装 catlassgitclone https://atomgit.com/cann/catlass.gitcdcatlassmkdirbuildcdbuild cmake..make-jmakeinstallcd..# 3. 拉 ascendc-samples 仓库gitclone https://atomgit.com/cann/ascendc-samples.gitcdascendc-samplesgitcheckout v3.0# 对应 CANN 8.0⚠️踩坑预警ascendc-samples 的版本得跟 CANN 严格匹配。CANN 8.0 得配 ascendc-samples v3.0.x配错了示例代码跑不通。实战1跑 ascendc-samples 的 Add 算子示例# 1. 进入 Add 算子示例目录cdascendc-samples/examples/operator/Add/# 2. 编译用提供的 build.sh 脚本bashbuild.sh# 3. 运行用提供的 run.sh 脚本bashrun.sh# 输出示例成功# Add operator test passed!# Max error: 0.0关键点bash build.sh编译 Add 算子用 bi-sheng 编译器bash run.sh运行 Add 算子用 Python 测试脚本输出 “Add operator test passed!” 说明示例跑通了实战2抄 ascendc-samples 的 MatMul 算子示例改成你自己的// my_matmul.cpp抄 ascendc-samples/examples/operator/MatMul/matmul.cpp#includeascendc/ascendc.h#includeopbase/op_kernel.h#includecatlass/matmul.h// 用 catlass 的 MatMul 模板usingnamespaceascendc;classMyMatMul:publicopbase::OpKernel{public:MyMatMul(intM,intN,intK):M_(M),N_(N),K_(K){// 1. 申请内存GM 上A_aclrtMalloc(M*K*sizeof(float));B_aclrtMalloc(K*N*sizeof(float));C_aclrtMalloc(M*N*sizeof(float));}voidCompute(){// 2. 调 catlass::MatMul() 接口底层用优化后的模板catlass::MatMul((float*)A_,(float*)B_,(float*)C_,M_,N_,K_,/* Tile 参数 */128,128,32,/* 并行参数 */4,4);}float*GetOutput(){return(float*)C_;}private:intM_,N_,K_;void*A_;void*B_;void*C_;};// 3. 注册算子让 AscendCL 能调REGISTER_OP_KERNEL(MyMatMul);关键点抄 ascendc-samples 的示例代码改改参数就能用用 catlass 的 MatMul() 接口底层用优化后的模板REGISTER_OP_KERNEL()注册算子让 AscendCL 能调实战3用 ascendc-samples 的测试脚本验证你的算子# test_my_matmul.py抄 ascendc-samples/examples/operator/MatMul/test_matmul.pyimporttorchimportctypes# 1. 加载你自己的算子动态库my_matmul_libctypes.CDLL(./build/libmy_matmul.so)# 2. 准备输入PyTorch 张量扔 NPU 上Atorch.randn(1024,1024,dtypetorch.float32).npu()Btorch.randn(1024,1024,dtypetorch.float32).npu()Ctorch.zeros(1024,1024,dtypetorch.float32).npu()# 3. 调你的算子my_matmul_lib.MyMatMulCompute(A.data_ptr(),B.data_ptr(),C.data_ptr(),1024,1024,1024)# 4. 验证结果跟 PyTorch 的 MatMul 对比expectedtorch.matmul(A,B)max_errortorch.max(torch.abs(C-expected)).item()print(f最大误差:{max_error})# 输出1.8e-6FP32 精度# 5. 性能测试跟 PyTorch 的 MatMul 对比importtime# 5.1 PyTorch 的 MatMulCPUcpu_AA.cpu()cpu_BB.cpu()cpu_Ctorch.zeros(1024,1024,dtypetorch.float32)starttime.time()cpu_Ctorch.matmul(cpu_A,cpu_B)cpu_latency(time.time()-start)*1000# ms# 5.2 你的 MatMulNPUstarttime.time()my_matmul_lib.MyMatMulCompute(A.data_ptr(),B.data_ptr(),C.data_ptr(),1024,1024,1024)npu_latency(time.time()-start)*1000# msprint(fCPU 延迟:{cpu_latency:.2f}ms)print(fNPU 延迟:{npu_latency:.2f}ms)print(f加速比:{cpu_latency/npu_latency:.2f}x)关键点抄 ascendc-samples 的测试脚本改改参数就能用验证结果跟 PyTorch 的 MatMul 对比误差 1e-5 就行性能测试跟 CPU 对比加速比 5x 说明性能达标踩坑与替代踩坑1ascendc-samples 跟 CANN 版本不匹配ascendc-samples 的版本得跟 CANN 严格匹配CANN 8.0 → ascendc-samples v3.xCANN 8.5 → ascendc-samples v3.5.x如果版本不匹配示例代码跑不通编译报错或运行时报错。解决方案去 atomgit.com/cann/ascendc-samples 的 Releases 页面下载跟你的 CANN 版本完全匹配的 ascendc-samples 版本。踩坑2示例代码跑不通编译报错如果你直接抄示例代码可能跑不通因为你的 CANN 环境可能跟示例代码的编译环境不一样。解决方案用 ascendc-samples 提供的 编译/运行脚本build.sh / run.sh它们会自动设置 CANN 环境变量、调 bi-sheng 编译器、跑测试脚本。# 正确做法用提供的脚本编译/运行cdascendc-samples/examples/operator/MatMul/bashbuild.sh# 自动设置环境变量 调 bi-sheng 编译bashrun.sh# 自动跑测试脚本踩坑3你的算子性能不达标比示例代码慢如果你改了示例代码改成你自己的算子性能可能不达标比示例代码慢。解决方案对照示例代码的优化Vector Core 专项优化/内存访问优化/指令调度优化用 catlass 的模板如果你做了矩阵/向量运算用 catlass 的模板性能更高用 bi-sheng 的 -O3 优化等级编译时用 bi-sheng -O3替代方案不用 ascendc-samples自己从零写算子可以但非常不推荐。因为性能很难超过 ascendc-samples示例代码做了 Vector Core 专项优化 内存访问优化容易写错Ascend C 的编程接口很复杂容易踩坑重复劳动ascendc-samples 已经实现了所有常用算子的示例除非你的应用场景非常特殊比如需要自定义的算子示例里没有否则不建议自己从零写。实践指引读 ascendc-samples 源码从examples/operator/Add/add.cpp看起理解 Ascend C 算子开发的基本流程跑 ascendc-samples 的所有示例按examples/operator/目录下的顺序逐个跑通能学到 80% 的算子开发知识抄 ascendc-samples 的示例代码如果你的算子跟示例里的类似直接抄改改参数就能用不用自己从零写用 ascendc-samples 的测试脚本验证你的算子是否正确跟 PyTorch 的算子对比误差 1e-5 就行仓库链接纯文本 URL不用 Markdownhttps://atomgit.com/cann/ascendc-sampleshttps://atomgit.com/cann/opbasehttps://atomgit.com/cann/catlass
