cann-robotdocs: 修复 README 编译运行样例链接失效及错误码 EZ9999

文件	最后提交记录	最后更新时间
arch22	统一所有算子获取内核数量方法为通过plantformC api获取，部分强行硬编码为8的算子需要自己手动调整，同时将部分重复在host侧定义的函数修改为使用公共库中的函数 Co-authored-by: Jett.chen<1519755291@qq.com> # message auto-generated for no-merge-commit merge: !163 merge plantformC into master 统一所有算子获取内核数量方法为通过plantformC api获取，部分强行硬编码为8的算子需要自己手动调整 Created-by: 2302_77046878 Commit-by: Jett.chen Merged-by: cann-robot Description: ## 描述 <!--在这里详细描述你的改动，包括改动的原因和所采取的方法。--> 统一所有算子获取内核数量方法为通过plantformC api获取，部分强行硬编码为8的算子需要自己手动调整，直接进行修改可能会出现预期外的错误 ## 关联的Issue <!-- 如果这个PR是为了解决特定的Issue，请在这里提供Issue链接。例如：关联Issue #000--> https://gitcode.com/cann/ops-blas/issues/136 ## 测试 <!--描述进行了哪些测试来验证你的改动。--> 修改后对于arch22，arch35架构下的算子进行了测试，全部通过 ![image.png](https://raw.gitcode.com/user-images/assets/8916851/1b4a1a18-b126-4957-929b-1fb9cd7be972/image.png 'image.png') ![image.png](https://raw.gitcode.com/user-images/assets/8916851/e6f2bbb5-5ec7-4052-803d-2882f2074957/image.png 'image.png') ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!163	1 天前
arch35	feat(nrm2): 新增 arch35 aclblasSnrm2接口 Co-authored-by: chensi79<chensi79@huawei.com> # message auto-generated for no-merge-commit merge: !181 merge aclblasSnrm2 into master feat(nrm2): 新增 arch35 aclblasSnrm2接口 Created-by: chensi79 Commit-by: chensi79@huawei.com;chensi79 Merged-by: cann-robot Description: ## 描述为 ascend950（arch35）新增 `aclblasSnrm2` 接口实现，并配套调整 arch22 命名与测试目录结构。 ### 改动原因 - 现有 `aclblasSnrm2` 仅有占位声明（`include/cann_ops_blas.h`），缺少面向 arch35（Ascend 950）的实现 - arch35 的 AIV 多核 + SIMT 编程模型与 arch22 不同，需要独立的双路径 kernel - 历史目录命名 `snrm2_` 统一为 `nrm2_` ### 改动方法 1. arch35 aclblasSnrm2 实现 - 接口签名（`include/cann_ops_blas.h`）：从占位声明 `(handle, const int64_t n, uint8_t* x, const int64_t incx, uint8_t* result)` 调整为 `(handle, int n, const float* x, int incx, float* result)` - Host（`blas/nrm2/arch35/nrm2_host.cpp`）：参数校验、`n≤0`/`incx≤0` 早返回 0、Tiling（useCoreNum = min(aivCoreNum, n, 64)）、workspace 复用 `aclblasGetEffectiveWorkspace` - Kernel（`blas/nrm2/arch35/nrm2_kernel.cpp` + `nrm2_kernel.h`）：双路径 + 汇总 - SIMD 路径（incx==1）：`snrm2_aiv_kernel` 多核 `Mul → ReduceSum` 后 atomic add 写入 workspace[blockIdx] - SIMT 路径（incx!=1）：`snrm2_simt_kernel` 多核 `asc_vf_call<Snrm2SimtCompute>`，warp-style 归约 - 汇总：`snrm2_reduce_kernel` 单核 `ReduceSum + Sqrt` - UB 预算：`SAFETY_MARGIN=32KB`、`UB_MAX_CHUNK_FLOATS=27392`（基于 BUFFER_NUM=2、248KB UB 推导） 2. arch22 命名规范化（纯重命名，内容基本不变） - `blas/nrm2/arch22/snrm2_host.cpp → nrm2_host.cpp` - `blas/nrm2/arch22/snrm2_kernel.cpp → nrm2_kernel.cpp` 3. 测试文件： - 新增 `test/nrm2/snrm2/arch35/`（`snrm2_test.cpp` / `snrm2_test.csv` / `snrm2_npu_wrapper.h`） - 新增 `test/nrm2/snrm2/snrm2_golden.h`（cblas_snrm2 参考）、`snrm2_param.h` 4. sasum_kernel.cpp（arch35）修正 - `maxCopyPadNum` 由 `(UINT16_MAX+1)/sizeof(float)` 改为 `UINT16_MAX/sizeof(float)`，避免溢出 - SIMT 归约改用 next-power-of-2 算法并补充 `(threadIdx.x + s) < blockDim.x` 边界检查，支持非 2 的幂 `blockDim.x` 5. 文档 - 更新 `blas/nrm2/README.md`（补充 arch22/arch35 双架构、双路径设计、硬件支持矩阵） ## 关联的Issue - https://gitcode.com/cann/ops-blas/issues/190 ## 测试 - 单元测试 `test/nrm2/snrm2/arch35/nrm2_test.cpp`，CSV 数据驱动，覆盖： - SIMD 路径：n=0/1/8/32/1024/8192（含 32B 对齐与非对齐） - SIMT 路径：incx=2/3/7（含素数 stride） - 错误路径：NullHandle、x/result nullptr、n≤0、incx≤0 - 参考实现 `test/nrm2/snrm2/nrm2_golden.h` 调用 `cblas_snrm2`做精度比对 ## 文档更新 - 更新 `blas/nrm2/README.md` ## 类型标签 - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!181	1 天前
README.md	docs: 修复 README 编译运行样例链接失效及错误码 EZ9999 Co-authored-by: Zhang Hua<zhanghua25@mails.ucas.ac.cn> # message auto-generated for no-merge-commit merge: !213 merge docs/issue-fix-readme-link-and-errorcode into master docs: 修复 README 编译运行样例链接失效及错误码 EZ9999 Created-by: zhanghua145 Commit-by: Zhang Hua Merged-by: cann-robot Description: ## 描述修复两个文档问题： 1. 链接修复：所有 `blas/${op}/README.md` 中指向 `compile_and_run_example.md` 的链接路径错误，无法跳转到实际文件 `docs/zh/develop/compile_and_run_example.md`。将 `compile_and_run_example.md` 替换为正确的相对路径 `../../docs/zh/develop/compile_and_run_example.md`。 2. 错误码修复：`compile_and_run_example.md` 中引用了内核层错误码 `EZ9999`，但该错误码在仓库 `include/` 目录及 CANN 官方文档中均不存在。SOC 不匹配时实际报错为 ACL 运行时层 `507035`（向量核异常）。删除 `EZ9999` 引用，仅保留 `507035`。 ## 关联的Issue [#217](https://gitcode.com/cann/ops-blas/issues/217) ## 测试 - 通过 `grep` 验证所有 42 个 blas README 链接路径已更新 - 通过 `grep` 验证仓库中不再包含 `EZ9999` 引用 ## 文档更新 - 42 个 `blas/${op}/README.md`：链接路径修正 - `docs/zh/develop/compile_and_run_example.md`：删除不存在的错误码 `EZ9999` ## 类型标签 - [x] 文档更新 See merge request: cann/ops-blas!213	22 小时前

Asum算子

算子概述

向量运算算子，计算向量元素绝对值之和（L1 范数 / 曼哈顿范数），常用于向量稀疏度度量和误差估计。

数学表达式：

result = sum(|x[i]|) for i = 0 to n-1

包含以下接口：

接口名	功能简述
aclblasSasum	实数向量绝对值之和

算子执行接口

aclblasSasum

产品支持情况

Ascend 950PR / Ascend 950DT：支持
Atlas A3 训练系列产品 / Atlas A3 推理系列产品：支持
Atlas A2 训练系列产品 / Atlas A2 推理系列产品：支持

函数原型

aclblasStatus_t aclblasSasum(aclblasHandle_t handle, int n, const float *x, int incx, float *result)

参数说明

参数名	输入/输出	参数类型	说明
handle	输入	aclblasHandle_t	ops-blas 库上下文句柄，携带 stream，Host 内存
n	输入	int	向量元素个数，Host 内存
x	输入	const float*（FP32）	float 向量，包含 n 个元素，Device 内存
incx	输入	int	x 中连续元素之间的步长，不可为 0，Host 内存
result	输出	float*（FP32）	向量元素绝对值之和，Device 内存

约束说明

n >= 0（n < 0 时返回错误）
incx != 0