| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
Feat：add default/user controlled workspace behavior in handle Co-authored-by: Zitao Wang<wangzitao4@huawei.com> # message auto-generated for no-merge-commit merge: !166 merge handle_workspace into master Feat：add default/user controlled workspace behavior in handle Created-by: wangzitao_leo Commit-by: Zitao Wang Merged-by: cann-robot Description: ## 描述 <!--在这里详细描述你的改动，包括改动的原因和所采取的方法。--> 本次 PR 为 aclblasHandle 引入了默认 workspace 与用户 workspace 的双轨机制，使库在 aclblasCreate 时自动预分配 4 MiB 的默认 workspace，同时保留用户通过 aclblasSetWorkspace 自行管理 device 内存的能力。核心数据结构 _aclblas_handle、创建/销毁/设置接口的行为均随之调整，版本宏也被拆分到独立的 aclblas_version.h 文件中。 主要改动: 1. workspace 双轨机制: _aclblas_handle 内部用 default_workspace（库管理）和 user_workspace（用户管理）替代了原先单一的 workspace 字段，并通过 use_user_workspace 标志位选择当前生效的 workspace，新增 aclblasGetEffectiveWorkspace / aclblasGetEffectiveWorkspaceSize 等内联函数来获取当前活跃的 workspace 指针和大小。 2. aclblasCreate / aclblasDestroy 行为变更: aclblasCreate 现在会自动调用 allocateDefaultWorkspace 预分配 4 MiB 默认 workspace；aclblasDestroy 增加了 aclrtSynchronizeStream 同步等待 GPU 完成后再释放默认 workspace，但不释放用户 workspace 内存。 3. aclblasSetWorkspace 语义调整: 传入 nullptr 时通过 aclblasResetToDefaultWorkspace 切回默认 workspace；传入有效指针时采用 grow-only 策略（仅在新的 size 大于当前用户 workspace size 时更新），且参数校验从"workspace 与 workspaceSize 必须同时空或同时非空"改为"workspace 非空时 workspaceSize 必须大于 0"。 4. aclblasSetStream 新增重置行为: 切换 stream 时自动调用 aclblasResetToDefaultWorkspace 将 workspace 重置为库默认，与 cuBLAS 行为保持一致。 ## 关联的Issue <!-- 如果这个PR是为了解决特定的Issue，请在这里提供Issue链接。例如：关联Issue #000--> 关联Issue [#156](https://gitcode.com/cann/ops-blas/issues/156) ## 测试 <!--描述进行了哪些测试来验证你的改动。--> A5：  ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> NA ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!166 | 12 天前 | |
Feat: 新增面向arch35的aclblasGemmBatchedEx接口 Co-authored-by: developer<developer@opencode.ai> Co-authored-by: yang-di52<yangdi52@huawei.com> # message auto-generated for no-merge-commit merge: !183 merge aclblasGemmBatchedEx into master Feat: 新增面向arch35的aclblasGemmBatchedEx接口 Created-by: yang-di52 Commit-by: yang-di52;developer Merged-by: cann-robot Description: ## 描述 新增面向 arch35 (Ascend950) 平台的 aclblasGemmBatchedEx 批量矩阵乘接口。 该接口是 aclblasGemmEx 的批量扩展版本，支持通过设备侧指针数组指定多组矩阵，在单次调用中完成所有批次的 GEMM 计算： C[i] = alpha * op(A[i]) * op(B[i]) + beta * C[i], for i ∈ [0, batchCount - 1] **支持的数据类型组合（7 种）**： | 序号 | Atype | Btype | Ctype | computeType | |------|-------|-------|-------|-------------| | 1 | FP16 | FP16 | FP16 | COMPUTE_16F | | 2 | FP16 | FP16 | FP16 | COMPUTE_32F | | 3 | BF16 | BF16 | BF16 | COMPUTE_32F | | 4 | FP8_E4M3 | FP8_E4M3 | FP16 | COMPUTE_32F | | 5 | FP8_E5M2 | FP8_E5M2 | FP16 | COMPUTE_32F | | 6 | FP8_E4M3 | FP8_E5M2 | FP16 | COMPUTE_32F | | 7 | FP8_E5M2 | FP8_E4M3 | FP16 | COMPUTE_32F | **核心设计**： - Batch × M × N 三维联合切分，grid-stride 循环多核分配 - 复用 gemm_ex 的 BlockMmad 低阶 API（Nd2Nz → LoadData → Mmad → Fixpipe） - Kernel 侧直接从 GM 读取设备侧指针数组，无需 Host 侧 D2H 拷贝 - Alpha/Beta Vector 后处理 + 256 MB 临时缓冲上限保护 + 分批处理策略 **交付物清单**： | 类别 | 文件 | |------|------| | Host 代码 | blas/gemm/gemm_batched_ex/arch35/gemm_batched_ex_host.cpp | | Kernel 代码 | blas/gemm/gemm_batched_ex/arch35/gemm_batched_ex_kernel.cpp | | TilingData | blas/gemm/gemm_batched_ex/arch35/gemm_batched_ex_tiling_data.h | | 测试代码 | test/gemm/gemm_batched_ex/ | | 算子文档 | blas/gemm/gemm_batched_ex/README.md | ## 关联的Issue [#178](https://gitcode.com/cann/ops-blas/issues/178) ## 测试 ### ST 精度测试 - **测试框架**：CSV 驱动 GTest，93 条用例（37 条 L0 + 56 条 L1） - **测试覆盖**：全部 7 种合法 dtype 组合、NN/NT/TN/TT/CC/NC/CN 转置组合、batchCount 1~1024、alpha/beta 各种取值、边界 shape（单行/单列/非对齐/极端维度）、非法参数校验 ### 性能测试 - **测试环境**：Ascend950 (PR/DT)，arch35，CANN 9.1.0，28 AI Cube Core，1650 MHz - **大规模性能**： - 2048×2048×2048 batch=2：3377 GFLOPS - 1024×1024×1024 batch=4：2226 GFLOPS - 512×512×512 batch=4：1113 GFLOPS - **多核利用率**：大 shape 下 28 核全部利用（100%），核间负载均衡 <2% - **满频运行**：Current Freq = Rated Freq = 1650 MHz ## 文档更新 - 新增 blas/gemm/gemm_batched_ex/README.md：接口说明、参数规格、支持的数据类型组合、已知限制、编译/测试命令、调用示例（含转置示例） ## 类型标签 - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!183 | 4 天前 | |
log接口支持 Co-authored-by: yang-di52<yangdi52@huawei.com> # message auto-generated for no-merge-commit merge: !106 merge fix into master log接口支持 Created-by: yang-di52 Commit-by: yang-di52 Merged-by: cann-robot Description: ## 描述 ### 变更摘要 本次 PR 为 ops-blas 新增了日志接口支持，核心是将 aclblasLoggerConfigure 中配置的日志级别通过 dlog_setlevel 实际生效到底层 dlog 日志系统。aclblas_logger_manager 模块从 C 链接方式改为 AclBlas 命名空间封装。 ### 主要改动 * 日志级别生效: 在 aclblasLoggerConfigure 中新增 dlog_setlevel 调用，根据传入的 aclblasLogLevel_t（INFO/ERROR/DEBUG）实际设置底层 dlog 日志等级，使日志配置不再只是存储参数。 * 命名空间重构: aclblas_logger_manager 模块从 extern "C" 包装改为 AclBlas 命名空间，头文件守卫从 #pragma once 改为传统的 #ifndef/#define 宏，并引入 log/log.h 头文件。 ## 关联的Issue [#96](https://gitcode.com/cann/ops-blas/issues/96) ## 测试 <!--描述进行了哪些测试来验证你的改动。--> ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!106 | 27 天前 | |
refactor: 统一头文件 include guard 为 #pragma once 并添加 CI lint 规则 Co-authored-by: Zhang Hua<1302896824@qq.com> # message auto-generated for no-merge-commit merge: !173 merge unify-include-guards into master refactor: 统一头文件 include guard 为 #pragma once 并添加 CI lint 规则 Created-by: zhanghua145 Commit-by: Zhang Hua Merged-by: cann-robot Description: ## 描述 统一仓库中所有头文件的 include guard 风格为 #pragma once，解决当前 #ifndef/#define/#endif 与 #pragma once 混用导致的风格不一致问题。 ## 关联 Issue 关联 Issue #162 ## 变更内容 ### 1. 统一现有头文件（125 个文件，+456/-698） | 类别 | 数量 | 说明 | |------|------|------| | #ifndef → #pragma once | 215 个 | 传统宏保护替换为 #pragma once | | 补充 #pragma once | 1 个 | cgerc_kernel_impl.h 原先缺少 include guard | | agent 模板同步 | 2 个 | 模板文件同步转换 | | **净减少行数** | **242 行** | 456 行增加，698 行删除 | ### 2. 添加 CI lint 检查脚本 新增 scripts/ci/check_pragma_once.sh，用于检查 PR 中新增/修改的 .h/.hpp 文件是否使用 #pragma once。 **调用方式**（与现有 scripts/ci/run_example.sh 一致，接收 PR 文件列表）： bash bash scripts/ci/check_pragma_once.sh <pr_filelist.txt> pr_filelist.txt 格式为每行一个文件路径（相对仓库根目录），仅检查 .h/.hpp 文件，其他文件类型自动跳过。 **检查规则**： - 必须使用 #pragma once，拒绝 #ifndef include guard 和无 guard 的文件 - 自动跳过第三方目录（.opencode/、asc-devkit/ 等） ## ⚠️ 需要维护者操作：接入 CI 流水线 本 PR 提供了 scripts/ci/check_pragma_once.sh 检查脚本，但**尚未接入 CI 流水线**。需要维护者在 GitCode 平台侧的流水线配置中新增一个步骤，在编译前调用该脚本。 建议的接入方式（参考现有 run_example.sh 的调用模式）： yaml # CI 流水线中新增 lint 步骤（在编译步骤之前） - name: check-pragma-once script: | # pr_filelist.txt 由平台自动生成，包含 PR 变更文件列表 bash scripts/ci/check_pragma_once.sh pr_filelist.txt 在维护者接入流水线之前，该脚本也可通过以下方式手动使用： - **本地验证**：bash scripts/ci/check_pragma_once.sh <文件列表> - **pre-commit hook**：已集成到 .pre-commit-config.yaml（需开发者自行安装 pre-commit） ## 关于 Issue 中提到的 3 个缺少 guard 的头文件 Issue #162 提到有 3 个头文件缺少标准 include guard。经逐 commit 追溯历史，原始 3 个文件为： | # | 文件 | 当前状态 | |---|------|----------| | 1 | blas/gerc/cgerc/arch22/cgerc_kernel_impl.h | 仍存在，**本次已修复** | | 2 | blas/cgerc/cgerc_kernel_impl.h | 已被删除（cgerc 目录重构，迁移至 blas/gerc/cgerc/arch22/） | | 3 | blas/common/kernel_launch/aclblas_kernel_do.h | 已被删除（commit cc4d43c "del common kernel launch header"） | 第 2、3 个文件已在其他 PR 的重构过程中被删除，因此本次只需修复第 1 个文件。 ## 未修改的 #ifndef 仓库中剩余约 114 处 #ifndef 均为**条件编译**（非 include guard），保持不变： - #ifndef __CCE_AICORE__ — 区分 Host/Device 编译目标 - #ifndef TEST_DEVICE_ID — 测试设备 ID 默认值 - #ifndef __force_inline__ — 编译器特性检测 - #ifndef ACLBLAS_OPERATION_DECLARED — 枚举声明保护 这些是功能性的编译条件判断，不属于 include guard，不在本次修改范围内。 ## 验证 - [x] 测试通过（bash build.sh --run --soc=ascend950，ops_blas + ops_blasLt 均成功）  - [x] 所有项目头文件（358 个）均已使用 #pragma once - [x] 无遗漏的 include guard 宏 - [x] CI lint 脚本测试通过（正确拒绝 #ifndef guard，放行 #pragma once，仅检查 PR 变更文件） ## 类型标签 - [ ] Bug 修复 - [ ] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [x] 其他：代码风格统一重构 See merge request: cann/ops-blas!173 | 12 天前 | |
Feat：add default/user controlled workspace behavior in handle Co-authored-by: Zitao Wang<wangzitao4@huawei.com> # message auto-generated for no-merge-commit merge: !166 merge handle_workspace into master Feat：add default/user controlled workspace behavior in handle Created-by: wangzitao_leo Commit-by: Zitao Wang Merged-by: cann-robot Description: ## 描述 <!--在这里详细描述你的改动，包括改动的原因和所采取的方法。--> 本次 PR 为 aclblasHandle 引入了默认 workspace 与用户 workspace 的双轨机制，使库在 aclblasCreate 时自动预分配 4 MiB 的默认 workspace，同时保留用户通过 aclblasSetWorkspace 自行管理 device 内存的能力。核心数据结构 _aclblas_handle、创建/销毁/设置接口的行为均随之调整，版本宏也被拆分到独立的 aclblas_version.h 文件中。 主要改动: 1. workspace 双轨机制: _aclblas_handle 内部用 default_workspace（库管理）和 user_workspace（用户管理）替代了原先单一的 workspace 字段，并通过 use_user_workspace 标志位选择当前生效的 workspace，新增 aclblasGetEffectiveWorkspace / aclblasGetEffectiveWorkspaceSize 等内联函数来获取当前活跃的 workspace 指针和大小。 2. aclblasCreate / aclblasDestroy 行为变更: aclblasCreate 现在会自动调用 allocateDefaultWorkspace 预分配 4 MiB 默认 workspace；aclblasDestroy 增加了 aclrtSynchronizeStream 同步等待 GPU 完成后再释放默认 workspace，但不释放用户 workspace 内存。 3. aclblasSetWorkspace 语义调整: 传入 nullptr 时通过 aclblasResetToDefaultWorkspace 切回默认 workspace；传入有效指针时采用 grow-only 策略（仅在新的 size 大于当前用户 workspace size 时更新），且参数校验从"workspace 与 workspaceSize 必须同时空或同时非空"改为"workspace 非空时 workspaceSize 必须大于 0"。 4. aclblasSetStream 新增重置行为: 切换 stream 时自动调用 aclblasResetToDefaultWorkspace 将 workspace 重置为库默认，与 cuBLAS 行为保持一致。 ## 关联的Issue <!-- 如果这个PR是为了解决特定的Issue，请在这里提供Issue链接。例如：关联Issue #000--> 关联Issue [#156](https://gitcode.com/cann/ops-blas/issues/156) ## 测试 <!--描述进行了哪些测试来验证你的改动。--> A5：  ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> NA ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!166 | 12 天前 | |
Feat: 新增面向arch35的aclblasGemmGroupedBatchedEx接口 Co-authored-by: Crrryyyy<chenruiyu4@huawei.com> # message auto-generated for no-merge-commit merge: !189 merge aclblasgemmgroupedbatchedEx into master Feat: 新增面向arch35的aclblasGemmGroupedBatchedEx接口 Created-by: Crrryyyy Commit-by: Crrryyyy Merged-by: cann-robot Description: ## 描述 新增面向 arch35 (Ascend950) 平台的 aclblasGemmGroupedBatchedEx 分组批量矩阵乘扩展接口。 该接口是 aclblasGemmBatchedEx（[PR #183](https://gitcode.com/cann/ops-blas/pull/183)）的分组扩展版本，支持通过设备侧指针数组指定多组矩阵，在单次调用中完成所有分组内批次的 GEMM 计算。相比 BatchedEx，新增 **分组（group）** 语义：每个分组可拥有独立的 (m, n, k, transa, transb, alpha, beta, lda, ldb, ldc, dtype) 参数及 groupSize： index = 0 for group in [0, groupCount): for problem in [0, groupSize[group]): C[index] = alpha[group] * op(A[index]) * op(B[index]) + beta[group] * C[index] index++ **支持的数据类型组合（7 种）**： | 序号 | Atype | Btype | Ctype | computeType | |------|-------|-------|-------|-------------| | 1 | FP16 | FP16 | FP16 | COMPUTE_16F | | 2 | FP16 | FP16 | FP16 | COMPUTE_32F | | 3 | BF16 | BF16 | BF16 | COMPUTE_32F | | 4 | FP8_E4M3 | FP8_E4M3 | FP16 | COMPUTE_32F | | 5 | FP8_E5M2 | FP8_E5M2 | FP16 | COMPUTE_32F | | 6 | FP8_E4M3 | FP8_E5M2 | FP16 | COMPUTE_32F | | 7 | FP8_E5M2 | FP8_E4M3 | FP16 | COMPUTE_32F | **核心设计**： - **双 Kernel 执行模型**：Cube kernel 完成全部矩阵乘法（FP32 workspace 输出）；AIV epilogue kernel 完成 alpha/beta 融合、类型转换和 k==0 路径 - **全设备侧计算**：Host 仅负责参数校验、tiling 和元数据上传，不逐组发起 GEMM，也不进行数值计算 - **多核切分**：problemCount × M × N 三维联合切分，grid-stride 循环多核分配 - **复用 BatchedEx 低阶 API**：Nd2Nz → LoadData → Mmad → Fixpipe 流水线 - Kernel 侧直接从 GM 读取设备侧指针数组，无需 Host 侧 D2H 拷贝 **交付物清单**： | 类别 | 文件 | |------|------| | Host 代码 | blas/gemm/gemm_grouped_batched_ex/arch35/gemm_grouped_batched_ex_host.cpp | | Cube Kernel | blas/gemm/gemm_grouped_batched_ex/arch35/gemm_grouped_batched_ex_kernel.cpp | | Epilogue Kernel | blas/gemm/gemm_grouped_batched_ex/arch35/gemm_grouped_batched_ex_epilogue_kernel.cpp | | TilingData | blas/gemm/gemm_grouped_batched_ex/arch35/gemm_grouped_batched_ex_tiling_data.h | | 公共工具 | blas/common/helper/device_utils.h、dtype_cast.h、math_utils.h、sync_utils.h | | 测试代码 | test/gemm_grouped_batched_ex/ | | 算子文档 | blas/gemm/gemm_grouped_batched_ex/README.md | | 接口声明 | include/cann_ops_blas.h | ## 关联的Issue [#188](https://gitcode.com/cann/ops-blas/issues/188) ## 测试 ### ST 精度测试 - **测试框架**：CSV 驱动 GTest + 手工 L1E 参数校验 - **CSV 用例**：62 条（8 条 L0 + 25 条 L1 + 10 条 L2 + 19 条 FP8） - **L1E 参数校验**：29 条手工用例，覆盖空指针、负维度、非法枚举、lda/ldb/ldc 约束、不支持 dtype、computeType 不匹配等 - **测试覆盖**：全部 7 种合法 dtype 组合、单组/多组、NN/NT/TN/TT 转置组合、groupCount/groupSize 变化、alpha/beta 各种取值、边界 shape、FP8 混合输入、非法参数校验 ### 编译测试 bash bash build.sh --ops=gemm_grouped_batched_ex --soc=ascend950 --run ## 文档更新 - 新增 blas/gemm/gemm_grouped_batched_ex/README.md：接口说明、Kernel 执行模型、支持范围、编译/测试命令 ## 类型标签 - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!189 | 4 天前 | |
统一所有算子获取内核数量方法为通过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架构下的算子进行了测试，全部通过   ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!163 | 3 天前 | |
Feat: 新增面向arch35的aclblasGemmEx接口 Co-authored-by: developer<developer@opencode.ai> Co-authored-by: yang-di52<yangdi52@huawei.com> # message auto-generated for no-merge-commit merge: !169 merge aclblasGemmEx into master Feat: 新增面向arch35的aclblasGemmEx接口 Created-by: yang-di52 Commit-by: yang-di52;developer Merged-by: cann-robot Description: ## 描述 新增 aclblasGemmEx 接口，实现通用矩阵乘法扩展（GEMM Ex），面向 arch35 (DAV_3510) 架构。 **支持的数据类型组合**： | 计算类型 | Scale 类型 | A/B 类型 | C 类型 | |---------|-----------|---------|--------| | ACLBLAS_COMPUTE_16F | ACL_FLOAT16 | ACL_FLOAT16 | ACL_FLOAT16 | | ACLBLAS_COMPUTE_32F | ACL_FLOAT | ACL_BF16 | ACL_BF16 | | ACLBLAS_COMPUTE_32F | ACL_FLOAT | ACL_FLOAT16 | ACL_FLOAT16 | | ACLBLAS_COMPUTE_32F | ACL_FLOAT | ACL_FLOAT | ACL_FLOAT | | ACLBLAS_COMPUTE_32F (FP8) | ACL_FLOAT | ACL_FLOAT8_E4M3FN | ACL_FLOAT16 | | ACLBLAS_COMPUTE_32F (FP8) | ACL_FLOAT | ACL_FLOAT8_E5M2 | ACL_FLOAT16 | ## 关联的Issue [#159](https://gitcode.com/cann/ops-blas/issues/159) ## 测试  ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-blas!169 | 5 天前 | |
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 | 3 天前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 12 天前 | ||
| 4 天前 | ||
| 27 天前 | ||
| 12 天前 | ||
| 12 天前 | ||
| 4 天前 | ||
| 3 天前 | ||
| 5 天前 | ||
| 3 天前 |