Star122
148
代码
代码
Issues30
Pull Requests28
Actions
项目讨论
Wiki
分析
Star122
148
已合并
feat(cdgmm): 迁移 ColwiseMul 为行主序 Cdgmm；新增 extensions/ 目录迁移 ComplexMatDot #314
2301_82146483创建于 7月28日
feat(cdgmm): 迁移 ColwiseMul 为行主序 Cdgmm；新增 extensions/ 目录迁移 ComplexMatDot #314
已合并
2301_82146483创建于 7月28日
从2301_82146483/ops-blas-yzyfeat/cdgmm-migration合入到cann/ops-blasmaster
共 30 个文件变更+1669-812
MCMakeLists.txt+3-0
@@ -158,6 +158,7 @@ set(ALL_BLAS_SRC_FILES "")
158set(OPS_BLAS ops_blas)158set(OPS_BLAS ops_blas)
159 159 
160add_subdirectory(blas)160add_subdirectory(blas)
161+add_subdirectory(extensions)
161 162 
162project(${OPS_BLAS} LANGUAGES ASC CXX)163project(${OPS_BLAS} LANGUAGES ASC CXX)
163 164 
@@ -175,6 +176,8 @@ target_include_directories(${OPS_BLAS} PRIVATE
175 ${ASCEND_CANN_PACKAGE_PATH}/include/op_common/176 ${ASCEND_CANN_PACKAGE_PATH}/include/op_common/
176 ./include177 ./include
177 ${CMAKE_CURRENT_LIST_DIR}/blas178 ${CMAKE_CURRENT_LIST_DIR}/blas
179+ ${CMAKE_CURRENT_LIST_DIR}/extensions
180+ ${CMAKE_CURRENT_LIST_DIR}/extensions/complexmatdot/arch22
178 "${_OPS_BLAS_ASC_ROOT}"181 "${_OPS_BLAS_ASC_ROOT}"
179 "${_OPS_BLAS_ASC_ROOT}/include"182 "${_OPS_BLAS_ASC_ROOT}/include"
180 "${_OPS_BLAS_ASC_ROOT}/include/utils"183 "${_OPS_BLAS_ASC_ROOT}/include/utils"
Mblas/CMakeLists.txt+6-4
@@ -68,12 +68,14 @@ if(NOT ENABLE_BLAS_TRMM)
68 STATUS68 STATUS
69 "Skipping blas trmm/strmm (requires asc-devkit >= 9.1, current ${ASC_DEVKIT_MAJOR}.${ASC_DEVKIT_MINOR})")69 "Skipping blas trmm/strmm (requires asc-devkit >= 9.1, current ${ASC_DEVKIT_MAJOR}.${ASC_DEVKIT_MINOR})")
70endif()70endif()
71-if(NOT ENABLE_BLAS_DGMM)71+# Only arch35 sdgmm uses tensor_api; arch22 cdgmm does not depend on it.
72- list(FILTER OP_SRC_FILES EXCLUDE REGEX "/dgmm/")72+# Filter only sdgmm sources so cdgmm (arch22) is not excluded.
73- list(FILTER ARCH_SRC_FILES EXCLUDE REGEX "/dgmm/")73+if(NOT ENABLE_BLAS_SDGMM)
74+ list(FILTER OP_SRC_FILES EXCLUDE REGEX "/dgmm/arch35/sdgmm")
75+ list(FILTER ARCH_SRC_FILES EXCLUDE REGEX "/dgmm/arch35/sdgmm")
74 message(76 message(
75 STATUS77 STATUS
76- "Skipping blas dgmm/sdgmm (requires asc-devkit >= 9.1, current ${ASC_DEVKIT_MAJOR}.${ASC_DEVKIT_MINOR})")78+ "Skipping blas dgmm/sdgmm arch35 (requires asc-devkit >= 9.1, current ${ASC_DEVKIT_MAJOR}.${ASC_DEVKIT_MINOR})")
77endif()79endif()
78if(NOT ENABLE_BLAS_GEMM_BATCHED)80if(NOT ENABLE_BLAS_GEMM_BATCHED)
79 list(FILTER OP_SRC_FILES EXCLUDE REGEX "/gemm_batched/")81 list(FILTER OP_SRC_FILES EXCLUDE REGEX "/gemm_batched/")
Dblas/colwise_mul/README.md+0-48
@@ -1,48 +0,0 @@
1-# ColwiseMul算子
2- 
3-## 算子概述
4- 
5-ColwiseMul（列方向复数乘法）算子实现了复数向量与矩阵每行的逐元素乘法运算，是 BLAS 基础线性代数库中的扩展算子之一。该算子针对复数运算特性进行了优化，高效完成向量与矩阵的列方向乘法操作。
6- 
7-数学表达式：
8- 
9-```
10-result[i, j] = vec[i] * mat[i, j]
11-```
12- 
13-包含以下接口：
14- 
15-| 接口名 | 功能简述 |
16-|--------|---------|
17-| aclblasColwiseMul | 复数向量与矩阵每行逐元素乘法 |
18- 
19-## 算子执行接口
20- 
21-### aclblasColwiseMul
22- 
23-#### 产品支持情况
24- 
25-- Ascend 950PR / Ascend 950DT：不支持
26-- Atlas A3 训练系列产品 / Atlas A3 推理系列产品：支持
27-- Atlas A2 训练系列产品 / Atlas A2 推理系列产品：支持
28- 
29-#### 函数原型
30- 
31-```cpp
32-int aclblasColwiseMul(const float *mat, const float *vec, float *result, const int64_t m, const int64_t n, void *stream)
33-```
34- 
35-#### 参数说明
36- 
37-| 参数名 | 输入/输出 | 参数类型 | 说明 |
38-|--------|----------|---------|------|
39-| mat | 输入 | const float*（FP32） | 复数矩阵，维度为 m x n，存储为 2*m*n 个 float，Device 内存 |
40-| vec | 输入 | const float*（FP32） | 复数向量，包含 m 个元素，存储为 2*m 个 float，Device 内存 |
41-| result | 输出 | float*（FP32） | 复数矩阵，维度为 m x n，存储为 2*m*n 个 float，Device 内存 |
42-| m | 输入 | int64_t | 矩阵的行数，Host 内存 |
43-| n | 输入 | int64_t | 矩阵的列数，Host 内存 |
44-| stream | 输入 | void* | 执行流，Host 内存 |
45- 
46-#### 约束说明
47- 
48-无
Dblas/colwise_mul/arch22/colwise_mul_host.cpp+0-158
@@ -1,158 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use the License for the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9- */
10- 
11-/* !
12- * \file colwise_mul_host.cpp
13- * \brief Host side implementation for colwise_mul operator
14- */
15- 
16-#include <cstdint>
17-#include <cstring>
18-#include <iostream>
19-#include <vector>
20-#include <algorithm>
21-#include "acl/acl.h"
22-#include "cann_ops_blas.h"
23-#include "common/helper/aclblas_handle_internal.h"
24-#include "common/helper/host_utils.h"
25- 
26-void colwise_mul_kernel_do(uint8_t* mat, uint8_t* vec, uint8_t* aug, uint8_t* result,
27- uint8_t* workSpace, uint8_t* tilingGm,
28- uint32_t numBlocks, void *stream);
29- 
30-constexpr uint32_t DEFAULT_VECTOR_NUM = 40;
31-constexpr uint32_t DEFAULT_CUBE_NUM = 20;
32- 
33-constexpr uint32_t COMPLEX_NUM = 2;
34-constexpr uint32_t FP32_BYTE_SIZE = 4;
35-constexpr uint32_t MAX_DATA_COUNT = 32 * 1024 / sizeof(float);
36- 
37-// Tiling data structure
38-struct ColwiseMulTilingData {
39- uint32_t m;
40- uint32_t n;
41- 
42- uint32_t startOffset[40];
43- uint32_t calRowNum[40];
44-};
45- 
46-// Tiling calculation
47-ColwiseMulTilingData CalColwiseMulTilingData(uint32_t m, uint32_t n, uint32_t vecCoreNum)
48-{
49- ColwiseMulTilingData tilingData;
50- memset(&tilingData, 0, sizeof(ColwiseMulTilingData));
51- 
52- if (vecCoreNum == 0) {
53- vecCoreNum = 1;
54- }
55- vecCoreNum = vecCoreNum > DEFAULT_VECTOR_NUM ? DEFAULT_VECTOR_NUM : vecCoreNum;
56- 
57- // n is already in float elements (for complex, it's 2 * num_complex)
58- uint32_t rowNumEachCore = m / vecCoreNum;
59- uint32_t remainRowNum = m % vecCoreNum;
60- 
61- if (rowNumEachCore == 0) {
62- for (uint32_t i = 0; i < remainRowNum; i++) {
63- tilingData.calRowNum[i] = 1;
64- tilingData.startOffset[i] = n * i; // each row has n FP32 elements
65- }
66- } else {
67- uint32_t currOffset = 0;
68- uint32_t currRowNum;
69- for (uint32_t i = 0; i < vecCoreNum; i++) {
70- if (i < remainRowNum) {
71- currRowNum = rowNumEachCore + 1;
72- } else {
73- currRowNum = rowNumEachCore;
74- }
75- tilingData.calRowNum[i] = currRowNum;
76- tilingData.startOffset[i] = currOffset;
77- currOffset += currRowNum * n;
78- }
79- }
80- 
81- tilingData.m = m; // num of rows
82- tilingData.n = n; // num of FP32 elements each row
83- 
84- return tilingData;
85-}
86- 
87-uint32_t* CreateAugColwiseMul()
88-{
89- uint32_t complexCount = MAX_DATA_COUNT / COMPLEX_NUM;
90- 
91- uint32_t* augData = nullptr;
92- 
93- augData = new uint32_t[MAX_DATA_COUNT];
94- 
95- for (uint32_t i = 0; i < complexCount; i++) {
96- augData[COMPLEX_NUM * i] = FP32_BYTE_SIZE * i;
97- augData[COMPLEX_NUM * i + 1] = FP32_BYTE_SIZE * (i + complexCount);
98- }
99- return augData;
100-}
101- 
102-aclblasStatus_t aclblasColwiseMul(
103- aclblasHandle_t handle, const int64_t m, const int64_t n, uint8_t* mat, uint8_t* vec, uint8_t* result)
104-{
105- auto* h = handle;
106- aclrtStream useStream = h->stream;
107- 
108- uint32_t nFloats = n * 2;
109- uint32_t numBlocks = 8;
110- 
111- ColwiseMulTilingData tiling = CalColwiseMulTilingData(m, nFloats, numBlocks);
112- 
113- uint32_t* aug = CreateAugColwiseMul();
114- size_t augByteSize = MAX_DATA_COUNT * sizeof(uint32_t);
115- size_t workspaceSize = 1024;
116- 
117- uint8_t* augDevice = nullptr;
118- uint8_t* workspaceDevice = nullptr;
119- uint8_t* tilingDevice = nullptr;
120- 
121- aclError aclRet = aclrtMalloc((void**)&augDevice, augByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
122- CHECK_RET(
123- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet);
124- return ACLBLAS_STATUS_ALLOC_FAILED);
125- 
126- aclRet = aclrtMalloc((void**)&workspaceDevice, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
127- CHECK_RET(
128- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet); aclrtFree(augDevice);
129- return ACLBLAS_STATUS_ALLOC_FAILED);
130- 
131- aclRet = aclrtMalloc((void**)&tilingDevice, sizeof(ColwiseMulTilingData), ACL_MEM_MALLOC_HUGE_FIRST);
132- CHECK_RET(
133- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet); aclrtFree(workspaceDevice);
134- aclrtFree(augDevice); return ACLBLAS_STATUS_ALLOC_FAILED);
135- 
136- aclRet = aclrtMemcpy(augDevice, augByteSize, aug, augByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
137- CHECK_RET(
138- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet); aclrtFree(tilingDevice);
139- aclrtFree(workspaceDevice); aclrtFree(augDevice); return ACLBLAS_STATUS_INTERNAL_ERROR);
140- 
141- aclRet = aclrtMemcpy(
142- tilingDevice, sizeof(ColwiseMulTilingData), &tiling, sizeof(ColwiseMulTilingData), ACL_MEMCPY_HOST_TO_DEVICE);
143- CHECK_RET(
144- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet); aclrtFree(tilingDevice);
145- aclrtFree(workspaceDevice); aclrtFree(augDevice); return ACLBLAS_STATUS_INTERNAL_ERROR);
146- 
147- colwise_mul_kernel_do(mat, vec, augDevice, result, workspaceDevice, tilingDevice, numBlocks, useStream);
148- aclRet = aclrtSynchronizeStream(useStream);
149- CHECK_RET(
150- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet); aclrtFree(tilingDevice);
151- aclrtFree(workspaceDevice); aclrtFree(augDevice); return ACLBLAS_STATUS_INTERNAL_ERROR);
152- 
153- aclrtFree(augDevice);
154- aclrtFree(workspaceDevice);
155- aclrtFree(tilingDevice);
156- 
157- return ACLBLAS_STATUS_SUCCESS;
158-}
Dblas/colwise_mul/arch22/colwise_mul_kernel.cpp+0-263
@@ -1,263 +0,0 @@
1-/**
2-* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-* CANN Open Software License Agreement Version 2.0 (the "License").
5-* Please refer to the License for details. You may not use the License for the License.
6-* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-* See LICENSE in the root of the software repository for the full text of the License.
9-*/
10- 
11-#include "kernel_operator.h"
12-#include "common/helper/kernel_utils.h"
13-#include "common/iterator/iterator.h"
14-#include "common/compute/simd.h"
15- 
16-__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub_uint32(
17- AscendC::LocalTensor<uint32_t> dst,
18- AscendC::GlobalTensor<uint32_t> src,
19- uint32_t len)
20-{
21- uint16_t nBurst = 1;
22- uint32_t lenBurst = len * sizeof(uint32_t);
23- uint8_t leftPaddingNum = 0;
24- uint8_t rightPaddingNum = 0;
25- uint32_t srcGap = 0;
26- uint32_t dstGap = 0;
27- gm_to_ub_align<ArchType::ASCEND_V220, uint32_t>(dst, src,
28- 0, // sid
29- nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
30-}
31- 
32-__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub(
33- AscendC::LocalTensor<float> dst,
34- AscendC::GlobalTensor<float> src,
35- uint32_t len)
36-{
37- uint16_t nBurst = 1;
38- uint32_t lenBurst = len * sizeof(float);
39- uint8_t leftPaddingNum = 0;
40- uint8_t rightPaddingNum = 0;
41- uint32_t srcGap = 0;
42- uint32_t dstGap = 0;
43- gm_to_ub_align<ArchType::ASCEND_V220, float>(dst, src,
44- 0, // sid
45- nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
46-}
47- 
48-__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_ub2gm(
49- AscendC::GlobalTensor<float> dst,
50- AscendC::LocalTensor<float> src,
51- uint32_t len)
52-{
53- uint16_t nBurst = 1;
54- uint32_t lenBurst = len * sizeof(float);
55- uint8_t leftPaddingNum = 0;
56- uint8_t rightPaddingNum = 0;
57- uint32_t srcGap = 0;
58- uint32_t dstGap = 0;
59- ub_to_gm_align<ArchType::ASCEND_V220, float>(dst, src,
60- 0, // sid
61- nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
62-}
63- 
64-__aicore__ __inline__ __attribute__((always_inline)) void colwise_mul_compute_aiv(
65- AscendC::GlobalTensor<float> gm_in,
66- AscendC::GlobalTensor<float> gm_out,
67- AscendC::LocalTensor<float> ub_in,
68- AscendC::LocalTensor<float> ub_out,
69- AscendC::LocalTensor<uint32_t> ub_offset,
70- float s_real, float s_imag, uint32_t copy_len, uint32_t len, uint32_t event_id)
71-{
72- uint32_t repeatTime = (len + 63) / 64; // 3
73- uint32_t computeRepeat = (len / 2 + 63) / 64; // 2
74- 
75- uint32_t real_offset = 0;
76- uint32_t imag_offset = 32 * 1024 / sizeof(float) / 2;
77- 
78- AscendC::LocalTensor<float> ub_out_real = ub_out;
79- AscendC::LocalTensor<float> ub_out_imag = ub_out[imag_offset];
80- 
81- copy_vec_gm2ub(ub_in, gm_in, copy_len);
82- 
83- SET_FLAG(MTE2, V, event_id);
84- WAIT_FLAG(MTE2, V, event_id);
85- 
86- uint32_t mask = 0;
87- uint64_t rsvdCnt = 0;
88- 
89- AscendC::GatherMask<float>(ub_out_real, ub_in, 1, false, mask,
90- {1, static_cast<uint16_t>(repeatTime), 8, 8}, rsvdCnt);
91- 
92- AscendC::GatherMask<float>(ub_out_imag, ub_in, 2, false, mask,
93- {1, static_cast<uint16_t>(repeatTime), 8, 8}, rsvdCnt);
94- 
95- PIPE_BARRIER(V);
96- 
97- // R * R
98- muls_v<ArchType::ASCEND_V220, float>(ub_in, ub_out_real, s_real, computeRepeat, 1, 1, 8, 8);
99- 
100- // R * I
101- muls_v<ArchType::ASCEND_V220, float>(ub_in[imag_offset], ub_out_real, s_imag, computeRepeat, 1, 1, 8, 8);
102- 
103- // I * I
104- muls_v<ArchType::ASCEND_V220, float>(ub_out_real, ub_out_imag, s_imag, computeRepeat, 1, 1, 8, 8);
105- 
106- PIPE_BARRIER(V);
107- // R * R - I * I
108- sub_v<ArchType::ASCEND_V220, float>(ub_in, ub_in, ub_out_real, computeRepeat, 1, 1, 1, 8, 8, 8);
109- 
110- // I * R
111- muls_v<ArchType::ASCEND_V220, float>(ub_out_imag, ub_out_imag, s_real, computeRepeat, 1, 1, 8, 8);
112- 
113- PIPE_BARRIER(V);
114- // R * I + I * R
115- add_v<ArchType::ASCEND_V220, float>(
116- ub_in[imag_offset], ub_out_imag, ub_in[imag_offset], computeRepeat, 1, 1, 1, 8, 8, 8);
117- 
118- PIPE_BARRIER(V);
119- 
120- AscendC::Gather(ub_out, ub_in, ub_offset, (uint32_t)0, repeatTime * 64);
121- PIPE_BARRIER(ALL);
122- 
123- SET_FLAG(V, MTE3, event_id);
124- WAIT_FLAG(V, MTE3, event_id);
125- 
126- copy_vec_ub2gm(gm_out, ub_out, copy_len);
127-}
128- 
129-__aicore__ __inline__ __attribute__((always_inline)) void colwise_mul_aiv(
130- AscendC::GlobalTensor<float> gm_in,
131- AscendC::GlobalTensor<float> gm_vec,
132- AscendC::GlobalTensor<uint32_t> gm_aug,
133- AscendC::GlobalTensor<float> gm_out,
134- uint32_t m, uint32_t cal_num, uint32_t offset, uint32_t row_num)
135-{
136- // ub 192kb
137- AsdopsBuffer<ArchType::ASCEND_V220> buf;
138- AscendC::LocalTensor<float> ub_out_ping = buf.GetBuffer<BufferType::ASCEND_UB, float>(0 * 1024);
139- AscendC::LocalTensor<float> ub_out_pong = buf.GetBuffer<BufferType::ASCEND_UB, float>(32 * 1024);
140- AscendC::LocalTensor<float> ub_in_ping = buf.GetBuffer<BufferType::ASCEND_UB, float>(64 * 1024);
141- AscendC::LocalTensor<float> ub_in_pong = buf.GetBuffer<BufferType::ASCEND_UB, float>(96 * 1024);
142- AscendC::LocalTensor<uint32_t> ub_offset = buf.GetBuffer<BufferType::ASCEND_UB, uint32_t>(128 * 1024);
143- 
144- uint32_t ping_flag = 1;
145- 
146- uint32_t maxDataCount = 32 * 1024 / sizeof(float);
147- 
148- uint32_t repeatTime = cal_num / maxDataCount;
149- uint32_t remainNum = cal_num % maxDataCount;
150- 
151- float s_real;
152- float s_imag;
153- 
154- // prepare offset
155- copy_vec_gm2ub_uint32(ub_offset, gm_aug, maxDataCount);
156- 
157- SET_FLAG(MTE2, V, EVENT_ID0);
158- WAIT_FLAG(MTE2, V, EVENT_ID0);
159- SET_FLAG(MTE2, V, EVENT_ID1);
160- WAIT_FLAG(MTE2, V, EVENT_ID1);
161- 
162- uint32_t curr_offset = offset;
163- if (cal_num == 0) {
164- return;
165- }
166- uint32_t curr_row = curr_offset / cal_num;
167- 
168- for (uint32_t row_idx = 0; row_idx < row_num; row_idx++) {
169- curr_offset = offset + cal_num * row_idx;
170- 
171- s_real = gm_vec.GetValue((curr_row + row_idx) * 2);
172- s_imag = gm_vec.GetValue((curr_row + row_idx) * 2 + 1);
173- 
174- SET_FLAG(S, V, EVENT_ID0);
175- WAIT_FLAG(S, V, EVENT_ID0);
176- SET_FLAG(S, V, EVENT_ID1);
177- WAIT_FLAG(S, V, EVENT_ID1);
178- 
179- if (repeatTime > 0) {
180- SET_FLAG(MTE3, MTE2, EVENT_ID0);
181- SET_FLAG(MTE3, MTE2, EVENT_ID1);
182- for (uint32_t i = 0; i < repeatTime; i++) {
183- auto ub_in = ping_flag ? ub_in_ping : ub_in_pong;
184- auto ub_out = ping_flag ? ub_out_ping : ub_out_pong;
185- 
186- auto event_id = ping_flag ? EVENT_ID0 : EVENT_ID1;
187- 
188- WAIT_FLAG(MTE3, MTE2, event_id);
189- 
190- colwise_mul_compute_aiv(gm_in[curr_offset], gm_out[curr_offset], ub_in, ub_out, ub_offset, s_real,
191- s_imag, maxDataCount, maxDataCount, event_id);
192- 
193- SET_FLAG(MTE3, MTE2, event_id);
194- 
195- curr_offset += maxDataCount;
196- ping_flag = 1 - ping_flag;
197- }
198- WAIT_FLAG(MTE3, MTE2, EVENT_ID0);
199- WAIT_FLAG(MTE3, MTE2, EVENT_ID1);
200- }
201- 
202- if (remainNum > 0) {
203- SET_FLAG(MTE3, MTE2, EVENT_ID0);
204- SET_FLAG(MTE3, MTE2, EVENT_ID1);
205- auto ub_in = ping_flag ? ub_in_ping : ub_in_pong;
206- auto ub_out = ping_flag ? ub_out_ping : ub_out_pong;
207- auto event_id = ping_flag ? EVENT_ID0 : EVENT_ID1;
208- WAIT_FLAG(MTE3, MTE2, event_id);
209- 
210- colwise_mul_compute_aiv(gm_in[curr_offset], gm_out[curr_offset], ub_in, ub_out, ub_offset, s_real, s_imag,
211- remainNum, remainNum, event_id);
212- 
213- SET_FLAG(MTE3, MTE2, event_id);
214- ping_flag = 1 - ping_flag;
215- curr_offset += remainNum;
216- WAIT_FLAG(MTE3, MTE2, EVENT_ID0);
217- WAIT_FLAG(MTE3, MTE2, EVENT_ID1);
218- }
219- }
220- PIPE_BARRIER(ALL);
221-}
222- 
223-__global__ __aicore__ __vector__ void colwise_mul(GM_ADDR mat, GM_ADDR vec,
224- GM_ADDR aug, GM_ADDR result,
225- GM_ADDR workSpace, GM_ADDR tilingGm)
226-{
227- AscendC::SetAtomicNone();
228- AscendC::SetMaskNorm();
229- // AscendC::SetVectorMask<float>((uint64_t)-1, (uint64_t)-1);
230- 
231- auto core_idx = AscendC::GetBlockIdx();
232- 
233- auto tiling_buf = reinterpret_cast<__gm__ uint8_t *>(tilingGm);
234- 
235- uint32_t m = (*(__gm__ uint32_t *)((__gm__ uint8_t *)tiling_buf)); // num of float elements
236- uint32_t n = (*(__gm__ uint32_t *)((__gm__ uint8_t *)tiling_buf + 4)); // num of float elements
237- 
238- uint32_t offset = (*(__gm__ uint32_t *)((__gm__ uint8_t *)tiling_buf + 8 + 4 * core_idx));
239- uint32_t row_num = (*(__gm__ uint32_t *)((__gm__ uint8_t *)tiling_buf + 8 + 40 * 4 + 4 * core_idx));
240- 
241- if (row_num <= 0)
242- return;
243- 
244- AscendC::GlobalTensor<float> mat_tensor;
245- AscendC::GlobalTensor<float> vec_tensor;
246- AscendC::GlobalTensor<uint32_t> aug_tensor;
247- AscendC::GlobalTensor<float> result_tensor;
248- 
249- mat_tensor.SetGlobalBuffer(reinterpret_cast<__gm__ float *>(mat));
250- vec_tensor.SetGlobalBuffer(reinterpret_cast<__gm__ float *>(vec));
251- aug_tensor.SetGlobalBuffer(reinterpret_cast<__gm__ uint32_t *>(aug));
252- result_tensor.SetGlobalBuffer(reinterpret_cast<__gm__ float *>(result));
253- colwise_mul_aiv(mat_tensor, vec_tensor, aug_tensor, result_tensor, m, n, offset, row_num);
254-}
255- 
256-// Wrapper function for host to call
257-void colwise_mul_kernel_do(GM_ADDR mat, GM_ADDR vec, GM_ADDR aug, GM_ADDR result,
258- GM_ADDR workSpace, GM_ADDR tilingGm,
259- uint32_t numBlocks, void *stream)
260-{
261- colwise_mul<<<numBlocks, nullptr, stream>>>(mat, vec, aug, result, workSpace, tilingGm);
262-}
263- 
Mblas/dgmm/README.md+52-2
@@ -2,7 +2,7 @@
2 2 
3## 算子概述3## 算子概述
4 4 
5-Dgmm（Diagonal Matrix-Matrix Multiplication）算子实现了对角矩阵与普通矩阵的乘法运算，使用一个向量构造对角矩阵，按行（LEFT）或按列（RIGHT）对输入矩阵进行缩放，核心运算为逐元素乘法。矩阵按列主序（BLAS 约定）存储。5+Dgmm（Diagonal Matrix-Matrix Multiplication）算子实现了对角矩阵与普通矩阵的乘法运算，使用一个向量构造对角矩阵，按行（LEFT）或按列（RIGHT）对输入矩阵进行缩放，核心运算为逐元素乘法。
6 6 
7数学表达式：7数学表达式：
8 8 
@@ -15,7 +15,8 @@ RIGHT 模式： C = A * diag(x)， C[i,j] = A[i,j] * x[j] （x 长度为 n
15 15 
16| 接口名 | 功能简述 |16| 接口名 | 功能简述 |
17|--------|---------|17|--------|---------|
18-| aclblasSdgmm | 单精度浮点对角矩阵乘法 |18+| aclblasSdgmm | 单精度实数对角矩阵乘法 |
19+| aclblasCdgmm | 单精度复数对角矩阵乘法，当前仅支持 LEFT 模式，RIGHT 模式暂未实现 |
19 20 
20## 算子执行接口21## 算子执行接口
21 22 
@@ -237,3 +238,52 @@ C[0,2] = 14.000000
237C[1,2] = 24.000000238C[1,2] = 24.000000
238C[2,2] = 36.000000239C[2,2] = 36.000000
239```240```
241+ 
242+### aclblasCdgmm
243+ 
244+#### 产品支持情况
245+ 
246+- Ascend 950PR / Ascend 950DT：不支持
247+- Atlas A3 训练系列产品 / Atlas A3 推理系列产品：支持
248+- Atlas A2 训练系列产品 / Atlas A2 推理系列产品：支持
249+ 
250+#### 函数原型
251+ 
252+```cpp
253+aclblasStatus_t aclblasCdgmm(aclblasHandle_t handle, aclblasSideMode_t mode, int m, int n, const aclblasComplex *A, int lda, const aclblasComplex *x, int incx, aclblasComplex *C, int ldc)
254+```
255+ 
256+#### 参数说明
257+ 
258+| 参数名 | 输入/输出 | 参数类型 | 说明 |
259+|--------|----------|---------|------|
260+| handle | 输入 | aclblasHandle_t | ops-blas 库上下文句柄，携带 stream，Host 内存 |
261+| mode | 输入 | aclblasSideMode_t | 缩放模式：当前仅支持 ACLBLAS_SIDE_LEFT（C = diag(x) * A，x 长度为 m）；ACLBLAS_SIDE_RIGHT 返回 ACLBLAS_STATUS_NOT_SUPPORTED，Host 内存 |
262+| m | 输入 | int | 矩阵 A/C 的行数，m >= 0，Host 内存 |
263+| n | 输入 | int | 矩阵 A/C 的列数，n >= 0，Host 内存 |
264+| A | 输入 | const aclblasComplex* | 输入矩阵，行主序存储，维度 m×n，Device 内存 |
265+| lda | 输入 | int | 矩阵 A 的行跨度（row-major），lda >= max(1, n)，Host 内存 |
266+| x | 输入 | const aclblasComplex* | 对角向量，LEFT 模式下长度为 m，Device 内存 |
267+| incx | 输入 | int | x 中相邻元素的步长，incx != 0，可为负数，Host 内存 |
268+| C | 输出 | aclblasComplex* | 输出矩阵，行主序存储，维度 m×n，Device 内存 |
269+| ldc | 输入 | int | 矩阵 C 的行跨度（row-major），ldc >= max(1, n)，Host 内存 |
270+ 
271+#### 约束说明
272+ 
273+- 当前仅支持 `ACLBLAS_SIDE_LEFT` 模式（C = diag(x) * A，C[i,j] = x[i] * A[i,j]）
274+- `ACLBLAS_SIDE_RIGHT` 是合法枚举，但当前返回 `ACLBLAS_STATUS_NOT_SUPPORTED`
275+- A 和 C 使用行主序（row-major）存储，与 cuBLAS `cublasCdgmm` 的列主序语义存在差异，不能把 cuBLAS 的列主序输入未经转换直接传入
276+- `lda` 和 `ldc` 表示行主序的行跨度（相邻两行起始位置之间的复数元素数），约束为不小于 n
277+- handle 不能为 nullptr，否则返回 `ACLBLAS_STATUS_HANDLE_IS_NULLPTR`
278+- mode 必须为 `ACLBLAS_SIDE_LEFT` 或 `ACLBLAS_SIDE_RIGHT`，否则返回 `ACLBLAS_STATUS_INVALID_ENUM`
279+- m >= 0, n >= 0，否则返回 `ACLBLAS_STATUS_INVALID_VALUE`
280+- incx != 0（可为负数，表示反向访问 x），否则返回 `ACLBLAS_STATUS_INVALID_VALUE`
281+- lda >= max(1, n)，否则返回 `ACLBLAS_STATUS_INVALID_VALUE`
282+- ldc >= max(1, n)，否则返回 `ACLBLAS_STATUS_INVALID_VALUE`
283+- 当 m > 0 且 n > 0 时，A、x、C 不能为 nullptr，否则返回 `ACLBLAS_STATUS_INVALID_VALUE`
284+- m == 0 或 n == 0 时为 no-op，直接返回 `ACLBLAS_STATUS_SUCCESS`
285+- 支持正负 `incx`
286+- 支持独立 `lda`、`ldc`（lda != ldc）
287+- C 的 padding 区域（C[i*ldc+n ... i*ldc+ldc-1]）不被修改
288+- 支持 `A == C && lda == ldc` 的原地执行；`A == C && lda != ldc` 返回 `ACLBLAS_STATUS_INVALID_VALUE`
289+- arch22 支持；arch35 暂不支持 Cdgmm
Ablas/dgmm/arch22/cdgmm_host.cpp+292-0
@@ -0,0 +1,292 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cdgmm_host.cpp
13+ * \brief Host-side implementation of aclblasCdgmm (arch22).
14+ * Complex diagonal matrix-matrix multiplication: C = diag(x) * A
15+ * (mode=LEFT). Row-major storage. RIGHT mode returns
16+ * ACLBLAS_STATUS_NOT_SUPPORTED.
17+ */
18+ 
19+#include <cstdint>
20+#include <vector>
21+#include <algorithm>
22+#include "acl/acl.h"
23+#include "log/log.h"
24+#include "cann_ops_blas.h"
25+#include "cdgmm_kernel.h"
26+#include "common/helper/aclblas_handle_internal.h"
27+#include "common/helper/host_utils.h"
28+ 
29+constexpr uint32_t CDGMM_DEFAULT_VECTOR_NUM = 40;
30+constexpr uint32_t CDGMM_COMPLEX_NUM = 2;
31+constexpr uint32_t CDGMM_FP32_BYTE_SIZE = 4;
32+constexpr uint32_t CDGMM_MAX_DATA_COUNT = 32 * 1024 / sizeof(float);
33+constexpr size_t CDGMM_WORKSPACE_SIZE = 1024;
34+ 
35+// ==========================================================================
36+// RAII wrapper for device buffers — guarantees aclrtFree on all paths.
37+// ==========================================================================
38+class AclDeviceBuffer {
39+public:
40+ AclDeviceBuffer() = default;
41+ ~AclDeviceBuffer()
42+ {
43+ if (data_ != nullptr) {
44+ aclrtFree(data_);
45+ }
46+ }
47+ 
48+ aclError Allocate(size_t size)
49+ {
50+ data_ = nullptr;
51+ return aclrtMalloc(reinterpret_cast<void**>(&data_), size, ACL_MEM_MALLOC_HUGE_FIRST);
52+ }
53+ 
54+ uint8_t* Get() const { return data_; }
55+ 
56+ AclDeviceBuffer(const AclDeviceBuffer&) = delete;
57+ AclDeviceBuffer& operator=(const AclDeviceBuffer&) = delete;
58+ 
59+private:
60+ uint8_t* data_ = nullptr;
61+};
62+ 
63+// ==========================================================================
64+// Parameter validation
65+// ==========================================================================
66+static aclblasStatus_t ValidateCdgmmParams(
67+ aclblasSideMode_t mode, int m, int n,
68+ const aclblasComplex* A, int lda, const aclblasComplex* x, int incx,
69+ aclblasComplex* C, int ldc)
70+{
71+ if (mode != ACLBLAS_SIDE_LEFT && mode != ACLBLAS_SIDE_RIGHT) {
72+ OP_LOGE("aclblasCdgmm", "mode must be SIDE_LEFT(141) or SIDE_RIGHT(142), got %d",
73+ static_cast<int>(mode));
74+ return ACLBLAS_STATUS_INVALID_ENUM;
75+ }
76+ if (mode == ACLBLAS_SIDE_RIGHT) {
77+ OP_LOGE("aclblasCdgmm",
78+ "ACLBLAS_SIDE_RIGHT is not supported in the current row-major implementation");
79+ return ACLBLAS_STATUS_NOT_SUPPORTED;
80+ }
81+ if (m < 0) {
82+ OP_LOGE("aclblasCdgmm", "m must be >= 0, got %d", m);
83+ return ACLBLAS_STATUS_INVALID_VALUE;
84+ }
85+ if (n < 0) {
86+ OP_LOGE("aclblasCdgmm", "n must be >= 0, got %d", n);
87+ return ACLBLAS_STATUS_INVALID_VALUE;
88+ }
89+ if (incx == 0) {
90+ OP_LOGE("aclblasCdgmm", "incx must not be zero");
91+ return ACLBLAS_STATUS_INVALID_VALUE;
92+ }
93+ if (lda < std::max(1, n)) {
94+ OP_LOGE("aclblasCdgmm", "lda must be >= max(1, n), got lda=%d, n=%d", lda, n);
95+ return ACLBLAS_STATUS_INVALID_VALUE;
96+ }
97+ if (ldc < std::max(1, n)) {
98+ OP_LOGE("aclblasCdgmm", "ldc must be >= max(1, n), got ldc=%d, n=%d", ldc, n);
99+ return ACLBLAS_STATUS_INVALID_VALUE;
100+ }
101+ if (m > 0 && n > 0 && (A == nullptr || x == nullptr || C == nullptr)) {
102+ OP_LOGE("aclblasCdgmm", "A/x/C must not be nullptr when m>0 and n>0");
103+ return ACLBLAS_STATUS_INVALID_VALUE;
104+ }
105+ if (A == C && lda != ldc) {
106+ OP_LOGE("aclblasCdgmm", "in-place execution (A==C) requires lda==ldc, got lda=%d, ldc=%d", lda, ldc);
107+ return ACLBLAS_STATUS_INVALID_VALUE;
108+ }
109+ return ACLBLAS_STATUS_SUCCESS;
110+}
111+ 
112+// ==========================================================================
113+// Unified numBlocks calculation — shared by tiling and host launch.
114+// ==========================================================================
115+static uint32_t CalcCdgmmNumBlocks(uint32_t m, uint32_t aivCoreNum)
116+{
117+ uint32_t clamped = std::max(1U, std::min(aivCoreNum, CDGMM_DEFAULT_VECTOR_NUM));
118+ return std::min(clamped, m);
119+}
120+ 
121+// ==========================================================================
122+// Tiling: row-split decomposition (Task 1 + Task 2)
123+// ==========================================================================
124+static CdgmmTilingData CalCdgmmTilingData(
125+ uint32_t mode, uint32_t m, uint32_t n, int32_t incx,
126+ uint32_t lda, uint32_t ldc, uint32_t aivCoreNum)
127+{
128+ CdgmmTilingData tiling{};
129+ 
130+ tiling.mode = mode;
131+ tiling.m = m;
132+ tiling.n = n;
133+ tiling.incx = incx;
134+ tiling.lda = lda;
135+ tiling.ldc = ldc;
136+ 
137+ // Guard: empty matrix — return zero-initialized tiling without division.
138+ if (m == 0) {
139+ return tiling;
140+ }
141+ 
142+ const uint32_t numBlocks = CalcCdgmmNumBlocks(m, aivCoreNum);
143+ if (numBlocks == 0) {
144+ return tiling;
145+ }
146+ const uint32_t rowsPerCore = m / numBlocks;
147+ const uint32_t remainder = m % numBlocks;
148+ 
149+ uint32_t currRow = 0;
150+ for (uint32_t i = 0; i < numBlocks; i++) {
151+ uint32_t currCount = (i < remainder) ? rowsPerCore + 1 : rowsPerCore;
152+ tiling.startRow[i] = currRow;
153+ tiling.rowCount[i] = currCount;
154+ currRow += currCount;
155+ }
156+ 
157+ return tiling;
158+}
159+ 
160+// ==========================================================================
161+// Aug offset table builder
162+// ==========================================================================
163+static std::vector<uint32_t> CreateAugCdgmm()
164+{
165+ uint32_t complexCount = CDGMM_MAX_DATA_COUNT / CDGMM_COMPLEX_NUM;
166+ std::vector<uint32_t> aug(CDGMM_MAX_DATA_COUNT);
167+ 
168+ for (uint32_t i = 0; i < complexCount; i++) {
169+ aug[CDGMM_COMPLEX_NUM * i] = CDGMM_FP32_BYTE_SIZE * i;
170+ aug[CDGMM_COMPLEX_NUM * i + 1] = CDGMM_FP32_BYTE_SIZE * (i + complexCount);
171+ }
172+ return aug;
173+}
174+ 
175+// ==========================================================================
176+// Build launch configuration (tiling, numBlocks, aug)
177+// ==========================================================================
178+struct CdgmmLaunchConfig {
179+ CdgmmTilingData tiling{};
180+ uint32_t numBlocks = 0;
181+ std::vector<uint32_t> aug;
182+};
183+ 
184+static aclblasStatus_t BuildCdgmmLaunchConfig(
185+ aclblasSideMode_t mode, int m, int n, int lda, int incx, int ldc,
186+ uint32_t aivCoreNum, CdgmmLaunchConfig& config)
187+{
188+ uint32_t modeNorm = (mode == ACLBLAS_SIDE_LEFT) ? CDGMM_MODE_LEFT : CDGMM_MODE_RIGHT;
189+ 
190+ config.tiling = CalCdgmmTilingData(
191+ modeNorm, static_cast<uint32_t>(m), static_cast<uint32_t>(n), incx,
192+ static_cast<uint32_t>(lda), static_cast<uint32_t>(ldc), aivCoreNum);
193+ 
194+ config.numBlocks = CalcCdgmmNumBlocks(static_cast<uint32_t>(m), aivCoreNum);
195+ config.aug = CreateAugCdgmm();
196+ 
197+ return ACLBLAS_STATUS_SUCCESS;
198+}
199+ 
200+// ==========================================================================
201+// Launch: allocate device buffers, copy, kernel, sync (Task 3)
202+// ==========================================================================
203+static aclblasStatus_t LaunchCdgmm(
204+ aclrtStream stream,
205+ const aclblasComplex* A, const aclblasComplex* x, aclblasComplex* C,
206+ const CdgmmLaunchConfig& config)
207+{
208+ AclDeviceBuffer augDevice;
209+ AclDeviceBuffer workspaceDevice;
210+ AclDeviceBuffer tilingDevice;
211+ 
212+ size_t augByteSize = config.aug.size() * sizeof(uint32_t);
213+ 
214+ aclError aclRet = augDevice.Allocate(augByteSize);
215+ if (aclRet != ACL_SUCCESS) {
216+ OP_LOGE("aclblasCdgmm", "aclrtMalloc aug failed. ERROR: %d", aclRet);
217+ return ACLBLAS_STATUS_ALLOC_FAILED;
218+ }
219+ 
220+ aclRet = workspaceDevice.Allocate(CDGMM_WORKSPACE_SIZE);
221+ if (aclRet != ACL_SUCCESS) {
222+ OP_LOGE("aclblasCdgmm", "aclrtMalloc workspace failed. ERROR: %d", aclRet);
223+ return ACLBLAS_STATUS_ALLOC_FAILED;
224+ }
225+ 
226+ aclRet = tilingDevice.Allocate(sizeof(CdgmmTilingData));
227+ if (aclRet != ACL_SUCCESS) {
228+ OP_LOGE("aclblasCdgmm", "aclrtMalloc tiling failed. ERROR: %d", aclRet);
229+ return ACLBLAS_STATUS_ALLOC_FAILED;
230+ }
231+ 
232+ aclRet = aclrtMemcpy(augDevice.Get(), augByteSize, config.aug.data(),
233+ augByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
234+ if (aclRet != ACL_SUCCESS) {
235+ OP_LOGE("aclblasCdgmm", "aclrtMemcpy aug failed. ERROR: %d", aclRet);
236+ return ACLBLAS_STATUS_INTERNAL_ERROR;
237+ }
238+ 
239+ aclRet = aclrtMemcpy(tilingDevice.Get(), sizeof(CdgmmTilingData), &config.tiling,
240+ sizeof(CdgmmTilingData), ACL_MEMCPY_HOST_TO_DEVICE);
241+ if (aclRet != ACL_SUCCESS) {
242+ OP_LOGE("aclblasCdgmm", "aclrtMemcpy tiling failed. ERROR: %d", aclRet);
243+ return ACLBLAS_STATUS_INTERNAL_ERROR;
244+ }
245+ 
246+ cdgmm_kernel_do(reinterpret_cast<GM_ADDR>(const_cast<aclblasComplex*>(A)),
247+ reinterpret_cast<GM_ADDR>(const_cast<aclblasComplex*>(x)),
248+ reinterpret_cast<GM_ADDR>(C),
249+ augDevice.Get(), workspaceDevice.Get(), tilingDevice.Get(),
250+ config.numBlocks, stream);
251+ 
252+ aclRet = aclrtSynchronizeStream(stream);
253+ if (aclRet != ACL_SUCCESS) {
254+ OP_LOGE("aclblasCdgmm", "aclrtSynchronizeStream failed. ERROR: %d", aclRet);
255+ return ACLBLAS_STATUS_INTERNAL_ERROR;
256+ }
257+ 
258+ return ACLBLAS_STATUS_SUCCESS;
259+}
260+ 
261+// ==========================================================================
262+// Public API — orchestrates validate → build config → launch
263+// ==========================================================================
264+extern "C" aclblasStatus_t aclblasCdgmm(
265+ aclblasHandle_t handle, aclblasSideMode_t mode,
266+ int m, int n, const aclblasComplex* A, int lda,
267+ const aclblasComplex* x, int incx, aclblasComplex* C, int ldc)
268+{
269+ if (handle == nullptr) {
270+ OP_LOGE("aclblasCdgmm", "handle is nullptr");
271+ return ACLBLAS_STATUS_HANDLE_IS_NULLPTR;
272+ }
273+ 
274+ aclblasStatus_t st = ValidateCdgmmParams(mode, m, n, A, lda, x, incx, C, ldc);
275+ if (st != ACLBLAS_STATUS_SUCCESS) {
276+ return st;
277+ }
278+ if (m == 0 || n == 0) {
279+ return ACLBLAS_STATUS_SUCCESS;
280+ }
281+ 
282+ uint32_t aivCoreNum = GetAivCoreCount();
283+ if (aivCoreNum == 0) {
284+ OP_LOGE("aclblasCdgmm", "GetAivCoreCount failed");
285+ return ACLBLAS_STATUS_INTERNAL_ERROR;
286+ }
287+ 
288+ CdgmmLaunchConfig config;
289+ BuildCdgmmLaunchConfig(mode, m, n, lda, incx, ldc, aivCoreNum, config);
290+ 
291+ return LaunchCdgmm(handle->stream, A, x, C, config);
292+}
Ablas/dgmm/arch22/cdgmm_kernel.cpp+374-0
@@ -0,0 +1,374 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cdgmm_kernel.cpp
13+ * \brief Device-side kernel for aclblasCdgmm (arch22).
14+ * Row-major storage. Complex elements are stored as interleaved
15+ * float pairs (real, imag).
16+ *
17+ * mode=LEFT: C[i,j] = x[i] * A[i,j]
18+ * Each row is contiguous and multiplied by a single complex
19+ * scalar x[i]. Reuses the original scalar-vector complex
20+ * multiply path (GatherMask deinterleave, muls_v, Gather
21+ * re-interleave) with ping-pong UB buffers.
22+ */
23+ 
24+#include "kernel_operator.h"
25+#include "common/helper/kernel_utils.h"
26+#include "common/iterator/iterator.h"
27+#include "common/compute/simd.h"
28+#include "cdgmm_kernel.h"
29+ 
30+// ==========================================================================
31+// UB copy helpers (kept from original colwise_mul, renamed)
32+// ==========================================================================
33+ 
34+__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub_uint32(
35+ AscendC::LocalTensor<uint32_t> dst,
36+ AscendC::GlobalTensor<uint32_t> src,
37+ uint32_t len)
38+{
39+ uint16_t nBurst = 1;
40+ uint32_t lenBurst = len * sizeof(uint32_t);
41+ uint8_t leftPaddingNum = 0;
42+ uint8_t rightPaddingNum = 0;
43+ uint32_t srcGap = 0;
44+ uint32_t dstGap = 0;
45+ gm_to_ub_align<ArchType::ASCEND_V220, uint32_t>(dst, src,
46+ 0, nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
47+}
48+ 
49+__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub(
50+ AscendC::LocalTensor<float> dst,
51+ AscendC::GlobalTensor<float> src,
52+ uint32_t len)
53+{
54+ uint16_t nBurst = 1;
55+ uint32_t lenBurst = len * sizeof(float);
56+ uint8_t leftPaddingNum = 0;
57+ uint8_t rightPaddingNum = 0;
58+ uint32_t srcGap = 0;
59+ uint32_t dstGap = 0;
60+ gm_to_ub_align<ArchType::ASCEND_V220, float>(dst, src,
61+ 0, nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
62+}
63+ 
64+__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_ub2gm(
65+ AscendC::GlobalTensor<float> dst,
66+ AscendC::LocalTensor<float> src,
67+ uint32_t len)
68+{
69+ uint16_t nBurst = 1;
70+ uint32_t lenBurst = len * sizeof(float);
71+ uint8_t leftPaddingNum = 0;
72+ uint8_t rightPaddingNum = 0;
73+ uint32_t srcGap = 0;
74+ uint32_t dstGap = 0;
75+ ub_to_gm_align<ArchType::ASCEND_V220, float>(dst, src,
76+ 0, nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
77+}
78+ 
79+// ==========================================================================
80+// Core compute: scalar complex multiply on a contiguous block.
81+// (Renamed from colwise_mul_compute_aiv, logic preserved.)
82+//
83+// Multiplies a contiguous block of interleaved complex floats by a complex
84+// scalar (s_real, s_imag):
85+// out = a * (s_real + s_imag*i)
86+// ==========================================================================
87+__aicore__ __inline__ __attribute__((always_inline)) void cdgmm_compute_aiv(
88+ AscendC::GlobalTensor<float> gm_in,
89+ AscendC::GlobalTensor<float> gm_out,
90+ AscendC::LocalTensor<float> ub_in,
91+ AscendC::LocalTensor<float> ub_out,
92+ AscendC::LocalTensor<uint32_t> ub_offset,
93+ float s_real, float s_imag, uint32_t copy_len, uint32_t event_id)
94+{
95+ uint32_t len = copy_len;
96+ uint32_t repeatTime = (len + 63) / 64;
97+ uint32_t computeRepeat = (len / 2 + 63) / 64;
98+ 
99+ uint32_t real_offset = 0;
100+ uint32_t imag_offset = 32 * 1024 / sizeof(float) / 2;
101+ 
102+ AscendC::LocalTensor<float> ub_out_real = ub_out;
103+ AscendC::LocalTensor<float> ub_out_imag = ub_out[imag_offset];
104+ 
105+ copy_vec_gm2ub(ub_in, gm_in, copy_len);
106+ 
107+ SET_FLAG(MTE2, V, event_id);
108+ WAIT_FLAG(MTE2, V, event_id);
109+ 
110+ uint32_t mask = 0;
111+ uint64_t rsvdCnt = 0;
112+ 
113+ AscendC::GatherMask<float>(ub_out_real, ub_in, 1, false, mask,
114+ {1, static_cast<uint16_t>(repeatTime), 8, 8}, rsvdCnt);
115+ 
116+ AscendC::GatherMask<float>(ub_out_imag, ub_in, 2, false, mask,
117+ {1, static_cast<uint16_t>(repeatTime), 8, 8}, rsvdCnt);
118+ 
119+ PIPE_BARRIER(V);
120+ 
121+ // R * R
122+ muls_v<ArchType::ASCEND_V220, float>(ub_in, ub_out_real, s_real, computeRepeat, 1, 1, 8, 8);
123+ 
124+ // R * I
125+ muls_v<ArchType::ASCEND_V220, float>(ub_in[imag_offset], ub_out_real, s_imag, computeRepeat, 1, 1, 8, 8);
126+ 
127+ // I * I
128+ muls_v<ArchType::ASCEND_V220, float>(ub_out_real, ub_out_imag, s_imag, computeRepeat, 1, 1, 8, 8);
129+ 
130+ PIPE_BARRIER(V);
131+ // R * R - I * I
132+ sub_v<ArchType::ASCEND_V220, float>(ub_in, ub_in, ub_out_real, computeRepeat, 1, 1, 1, 8, 8, 8);
133+ 
134+ // I * R
135+ muls_v<ArchType::ASCEND_V220, float>(ub_out_imag, ub_out_imag, s_real, computeRepeat, 1, 1, 8, 8);
136+ 
137+ PIPE_BARRIER(V);
138+ // R * I + I * R
139+ add_v<ArchType::ASCEND_V220, float>(
140+ ub_in[imag_offset], ub_out_imag, ub_in[imag_offset], computeRepeat, 1, 1, 1, 8, 8, 8);
141+ 
142+ PIPE_BARRIER(V);
143+ 
144+ AscendC::Gather(ub_out, ub_in, ub_offset, (uint32_t)0, repeatTime * 64);
145+ PIPE_BARRIER(ALL);
146+ 
147+ SET_FLAG(V, MTE3, event_id);
148+ WAIT_FLAG(V, MTE3, event_id);
149+ 
150+ copy_vec_ub2gm(gm_out, ub_out, copy_len);
151+}
152+ 
153+// ==========================================================================
154+// Pure computation: x index for given row with signed incx (Task 4)
155+// ==========================================================================
156+__aicore__ __inline__ __attribute__((always_inline)) int64_t CalcCdgmmXIndex(
157+ uint32_t row, uint32_t m, int32_t incx)
158+{
159+ if (incx >= 0) {
160+ return static_cast<int64_t>(row) * static_cast<int64_t>(incx);
161+ }
162+ int64_t absIncx = -static_cast<int64_t>(incx);
163+ return (static_cast<int64_t>(m) - 1 - static_cast<int64_t>(row)) * absIncx;
164+}
165+ 
166+// ==========================================================================
167+// UB buffer context for LEFT-mode ping-pong processing
168+// ==========================================================================
169+struct CdgmmLeftUbCtx {
170+ AscendC::LocalTensor<float> ubOutPing;
171+ AscendC::LocalTensor<float> ubOutPong;
172+ AscendC::LocalTensor<float> ubInPing;
173+ AscendC::LocalTensor<float> ubInPong;
174+ AscendC::LocalTensor<uint32_t> ubOffset;
175+ uint32_t pingFlag;
176+ uint32_t maxDataCount;
177+ uint32_t repeatTime;
178+ uint32_t remainNum;
179+};
180+ 
181+// ==========================================================================
182+// Process full chunks for a single row (Task 4)
183+// ==========================================================================
184+__aicore__ __inline__ __attribute__((always_inline)) void ProcessCdgmmFullChunks(
185+ AscendC::GlobalTensor<float> gm_a,
186+ AscendC::GlobalTensor<float> gm_c,
187+ AscendC::GlobalTensor<uint32_t> gm_aug,
188+ float s_real, float s_imag,
189+ uint64_t aBase, uint64_t cBase,
190+ CdgmmLeftUbCtx& ctx)
191+{
192+ if (ctx.repeatTime == 0) {
193+ return;
194+ }
195+ 
196+ uint64_t currOffset = 0;
197+ SET_FLAG(MTE3, MTE2, EVENT_ID0);
198+ SET_FLAG(MTE3, MTE2, EVENT_ID1);
199+ 
200+ for (uint32_t i = 0; i < ctx.repeatTime; i++) {
201+ auto ubIn = ctx.pingFlag ? ctx.ubInPing : ctx.ubInPong;
202+ auto ubOut = ctx.pingFlag ? ctx.ubOutPing : ctx.ubOutPong;
203+ auto eventId = ctx.pingFlag ? EVENT_ID0 : EVENT_ID1;
204+ 
205+ WAIT_FLAG(MTE3, MTE2, eventId);
206+ 
207+ cdgmm_compute_aiv(gm_a[aBase + currOffset], gm_c[cBase + currOffset],
208+ ubIn, ubOut, ctx.ubOffset, s_real, s_imag,
209+ ctx.maxDataCount, eventId);
210+ 
211+ SET_FLAG(MTE3, MTE2, eventId);
212+ currOffset += ctx.maxDataCount;
213+ ctx.pingFlag = 1 - ctx.pingFlag;
214+ }
215+ WAIT_FLAG(MTE3, MTE2, EVENT_ID0);
216+ WAIT_FLAG(MTE3, MTE2, EVENT_ID1);
217+}
218+ 
219+// ==========================================================================
220+// Process tail chunk for a single row (Task 4)
221+// ==========================================================================
222+__aicore__ __inline__ __attribute__((always_inline)) void ProcessCdgmmTailChunk(
223+ AscendC::GlobalTensor<float> gm_a,
224+ AscendC::GlobalTensor<float> gm_c,
225+ AscendC::GlobalTensor<uint32_t> gm_aug,
226+ float s_real, float s_imag,
227+ uint64_t aBase, uint64_t cBase, uint64_t prevOffset,
228+ CdgmmLeftUbCtx& ctx)
229+{
230+ if (ctx.remainNum == 0) {
231+ return;
232+ }
233+ 
234+ uint64_t currOffset = prevOffset;
235+ SET_FLAG(MTE3, MTE2, EVENT_ID0);
236+ SET_FLAG(MTE3, MTE2, EVENT_ID1);
237+ 
238+ auto ubIn = ctx.pingFlag ? ctx.ubInPing : ctx.ubInPong;
239+ auto ubOut = ctx.pingFlag ? ctx.ubOutPing : ctx.ubOutPong;
240+ auto eventId = ctx.pingFlag ? EVENT_ID0 : EVENT_ID1;
241+ WAIT_FLAG(MTE3, MTE2, eventId);
242+ 
243+ cdgmm_compute_aiv(gm_a[aBase + currOffset], gm_c[cBase + currOffset],
244+ ubIn, ubOut, ctx.ubOffset, s_real, s_imag,
245+ ctx.remainNum, eventId);
246+ 
247+ SET_FLAG(MTE3, MTE2, eventId);
248+ ctx.pingFlag = 1 - ctx.pingFlag;
249+ WAIT_FLAG(MTE3, MTE2, EVENT_ID0);
250+ WAIT_FLAG(MTE3, MTE2, EVENT_ID1);
251+}
252+ 
253+// ==========================================================================
254+// LEFT mode row processing (ping-pong, migrated from colwise_mul_aiv)
255+//
256+// Row-major: row i starts at A_float[2*i*lda] and C_float[2*i*ldc].
257+// Each row has n complex elements = 2*n float elements.
258+// x[i] is read using incx stride: xIndex = (incx>0) ? i*incx : (m-1-i)*|incx|
259+// ==========================================================================
260+__aicore__ __inline__ __attribute__((always_inline)) void cdgmm_left_aiv(
261+ AscendC::GlobalTensor<float> gm_a,
262+ AscendC::GlobalTensor<float> gm_x,
263+ AscendC::GlobalTensor<uint32_t> gm_aug,
264+ AscendC::GlobalTensor<float> gm_c,
265+ uint32_t m, uint32_t n, int32_t incx,
266+ uint32_t lda, uint32_t ldc,
267+ uint32_t startRow, uint32_t rowCount)
268+{
269+ // ub 192kb
270+ AsdopsBuffer<ArchType::ASCEND_V220> buf;
271+ 
272+ CdgmmLeftUbCtx ctx;
273+ ctx.ubOutPing = buf.GetBuffer<BufferType::ASCEND_UB, float>(0 * 1024);
274+ ctx.ubOutPong = buf.GetBuffer<BufferType::ASCEND_UB, float>(32 * 1024);
275+ ctx.ubInPing = buf.GetBuffer<BufferType::ASCEND_UB, float>(64 * 1024);
276+ ctx.ubInPong = buf.GetBuffer<BufferType::ASCEND_UB, float>(96 * 1024);
277+ ctx.ubOffset = buf.GetBuffer<BufferType::ASCEND_UB, uint32_t>(128 * 1024);
278+ ctx.pingFlag = 1;
279+ ctx.maxDataCount = 32 * 1024 / sizeof(float);
280+ 
281+ uint32_t rowFloats = 2 * n;
282+ ctx.repeatTime = rowFloats / ctx.maxDataCount;
283+ ctx.remainNum = rowFloats % ctx.maxDataCount;
284+ 
285+ // prepare offset
286+ copy_vec_gm2ub_uint32(ctx.ubOffset, gm_aug, ctx.maxDataCount);
287+ 
288+ SET_FLAG(MTE2, V, EVENT_ID0);
289+ WAIT_FLAG(MTE2, V, EVENT_ID0);
290+ SET_FLAG(MTE2, V, EVENT_ID1);
291+ WAIT_FLAG(MTE2, V, EVENT_ID1);
292+ 
293+ if (rowCount == 0) {
294+ return;
295+ }
296+ 
297+ for (uint32_t localRow = 0; localRow < rowCount; localRow++) {
298+ uint32_t row = startRow + localRow;
299+ 
300+ uint64_t aFloatOffset = 2ULL * static_cast<uint64_t>(row) * lda;
301+ uint64_t cFloatOffset = 2ULL * static_cast<uint64_t>(row) * ldc;
302+ 
303+ int64_t xIndex = CalcCdgmmXIndex(row, m, incx);
304+ 
305+ float s_real = gm_x.GetValue(2ULL * static_cast<uint64_t>(xIndex));
306+ float s_imag = gm_x.GetValue(2ULL * static_cast<uint64_t>(xIndex) + 1);
307+ 
308+ SET_FLAG(S, V, EVENT_ID0);
309+ WAIT_FLAG(S, V, EVENT_ID0);
310+ SET_FLAG(S, V, EVENT_ID1);
311+ WAIT_FLAG(S, V, EVENT_ID1);
312+ 
313+ ProcessCdgmmFullChunks(gm_a, gm_c, gm_aug, s_real, s_imag,
314+ aFloatOffset, cFloatOffset, ctx);
315+ 
316+ uint64_t prevOffset = static_cast<uint64_t>(ctx.repeatTime) * ctx.maxDataCount;
317+ ProcessCdgmmTailChunk(gm_a, gm_c, gm_aug, s_real, s_imag,
318+ aFloatOffset, cFloatOffset, prevOffset, ctx);
319+ }
320+ PIPE_BARRIER(ALL);
321+}
322+ 
323+// ==========================================================================
324+// Kernel entry point
325+// ==========================================================================
326+__global__ __aicore__ __vector__ void cdgmm(GM_ADDR A, GM_ADDR x,
327+ GM_ADDR C, GM_ADDR aug,
328+ GM_ADDR workSpace, GM_ADDR tilingGm)
329+{
330+ AscendC::SetAtomicNone();
331+ AscendC::SetMaskNorm();
332+ 
333+ auto core_idx = AscendC::GetBlockIdx();
334+ 
335+ auto* tiling = reinterpret_cast<__gm__ CdgmmTilingData*>(tilingGm);
336+ 
337+ uint32_t mode = tiling->mode;
338+ uint32_t m = tiling->m;
339+ uint32_t n = tiling->n;
340+ int32_t incx = tiling->incx;
341+ uint32_t lda = tiling->lda;
342+ uint32_t ldc = tiling->ldc;
343+ 
344+ uint32_t startRow = tiling->startRow[core_idx];
345+ uint32_t rowCount = tiling->rowCount[core_idx];
346+ 
347+ if (rowCount == 0) {
348+ return;
349+ }
350+ 
351+ AscendC::GlobalTensor<float> a_tensor;
352+ AscendC::GlobalTensor<float> x_tensor;
353+ AscendC::GlobalTensor<uint32_t> aug_tensor;
354+ AscendC::GlobalTensor<float> c_tensor;
355+ 
356+ a_tensor.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(A));
357+ x_tensor.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(x));
358+ aug_tensor.SetGlobalBuffer(reinterpret_cast<__gm__ uint32_t*>(aug));
359+ c_tensor.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(C));
360+ 
361+ // Only LEFT is supported; RIGHT is rejected by Host before kernel launch.
362+ if (mode == CDGMM_MODE_LEFT) {
363+ cdgmm_left_aiv(a_tensor, x_tensor, aug_tensor, c_tensor,
364+ m, n, incx, lda, ldc, startRow, rowCount);
365+ }
366+}
367+ 
368+// Wrapper function for host to call
369+void cdgmm_kernel_do(GM_ADDR A, GM_ADDR x, GM_ADDR C,
370+ GM_ADDR aug, GM_ADDR workSpace, GM_ADDR tilingGm,
371+ uint32_t numBlocks, void* stream)
372+{
373+ cdgmm<<<numBlocks, nullptr, stream>>>(A, x, C, aug, workSpace, tilingGm);
374+}
Ablas/dgmm/arch22/cdgmm_kernel.h+42-0
@@ -0,0 +1,42 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cdgmm_kernel.h
13+ * \brief Declaration of the kernel launcher for aclblasCdgmm (arch22).
14+ * Shared by host.cpp and kernel.cpp.
15+ *
16+ * Row-major storage. Only LEFT mode is implemented (C[i,j] = x[i] * A[i,j]).
17+ */
18+ 
19+#pragma once
20+ 
21+#include <cstdint>
22+#include "cdgmm_tiling_data.h"
23+ 
24+#ifndef GM_ADDR
25+#define GM_ADDR uint8_t*
26+#endif
27+ 
28+/*!
29+ * \brief Kernel launcher: synchronously launches the Cdgmm kernel.
30+ *
31+ * \param A GM address of matrix A (row-major complex, m x n, stride lda)
32+ * \param x GM address of vector x (interleaved complex)
33+ * \param C GM address of output matrix C (row-major complex, m x n, stride ldc)
34+ * \param aug GM address of gather offset table (pre-computed by Host)
35+ * \param workSpace GM address of workspace (currently unused, reserved)
36+ * \param tilingGm GM address of tiling data
37+ * \param numBlocks block count for the <<<>>> launch
38+ * \param stream aclrtStream handle
39+ */
40+void cdgmm_kernel_do(GM_ADDR A, GM_ADDR x, GM_ADDR C,
41+ GM_ADDR aug, GM_ADDR workSpace, GM_ADDR tilingGm,
42+ uint32_t numBlocks, void* stream);
Ablas/dgmm/arch22/cdgmm_tiling_data.h+51-0
@@ -0,0 +1,51 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cdgmm_tiling_data.h
13+ * \brief Tiling data structure for aclblasCdgmm (arch22).
14+ * Shared by host side and kernel side to avoid layout mismatch.
15+ *
16+ * Row-major storage: A and C are m x n row-major complex matrices.
17+ * Complex elements are stored as interleaved float pairs (real, imag).
18+ * Only LEFT mode is implemented; RIGHT is rejected by the Host.
19+ */
20+ 
21+#pragma once
22+ 
23+#include <cstdint>
24+ 
25+static constexpr uint32_t CDGMM_MODE_LEFT = 0;
26+static constexpr uint32_t CDGMM_MODE_RIGHT = 1;
27+ 
28+static constexpr uint32_t CDGMM_MAX_CORES = 40;
29+ 
30+/*!
31+ * \brief Tiling data for Cdgmm.
32+ *
33+ * Row-split decomposition: each core handles a contiguous range of
34+ * rows [startRow[i], startRow[i] + rowCount[i]). The Host balances
35+ * rows across at most CDGMM_MAX_CORES cores (capped by aivCoreNum and m).
36+ *
37+ * Matrices A and C are row-major. lda/ldc are row strides (number of
38+ * complex elements between consecutive rows). Complex elements are
39+ * stored as interleaved float pairs (real, imag).
40+ */
41+struct CdgmmTilingData {
42+ uint32_t mode; // normalized: CDGMM_MODE_LEFT or CDGMM_MODE_RIGHT
43+ uint32_t m; // number of rows of matrix A/C
44+ uint32_t n; // number of columns of matrix A/C
45+ int32_t incx; // stride of vector x (may be negative)
46+ uint32_t lda; // row stride of A (row-major, in complex elements)
47+ uint32_t ldc; // row stride of C (row-major, in complex elements)
48+ 
49+ uint32_t startRow[CDGMM_MAX_CORES]; // start row per core
50+ uint32_t rowCount[CDGMM_MAX_CORES]; // row count per core
51+};
Mcmake/asc_devkit_version.cmake+8-8
@@ -12,7 +12,8 @@
12# - MXFP8/MXFP4：blasLt 矩阵乘法，需 asc-devkit >= 9.112# - MXFP8/MXFP4：blasLt 矩阵乘法，需 asc-devkit >= 9.1
13# - TRMM/STRMM：仅 arch35(ascend950) 的 strmm 使用 tensor_api，需 asc-devkit >= 9.1；13# - TRMM/STRMM：仅 arch35(ascend950) 的 strmm 使用 tensor_api，需 asc-devkit >= 9.1；
14# 其他架构的 strmm 不依赖 tensor_api，不受此版本限制14# 其他架构的 strmm 不依赖 tensor_api，不受此版本限制
15-# - DGMM/SDGMM：仅 arch35(ascend950) 的 sdgmm 使用 tensor_api，需 asc-devkit >= 9.115+# - SDGMM：仅 arch35(ascend950) 的 sdgmm 使用 tensor_api，需 asc-devkit >= 9.1；
16+# arch22 的 cdgmm 不依赖 tensor_api，不受此版本限制
16# - GEMM_BATCHED：仅 arch35(ascend950) 的 gemm_batched 使用 tensor_api，需 asc-devkit >= 9.117# - GEMM_BATCHED：仅 arch35(ascend950) 的 gemm_batched 使用 tensor_api，需 asc-devkit >= 9.1
17# - GEMM3M/SGEMM3M：仅 arch35(ascend950) 的 sgemm3m 使用 tensor_api，需 asc-devkit >= 9.118# - GEMM3M/SGEMM3M：仅 arch35(ascend950) 的 sgemm3m 使用 tensor_api，需 asc-devkit >= 9.1
18# - SYRK/SSYRK：仅 arch35(ascend950) 的 ssyrk 使用 tensor_api，需 asc-devkit >= 9.119# - SYRK/SSYRK：仅 arch35(ascend950) 的 ssyrk 使用 tensor_api，需 asc-devkit >= 9.1
@@ -23,7 +24,7 @@ function(ops_blas_detect_asc_devkit_version)
23 set(ASC_DEVKIT_MINOR 0)24 set(ASC_DEVKIT_MINOR 0)
24 set(ENABLE_BLASLT_MXFP8 FALSE)25 set(ENABLE_BLASLT_MXFP8 FALSE)
25 set(ENABLE_BLAS_TRMM TRUE)26 set(ENABLE_BLAS_TRMM TRUE)
26- set(ENABLE_BLAS_DGMM TRUE)27+ set(ENABLE_BLAS_SDGMM TRUE)
27 set(ENABLE_BLAS_GEMM_BATCHED TRUE)28 set(ENABLE_BLAS_GEMM_BATCHED TRUE)
28 set(ENABLE_BLAS_GEMM3M TRUE)29 set(ENABLE_BLAS_GEMM3M TRUE)
29 set(ENABLE_BLAS_SYRK TRUE)30 set(ENABLE_BLAS_SYRK TRUE)
@@ -42,30 +43,29 @@ function(ops_blas_detect_asc_devkit_version)
42 # arch35 的 strmm/sdgmm/gemm_batched/sgemm3m/ssyrk 使用 tensor_api，需 devkit >= 9.1；其他架构不受限43 # arch35 的 strmm/sdgmm/gemm_batched/sgemm3m/ssyrk 使用 tensor_api，需 devkit >= 9.1；其他架构不受限
43 if("arch35" IN_LIST SOC_ARCH_DIRS AND NOT (ASC_DEVKIT_MAJOR GREATER_EQUAL 9 AND ASC_DEVKIT_MINOR GREATER 0))44 if("arch35" IN_LIST SOC_ARCH_DIRS AND NOT (ASC_DEVKIT_MAJOR GREATER_EQUAL 9 AND ASC_DEVKIT_MINOR GREATER 0))
44 set(ENABLE_BLAS_TRMM FALSE)45 set(ENABLE_BLAS_TRMM FALSE)
45- set(ENABLE_BLAS_DGMM FALSE)46+ set(ENABLE_BLAS_SDGMM FALSE)
46 set(ENABLE_BLAS_GEMM_BATCHED FALSE)47 set(ENABLE_BLAS_GEMM_BATCHED FALSE)
47 set(ENABLE_BLAS_GEMM3M FALSE)48 set(ENABLE_BLAS_GEMM3M FALSE)
48 set(ENABLE_BLAS_SYRK FALSE)49 set(ENABLE_BLAS_SYRK FALSE)
49 endif()50 endif()
50 else()51 else()
51 set(ENABLE_BLAS_TRMM FALSE)52 set(ENABLE_BLAS_TRMM FALSE)
52- set(ENABLE_BLAS_DGMM FALSE)53+ set(ENABLE_BLAS_SDGMM FALSE)
53 set(ENABLE_BLAS_GEMM_BATCHED FALSE)54 set(ENABLE_BLAS_GEMM_BATCHED FALSE)
54 set(ENABLE_BLAS_GEMM3M FALSE)55 set(ENABLE_BLAS_GEMM3M FALSE)
55- set(ENABLE_BLAS_SYRK FALSE)56+ message(WARNING "asc_devkit_version.h not found: ${_header}, MXFP8/TRMM/SDGMM/GEMM_BATCHED/GEMM3M will be skipped")
56- message(WARNING "asc_devkit_version.h not found: ${_header}, MXFP8/TRMM/DGMM/GEMM_BATCHED/GEMM3M/SYRK will be skipped")
57 endif()57 endif()
58 58 
59 set(ASC_DEVKIT_MAJOR ${ASC_DEVKIT_MAJOR} PARENT_SCOPE)59 set(ASC_DEVKIT_MAJOR ${ASC_DEVKIT_MAJOR} PARENT_SCOPE)
60 set(ASC_DEVKIT_MINOR ${ASC_DEVKIT_MINOR} PARENT_SCOPE)60 set(ASC_DEVKIT_MINOR ${ASC_DEVKIT_MINOR} PARENT_SCOPE)
61 set(ENABLE_BLASLT_MXFP8 ${ENABLE_BLASLT_MXFP8} PARENT_SCOPE)61 set(ENABLE_BLASLT_MXFP8 ${ENABLE_BLASLT_MXFP8} PARENT_SCOPE)
62 set(ENABLE_BLAS_TRMM ${ENABLE_BLAS_TRMM} PARENT_SCOPE)62 set(ENABLE_BLAS_TRMM ${ENABLE_BLAS_TRMM} PARENT_SCOPE)
63- set(ENABLE_BLAS_DGMM ${ENABLE_BLAS_DGMM} PARENT_SCOPE)63+ set(ENABLE_BLAS_SDGMM ${ENABLE_BLAS_SDGMM} PARENT_SCOPE)
64 set(ENABLE_BLAS_GEMM_BATCHED ${ENABLE_BLAS_GEMM_BATCHED} PARENT_SCOPE)64 set(ENABLE_BLAS_GEMM_BATCHED ${ENABLE_BLAS_GEMM_BATCHED} PARENT_SCOPE)
65 set(ENABLE_BLAS_GEMM3M ${ENABLE_BLAS_GEMM3M} PARENT_SCOPE)65 set(ENABLE_BLAS_GEMM3M ${ENABLE_BLAS_GEMM3M} PARENT_SCOPE)
66 set(ENABLE_BLAS_SYRK ${ENABLE_BLAS_SYRK} PARENT_SCOPE)66 set(ENABLE_BLAS_SYRK ${ENABLE_BLAS_SYRK} PARENT_SCOPE)
67 message(67 message(
68 STATUS68 STATUS
69- "ASC_DEVKIT_MAJOR=${ASC_DEVKIT_MAJOR}, ASC_DEVKIT_MINOR=${ASC_DEVKIT_MINOR}, ENABLE_BLASLT_MXFP8=${ENABLE_BLASLT_MXFP8}, ENABLE_BLAS_TRMM=${ENABLE_BLAS_TRMM}, ENABLE_BLAS_DGMM=${ENABLE_BLAS_DGMM}, ENABLE_BLAS_GEMM_BATCHED=${ENABLE_BLAS_GEMM_BATCHED}, ENABLE_BLAS_GEMM3M=${ENABLE_BLAS_GEMM3M}, ENABLE_BLAS_SYRK=${ENABLE_BLAS_SYRK}"69+ "ASC_DEVKIT_MAJOR=${ASC_DEVKIT_MAJOR}, ASC_DEVKIT_MINOR=${ASC_DEVKIT_MINOR}, ENABLE_BLASLT_MXFP8=${ENABLE_BLASLT_MXFP8}, ENABLE_BLAS_TRMM=${ENABLE_BLAS_TRMM}, ENABLE_BLAS_SDGMM=${ENABLE_BLAS_SDGMM}, ENABLE_BLAS_GEMM_BATCHED=${ENABLE_BLAS_GEMM_BATCHED}, ENABLE_BLAS_GEMM3M=${ENABLE_BLAS_GEMM3M}"
70 )70 )
71endfunction()71endfunction()
Mcmake/test.cmake+9-4
@@ -17,7 +17,7 @@ function(_ops_blas_has_blas_op_sources op_name out_var)
17 set(${out_var} FALSE PARENT_SCOPE)17 set(${out_var} FALSE PARENT_SCOPE)
18 return()18 return()
19 endif()19 endif()
20- if(NOT ENABLE_BLAS_DGMM AND (op_name STREQUAL "sdgmm" OR op_name STREQUAL "dgmm"))20+ if(NOT ENABLE_BLAS_SDGMM AND (op_name STREQUAL "sdgmm"))
21 set(${out_var} FALSE PARENT_SCOPE)21 set(${out_var} FALSE PARENT_SCOPE)
22 return()22 return()
23 endif()23 endif()
@@ -37,7 +37,9 @@ function(_ops_blas_has_blas_op_sources op_name out_var)
37 foreach(arch_dir ${SOC_ARCH_DIRS})37 foreach(arch_dir ${SOC_ARCH_DIRS})
38 file(GLOB arch_dir_srcs ${CMAKE_SOURCE_DIR}/blas/${op_name}/${arch_dir}/*.cpp38 file(GLOB arch_dir_srcs ${CMAKE_SOURCE_DIR}/blas/${op_name}/${arch_dir}/*.cpp
39 ${CMAKE_SOURCE_DIR}/blas/*/${op_name}/${arch_dir}/*.cpp39 ${CMAKE_SOURCE_DIR}/blas/*/${op_name}/${arch_dir}/*.cpp
40- ${CMAKE_SOURCE_DIR}/blas/*/${arch_dir}/${op_name}_*.cpp)40+ ${CMAKE_SOURCE_DIR}/blas/*/${arch_dir}/${op_name}_*.cpp
41+ ${CMAKE_SOURCE_DIR}/extensions/*/${op_name}/${arch_dir}/*.cpp
42+ ${CMAKE_SOURCE_DIR}/extensions/${op_name}/${arch_dir}/*.cpp)
41 if(arch_dir_srcs)43 if(arch_dir_srcs)
42 set(has_sources TRUE)44 set(has_sources TRUE)
43 break()45 break()
@@ -47,7 +49,8 @@ function(_ops_blas_has_blas_op_sources op_name out_var)
47 if(NOT has_sources AND op_name MATCHES "^[a-zA-Z]")49 if(NOT has_sources AND op_name MATCHES "^[a-zA-Z]")
48 string(SUBSTRING "${op_name}" 1 -1 _stripped)50 string(SUBSTRING "${op_name}" 1 -1 _stripped)
49 foreach(arch_dir ${SOC_ARCH_DIRS})51 foreach(arch_dir ${SOC_ARCH_DIRS})
50- file(GLOB arch_dir_srcs ${CMAKE_SOURCE_DIR}/blas/${_stripped}/${arch_dir}/*.cpp)52+ file(GLOB arch_dir_srcs ${CMAKE_SOURCE_DIR}/blas/${_stripped}/${arch_dir}/*.cpp
53+ ${CMAKE_SOURCE_DIR}/extensions/${_stripped}/${arch_dir}/*.cpp)
51 if(arch_dir_srcs)54 if(arch_dir_srcs)
52 set(has_sources TRUE)55 set(has_sources TRUE)
53 break()56 break()
@@ -56,7 +59,9 @@ function(_ops_blas_has_blas_op_sources op_name out_var)
56 endif()59 endif()
57 if(NOT has_sources)60 if(NOT has_sources)
58 file(GLOB base_srcs ${CMAKE_SOURCE_DIR}/blas/${op_name}/*.cpp61 file(GLOB base_srcs ${CMAKE_SOURCE_DIR}/blas/${op_name}/*.cpp
59- ${CMAKE_SOURCE_DIR}/blas/*/${op_name}/*.cpp)62+ ${CMAKE_SOURCE_DIR}/blas/*/${op_name}/*.cpp
63+ ${CMAKE_SOURCE_DIR}/extensions/${op_name}/*.cpp
64+ ${CMAKE_SOURCE_DIR}/extensions/*/${op_name}/*.cpp)
60 if(base_srcs)65 if(base_srcs)
61 set(has_sources TRUE)66 set(has_sources TRUE)
62 endif()67 endif()
Mdocs/zh/api_list.md+2-2
@@ -812,8 +812,8 @@ Level 1 接口在向量之间进行运算，典型操作包括向量缩放（sca
812| [aclblasSasum](../../blas/asum/README.md) | 实数向量绝对值之和 |812| [aclblasSasum](../../blas/asum/README.md) | 实数向量绝对值之和 |
813| [aclblasSaxpy](../../blas/axpy/README.md) | 单精度浮点 AXPY（y = αx + y） |813| [aclblasSaxpy](../../blas/axpy/README.md) | 单精度浮点 AXPY（y = αx + y） |
814| [aclblasCaxpy](../../blas/axpy/README.md) | 复数 AXPY |814| [aclblasCaxpy](../../blas/axpy/README.md) | 复数 AXPY |
815-| [aclblasColwiseMul](../../blas/colwise_mul/README.md) | 复数向量与矩阵每行逐元素乘法 |815+| [aclblasCdgmm](../../blas/dgmm/README.md) | 行主序单精度复数对角矩阵左乘 |
816-| [aclblasComplexMatDot](../../blas/complex_mat_dot/README.md) | 复数矩阵逐元素点乘 |816+| [aclblasComplexMatDot](../../extensions/complexmatdot/README.md) | 复数矩阵逐元素点乘 |
817| [aclblasScopy](../../blas/copy/README.md) | 单精度浮点向量拷贝 |817| [aclblasScopy](../../blas/copy/README.md) | 单精度浮点向量拷贝 |
818| [aclblasCcopy](../../blas/copy/README.md) | 复数向量拷贝 |818| [aclblasCcopy](../../blas/copy/README.md) | 复数向量拷贝 |
819| [aclblasSdot](../../blas/dot/README.md) | 实数向量点积 |819| [aclblasSdot](../../blas/dot/README.md) | 实数向量点积 |
Aextensions/CMakeLists.txt+64-0
@@ -0,0 +1,64 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use the License for the License.
6+# THIS FILE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+cmake_minimum_required(VERSION 3.16)
12+ 
13+# Collect base source files (non-arch-specific)
14+file(GLOB BASE_SRC_FILES ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp)
15+ 
16+# Auto-collect operator source files from subdirectories
17+set(OP_SRC_FILES "")
18+file(GLOB children RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
19+foreach(child ${children})
20+ if(IS_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/${child})
21+ # Recursively collect non-arch-specific source files
22+ file(GLOB_RECURSE dir_srcs ${CMAKE_CURRENT_SOURCE_DIR}/${child}/*.cpp)
23+ foreach(src_file ${dir_srcs})
24+ set(is_arch_specific FALSE)
25+ foreach(arch_dir ${ARCH_SPECIFIC_DIRS})
26+ if(src_file MATCHES "/${arch_dir}/")
27+ set(is_arch_specific TRUE)
28+ break()
29+ endif()
30+ endforeach()
31+ if(NOT is_arch_specific)
32+ get_filename_component(_src_fname ${src_file} NAME)
33+ get_filename_component(_src_dir ${src_file} DIRECTORY)
34+ set(_has_arch_counterpart FALSE)
35+ foreach(_arch_dir ${SOC_ARCH_DIRS})
36+ if(EXISTS ${_src_dir}/${_arch_dir}/${_src_fname})
37+ set(_has_arch_counterpart TRUE)
38+ break()
39+ endif()
40+ endforeach()
41+ if(NOT _has_arch_counterpart)
42+ list(APPEND OP_SRC_FILES ${src_file})
43+ endif()
44+ endif()
45+ endforeach()
46+ endif()
47+endforeach()
48+ 
49+# Collect arch-specific source files for the target SOC
50+set(ARCH_SRC_FILES "")
51+foreach(arch_dir ${SOC_ARCH_DIRS})
52+ file(GLOB_RECURSE arch_dir_srcs
53+ ${CMAKE_CURRENT_SOURCE_DIR}/*/${arch_dir}/*.cpp
54+ )
55+ list(APPEND ARCH_SRC_FILES ${arch_dir_srcs})
56+endforeach()
57+ 
58+set(ALL_BLAS_SRC_FILES
59+ ${ALL_BLAS_SRC_FILES}
60+ ${BASE_SRC_FILES}
61+ ${OP_SRC_FILES}
62+ ${ARCH_SRC_FILES}
63+ PARENT_SCOPE
64+)
Rblas/complex_mat_dot/README.md→extensions/complexmatdot/README.md+1-1
@@ -45,4 +45,4 @@ int aclblasComplexMatDot(const float *matx, const float *maty, float *result, co
45 45 
46#### 约束说明46#### 约束说明
47 47 
48-无48+无
Rblas/complex_mat_dot/arch22/complex_mat_dot_host.cpp→extensions/complexmatdot/arch22/complexmatdot_host.cpp+9-13
@@ -2,15 +2,15 @@
2 * Copyright (c) 2026 Huawei Technologies Co., Ltd.2 * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.5+ * Please refer to the License for details. You may not use the License for the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11/* !11/* !
12- * \file complex_mat_dot_host.cpp12+ * \file complexmatdot_host.cpp
13- * \brief Complex matrix dot product host implementation13+ * \brief Complex matrix dot product host implementation (extensions)
14 */14 */
15 15 
16#include <cstdint>16#include <cstdint>
@@ -20,12 +20,10 @@
20#include <iterator>20#include <iterator>
21#include "acl/acl.h"21#include "acl/acl.h"
22#include "cann_ops_blas.h"22#include "cann_ops_blas.h"
23+#include "complexmatdot_kernel.h"
23#include "common/helper/aclblas_handle_internal.h"24#include "common/helper/aclblas_handle_internal.h"
24#include "common/helper/host_utils.h"25#include "common/helper/host_utils.h"
25 26 
26-void complex_mat_dot_kernel_do(uint8_t* matx, uint8_t* maty, uint8_t* aug, uint8_t* result,
27- uint8_t* tilingGm, uint32_t numBlocks, void *stream);
28- 
29constexpr uint32_t COMPLEX_NUM = 2;27constexpr uint32_t COMPLEX_NUM = 2;
30 28 
31constexpr uint32_t MAX_DATA_COUNT = 27 * 1024 / sizeof(float);29constexpr uint32_t MAX_DATA_COUNT = 27 * 1024 / sizeof(float);
@@ -79,13 +77,11 @@ static void CalTilingData(ComplexMatDotTilingData& tilingData, uint32_t m, uint3
79 }77 }
80}78}
81 79 
82-uint32_t* CreateAugComplexMatDot()80+static std::vector<uint32_t> CreateAugComplexMatDot()
83{81{
84 uint32_t complexCount = MAX_DATA_COUNT / 2;82 uint32_t complexCount = MAX_DATA_COUNT / 2;
85 83 
86- uint32_t* augData = nullptr;84+ std::vector<uint32_t> augData(MAX_DATA_COUNT);
87- 
88- augData = new uint32_t[MAX_DATA_COUNT];
89 85 
90 for (uint32_t i = 0; i < complexCount; i++) {86 for (uint32_t i = 0; i < complexCount; i++) {
91 augData[MUL_NUM * i] = FOUR_NUM * i;87 augData[MUL_NUM * i] = FOUR_NUM * i;
@@ -105,7 +101,7 @@ aclblasStatus_t aclblasComplexMatDot(
105 101 
106 ComplexMatDotTilingData tiling;102 ComplexMatDotTilingData tiling;
107 CalTilingData(tiling, m, n, numBlocks);103 CalTilingData(tiling, m, n, numBlocks);
108- uint32_t* aug = CreateAugComplexMatDot();104+ std::vector<uint32_t> aug = CreateAugComplexMatDot();
109 105 
110 size_t augByteSize = MAX_DATA_COUNT * sizeof(uint32_t);106 size_t augByteSize = MAX_DATA_COUNT * sizeof(uint32_t);
111 107 
@@ -122,7 +118,7 @@ aclblasStatus_t aclblasComplexMatDot(
122 aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet); aclrtFree(augDevice);118 aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet); aclrtFree(augDevice);
123 return ACLBLAS_STATUS_ALLOC_FAILED);119 return ACLBLAS_STATUS_ALLOC_FAILED);
124 120 
125- aclRet = aclrtMemcpy(augDevice, augByteSize, aug, augByteSize, ACL_MEMCPY_HOST_TO_DEVICE);121+ aclRet = aclrtMemcpy(augDevice, augByteSize, aug.data(), augByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
126 CHECK_RET(122 CHECK_RET(
127 aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet); aclrtFree(tilingDevice);123 aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet); aclrtFree(tilingDevice);
128 aclrtFree(augDevice); return ACLBLAS_STATUS_INTERNAL_ERROR);124 aclrtFree(augDevice); return ACLBLAS_STATUS_INTERNAL_ERROR);
Rblas/complex_mat_dot/arch22/complex_mat_dot_kernel.cpp→extensions/complexmatdot/arch22/complexmatdot_kernel.cpp+3-7
@@ -2,19 +2,17 @@
2* Copyright (c) 2026 Huawei Technologies Co., Ltd.2* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3* This program is free software, you can redistribute it and/or modify it under the terms and conditions of3* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4* CANN Open Software License Agreement Version 2.0 (the "License").4* CANN Open Software License Agreement Version 2.0 (the "License").
5-* Please refer to the License for details. You may not use this file except in compliance with the License.5+* Please refer to the License for details. You may not use the License for the License.
6-* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8* See LICENSE in the root of the software repository for the full text of the License.8* See LICENSE in the root of the software repository for the full text of the License.
9*/9*/
10 10 
11-#ifndef COMPLEX_MAT_DOT_KERNEL_H
12-#define COMPLEX_MAT_DOT_KERNEL_H
13- 
14#include <type_traits>11#include <type_traits>
15#include <cstdint>12#include <cstdint>
16#include "kernel_operator.h"13#include "kernel_operator.h"
17#include "lib/matmul_intf.h"14#include "lib/matmul_intf.h"
15+#include "complexmatdot_kernel.h"
18 16 
19using namespace AscendC;17using namespace AscendC;
20using namespace matmul;18using namespace matmul;
@@ -257,5 +255,3 @@ void complex_mat_dot_kernel_do(GM_ADDR matx, GM_ADDR maty, GM_ADDR aug, GM_ADDR
257{255{
258 complex_mat_dot_kernel<<<numBlocks, nullptr, stream>>>(matx, maty, aug, result, tilingGm);256 complex_mat_dot_kernel<<<numBlocks, nullptr, stream>>>(matx, maty, aug, result, tilingGm);
259}257}
260- 
261-#endif // COMPLEX_MAT_DOT_KERNEL_H
Aextensions/complexmatdot/arch22/complexmatdot_kernel.h+26-0
@@ -0,0 +1,26 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file complexmatdot_kernel.h
13+ * \brief Declaration of the kernel launcher for aclblasComplexMatDot (arch22).
14+ * Shared by host.cpp and kernel.cpp.
15+ */
16+ 
17+#pragma once
18+ 
19+#include <cstdint>
20+ 
21+#ifndef GM_ADDR
22+#define GM_ADDR uint8_t*
23+#endif
24+ 
25+void complex_mat_dot_kernel_do(GM_ADDR matx, GM_ADDR maty, GM_ADDR aug, GM_ADDR result,
26+ GM_ADDR tilingGm, uint32_t numBlocks, void *stream);
Minclude/cann_ops_blas.h+5-3
@@ -198,9 +198,6 @@ aclblasStatus_t aclblasCsrot(
198 aclblasHandle_t handle, const int64_t n, aclblasComplex* x, const int64_t incx, aclblasComplex* y,198 aclblasHandle_t handle, const int64_t n, aclblasComplex* x, const int64_t incx, aclblasComplex* y,
199 const int64_t incy, const float c, const float s);199 const int64_t incy, const float c, const float s);
200 200 
201-aclblasStatus_t aclblasColwiseMul(
202- aclblasHandle_t handle, const int64_t m, const int64_t n, uint8_t* mat, uint8_t* vec, uint8_t* result);
203- 
204aclblasStatus_t aclblasComplexMatDot(201aclblasStatus_t aclblasComplexMatDot(
205 aclblasHandle_t handle, const int64_t m, const int64_t n, aclblasComplex* matx, aclblasComplex* maty,202 aclblasHandle_t handle, const int64_t m, const int64_t n, aclblasComplex* matx, aclblasComplex* maty,
206 aclblasComplex* result);203 aclblasComplex* result);
@@ -466,6 +463,11 @@ aclblasStatus_t aclblasSdgmm(
466 const float* A, int lda, const float* x, int incx,463 const float* A, int lda, const float* x, int incx,
467 float* C, int ldc);464 float* C, int ldc);
468 465 
466+aclblasStatus_t aclblasCdgmm(
467+ aclblasHandle_t handle, aclblasSideMode_t mode, int m, int n,
468+ const aclblasComplex* A, int lda, const aclblasComplex* x, int incx,
469+ aclblasComplex* C, int ldc);
470+ 
469aclblasStatus_t aclblasSgemm3m(471aclblasStatus_t aclblasSgemm3m(
470 aclblasHandle handle, aclblasOperation_t transA, aclblasOperation_t transB, int m, int n, int k,472 aclblasHandle handle, aclblasOperation_t transA, aclblasOperation_t transB, int m, int n, int k,
471 const float* alpha, const float* A, int lda,473 const float* alpha, const float* A, int lda,
Mtest/CMakeLists.txt+1-1
@@ -66,7 +66,7 @@ if(BUILD_TEST)
66 if(NOT ENABLE_BLAS_TRMM AND (TEST_NAME STREQUAL "strmm" OR TEST_NAME STREQUAL "trmm"))66 if(NOT ENABLE_BLAS_TRMM AND (TEST_NAME STREQUAL "strmm" OR TEST_NAME STREQUAL "trmm"))
67 list(APPEND SKIP_REASONS67 list(APPEND SKIP_REASONS
68 "requires asc-devkit >= 9.1, current ${ASC_DEVKIT_MAJOR}.${ASC_DEVKIT_MINOR}")68 "requires asc-devkit >= 9.1, current ${ASC_DEVKIT_MAJOR}.${ASC_DEVKIT_MINOR}")
69- elseif(NOT ENABLE_BLAS_DGMM AND (TEST_NAME STREQUAL "sdgmm" OR TEST_NAME STREQUAL "dgmm"))69+ elseif(NOT ENABLE_BLAS_SDGMM AND TEST_NAME STREQUAL "sdgmm")
70 list(APPEND SKIP_REASONS70 list(APPEND SKIP_REASONS
71 "requires asc-devkit >= 9.1, current ${ASC_DEVKIT_MAJOR}.${ASC_DEVKIT_MINOR}")71 "requires asc-devkit >= 9.1, current ${ASC_DEVKIT_MAJOR}.${ASC_DEVKIT_MINOR}")
72 elseif(NOT ENABLE_BLAS_GEMM_BATCHED AND (TEST_NAME STREQUAL "sgemm_batched" OR TEST_NAME STREQUAL "gemm_batched"))72 elseif(NOT ENABLE_BLAS_GEMM_BATCHED AND (TEST_NAME STREQUAL "sgemm_batched" OR TEST_NAME STREQUAL "gemm_batched"))
Dtest/colwise_mul/colwise_mul/arch22/colwise_mul_test.cpp+0-149
@@ -1,149 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9- */
10- 
11-/* !
12- * \file colwise_mul_test.cpp
13- * \brief Test for colwise_mul operator
14- */
15- 
16-#include <cstdint>
17-#include <iostream>
18-#include <vector>
19-#include <cmath>
20-#include "acl/acl.h"
21-#include "cann_ops_blas.h"
22- 
23-#define CHECK_RET(cond, return_expr) \
24- do { \
25- if (!(cond)) { \
26- return_expr; \
27- } \
28- } while (0)
29- 
30-#define LOG_PRINT(message, ...) \
31- do { \
32- printf(message, ##__VA_ARGS__); \
33- } while (0)
34- 
35-uint32_t VerifyResult(const float* output, const float* golden, size_t size, const char* test_name)
36-{
37- std::cout << "\n========== " << test_name << " ==========" << std::endl;
38- 
39- const float epsilon = 1e-4;
40- uint32_t errors = 0;
41- 
42- for (size_t i = 0; i < size; i++) {
43- if (std::abs(output[i] - golden[i]) > epsilon) {
44- if (errors < 5) {
45- std::cout << "Mismatch at index " << i << ": output=" << output[i] << ", golden=" << golden[i]
46- << std::endl;
47- }
48- errors++;
49- }
50- }
51- 
52- if (errors == 0) {
53- std::cout << "[Success] " << test_name << " verification passed." << std::endl;
54- return 0;
55- } else {
56- std::cout << "[Failed] " << test_name << " verification failed with " << errors << " errors!" << std::endl;
57- return 1;
58- }
59-}
60- 
61-int32_t main(int32_t argc, char* argv[])
62-{
63- int32_t deviceId = 0;
64- 
65- aclInit(nullptr);
66- aclrtSetDevice(deviceId);
67- 
68- aclblasHandle_t handle = nullptr;
69- auto ret = aclblasCreate(&handle);
70- CHECK_RET(ret == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", ret); return ret);
71- 
72- aclrtStream stream = nullptr;
73- aclrtCreateStream(&stream);
74- ret = aclblasSetStream(handle, stream);
75- CHECK_RET(ret == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", ret); return ret);
76- 
77- // Test: 2x3 complex matrix multiplied by 2-element complex vector
78- // Matrix (complex):
79- // Row 0: (1+2i, 3+4i, 5+6i)
80- // Row 1: (7+8i, 9+10i, 11+12i)
81- // Vector (complex):
82- // vec[0] = 2+3i
83- // vec[1] = 4+5i
84- // Result:
85- // Row 0: (2+3i) * (1+2i, 3+4i, 5+6i) = (-4+7i, -6+17i, -8+27i)
86- // Row 1: (4+5i) * (7+8i, 9+10i, 11+12i) = (-12+67i, -14+85i, -16+103i)
87- 
88- constexpr int64_t m = 2;
89- constexpr int64_t n = 3;
90- 
91- std::vector<float> mat = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f, 10.0f, 11.0f, 12.0f};
92- 
93- std::vector<float> vec = {2.0f, 3.0f, 4.0f, 5.0f};
94- 
95- std::vector<float> result(m * n * 2);
96- 
97- std::vector<float> golden = {-4.0f, 7.0f, -6.0f, 17.0f, -8.0f, 27.0f, -12.0f, 67.0f, -14.0f, 85.0f, -16.0f, 103.0f};
98- 
99- uint8_t* matDevice = nullptr;
100- uint8_t* vecDevice = nullptr;
101- uint8_t* resultDevice = nullptr;
102- size_t matByteSize = mat.size() * sizeof(float);
103- size_t vecByteSize = vec.size() * sizeof(float);
104- size_t resultByteSize = result.size() * sizeof(float);
105- 
106- aclError aclRet = aclrtMalloc((void**)&matDevice, matByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
107- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc matDevice failed. ERROR: %d\n", aclRet); return aclRet);
108- aclRet = aclrtMalloc((void**)&vecDevice, vecByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
109- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc vecDevice failed. ERROR: %d\n", aclRet); return aclRet);
110- aclRet = aclrtMalloc((void**)&resultDevice, resultByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
111- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc resultDevice failed. ERROR: %d\n", aclRet); return aclRet);
112- aclRet = aclrtMemcpy(matDevice, matByteSize, mat.data(), matByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
113- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy matDevice failed. ERROR: %d\n", aclRet); return aclRet);
114- aclRet = aclrtMemcpy(vecDevice, vecByteSize, vec.data(), vecByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
115- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy vecDevice failed. ERROR: %d\n", aclRet); return aclRet);
116- 
117- aclblasStatus_t statusRet = aclblasColwiseMul(handle, m, n, matDevice, vecDevice, resultDevice);
118- CHECK_RET(
119- statusRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasColwiseMul failed. ERROR: %d\n", statusRet);
120- return statusRet);
121- 
122- aclRet = aclrtSynchronizeStream(stream);
123- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet); return aclRet);
124- aclRet = aclrtMemcpy(result.data(), resultByteSize, resultDevice, resultByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
125- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy result failed. ERROR: %d\n", aclRet); return aclRet);
126- 
127- int32_t testRet = VerifyResult(result.data(), golden.data(), m * n * 2, "ColwiseMul Complex Test");
128- 
129- aclrtFree(matDevice);
130- aclrtFree(vecDevice);
131- aclrtFree(resultDevice);
132- 
133- aclrtDestroyStream(stream);
134- aclblasDestroy(handle);
135- aclrtResetDevice(deviceId);
136- aclFinalize();
137- 
138- if (testRet == 0) {
139- std::cout << "\n========================================" << std::endl;
140- std::cout << "Test passed successfully!" << std::endl;
141- std::cout << "========================================" << std::endl;
142- } else {
143- std::cout << "\n========================================" << std::endl;
144- std::cout << "Test failed!" << std::endl;
145- std::cout << "========================================" << std::endl;
146- }
147- 
148- return testRet;
149-}
Dtest/complex_mat_dot/complex_mat_dot/arch22/complex_mat_dot_test.cpp+0-144
@@ -1,144 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9- */
10- 
11-/* !
12- * \file complex_mat_dot_test.cpp
13- * \brief Test for complex matrix dot product
14- */
15- 
16-#include <cstdint>
17-#include <iostream>
18-#include <vector>
19-#include <algorithm>
20-#include <iterator>
21-#include <cmath>
22-#include "acl/acl.h"
23-#include "cann_ops_blas.h"
24- 
25-#define CHECK_RET(cond, return_expr) \
26- do { \
27- if (!(cond)) { \
28- return_expr; \
29- } \
30- } while (0)
31- 
32-#define LOG_PRINT(message, ...) \
33- do { \
34- printf(message, ##__VA_ARGS__); \
35- } while (0)
36- 
37-uint32_t VerifyResult(std::vector<float>& output, std::vector<float>& golden)
38-{
39- auto printTensor = [](std::vector<float>& tensor, const char* name) {
40- constexpr size_t maxPrintSize = 20;
41- std::cout << name << ": ";
42- std::copy(
43- tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
44- std::ostream_iterator<float>(std::cout, " "));
45- if (tensor.size() > maxPrintSize) {
46- std::cout << "...";
47- }
48- std::cout << std::endl;
49- };
50- printTensor(output, "Output");
51- printTensor(golden, "Golden");
52- 
53- // Use relative error for floating point comparison
54- constexpr float epsilon = 1e-5f;
55- size_t errorCount = 0;
56- for (size_t i = 0; i < output.size(); i++) {
57- float relError = std::abs(output[i] - golden[i]) / (std::abs(golden[i]) + 1e-10f);
58- if (relError > epsilon) {
59- errorCount++;
60- }
61- }
62- 
63- if (errorCount == 0) {
64- std::cout << "[Success] Case accuracy is verification passed." << std::endl;
65- return 0;
66- } else {
67- std::cout << "[Failed] Case accuracy is verification failed! Error count: " << errorCount << std::endl;
68- return 1;
69- }
70-}
71- 
72-int32_t main(int32_t argc, char* argv[])
73-{
74- int32_t deviceId = 0;
75- aclrtStream stream = nullptr;
76- aclblasHandle handle = nullptr;
77- 
78- constexpr uint32_t m = 4;
79- constexpr uint32_t n = 4;
80- constexpr uint32_t complexSize = m * n * 2;
81- 
82- std::vector<float> matx(complexSize);
83- std::vector<float> maty(complexSize);
84- std::vector<float> result(complexSize, 0.0f);
85- 
86- for (uint32_t i = 0; i < m * n; i++) {
87- matx[i * 2] = 1.0f;
88- matx[i * 2 + 1] = 2.0f;
89- maty[i * 2] = 3.0f;
90- maty[i * 2 + 1] = 4.0f;
91- }
92- 
93- aclInit(nullptr);
94- aclrtSetDevice(deviceId);
95- aclrtCreateStream(&stream);
96- aclblasCreate(&handle);
97- aclblasSetStream(handle, stream);
98- 
99- size_t dataSize = complexSize * sizeof(float);
100- 
101- aclblasComplex* matxDevice = nullptr;
102- aclblasComplex* matyDevice = nullptr;
103- aclblasComplex* resultDevice = nullptr;
104- 
105- aclError aclRet = aclrtMalloc((void**)&matxDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
106- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc matxDevice failed. ERROR: %d\n", aclRet); return aclRet);
107- aclRet = aclrtMalloc((void**)&matyDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
108- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc matyDevice failed. ERROR: %d\n", aclRet); return aclRet);
109- aclRet = aclrtMalloc((void**)&resultDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
110- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc resultDevice failed. ERROR: %d\n", aclRet); return aclRet);
111- 
112- aclRet = aclrtMemcpy(matxDevice, dataSize, matx.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
113- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy matxDevice failed. ERROR: %d\n", aclRet); return aclRet);
114- aclRet = aclrtMemcpy(matyDevice, dataSize, maty.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
115- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy matyDevice failed. ERROR: %d\n", aclRet); return aclRet);
116- aclRet = aclrtMemcpy(resultDevice, dataSize, result.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
117- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy resultDevice failed. ERROR: %d\n", aclRet); return aclRet);
118- 
119- auto ret = aclblasComplexMatDot(handle, m, n, matxDevice, matyDevice, resultDevice);
120- CHECK_RET(ret == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasComplexMatDot failed. ERROR: %d\n", ret); return ret);
121- 
122- aclRet = aclrtSynchronizeStream(stream);
123- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet); return aclRet);
124- 
125- aclRet = aclrtMemcpy(result.data(), dataSize, resultDevice, dataSize, ACL_MEMCPY_DEVICE_TO_HOST);
126- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy result failed. ERROR: %d\n", aclRet); return aclRet);
127- 
128- aclrtFree(matxDevice);
129- aclrtFree(matyDevice);
130- aclrtFree(resultDevice);
131- 
132- aclblasDestroy(handle);
133- aclrtDestroyStream(stream);
134- aclrtResetDevice(deviceId);
135- aclFinalize();
136- 
137- std::vector<float> golden(complexSize);
138- for (uint32_t i = 0; i < m * n; i++) {
139- golden[i * 2] = matx[i * 2] * maty[i * 2] - matx[i * 2 + 1] * maty[i * 2 + 1];
140- golden[i * 2 + 1] = matx[i * 2] * maty[i * 2 + 1] + matx[i * 2 + 1] * maty[i * 2];
141- }
142- 
143- return VerifyResult(result, golden);
144-}
Rtest/complex_mat_dot/complex_mat_dot/CMakeLists.txt→test/dgmm/cdgmm/CMakeLists.txt+3-3
@@ -2,10 +2,10 @@
2# Copyright (c) 2026 Huawei Technologies Co., Ltd.2# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.5+# Please refer to the License for details. You may not use the License for the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+# THIS FILE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------------------------------------9# ----------------------------------------------------------------------------------------------------------
10 10 
11-ops_blas_add_tests(${OPS_BLAS})11+ops_blas_add_gtest_tests(${OPS_BLAS} cdgmm_test)
Atest/dgmm/cdgmm/arch22/cdgmm_npu_wrapper.h+122-0
@@ -0,0 +1,122 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#pragma once
12+ 
13+#include <cstdint>
14+#include <vector>
15+ 
16+#include "acl/acl.h"
17+#include "cann_ops_blas.h"
18+#include "fill.h"
19+ 
20+static inline bool CdgmmNeedPassThrough(
21+ aclblasHandle_t handle, aclblasSideMode_t mode, int m, int n)
22+{
23+ return handle == nullptr ||
24+ m <= 0 || n <= 0 ||
25+ (mode != ACLBLAS_SIDE_LEFT && mode != ACLBLAS_SIDE_RIGHT);
26+}
27+ 
28+static inline aclError CdgmmAllocCopyH2D(void*& dPtr, const void* hPtr, size_t bytes)
29+{
30+ dPtr = nullptr;
31+ if (hPtr == nullptr) return ACL_SUCCESS;
32+ aclError ret = aclrtMalloc(&dPtr, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
33+ if (ret != ACL_SUCCESS) return ret;
34+ ret = aclrtMemcpy(dPtr, bytes, hPtr, bytes, ACL_MEMCPY_HOST_TO_DEVICE);
35+ if (ret != ACL_SUCCESS) {
36+ aclrtFree(dPtr);
37+ dPtr = nullptr;
38+ }
39+ return ret;
40+}
41+ 
42+static inline void CdgmmFreeAll(void* dX, void* dA, void* dC)
43+{
44+ if (dX) aclrtFree(dX);
45+ if (dA) aclrtFree(dA);
46+ if (dC) aclrtFree(dC);
47+}
48+ 
49+static inline aclError CdgmmAllocAndFillC(void*& dC, aclblasComplex* C, size_t cBytes)
50+{
51+ dC = nullptr;
52+ if (C == nullptr) return ACL_SUCCESS;
53+ aclError ret = aclrtMalloc(&dC, cBytes, ACL_MEM_MALLOC_HUGE_FIRST);
54+ if (ret != ACL_SUCCESS) return ret;
55+ // Fill with sentinel so unmodified padding matches the golden's untouched region.
56+ std::vector<float> sentinelBuf(cBytes / sizeof(float), kBlasSentinel);
57+ ret = aclrtMemcpy(dC, cBytes, sentinelBuf.data(), cBytes, ACL_MEMCPY_HOST_TO_DEVICE);
58+ if (ret != ACL_SUCCESS) {
59+ aclrtFree(dC);
60+ dC = nullptr;
61+ }
62+ return ret;
63+}
64+ 
65+inline aclblasStatus_t aclblasCdgmm_npu(
66+ aclblasHandle_t handle,
67+ aclblasSideMode_t mode,
68+ int m, int n,
69+ const aclblasComplex* A, int lda,
70+ const aclblasComplex* x, int incx,
71+ aclblasComplex* C, int ldc)
72+{
73+ if (CdgmmNeedPassThrough(handle, mode, m, n)) {
74+ return aclblasCdgmm(handle, mode, m, n, A, lda, x, incx, C, ldc);
75+ }
76+ 
77+ // Row-major: x length is m (LEFT only). Storage = lda * m complex elements.
78+ const int xLen = m;
79+ const int64_t absIncx = (incx >= 0) ? static_cast<int64_t>(incx)
80+ : -static_cast<int64_t>(incx);
81+ const size_t xTotalEl = static_cast<size_t>(xLen - 1) * static_cast<size_t>(absIncx) + 1;
82+ const size_t xBytes = xTotalEl * sizeof(aclblasComplex);
83+ const size_t aBytes = static_cast<size_t>(lda) * static_cast<size_t>(m) * sizeof(aclblasComplex);
84+ const size_t cBytes = static_cast<size_t>(ldc) * static_cast<size_t>(m) * sizeof(aclblasComplex);
85+ 
86+ void* dX = nullptr;
87+ void* dA = nullptr;
88+ void* dC = nullptr;
89+ 
90+ if (CdgmmAllocCopyH2D(dX, x, xBytes) != ACL_SUCCESS) {
91+ return ACLBLAS_STATUS_ALLOC_FAILED;
92+ }
93+ if (CdgmmAllocCopyH2D(dA, A, aBytes) != ACL_SUCCESS) {
94+ CdgmmFreeAll(dX, dA, dC);
95+ return ACLBLAS_STATUS_ALLOC_FAILED;
96+ }
97+ if (CdgmmAllocAndFillC(dC, C, cBytes) != ACL_SUCCESS) {
98+ CdgmmFreeAll(dX, dA, dC);
99+ return ACLBLAS_STATUS_ALLOC_FAILED;
100+ }
101+ 
102+ aclblasStatus_t ret = aclblasCdgmm(
103+ handle, mode, m, n,
104+ static_cast<const aclblasComplex*>(dA), lda,
105+ static_cast<const aclblasComplex*>(dX), incx,
106+ static_cast<aclblasComplex*>(dC), ldc);
107+ 
108+ if (aclrtSynchronizeDevice() != ACL_SUCCESS) {
109+ CdgmmFreeAll(dX, dA, dC);
110+ return ACLBLAS_STATUS_INTERNAL_ERROR;
111+ }
112+ 
113+ if (ret == ACLBLAS_STATUS_SUCCESS && C != nullptr && dC != nullptr) {
114+ if (aclrtMemcpy(C, cBytes, dC, cBytes, ACL_MEMCPY_DEVICE_TO_HOST) != ACL_SUCCESS) {
115+ CdgmmFreeAll(dX, dA, dC);
116+ return ACLBLAS_STATUS_INTERNAL_ERROR;
117+ }
118+ }
119+ 
120+ CdgmmFreeAll(dX, dA, dC);
121+ return ret;
122+}
Atest/dgmm/cdgmm/arch22/cdgmm_test.cpp+157-0
@@ -0,0 +1,157 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <algorithm>
12+#include <vector>
13+ 
14+#include "verify.h"
15+#include "blas_test.h"
16+#include "csv_loader.h"
17+#include "fill.h"
18+#include "cdgmm_param.h"
19+#include "cdgmm_golden.h"
20+#include "cdgmm_npu_wrapper.h"
21+ 
22+// ── Helpers for complex (interleaved float) data generation ──────────────────
23+ 
24+// Generate `count` complex elements (= 2*count floats) with stride `inc`
25+// (in complex elements). Real and imag use the same generator with a seed
26+// offset so both components vary. Returns interleaved float vector.
27+static inline std::vector<float> makeComplexStrided(
28+ int count, int inc, const BlasFillMode& fill, uint32_t seed)
29+{
30+ if (fill.method == BlasFillMode::M_NULLPTR || count <= 0)
31+ return {};
32+ 
33+ const int absInc = std::abs(inc);
34+ const size_t storageEl = static_cast<size_t>((count - 1) * absInc + 1);
35+ std::vector<float> data(storageEl * 2, 0.0f);
36+ 
37+ std::mt19937 rngReal(seed ? seed : 42);
38+ std::mt19937 rngImag((seed ? seed : 42) + 1000);
39+ auto genReal = createGenerator(fill, rngReal);
40+ auto genImag = createGenerator(fill, rngImag);
41+ 
42+ for (int i = 0; i < count; i++) {
43+ int idx = (inc > 0) ? (i * inc) : ((count - 1 - i) * absInc);
44+ data[static_cast<size_t>(idx) * 2] = genReal->at(i);
45+ data[static_cast<size_t>(idx) * 2 + 1] = genImag->at(i);
46+ }
47+ return data;
48+}
49+ 
50+// Generate a row-major complex matrix: m rows, n cols, row stride lda.
51+// Returns interleaved float vector of size m * lda * 2.
52+static inline std::vector<float> makeComplexMatrixRM(
53+ int m, int n, int lda, const BlasFillMode& fill, uint32_t seed)
54+{
55+ if (fill.method == BlasFillMode::M_NULLPTR || m <= 0 || n <= 0 || lda <= 0)
56+ return {};
57+ 
58+ const size_t storageSize = static_cast<size_t>(m) * lda * 2;
59+ std::vector<float> data(storageSize, 0.0f);
60+ 
61+ if (fill.method == BlasFillMode::M_VALUE) {
62+ for (size_t i = 0; i < storageSize; i++)
63+ data[i] = fill.val1;
64+ return data;
65+ }
66+ 
67+ std::mt19937 rngReal(seed ? seed : 42);
68+ std::mt19937 rngImag((seed ? seed : 42) + 2000);
69+ auto genReal = createGenerator(fill, rngReal);
70+ auto genImag = createGenerator(fill, rngImag);
71+ 
72+ for (int i = 0; i < m; i++) {
73+ for (int j = 0; j < n; j++) {
74+ size_t idx = (static_cast<size_t>(i) * lda + j) * 2;
75+ data[idx] = genReal->at(i * n + j);
76+ data[idx + 1] = genImag->at(i * n + j);
77+ }
78+ }
79+ return data;
80+}
81+ 
82+// ── Test fixture ─────────────────────────────────────────────────────────────
83+class CdgmmArch22Test : public BlasTest<CdgmmParam> { };
84+ 
85+// ── TEST_F: null handle (not in CSV) ─────────────────────────────────────────
86+TEST_F(CdgmmArch22Test, NullHandle) {
87+ aclblasStatus_t ret = aclblasCdgmm_npu(
88+ nullptr, ACLBLAS_SIDE_LEFT, 4, 4,
89+ nullptr, 4, nullptr, 1, nullptr, 4);
90+ EXPECT_EQ(static_cast<int>(ret), static_cast<int>(ACLBLAS_STATUS_HANDLE_IS_NULLPTR));
91+}
92+ 
93+// ── CSV parameterised test suite ─────────────────────────────────────────────
94+INSTANTIATE_TEST_SUITE_P(
95+ Cdgmm, CdgmmArch22Test,
96+ ::testing::ValuesIn(GetCasesFromCsv<CdgmmParam>(ReplaceFileExtension2Csv(__FILE__))),
97+ PrintCaseInfoString<CdgmmParam>);
98+ 
99+// ── TEST_P: 5-step CSV-driven flow ───────────────────────────────────────────
100+TEST_P(CdgmmArch22Test, CsvDriven) {
101+ const auto& p = GetParam();
102+ 
103+ // Step 1: Generate host data (interleaved complex floats)
104+ // Row-major: x length is m (LEFT only).
105+ const int xLen = p.m;
106+ std::vector<float> xHost;
107+ if (p.nullx == 0 && p.m > 0 && p.n > 0) {
108+ xHost = makeComplexStrided(xLen, p.incx, p.xFill, p.randomSeed);
109+ }
110+ 
111+ std::vector<float> aHost;
112+ if (p.nullA == 0 && p.m > 0 && p.n > 0) {
113+ aHost = makeComplexMatrixRM(p.m, p.n, p.lda, p.aFill, p.randomSeed);
114+ }
115+ 
116+ // C is the output buffer; initialise with sentinel so unmodified padding
117+ // (if any) matches the golden's untouched region.
118+ std::vector<float> cHost;
119+ if (p.nullC == 0 && p.m > 0 && p.n > 0) {
120+ cHost.assign(static_cast<size_t>(p.ldc) * static_cast<size_t>(p.m) * 2, kBlasSentinel);
121+ }
122+ 
123+ const aclblasComplex* xPtr = xHost.empty() ? nullptr : reinterpret_cast<const aclblasComplex*>(xHost.data());
124+ const aclblasComplex* aPtr = aHost.empty() ? nullptr : reinterpret_cast<const aclblasComplex*>(aHost.data());
125+ aclblasComplex* cPtr = cHost.empty() ? nullptr : reinterpret_cast<aclblasComplex*>(cHost.data());
126+ 
127+ // Step 2: Execute on NPU (wrapper handles nullptr passthrough, device memory)
128+ aclblasStatus_t ret = aclblasCdgmm_npu(
129+ CdgmmArch22Test::handle_, p.mode, p.m, p.n,
130+ aPtr, p.lda, xPtr, p.incx, cPtr, p.ldc);
131+ 
132+ // Step 3: Verify expected return code
133+ EXPECT_EQ(static_cast<int>(ret), static_cast<int>(p.expectResult));
134+ if (p.expectResult != ACLBLAS_STATUS_SUCCESS) return;
135+ 
136+ // m==0 or n==0: operator returns SUCCESS without computing; no output to verify
137+ if (p.m == 0 || p.n == 0) return;
138+ 
139+ // Step 4: Compute golden on CPU (row-major LEFT)
140+ std::vector<float> goldenC(cHost.size(), kBlasSentinel);
141+ aclblasStatus_t cpuRet = aclblasCdgmm_cpu(
142+ CdgmmArch22Test::handle_, p.mode, p.m, p.n,
143+ aPtr, p.lda, xPtr, p.incx,
144+ reinterpret_cast<aclblasComplex*>(goldenC.data()), p.ldc);
145+ EXPECT_EQ(static_cast<int>(cpuRet), static_cast<int>(ACLBLAS_STATUS_SUCCESS));
146+ 
147+ // Step 5: Precision verification — compare entire C storage as floats.
148+ // C storage = ldc * m * 2 floats (interleaved complex).
149+ // Padding columns (if ldc > n) are sentinel in both cHost and goldenC.
150+ VerifyConfig cfg;
151+ cfg.mode = PrecisionMode::MERE_MARE;
152+ cfg.mereThreshold = p.mereThreshold;
153+ cfg.mareMultiplier = p.mareMultiplier;
154+ 
155+ EXPECT_TRUE(Verifier::verifyVector(
156+ cHost.data(), goldenC.data(), cHost.size(), 1, cfg, p.caseName));
157+}
Atest/dgmm/cdgmm/arch22/cdgmm_test.csv+65-0
@@ -0,0 +1,65 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use the License for the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+# Row-major LEFT-only Cdgmm. lda/ldc are row strides (>= n).
11+# x length = m. incx may be negative.
12+case_name,description,mode,m,n,incx,lda,ldc,x_fill,a_fill,nullx,nullA,nullC,expect_result,mere_threshold,mare_multiplier,random_seed
13+TC_L0_01,mode=L 2x3 basic (original ColwiseMul case),LEFT,2,3,1,3,3,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
14+TC_L0_02,mode=L 1x1 minimal,LEFT,1,1,1,1,1,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
15+TC_L0_03,mode=L single row,LEFT,1,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
16+TC_L0_04,mode=L single col,LEFT,4,1,1,1,1,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
17+TC_L0_05,mode=L m>n,LEFT,8,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
18+TC_L0_06,mode=L m<n,LEFT,4,8,1,8,8,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
19+TC_L0_07,m=0 early return,LEFT,0,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0,0,42
20+TC_L0_08,n=0 early return,LEFT,4,0,1,1,1,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0,0,42
21+TC_L0_09,x nullptr,LEFT,4,4,1,4,4,NULLPTR,RANDOM_1_1,1,0,0,ACLBLAS_STATUS_INVALID_VALUE,0,0,42
22+TC_L0_10,A nullptr,LEFT,4,4,1,4,4,RANDOM_1_1,NULLPTR,0,1,0,ACLBLAS_STATUS_INVALID_VALUE,0,0,42
23+TC_L0_11,C nullptr,LEFT,4,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,1,ACLBLAS_STATUS_INVALID_VALUE,0,0,42
24+TC_L0_12,mode invalid 0xFF,255,4,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_INVALID_ENUM,0,0,42
25+TC_L0_13,m negative,LEFT,-1,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_INVALID_VALUE,0,0,42
26+TC_L0_14,n negative,LEFT,4,-1,1,1,1,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_INVALID_VALUE,0,0,42
27+TC_L0_15,incx zero,LEFT,4,4,0,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_INVALID_VALUE,0,0,42
28+TC_L0_16,lda too small,LEFT,4,4,1,3,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_INVALID_VALUE,0,0,42
29+TC_L0_17,ldc too small,LEFT,4,4,1,4,3,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_INVALID_VALUE,0,0,42
30+TC_L0_18,mode=R not supported,RIGHT,4,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_NOT_SUPPORTED,0,0,42
31+TC_L1_01,mode=L incx=2,LEFT,4,4,2,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
32+TC_L1_02,mode=L incx=3,LEFT,8,8,3,8,8,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
33+TC_L1_03,mode=L incx=-1 reverse,LEFT,4,4,-1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
34+TC_L1_04,mode=L incx=-2 reverse,LEFT,8,8,-2,8,8,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
35+TC_L1_05,mode=L lda padding,LEFT,4,8,1,16,8,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
36+TC_L1_06,mode=L ldc padding,LEFT,4,8,1,8,16,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
37+TC_L1_07,mode=L lda!=ldc,LEFT,4,4,1,8,6,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
38+TC_L1_08,mode=L lda+ldc padding,LEFT,8,4,1,16,16,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
39+TC_L1_09,mode=L 4x4 basic,LEFT,4,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
40+TC_L1_10,mode=L 8x8 basic,LEFT,8,8,1,8,8,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
41+TC_L1_11,mode=L 16x16 basic,LEFT,16,16,1,16,16,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
42+TC_L1_12,mode=L 7x7 unaligned,LEFT,7,7,1,7,7,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
43+TC_L1_13,mode=L 33x33 unaligned,LEFT,33,33,1,33,33,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
44+TC_L1_14,mode=L 64x64,LEFT,64,64,1,64,64,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
45+TC_L1_15,mode=L 128x128,LEFT,128,128,1,128,128,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
46+TC_L1_16,mode=L 256x256,LEFT,256,256,1,256,256,RANDOM_5_5,RANDOM_5_5,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
47+TC_L1_17,mode=L 512x512,LEFT,512,512,1,512,512,RANDOM_5_5,RANDOM_5_5,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
48+TC_L1_18,mode=L 1024x1024,LEFT,1024,1024,1,1024,1024,RANDOM_5_5,RANDOM_5_5,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
49+TC_L1_19,mode=L 1024x2048,LEFT,1024,2048,1,2048,2048,RANDOM_5_5,RANDOM_5_5,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
50+TC_L1_20,mode=L 2048x1024,LEFT,2048,1024,1,1024,1024,RANDOM_5_5,RANDOM_5_5,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
51+TC_L1_21,mode=L incx=2 64x64,LEFT,64,64,2,64,64,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
52+TC_L1_22,mode=L incx=-1 64x64,LEFT,64,64,-1,64,64,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
53+TC_L1_23,mode=L incx=3 128x128,LEFT,128,128,3,128,128,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
54+TC_L1_24,x all 0,LEFT,8,8,1,8,8,VALUE_NORM_0,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
55+TC_L1_25,x all 1,LEFT,8,8,1,8,8,VALUE_NORM_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
56+TC_L1_26,x all -1,LEFT,8,8,1,8,8,VALUE_NORM_N1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
57+TC_L1_27,x INF,LEFT,8,8,1,8,8,VALUE_NORM_INF,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
58+TC_L1_28,x NAN,LEFT,8,8,1,8,8,VALUE_NORM_NAN,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
59+TC_L1_29,A all 0,LEFT,8,8,1,8,8,RANDOM_1_1,VALUE_NORM_0,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
60+TC_L1_30,A extreme,LEFT,8,8,1,8,8,RANDOM_1_1,RANDOM_EXTREME,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
61+TC_L1_31,mode=L long row exceeds UB tile,LEFT,4,8192,1,8192,8192,RANDOM_5_5,RANDOM_5_5,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
62+TC_L1_32,mode=L large m multi-core,LEFT,4096,4,1,4,4,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
63+TC_L1_33,mode=L in-place A==C lda==ldc,LEFT,8,8,1,8,8,RANDOM_1_1,RANDOM_1_1,0,0,0,ACLBLAS_STATUS_SUCCESS,0.000244140625,10.0,42
64+TC_L1_34,m=0 ptrs null early return,LEFT,0,4,1,4,4,NULLPTR,NULLPTR,1,1,1,ACLBLAS_STATUS_SUCCESS,0,0,42
65+TC_L1_35,n=0 ptrs null early return,LEFT,4,0,1,1,1,NULLPTR,NULLPTR,1,1,1,ACLBLAS_STATUS_SUCCESS,0,0,42
Atest/dgmm/cdgmm/cdgmm_golden.h+90-0
@@ -0,0 +1,90 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#pragma once
12+ 
13+#include <algorithm>
14+#include <cstdint>
15+ 
16+#include "cann_ops_blas.h"
17+ 
18+static inline aclblasStatus_t CdgmmValidateParams(
19+ aclblasHandle_t handle, aclblasSideMode_t mode,
20+ int m, int n, const aclblasComplex* A, int lda,
21+ const aclblasComplex* x, int incx, aclblasComplex* C, int ldc)
22+{
23+ if (handle == nullptr) {
24+ return ACLBLAS_STATUS_HANDLE_IS_NULLPTR;
25+ }
26+ if (mode != ACLBLAS_SIDE_LEFT && mode != ACLBLAS_SIDE_RIGHT) {
27+ return ACLBLAS_STATUS_INVALID_ENUM;
28+ }
29+ if (mode == ACLBLAS_SIDE_RIGHT) {
30+ return ACLBLAS_STATUS_NOT_SUPPORTED;
31+ }
32+ if (m < 0 || n < 0 || incx == 0) {
33+ return ACLBLAS_STATUS_INVALID_VALUE;
34+ }
35+ if (lda < std::max(1, n) || ldc < std::max(1, n)) {
36+ return ACLBLAS_STATUS_INVALID_VALUE;
37+ }
38+ if (m > 0 && n > 0 && (x == nullptr || A == nullptr || C == nullptr)) {
39+ return ACLBLAS_STATUS_INVALID_VALUE;
40+ }
41+ if (A == C && lda != ldc) {
42+ return ACLBLAS_STATUS_INVALID_VALUE;
43+ }
44+ return ACLBLAS_STATUS_SUCCESS;
45+}
46+ 
47+// Row-major complex multiply
48+inline aclblasComplex ComplexMul(aclblasComplex a, aclblasComplex b)
49+{
50+ return {a.real * b.real - a.imag * b.imag,
51+ a.real * b.imag + a.imag * b.real};
52+}
53+ 
54+// Row-major LEFT golden: C[i,j] = x[i] * A[i,j]
55+// A and C are row-major with row strides lda and ldc (in complex elements).
56+// x has logical length m with stride incx.
57+// Data is stored as aclblasComplex (interleaved float pairs).
58+inline aclblasStatus_t aclblasCdgmm_cpu(
59+ aclblasHandle_t handle,
60+ aclblasSideMode_t mode,
61+ int m, int n,
62+ const aclblasComplex* A, int lda,
63+ const aclblasComplex* x, int incx,
64+ aclblasComplex* C, int ldc)
65+{
66+ aclblasStatus_t st = CdgmmValidateParams(handle, mode, m, n, A, lda, x, incx, C, ldc);
67+ if (st != ACLBLAS_STATUS_SUCCESS) {
68+ return st;
69+ }
70+ if (m == 0 || n == 0) {
71+ return ACLBLAS_STATUS_SUCCESS;
72+ }
73+ 
74+ const int xLen = m;
75+ const int64_t absIncx = (incx >= 0) ? static_cast<int64_t>(incx)
76+ : -static_cast<int64_t>(incx);
77+ 
78+ for (int i = 0; i < m; i++) {
79+ int64_t xIdx = (incx >= 0)
80+ ? static_cast<int64_t>(i) * incx
81+ : static_cast<int64_t>(xLen - 1 - i) * absIncx;
82+ aclblasComplex xVal = x[xIdx];
83+ 
84+ for (int j = 0; j < n; j++) {
85+ C[static_cast<int64_t>(i) * ldc + j] =
86+ ComplexMul(A[static_cast<int64_t>(i) * lda + j], xVal);
87+ }
88+ }
89+ return ACLBLAS_STATUS_SUCCESS;
90+}
Atest/dgmm/cdgmm/cdgmm_param.h+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#pragma once
12+ 
13+#include <string>
14+#include "cann_ops_blas.h"
15+#include "csv_loader.h"
16+ 
17+struct CdgmmParam : public BlasTestParamBase {
18+ aclblasSideMode_t mode = ACLBLAS_SIDE_LEFT;
19+ int m = 0;
20+ int n = 0;
21+ int incx = 1;
22+ int lda = 0;
23+ int ldc = 0;
24+ BlasFillMode xFill = BlasFillMode("RANDOM_1_1");
25+ BlasFillMode aFill = BlasFillMode("RANDOM_1_1");
26+ std::string xFillRaw;
27+ int nullx = 0;
28+ int nullA = 0;
29+ int nullC = 0;
30+ 
31+ CdgmmParam(const csv_map& csv) : BlasTestParamBase(csv)
32+ {
33+ mode = parseSideMode(ReadMap(csv, "mode", "LEFT"));
34+ m = parseInt(ReadMap(csv, "m", "0"));
35+ n = parseInt(ReadMap(csv, "n", "0"));
36+ incx = parseInt(ReadMap(csv, "incx", "1"));
37+ // Row-major: lda/ldc are row strides, constrained by n (not m).
38+ lda = parseInt(ReadMap(csv, "lda", std::to_string(std::max(1, n))));
39+ ldc = parseInt(ReadMap(csv, "ldc", std::to_string(std::max(1, n))));
40+ xFill = BlasFillMode(ReadMap(csv, "x_fill", "RANDOM_1_1"));
41+ xFillRaw = ReadMap(csv, "x_fill", "RANDOM_1_1");
42+ aFill = BlasFillMode(ReadMap(csv, "a_fill", "RANDOM_1_1"));
43+ nullx = parseInt(ReadMap(csv, "nullx", "0"));
44+ nullA = parseInt(ReadMap(csv, "nullA", "0"));
45+ nullC = parseInt(ReadMap(csv, "nullC", "0"));
46+ }
47+};
Rtest/colwise_mul/colwise_mul/CMakeLists.txt→test/extensions/complexmatdot/CMakeLists.txt+2-2
@@ -2,8 +2,8 @@
2# Copyright (c) 2026 Huawei Technologies Co., Ltd.2# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.5+# Please refer to the License for details. You may not use the License for the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+# THIS FILE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------------------------------------9# ----------------------------------------------------------------------------------------------------------
Atest/extensions/complexmatdot/arch22/complexmatdot_test.cpp+235-0
@@ -0,0 +1,235 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use the License for the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file complexmatdot_test.cpp
13+ * \brief Test for complex matrix dot product
14+ */
15+ 
16+#include <cstdint>
17+#include <iostream>
18+#include <vector>
19+#include <algorithm>
20+#include <iterator>
21+#include <cmath>
22+#include "acl/acl.h"
23+#include "cann_ops_blas.h"
24+ 
25+#define LOG_PRINT(message, ...) \
26+ do { \
27+ printf(message, ##__VA_ARGS__); \
28+ } while (0)
29+ 
30+// ── Test data ────────────────────────────────────────────────────────────────
31+ 
32+constexpr uint32_t M = 4;
33+constexpr uint32_t N = 4;
34+constexpr uint32_t COMPLEX_SIZE = M * N * 2;
35+ 
36+struct ComplexMatDotTestData {
37+ std::vector<float> matx;
38+ std::vector<float> maty;
39+ std::vector<float> result;
40+};
41+ 
42+static ComplexMatDotTestData BuildComplexMatDotInput()
43+{
44+ ComplexMatDotTestData data;
45+ data.matx.resize(COMPLEX_SIZE);
46+ data.maty.resize(COMPLEX_SIZE);
47+ data.result.assign(COMPLEX_SIZE, 0.0f);
48+ 
49+ for (uint32_t i = 0; i < M * N; i++) {
50+ data.matx[i * 2] = 1.0f;
51+ data.matx[i * 2 + 1] = 2.0f;
52+ data.maty[i * 2] = 3.0f;
53+ data.maty[i * 2 + 1] = 4.0f;
54+ }
55+ return data;
56+}
57+ 
58+static std::vector<float> BuildComplexMatDotGolden(const ComplexMatDotTestData& data)
59+{
60+ std::vector<float> golden(COMPLEX_SIZE);
61+ for (uint32_t i = 0; i < M * N; i++) {
62+ golden[i * 2] = data.matx[i * 2] * data.maty[i * 2] - data.matx[i * 2 + 1] * data.maty[i * 2 + 1];
63+ golden[i * 2 + 1] = data.matx[i * 2] * data.maty[i * 2 + 1] + data.matx[i * 2 + 1] * data.maty[i * 2];
64+ }
65+ return golden;
66+}
67+ 
68+// ── Test runtime context with RAII cleanup ───────────────────────────────────
69+ 
70+struct ComplexMatDotTestContext {
71+ int32_t deviceId = 0;
72+ aclrtStream stream = nullptr;
73+ aclblasHandle handle = nullptr;
74+ aclblasComplex* matxDevice = nullptr;
75+ aclblasComplex* matyDevice = nullptr;
76+ aclblasComplex* resultDevice = nullptr;
77+};
78+ 
79+static void CleanupTestRuntime(ComplexMatDotTestContext& ctx)
80+{
81+ if (ctx.resultDevice != nullptr) {
82+ aclrtFree(ctx.resultDevice);
83+ ctx.resultDevice = nullptr;
84+ }
85+ if (ctx.matyDevice != nullptr) {
86+ aclrtFree(ctx.matyDevice);
87+ ctx.matyDevice = nullptr;
88+ }
89+ if (ctx.matxDevice != nullptr) {
90+ aclrtFree(ctx.matxDevice);
91+ ctx.matxDevice = nullptr;
92+ }
93+ if (ctx.handle != nullptr) {
94+ aclblasDestroy(ctx.handle);
95+ ctx.handle = nullptr;
96+ }
97+ if (ctx.stream != nullptr) {
98+ aclrtDestroyStream(ctx.stream);
99+ ctx.stream = nullptr;
100+ }
101+ aclrtResetDevice(ctx.deviceId);
102+ aclFinalize();
103+}
104+ 
105+static int32_t InitTestRuntime(ComplexMatDotTestContext& ctx)
106+{
107+ aclInit(nullptr);
108+ aclrtSetDevice(ctx.deviceId);
109+ aclrtCreateStream(&ctx.stream);
110+ aclblasCreate(&ctx.handle);
111+ aclblasSetStream(ctx.handle, ctx.stream);
112+ return 0;
113+}
114+ 
115+// ── Run operator on NPU ──────────────────────────────────────────────────────
116+ 
117+static int32_t RunComplexMatDot(ComplexMatDotTestContext& ctx, const ComplexMatDotTestData& data)
118+{
119+ size_t dataSize = COMPLEX_SIZE * sizeof(float);
120+ 
121+ aclError aclRet = aclrtMalloc((void**)&ctx.matxDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
122+ if (aclRet != ACL_SUCCESS) {
123+ LOG_PRINT("aclrtMalloc matxDevice failed. ERROR: %d\n", aclRet);
124+ return static_cast<int32_t>(aclRet);
125+ }
126+ aclRet = aclrtMalloc((void**)&ctx.matyDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
127+ if (aclRet != ACL_SUCCESS) {
128+ LOG_PRINT("aclrtMalloc matyDevice failed. ERROR: %d\n", aclRet);
129+ return static_cast<int32_t>(aclRet);
130+ }
131+ aclRet = aclrtMalloc((void**)&ctx.resultDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
132+ if (aclRet != ACL_SUCCESS) {
133+ LOG_PRINT("aclrtMalloc resultDevice failed. ERROR: %d\n", aclRet);
134+ return static_cast<int32_t>(aclRet);
135+ }
136+ 
137+ aclRet = aclrtMemcpy(ctx.matxDevice, dataSize, data.matx.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
138+ if (aclRet != ACL_SUCCESS) {
139+ LOG_PRINT("aclrtMemcpy matxDevice failed. ERROR: %d\n", aclRet);
140+ return static_cast<int32_t>(aclRet);
141+ }
142+ aclRet = aclrtMemcpy(ctx.matyDevice, dataSize, data.maty.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
143+ if (aclRet != ACL_SUCCESS) {
144+ LOG_PRINT("aclrtMemcpy matyDevice failed. ERROR: %d\n", aclRet);
145+ return static_cast<int32_t>(aclRet);
146+ }
147+ aclRet = aclrtMemcpy(ctx.resultDevice, dataSize, data.result.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
148+ if (aclRet != ACL_SUCCESS) {
149+ LOG_PRINT("aclrtMemcpy resultDevice failed. ERROR: %d\n", aclRet);
150+ return static_cast<int32_t>(aclRet);
151+ }
152+ 
153+ auto ret = aclblasComplexMatDot(ctx.handle, M, N, ctx.matxDevice, ctx.matyDevice, ctx.resultDevice);
154+ if (ret != ACLBLAS_STATUS_SUCCESS) {
155+ LOG_PRINT("aclblasComplexMatDot failed. ERROR: %d\n", ret);
156+ return static_cast<int32_t>(ret);
157+ }
158+ 
159+ aclRet = aclrtSynchronizeStream(ctx.stream);
160+ if (aclRet != ACL_SUCCESS) {
161+ LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet);
162+ return static_cast<int32_t>(aclRet);
163+ }
164+ 
165+ return 0;
166+}
167+ 
168+// ── Verify ───────────────────────────────────────────────────────────────────
169+ 
170+static uint32_t VerifyResult(const std::vector<float>& output, const std::vector<float>& golden)
171+{
172+ if (output.size() != golden.size()) {
173+ std::cout << "[Failed] Size mismatch: output=" << output.size() << " golden=" << golden.size() << std::endl;
174+ return 1;
175+ }
176+ 
177+ auto printTensor = [](const std::vector<float>& tensor, const char* name) {
178+ constexpr size_t maxPrintSize = 20;
179+ std::cout << name << ": ";
180+ std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
181+ std::ostream_iterator<float>(std::cout, " "));
182+ if (tensor.size() > maxPrintSize) {
183+ std::cout << "...";
184+ }
185+ std::cout << std::endl;
186+ };
187+ printTensor(output, "Output");
188+ printTensor(golden, "Golden");
189+ 
190+ constexpr float epsilon = 1e-5f;
191+ size_t errorCount = 0;
192+ for (size_t i = 0; i < output.size(); i++) {
193+ float relError = std::abs(output[i] - golden[i]) / (std::abs(golden[i]) + 1e-10f);
194+ if (relError > epsilon) {
195+ errorCount++;
196+ }
197+ }
198+ 
199+ if (errorCount == 0) {
200+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
201+ return 0;
202+ }
203+ std::cout << "[Failed] Case accuracy is verification failed! Error count: " << errorCount << std::endl;
204+ return 1;
205+}
206+ 
207+// ── Main — orchestration only ────────────────────────────────────────────────
208+ 
209+int32_t main()
210+{
211+ ComplexMatDotTestData data = BuildComplexMatDotInput();
212+ 
213+ ComplexMatDotTestContext ctx;
214+ InitTestRuntime(ctx);
215+ 
216+ int32_t runStatus = RunComplexMatDot(ctx, data);
217+ if (runStatus != 0) {
218+ CleanupTestRuntime(ctx);
219+ return runStatus;
220+ }
221+ 
222+ size_t dataSize = COMPLEX_SIZE * sizeof(float);
223+ aclError aclRet = aclrtMemcpy(data.result.data(), dataSize, ctx.resultDevice, dataSize,
224+ ACL_MEMCPY_DEVICE_TO_HOST);
225+ if (aclRet != ACL_SUCCESS) {
226+ LOG_PRINT("aclrtMemcpy result failed. ERROR: %d\n", aclRet);
227+ CleanupTestRuntime(ctx);
228+ return static_cast<int32_t>(aclRet);
229+ }
230+ 
231+ CleanupTestRuntime(ctx);
232+ 
233+ std::vector<float> golden = BuildComplexMatDotGolden(data);
234+ return VerifyResult(data.result, golden);
235+}