| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
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 小时前 | |
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 天前 | |
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 天前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 8 小时前 | ||
| 14 天前 | ||
| 2 天前 |