master

cann-robotfeat(nrm2): 新增 arch35 aclblasSnrm2接口

cd498f8a创建于 8 小时前历史提交

文件	最后提交记录	最后更新时间
cann_ops_blas.h	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	8 小时前
cann_ops_blasLt.h	feat: 新增 arch35 矩阵变换接口 aclblasLtMatrixTransform Co-authored-by: tangpingchuan<tangpingchuan@huawei.com> # message auto-generated for no-merge-commit merge: !150 merge aclblasLtMatrixTransform into master feat: 新增 arch35 矩阵变换接口 aclblasLtMatrixTransform Created-by: pingchuantang Commit-by: tangpingchuan Merged-by: cann-robot Description: ## 描述 新增 aclBLASLt 矩阵变换接口 aclblasLtMatrixTransform，实现 C = alpha·op(A) + beta·op(B) 的矩阵转置、缩放、加法及内存布局/数据类型转换，目标芯片 Ascend950，目标架构 arch35（DAV_3510）。 **新增接口**（include/cann_ops_blasLt.h）： - aclblasLtMatrixTransformDescCreate / aclblasLtMatrixTransformDescDestroy / aclblasLtMatrixTransformDescSetAttribute / aclblasLtMatrixTransformDescGetAttribute - aclblasLtMatrixTransform 主计算接口 - 新增类型 aclblasLtMatrixTransformDesc_t、aclblasLtMatrixTransformDescAttribute_t（SCALE_TYPE / POINTER_MODE / TRANSA / TRANSB） - aclblasLtOrder_t 扩展 COL32 / COL4_4R2_8C / COL32_2R_4R4 三种布局 **支持规格**： - dtype（A/B/C 独立）：FP32 / FP16 / BF16 / INT8 / INT32 / FP8_E4M3FN / FP8_E5M2 / FP4_E2M1；scaleType 按浮点 FP32 / 整数 INT32 / FP8 FP32 / FP4 BF16 分路径 - op：N / T（C 等价 T）；order：COL / ROW / COL32 / COL4_4R2_8C / COL32_2R_4R4 - 复杂量化布局（COL4_4R2_8C / COL32_2R_4R4）仅支持量化路径 dtype（INT8/INT32/FP8/FP4）；FP4 packed ld 经 ceil 字节换算，奇数维合法 - 不支持 batchCount>1、跨 scaleType 路径转换、FP64/复数/HiF8 等 **实现**：host 侧 tiling（blasLt/matrixtransform/arch35/matrix_transform_host.cpp 及 tiling data / perm table 头文件），arch35 kernel（matrix_transform_kernel.cpp），ACL 入口在 blasLt/aclblasLt.cpp 完成参数校验与任务下发。 ## 关联的Issue 关联Issue #147 ：https://gitcode.com/cann/ops-blas/issues/147 ## 测试 CSV 驱动 159 个用例 + 4 个复杂布局 anchor 用例（test/blasLtMatrixTransform/），覆盖全 dtype/order/op 组合、空矩阵、异常入参；整数路径位精确（MARE=MERE=0），浮点/FP8/FP4 按 MARE/MERE 单标杆标准比对。 ## 文档更新 新增 blasLt/matrixtransform/README.md（计算语义、支持规格、精度标准、接口说明）。 ## 类型标签 - [x] 新特性 See merge request: cann/ops-blas!150	14 天前
cann_ops_blas_common.h	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	2 天前