已合并
feat(arch35): 新增 geam 算子，包含 aclblasSgeam 和 aclblasCgeam 接口 #328
iuyi创建于 10 天前
feat(arch35): 新增 geam 算子，包含 aclblasSgeam 和 aclblasCgeam 接口 #328
已合并
共 26 个文件变更+3377-3
| @@ -0,0 +1,532 @@ | |||
| 1 | +# Geam算子 | ||
| 2 | + | ||
| 3 | +## 算子概述 | ||
| 4 | + | ||
| 5 | +Geam（General Matrix Add）算子执行带标量缩放的矩阵加法运算，支持对输入矩阵 A、B 分别施加可选的转置（Transpose）或共轭转置（ConjTrans）操作。数学定义为： | ||
| 6 | + | ||
| 7 | +```txt | ||
| 8 | +C[i, j] = alpha * op(A)[i, j] + beta * op(B)[i, j], i ∈ [0, m), j ∈ [0, n) | ||
| 9 | +``` | ||
| 10 | + | ||
| 11 | +其中 `op(X)` 根据转置参数取 `X`（NoTrans）、`X^T`（Trans）或 `X^H`（ConjTrans，仅复数有意义，浮点数等价于 Trans）。矩阵采用列主序（Column-Major）存储，`X[i, j]` 的物理地址为 `X[i + j * ldX]`。 | ||
| 12 | + | ||
| 13 | +本算子提供两个接口，覆盖实数与复数两类矩阵加法场景： | ||
| 14 | + | ||
| 15 | +| 接口名 | 功能简述 | | ||
| 16 | +|--------|----------| | ||
| 17 | +| aclblasSgeam | 单精度浮点 GEAM（General Matrix Add） | | ||
| 18 | +| aclblasCgeam | 单精度复数 GEAM（General Matrix Add） | | ||
| 19 | + | ||
| 20 | +## 算子执行接口 | ||
| 21 | + | ||
| 22 | +### aclblasSgeam | ||
| 23 | + | ||
| 24 | +#### 产品支持情况 | ||
| 25 | + | ||
| 26 | +- Ascend 950PR / Ascend 950DT：支持 | ||
| 27 | +- Atlas A3 训练系列产品 / Atlas A3 推理系列产品：不支持 | ||
| 28 | +- Atlas A2 训练系列产品 / Atlas A2 推理系列产品：不支持 | ||
| 29 | + | ||
| 30 | +#### 函数原型 | ||
| 31 | + | ||
| 32 | +```cpp | ||
| 33 | +aclblasStatus_t aclblasSgeam(aclblasHandle_t handle, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, const float* alpha, const float* A, int lda, const float* beta, const float* B, int ldb, float* C, int ldc) | ||
| 34 | +``` | ||
| 35 | + | ||
| 36 | +#### 参数说明 | ||
| 37 | + | ||
| 38 | +| 参数名 | 输入/输出 | 参数类型 | 说明 | | ||
| 39 | +|--------|----------|---------|------| | ||
| 40 | +| handle | 输入 | aclblasHandle_t | ops-blas 库上下文句柄，携带 stream，Host 内存 | | ||
| 41 | +| transa | 输入 | aclblasOperation_t | 矩阵 A 的转置操作：ACLBLAS_OP_N（不转置）、ACLBLAS_OP_T（转置）、ACLBLAS_OP_C（共轭转置，FP32 实数等价于转置），Host 内存 | | ||
| 42 | +| transb | 输入 | aclblasOperation_t | 矩阵 B 的转置操作（同 transa），Host 内存 | | ||
| 43 | +| m | 输入 | int | 输出矩阵 C（及 op(A)）的行数，m >= 0，Host 内存 | | ||
| 44 | +| n | 输入 | int | 输出矩阵 C（及 op(B)）的列数，n >= 0，Host 内存 | | ||
| 45 | +| alpha | 输入 | const float*（FP32） | A 的缩放因子指针，指向 Host 侧单个标量；不可为 nullptr，Host 内存 | | ||
| 46 | +| A | 输入 | const float*（FP32） | 输入矩阵 A，列主序存储；当 *alpha != 0 时不可为 nullptr，Device 内存 | | ||
| 47 | +| lda | 输入 | int | 矩阵 A 的主维度，transa=N 时 lda >= max(1, m)，transa=T/C 时 lda >= max(1, n)，Host 内存 | | ||
| 48 | +| beta | 输入 | const float*（FP32） | B 的缩放因子指针，指向 Host 侧单个标量；不可为 nullptr，Host 内存 | | ||
| 49 | +| B | 输入 | const float*（FP32） | 输入矩阵 B，列主序存储；当 *beta != 0 时不可为 nullptr，Device 内存 | | ||
| 50 | +| ldb | 输入 | int | 矩阵 B 的主维度，transb=N 时 ldb >= max(1, m)，transb=T/C 时 ldb >= max(1, n)，Host 内存 | | ||
| 51 | +| C | 输出 | float*（FP32） | 输出矩阵 C，列主序存储；m > 0 且 n > 0 时不可为 nullptr，Device 内存 | | ||
| 52 | +| ldc | 输入 | int | 矩阵 C 的主维度，ldc >= max(1, m)，Host 内存 | | ||
| 53 | + | ||
| 54 | +#### 约束说明 | ||
| 55 | + | ||
| 56 | +- handle 不能为 nullptr，否则返回 ACLBLAS_STATUS_HANDLE_IS_NULLPTR | ||
| 57 | +- transa / transb 必须为 ACLBLAS_OP_N、ACLBLAS_OP_T 或 ACLBLAS_OP_C，否则返回 ACLBLAS_STATUS_INVALID_ENUM | ||
| 58 | +- m >= 0，n >= 0，否则返回 ACLBLAS_STATUS_INVALID_VALUE | ||
| 59 | +- transa = N 时 lda >= max(1, m)；transa = T/C 时 lda >= max(1, n) | ||
| 60 | +- transb = N 时 ldb >= max(1, m)；transb = T/C 时 ldb >= max(1, n) | ||
| 61 | +- ldc >= max(1, m) | ||
| 62 | +- alpha 不允许为 nullptr，否则返回 ACLBLAS_STATUS_INVALID_VALUE | ||
| 63 | +- beta 不允许为 nullptr，否则返回 ACLBLAS_STATUS_INVALID_VALUE | ||
| 64 | +- *alpha == 0 时，A 不被引用，可为 nullptr | ||
| 65 | +- *beta == 0 时，B 不被引用，可为 nullptr | ||
| 66 | +- *alpha != 0 时，A 不能为 nullptr | ||
| 67 | +- *beta != 0 时，B 不能为 nullptr | ||
| 68 | +- m > 0 且 n > 0 时，C 不能为 nullptr | ||
| 69 | +- 支持 in-place：C == A 时要求 transa == N 且 lda == ldc；C == B 时要求 transb == N 且 ldb == ldc | ||
| 70 | +- m == 0 或 n == 0 时直接返回 ACLBLAS_STATUS_SUCCESS，不执行计算 | ||
| 71 | + | ||
| 72 | +#### 调用示例 | ||
| 73 | + | ||
| 74 | +示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](https://www.hiascend.com/document/detail/zh/CANN/community/8.2.RC1/quickstart/quickstart_18_0041.html)。 | ||
| 75 | + | ||
| 76 | +以下示例演示 NN 模式（NoTrans/NoTrans），取 m=4, n=3, alpha=1.0, beta=0.5，计算 C(4x3) = 1.0 * A(4x3) + 0.5 * B(4x3)： | ||
| 77 | + | ||
| 78 | +```cpp | ||
| 79 | +#include <cstdio> | ||
| 80 | +#include <memory> | ||
| 81 | +#include <vector> | ||
| 82 | + | ||
| 83 | +#include "acl/acl.h" | ||
| 84 | +#include "cann_ops_blas.h" | ||
| 85 | + | ||
| 86 | +#define CHECK_RET(cond, return_expr) \ | ||
| 87 | + do { \ | ||
| 88 | + if (!(cond)) { \ | ||
| 89 | + return_expr; \ | ||
| 90 | + } \ | ||
| 91 | + } while (0) | ||
| 92 | + | ||
| 93 | +#define LOG_PRINT(message, ...) \ | ||
| 94 | + do { \ | ||
| 95 | + printf(message, ##__VA_ARGS__); \ | ||
| 96 | + } while (0) | ||
| 97 | + | ||
| 98 | +class AclContext { | ||
| 99 | +public: | ||
| 100 | + explicit AclContext(int32_t deviceId) : deviceId_(deviceId) {} | ||
| 101 | + | ||
| 102 | + ~AclContext() | ||
| 103 | + { | ||
| 104 | + if (stream_ != nullptr) { | ||
| 105 | + aclrtDestroyStream(stream_); | ||
| 106 | + stream_ = nullptr; | ||
| 107 | + } | ||
| 108 | + if (deviceSet_) { | ||
| 109 | + aclrtResetDevice(deviceId_); | ||
| 110 | + deviceSet_ = false; | ||
| 111 | + } | ||
| 112 | + if (aclInited_) { | ||
| 113 | + aclFinalize(); | ||
| 114 | + aclInited_ = false; | ||
| 115 | + } | ||
| 116 | + } | ||
| 117 | + | ||
| 118 | + int Init() | ||
| 119 | + { | ||
| 120 | + auto ret = aclInit(nullptr); | ||
| 121 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 122 | + aclInited_ = true; | ||
| 123 | + | ||
| 124 | + ret = aclrtSetDevice(deviceId_); | ||
| 125 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 126 | + deviceSet_ = true; | ||
| 127 | + | ||
| 128 | + ret = aclrtCreateStream(&stream_); | ||
| 129 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 130 | + return ACL_SUCCESS; | ||
| 131 | + } | ||
| 132 | + | ||
| 133 | + aclrtStream Stream() const { return stream_; } | ||
| 134 | + | ||
| 135 | +private: | ||
| 136 | + int32_t deviceId_; | ||
| 137 | + aclrtStream stream_ = nullptr; | ||
| 138 | + bool aclInited_ = false; | ||
| 139 | + bool deviceSet_ = false; | ||
| 140 | +}; | ||
| 141 | + | ||
| 142 | +struct AclMemDeleter { | ||
| 143 | + void operator()(void* p) const { aclrtFree(p); } | ||
| 144 | +}; | ||
| 145 | + | ||
| 146 | +int aclblasSgeamTest(AclContext& ctx) | ||
| 147 | +{ | ||
| 148 | + aclrtStream stream = ctx.Stream(); | ||
| 149 | + | ||
| 150 | + // 1. 创建 ops-blas 句柄 | ||
| 151 | + aclblasHandle_t rawHandle = nullptr; | ||
| 152 | + auto blasRet = aclblasCreate(&rawHandle); | ||
| 153 | + CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet); | ||
| 154 | + return blasRet); | ||
| 155 | + std::unique_ptr<void, aclblasStatus_t (*)(void*)> handlePtr(rawHandle, aclblasDestroy); | ||
| 156 | + | ||
| 157 | + blasRet = aclblasSetStream(static_cast<aclblasHandle_t>(handlePtr.get()), stream); | ||
| 158 | + CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet); | ||
| 159 | + return blasRet); | ||
| 160 | + | ||
| 161 | + // 2. 准备 Host 数据 | ||
| 162 | + // C(4x3) = 1.0 * A(4x3) + 0.5 * B(4x3)，NN 模式 | ||
| 163 | + int m = 4, n = 3; | ||
| 164 | + int lda = m, ldb = m, ldc = m; | ||
| 165 | + float alpha = 1.0f; | ||
| 166 | + float beta = 0.5f; | ||
| 167 | + | ||
| 168 | + // A (m=4, n=3, column-major): 列0=[1,2,3,4], 列1=[5,6,7,8], 列2=[9,10,11,12] | ||
| 169 | + std::vector<float> hA = {1.0f, 2.0f, 3.0f, 4.0f, | ||
| 170 | + 5.0f, 6.0f, 7.0f, 8.0f, | ||
| 171 | + 9.0f, 10.0f, 11.0f, 12.0f}; | ||
| 172 | + // B (m=4, n=3, column-major): 列0=[2,4,6,8], 列1=[10,12,14,16], 列2=[18,20,22,24] | ||
| 173 | + std::vector<float> hB = {2.0f, 4.0f, 6.0f, 8.0f, | ||
| 174 | + 10.0f, 12.0f, 14.0f, 16.0f, | ||
| 175 | + 18.0f, 20.0f, 22.0f, 24.0f}; | ||
| 176 | + | ||
| 177 | + size_t aBytes = hA.size() * sizeof(float); | ||
| 178 | + size_t bBytes = hB.size() * sizeof(float); | ||
| 179 | + size_t cBytes = static_cast<size_t>(ldc) * static_cast<size_t>(n) * sizeof(float); | ||
| 180 | + | ||
| 181 | + // 3. 申请 Device 内存并拷贝数据 | ||
| 182 | + void* rawA = nullptr; | ||
| 183 | + auto aclRet = aclrtMalloc(&rawA, aBytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 184 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for A failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 185 | + std::unique_ptr<float, AclMemDeleter> aDevicePtr(static_cast<float*>(rawA)); | ||
| 186 | + | ||
| 187 | + void* rawB = nullptr; | ||
| 188 | + aclRet = aclrtMalloc(&rawB, bBytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 189 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for B failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 190 | + std::unique_ptr<float, AclMemDeleter> bDevicePtr(static_cast<float*>(rawB)); | ||
| 191 | + | ||
| 192 | + void* rawC = nullptr; | ||
| 193 | + aclRet = aclrtMalloc(&rawC, cBytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 194 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for C failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 195 | + std::unique_ptr<float, AclMemDeleter> cDevicePtr(static_cast<float*>(rawC)); | ||
| 196 | + | ||
| 197 | + aclRet = aclrtMemcpy(aDevicePtr.get(), aBytes, hA.data(), aBytes, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 198 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for A failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 199 | + | ||
| 200 | + aclRet = aclrtMemcpy(bDevicePtr.get(), bBytes, hB.data(), bBytes, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 201 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for B failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 202 | + | ||
| 203 | + // 4. 调用 aclblasSgeam | ||
| 204 | + blasRet = aclblasSgeam(static_cast<aclblasHandle_t>(handlePtr.get()), | ||
| 205 | + ACLBLAS_OP_N, ACLBLAS_OP_N, m, n, | ||
| 206 | + &alpha, aDevicePtr.get(), lda, | ||
| 207 | + &beta, bDevicePtr.get(), ldb, | ||
| 208 | + static_cast<float*>(cDevicePtr.get()), ldc); | ||
| 209 | + CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSgeam failed. ERROR: %d\n", blasRet); | ||
| 210 | + return blasRet); | ||
| 211 | + | ||
| 212 | + // 5. 同步等待任务执行结束 | ||
| 213 | + aclRet = aclrtSynchronizeStream(stream); | ||
| 214 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 215 | + | ||
| 216 | + // 6. 将结果从 Device 拷贝回 Host 并打印 | ||
| 217 | + // 预期 C = alpha*A + beta*B (列主序) | ||
| 218 | + // 列0=[1+1, 2+2, 3+3, 4+4]=[2,4,6,8] | ||
| 219 | + // 列1=[5+5, 6+6, 7+7, 8+8]=[10,12,14,16] | ||
| 220 | + // 列2=[9+9, 10+10, 11+11, 12+12]=[18,20,22,24] | ||
| 221 | + std::vector<float> hC(static_cast<size_t>(ldc) * static_cast<size_t>(n), 0.0f); | ||
| 222 | + aclRet = aclrtMemcpy(hC.data(), cBytes, cDevicePtr.get(), cBytes, ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 223 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 224 | + | ||
| 225 | + LOG_PRINT("result C (column-major):\n"); | ||
| 226 | + for (int col = 0; col < n; col++) { | ||
| 227 | + for (int row = 0; row < m; row++) { | ||
| 228 | + LOG_PRINT(" C[%d][%d] = %f\n", row, col, hC[static_cast<size_t>(col) * ldc + row]); | ||
| 229 | + } | ||
| 230 | + } | ||
| 231 | + | ||
| 232 | + return ACL_SUCCESS; | ||
| 233 | +} | ||
| 234 | + | ||
| 235 | +int main() | ||
| 236 | +{ | ||
| 237 | + AclContext ctx(0); | ||
| 238 | + auto ret = ctx.Init(); | ||
| 239 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 240 | + | ||
| 241 | + ret = aclblasSgeamTest(ctx); | ||
| 242 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclblasSgeamTest failed. ERROR: %d\n", ret); return ret); | ||
| 243 | + return 0; | ||
| 244 | +} | ||
| 245 | +``` | ||
| 246 | + | ||
| 247 | +预期输出： | ||
| 248 | + | ||
| 249 | +``` | ||
| 250 | +result C (column-major): | ||
| 251 | + C[0][0] = 2.000000 | ||
| 252 | + C[1][0] = 4.000000 | ||
| 253 | + C[2][0] = 6.000000 | ||
| 254 | + C[3][0] = 8.000000 | ||
| 255 | + C[0][1] = 10.000000 | ||
| 256 | + C[1][1] = 12.000000 | ||
| 257 | + C[2][1] = 14.000000 | ||
| 258 | + C[3][1] = 16.000000 | ||
| 259 | + C[0][2] = 18.000000 | ||
| 260 | + C[1][2] = 20.000000 | ||
| 261 | + C[2][2] = 22.000000 | ||
| 262 | + C[3][2] = 24.000000 | ||
| 263 | +``` | ||
| 264 | + | ||
| 265 | +### aclblasCgeam | ||
| 266 | + | ||
| 267 | +#### 产品支持情况 | ||
| 268 | + | ||
| 269 | +- Ascend 950PR / Ascend 950DT：支持 | ||
| 270 | +- Atlas A3 训练系列产品 / Atlas A3 推理系列产品：不支持 | ||
| 271 | +- Atlas A2 训练系列产品 / Atlas A2 推理系列产品：不支持 | ||
| 272 | + | ||
| 273 | +#### 函数原型 | ||
| 274 | + | ||
| 275 | +```cpp | ||
| 276 | +aclblasStatus_t aclblasCgeam(aclblasHandle_t handle, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, const aclblasComplex* alpha, const aclblasComplex* A, int lda, const aclblasComplex* beta, const aclblasComplex* B, int ldb, aclblasComplex* C, int ldc) | ||
| 277 | +``` | ||
| 278 | + | ||
| 279 | +#### 参数说明 | ||
| 280 | + | ||
| 281 | +| 参数名 | 输入/输出 | 参数类型 | 说明 | | ||
| 282 | +|--------|----------|---------|------| | ||
| 283 | +| handle | 输入 | aclblasHandle_t | ops-blas 库上下文句柄，携带 stream，Host 内存 | | ||
| 284 | +| transa | 输入 | aclblasOperation_t | 矩阵 A 的转置操作：ACLBLAS_OP_N（不转置）、ACLBLAS_OP_T（转置）、ACLBLAS_OP_C（共轭转置），Host 内存 | | ||
| 285 | +| transb | 输入 | aclblasOperation_t | 矩阵 B 的转置操作（同 transa），Host 内存 | | ||
| 286 | +| m | 输入 | int | 输出矩阵 C（及 op(A)）的行数，m >= 0，Host 内存 | | ||
| 287 | +| n | 输入 | int | 输出矩阵 C（及 op(B)）的列数，n >= 0，Host 内存 | | ||
| 288 | +| alpha | 输入 | const aclblasComplex* | A 的缩放因子指针，指向 Host 侧单个标量；不可为 nullptr，Host 内存 | | ||
| 289 | +| A | 输入 | const aclblasComplex* | 输入矩阵 A，列主序存储；当 *alpha != 0 时不可为 nullptr，Device 内存 | | ||
| 290 | +| lda | 输入 | int | 矩阵 A 的主维度，transa=N 时 lda >= max(1, m)，transa=T/C 时 lda >= max(1, n)，Host 内存 | | ||
| 291 | +| beta | 输入 | const aclblasComplex* | B 的缩放因子指针，指向 Host 侧单个标量；不可为 nullptr，Host 内存 | | ||
| 292 | +| B | 输入 | const aclblasComplex* | 输入矩阵 B，列主序存储；当 *beta != 0 时不可为 nullptr，Device 内存 | | ||
| 293 | +| ldb | 输入 | int | 矩阵 B 的主维度，transb=N 时 ldb >= max(1, m)，transb=T/C 时 ldb >= max(1, n)，Host 内存 | | ||
| 294 | +| C | 输出 | aclblasComplex* | 输出矩阵 C，列主序存储；m > 0 且 n > 0 时不可为 nullptr，Device 内存 | | ||
| 295 | +| ldc | 输入 | int | 矩阵 C 的主维度，ldc >= max(1, m)，Host 内存 | | ||
| 296 | + | ||
| 297 | +#### 约束说明 | ||
| 298 | + | ||
| 299 | +- handle 不能为 nullptr，否则返回 ACLBLAS_STATUS_HANDLE_IS_NULLPTR | ||
| 300 | +- transa / transb 必须为 ACLBLAS_OP_N、ACLBLAS_OP_T 或 ACLBLAS_OP_C，否则返回 ACLBLAS_STATUS_INVALID_ENUM | ||
| 301 | +- m >= 0，n >= 0，否则返回 ACLBLAS_STATUS_INVALID_VALUE | ||
| 302 | +- transa = N 时 lda >= max(1, m)；transa = T/C 时 lda >= max(1, n) | ||
| 303 | +- transb = N 时 ldb >= max(1, m)；transb = T/C 时 ldb >= max(1, n) | ||
| 304 | +- ldc >= max(1, m) | ||
| 305 | +- alpha 不允许为 nullptr，否则返回 ACLBLAS_STATUS_INVALID_VALUE | ||
| 306 | +- beta 不允许为 nullptr，否则返回 ACLBLAS_STATUS_INVALID_VALUE | ||
| 307 | +- *alpha == 0 时，A 不被引用，可为 nullptr | ||
| 308 | +- *beta == 0 时，B 不被引用，可为 nullptr | ||
| 309 | +- *alpha != 0 时，A 不能为 nullptr | ||
| 310 | +- *beta != 0 时，B 不能为 nullptr | ||
| 311 | +- m > 0 且 n > 0 时，C 不能为 nullptr | ||
| 312 | +- 支持 in-place：C == A 时要求 transa == N 且 lda == ldc；C == B 时要求 transb == N 且 ldb == ldc | ||
| 313 | +- m == 0 或 n == 0 时直接返回 ACLBLAS_STATUS_SUCCESS，不执行计算 | ||
| 314 | +- ACLBLAS_OP_C 对复数执行共轭转置，实部不变，虚部取反 | ||
| 315 | + | ||
| 316 | +#### 调用示例 | ||
| 317 | + | ||
| 318 | +示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](https://www.hiascend.com/document/detail/zh/CANN/community/8.2.RC1/quickstart/quickstart_18_0041.html)。 | ||
| 319 | + | ||
| 320 | +以下示例演示 NC 模式（NoTrans/ConjTrans），取 m=2, n=3, alpha=(1+1i), beta=(0.5+0.5i)，计算 C(2x3) = (1+i) * A(2x3) + (0.5+0.5i) * B^H(2x3)： | ||
| 321 | + | ||
| 322 | +```cpp | ||
| 323 | +#include <cstdio> | ||
| 324 | +#include <memory> | ||
| 325 | +#include <vector> | ||
| 326 | + | ||
| 327 | +#include "acl/acl.h" | ||
| 328 | +#include "cann_ops_blas.h" | ||
| 329 | + | ||
| 330 | +#define CHECK_RET(cond, return_expr) \ | ||
| 331 | + do { \ | ||
| 332 | + if (!(cond)) { \ | ||
| 333 | + return_expr; \ | ||
| 334 | + } \ | ||
| 335 | + } while (0) | ||
| 336 | + | ||
| 337 | +#define LOG_PRINT(message, ...) \ | ||
| 338 | + do { \ | ||
| 339 | + printf(message, ##__VA_ARGS__); \ | ||
| 340 | + } while (0) | ||
| 341 | + | ||
| 342 | +class AclContext { | ||
| 343 | +public: | ||
| 344 | + explicit AclContext(int32_t deviceId) : deviceId_(deviceId) {} | ||
| 345 | + | ||
| 346 | + ~AclContext() | ||
| 347 | + { | ||
| 348 | + if (stream_ != nullptr) { | ||
| 349 | + aclrtDestroyStream(stream_); | ||
| 350 | + stream_ = nullptr; | ||
| 351 | + } | ||
| 352 | + if (deviceSet_) { | ||
| 353 | + aclrtResetDevice(deviceId_); | ||
| 354 | + deviceSet_ = false; | ||
| 355 | + } | ||
| 356 | + if (aclInited_) { | ||
| 357 | + aclFinalize(); | ||
| 358 | + aclInited_ = false; | ||
| 359 | + } | ||
| 360 | + } | ||
| 361 | + | ||
| 362 | + int Init() | ||
| 363 | + { | ||
| 364 | + auto ret = aclInit(nullptr); | ||
| 365 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 366 | + aclInited_ = true; | ||
| 367 | + | ||
| 368 | + ret = aclrtSetDevice(deviceId_); | ||
| 369 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 370 | + deviceSet_ = true; | ||
| 371 | + | ||
| 372 | + ret = aclrtCreateStream(&stream_); | ||
| 373 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 374 | + return ACL_SUCCESS; | ||
| 375 | + } | ||
| 376 | + | ||
| 377 | + aclrtStream Stream() const { return stream_; } | ||
| 378 | + | ||
| 379 | +private: | ||
| 380 | + int32_t deviceId_; | ||
| 381 | + aclrtStream stream_ = nullptr; | ||
| 382 | + bool aclInited_ = false; | ||
| 383 | + bool deviceSet_ = false; | ||
| 384 | +}; | ||
| 385 | + | ||
| 386 | +struct AclMemDeleter { | ||
| 387 | + void operator()(void* p) const { aclrtFree(p); } | ||
| 388 | +}; | ||
| 389 | + | ||
| 390 | +int aclblasCgeamTest(AclContext& ctx) | ||
| 391 | +{ | ||
| 392 | + aclrtStream stream = ctx.Stream(); | ||
| 393 | + | ||
| 394 | + // 1. 创建 ops-blas 句柄 | ||
| 395 | + aclblasHandle_t rawHandle = nullptr; | ||
| 396 | + auto blasRet = aclblasCreate(&rawHandle); | ||
| 397 | + CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet); | ||
| 398 | + return blasRet); | ||
| 399 | + std::unique_ptr<void, aclblasStatus_t (*)(void*)> handlePtr(rawHandle, aclblasDestroy); | ||
| 400 | + | ||
| 401 | + blasRet = aclblasSetStream(static_cast<aclblasHandle_t>(handlePtr.get()), stream); | ||
| 402 | + CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet); | ||
| 403 | + return blasRet); | ||
| 404 | + | ||
| 405 | + // 2. 准备 Host 数据 | ||
| 406 | + // C(2x3) = (1+i) * A(2x3) + (0.5+0.5i) * B^H(2x3)，NC 模式 | ||
| 407 | + int m = 2, n = 3; | ||
| 408 | + int lda = m, ldb = n, ldc = m; // transb=C 时 ldb >= max(1, n=3) | ||
| 409 | + aclblasComplex alpha = {1.0f, 1.0f}; // 1+i | ||
| 410 | + aclblasComplex beta = {0.5f, 0.5f}; // 0.5+0.5i | ||
| 411 | + | ||
| 412 | + // A (m=2, n=3, column-major): 列0=[(1+i),(2+0i)], 列1=[(3+0i),(4+i)], 列2=[(0+2i),(1+1i)] | ||
| 413 | + std::vector<aclblasComplex> hA = {{{1.0f, 1.0f}, {2.0f, 0.0f}, | ||
| 414 | + {3.0f, 0.0f}, {4.0f, 1.0f}, | ||
| 415 | + {0.0f, 2.0f}, {1.0f, 1.0f}}}; | ||
| 416 | + // B 原始 (n=3, m=2, column-major)，转置前 B 形状为 (3,2) | ||
| 417 | + // 列0=[(1+0i),(0+1i),(1+1i)], 列1=[(2+0i),(0+0i),(1+0i)] | ||
| 418 | + // B^H 将 B 共轭转置为 (m=2, n=3) | ||
| 419 | + std::vector<aclblasComplex> hB = {{{1.0f, 0.0f}, {0.0f, 1.0f}, {1.0f, 1.0f}, | ||
| 420 | + {2.0f, 0.0f}, {0.0f, 0.0f}, {1.0f, 0.0f}}}; | ||
| 421 | + | ||
| 422 | + size_t aBytes = hA.size() * sizeof(aclblasComplex); | ||
| 423 | + size_t bBytes = hB.size() * sizeof(aclblasComplex); | ||
| 424 | + size_t cBytes = static_cast<size_t>(ldc) * static_cast<size_t>(n) * sizeof(aclblasComplex); | ||
| 425 | + | ||
| 426 | + // 3. 申请 Device 内存并拷贝数据 | ||
| 427 | + aclblasComplex* rawA = nullptr; | ||
| 428 | + auto aclRet = aclrtMalloc(reinterpret_cast<void**>(&rawA), aBytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 429 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for A failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 430 | + std::unique_ptr<aclblasComplex, AclMemDeleter> aDevicePtr(rawA); | ||
| 431 | + | ||
| 432 | + aclblasComplex* rawB = nullptr; | ||
| 433 | + aclRet = aclrtMalloc(reinterpret_cast<void**>(&rawB), bBytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 434 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for B failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 435 | + std::unique_ptr<aclblasComplex, AclMemDeleter> bDevicePtr(rawB); | ||
| 436 | + | ||
| 437 | + aclblasComplex* rawC = nullptr; | ||
| 438 | + aclRet = aclrtMalloc(reinterpret_cast<void**>(&rawC), cBytes, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 439 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for C failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 440 | + std::unique_ptr<aclblasComplex, AclMemDeleter> cDevicePtr(rawC); | ||
| 441 | + | ||
| 442 | + aclRet = aclrtMemcpy(aDevicePtr.get(), aBytes, hA.data(), aBytes, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 443 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for A failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 444 | + | ||
| 445 | + aclRet = aclrtMemcpy(bDevicePtr.get(), bBytes, hB.data(), bBytes, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 446 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for B failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 447 | + | ||
| 448 | + // 4. 调用 aclblasCgeam | ||
| 449 | + blasRet = aclblasCgeam(static_cast<aclblasHandle_t>(handlePtr.get()), | ||
| 450 | + ACLBLAS_OP_N, ACLBLAS_OP_C, m, n, | ||
| 451 | + &alpha, aDevicePtr.get(), lda, | ||
| 452 | + &beta, bDevicePtr.get(), ldb, | ||
| 453 | + cDevicePtr.get(), ldc); | ||
| 454 | + CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCgeam failed. ERROR: %d\n", blasRet); | ||
| 455 | + return blasRet); | ||
| 456 | + | ||
| 457 | + // 5. 同步等待任务执行结束 | ||
| 458 | + aclRet = aclrtSynchronizeStream(stream); | ||
| 459 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 460 | + | ||
| 461 | + // 6. 将结果从 Device 拷贝回 Host 并打印 | ||
| 462 | + std::vector<aclblasComplex> hC(static_cast<size_t>(ldc) * static_cast<size_t>(n), {0.0f, 0.0f}); | ||
| 463 | + aclRet = aclrtMemcpy(hC.data(), cBytes, cDevicePtr.get(), cBytes, ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 464 | + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", aclRet); return aclRet); | ||
| 465 | + | ||
| 466 | + LOG_PRINT("result C (column-major):\n"); | ||
| 467 | + for (int col = 0; col < n; col++) { | ||
| 468 | + for (int row = 0; row < m; row++) { | ||
| 469 | + const auto& val = hC[static_cast<size_t>(col) * ldc + row]; | ||
| 470 | + LOG_PRINT(" C[%d][%d] = %f + %fi\n", row, col, val.real, val.imag); | ||
| 471 | + } | ||
| 472 | + } | ||
| 473 | + | ||
| 474 | + return ACL_SUCCESS; | ||
| 475 | +} | ||
| 476 | + | ||
| 477 | +int main() | ||
| 478 | +{ | ||
| 479 | + AclContext ctx(0); | ||
| 480 | + auto ret = ctx.Init(); | ||
| 481 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 482 | + | ||
| 483 | + ret = aclblasCgeamTest(ctx); | ||
| 484 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclblasCgeamTest failed. ERROR: %d\n", ret); return ret); | ||
| 485 | + return 0; | ||
| 486 | +} | ||
| 487 | +``` | ||
| 488 | + | ||
| 489 | +预期输出： | ||
| 490 | + | ||
| 491 | +``` | ||
| 492 | +result C (column-major): | ||
| 493 | + C[0][0] = 0.500000 + 2.500000i | ||
| 494 | + C[1][0] = 2.500000 + 2.500000i | ||
| 495 | + C[0][1] = 3.500000 + 3.500000i | ||
| 496 | + C[1][1] = 4.500000 + 5.500000i | ||
| 497 | + C[0][2] = -1.500000 + 3.000000i | ||
| 498 | + C[1][2] = 1.000000 + 2.000000i | ||
| 499 | +``` | ||
| 500 | + | ||
| 501 | +## 支持的数据类型 | ||
| 502 | + | ||
| 503 | +| 接口 | 数据类型 | 说明 | | ||
| 504 | +|------|----------|------| | ||
| 505 | +| `aclblasSgeam` | `float` (FP32) | 单精度实数 | | ||
| 506 | +| `aclblasCgeam` | `aclblasComplex` | 单精度复数（实部 + 虚部各为 float） | | ||
| 507 | + | ||
| 508 | +## 返回值 / 错误码 | ||
| 509 | + | ||
| 510 | +| 返回值 | 含义 | | ||
| 511 | +|--------|------| | ||
| 512 | +| `ACLBLAS_STATUS_SUCCESS` | 执行成功 | | ||
| 513 | +| `ACLBLAS_STATUS_HANDLE_IS_NULLPTR` | handle 为空指针 | | ||
| 514 | +| `ACLBLAS_STATUS_INVALID_VALUE` | 维度非法（m < 0、n < 0）、leading dimension 不满足约束、必需指针为空（alpha、beta、C、A/B 在对应标量非零时为 nullptr） | | ||
| 515 | +| `ACLBLAS_STATUS_INVALID_ENUM` | transa / transb 取值不在 N/T/C 范围内 | | ||
| 516 | +| `ACLBLAS_STATUS_INTERNAL_ERROR` | 内部错误（如获取核心数失败） | | ||
| 517 | + | ||
| 518 | +## 精度标准 | ||
| 519 | + | ||
| 520 | +| 数据类型 | rtol | atol | required_matched_ratio | max_abs_error_limit | | ||
| 521 | +|----------|------|------|----------------------|-------------------| | ||
| 522 | +| `float` (FP32) | 2^-10 (≈ 9.77e-4) | 2^-16 (≈ 1.53e-5) | 0.99 | 1e-2 或 32×ULP | | ||
| 523 | +| `aclblasComplex` | 实部 / 虚部分别按 float 标准校验 | 同上 | 同上 | 同上 | | ||
| 524 | + | ||
| 525 | +## 支持芯片 | ||
| 526 | + | ||
| 527 | +| 芯片 | 架构 | 支持状态 | | ||
| 528 | +|------|------|----------| | ||
| 529 | +| Ascend 950PR | arch35 (DAV_3510) | 支持 | | ||
| 530 | +| Ascend 950DT | arch35 (DAV_3510) | 支持 | | ||
| 531 | +| Atlas A3 | — | 不支持 | | ||
| 532 | +| Atlas A2 | — | 不支持 | | ||
| @@ -0,0 +1,165 @@ | |||
| 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 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 8 | + */ | ||
| 9 | + | ||
| 10 | +/** | ||
| 11 | + * @file cgeam_host.cpp | ||
| 12 | + * @brief Host-side implementation for Cgeam operator (complex) | ||
| 13 | + * Handles parameter validation, tiling calculation, and kernel launch | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | +namespace { | ||
| 26 | + | ||
| 27 | +constexpr uint32_t UB_SIZE = 248 * 1024; // 248 KB | ||
| 28 | +constexpr uint32_t UB_RESERVE = 256; | ||
| 29 | +constexpr uint32_t ALIGN_UNIT = 8; // 32 bytes / sizeof(float) | ||
| 30 | +constexpr uint32_t CGEAM_BUFFERS = 12; // 4 TQue with num=2 (A_r*2, A_i*2, C_r*2, C_i*2) + 4 TBuf (B_r, B_i, calcBufR, calcBufI) | ||
| 31 | +constexpr uint32_t CGEAM_BUFFERS_BETA_ZERO = 10; // Same minus B_r, B_i (not allocated when beta=0) | ||
| 32 | +// Hardware limit for DataCopyPadExt blockCount parameter. | ||
| 33 | +// When curM exceeds this, DataCopyPad blockCount exceeds hardware limit | ||
| 34 | +// and behavior is undefined. Each complex element maps to one blockCount | ||
| 35 | +// unit (real/imag each have curM floats read separately). | ||
| 36 | +constexpr uint32_t MAX_BLOCK_COUNT = 4095; | ||
| 37 | + | ||
| 38 | +CgeamTilingData CalCgeamTilingData( | ||
| 39 | + aclblasOperation_t transa, aclblasOperation_t transb, uint32_t m, uint32_t n, float alphaR, float alphaI, | ||
| 40 | + float betaR, float betaI, uint32_t lda, uint32_t ldb, uint32_t ldc, uint32_t alphaIsZero, uint32_t betaIsZero, | ||
| 41 | + uint32_t aivCoreNum) | ||
| 42 | +{ | ||
| 43 | + CgeamTilingData tiling{}; | ||
| 44 | + tiling.m = m; | ||
| 45 | + tiling.n = n; | ||
| 46 | + tiling.lda = lda; | ||
| 47 | + tiling.ldb = ldb; | ||
| 48 | + tiling.ldc = ldc; | ||
| 49 | + tiling.alphaR = alphaR; | ||
| 50 | + tiling.alphaI = alphaI; | ||
| 51 | + tiling.betaR = betaR; | ||
| 52 | + tiling.betaI = betaI; | ||
| 53 | + tiling.alphaIsZero = alphaIsZero; | ||
| 54 | + tiling.betaIsZero = betaIsZero; | ||
| 55 | + tiling.opA = transa; | ||
| 56 | + tiling.opB = transb; | ||
| 57 | + | ||
| 58 | + // Calculate tileM based on UB capacity, capped to MAX_BLOCK_COUNT | ||
| 59 | + // to avoid triggering the kernel batch-loop error path when curM > 4095. | ||
| 60 | + // For complex with beta!=0: 8 buffers (A_r, A_i, B_r, B_i, C_r, C_i, calcBufR, calcBufI) | ||
| 61 | + // For complex with beta=0: 6 buffers (A_r, A_i, C_r, C_i, calcBufR, calcBufI) | ||
| 62 | + uint32_t buffersPerTile = betaIsZero ? CGEAM_BUFFERS_BETA_ZERO : CGEAM_BUFFERS; | ||
 | |||
| 63 | + uint32_t maxTileM = (UB_SIZE - UB_RESERVE) / (buffersPerTile * sizeof(float)); | ||
| 64 | + uint32_t tileM = (maxTileM / ALIGN_UNIT) * ALIGN_UNIT; | ||
| 65 | + tileM = std::min(tileM, MAX_BLOCK_COUNT); | ||
| 66 | + tiling.tileM = tileM; | ||
| 67 | + | ||
| 68 | + // 2D block decomposition: colBlocks x mBlocks | ||
| 69 | + uint32_t colBlocks = std::min(n, aivCoreNum); | ||
| 70 | + if (colBlocks == 0) { | ||
| 71 | + colBlocks = 1; | ||
| 72 | + } | ||
| 73 | + tiling.colBlocks = colBlocks; | ||
| 74 | + tiling.perCoreN = n / colBlocks; | ||
| 75 | + tiling.remainder = n % colBlocks; | ||
| 76 | + | ||
| 77 | + // Calculate m-blocks | ||
| 78 | + uint32_t totalMTiles = (tiling.tileM > 0) ? ((m + tiling.tileM - 1) / tiling.tileM) : 1; | ||
| 79 | + uint32_t mBlocks = 1; | ||
| 80 | + if (colBlocks < aivCoreNum && totalMTiles > 1) { | ||
| 81 | + uint32_t maxMBlocks = aivCoreNum / colBlocks; | ||
| 82 | + mBlocks = std::min(maxMBlocks, totalMTiles); | ||
| 83 | + } | ||
| 84 | + if (mBlocks == 0) // to suppress 'maybe divide by zero' warning | ||
| 85 | + __builtin_unreachable(); | ||
| 86 | + tiling.mBlocks = mBlocks; | ||
| 87 | + tiling.perCoreMTile = totalMTiles / mBlocks; | ||
| 88 | + tiling.mTileRemainder = totalMTiles % mBlocks; | ||
| 89 | + | ||
| 90 | + return tiling; | ||
| 91 | +} | ||
| 92 | + | ||
| 93 | +aclblasStatus_t LaunchCgeamKernel( | ||
| 94 | + aclblasHandle_t handle, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, | ||
| 95 | + const aclblasComplex* alpha, const aclblasComplex* A, int lda, const aclblasComplex* beta, const aclblasComplex* B, | ||
| 96 | + int ldb, aclblasComplex* C, int ldc) | ||
| 97 | +{ | ||
| 98 | + // Get core count | ||
| 99 | + uint32_t aivCoreNum = GetAivCoreCount(); | ||
| 100 | + if (aivCoreNum == 0) { | ||
| 101 | + OP_LOGE("aclblasCgeam", "Failed to get AIV core count"); | ||
| 102 | + return ACLBLAS_STATUS_INTERNAL_ERROR; | ||
| 103 | + } | ||
| 104 | + | ||
| 105 | + float alphaR = alpha->real; | ||
| 106 | + float alphaI = alpha->imag; | ||
| 107 | + float betaR = beta->real; | ||
| 108 | + float betaI = beta->imag; | ||
| 109 | + uint32_t alphaIsZero = (alphaR == 0.0f && alphaI == 0.0f) ? 1u : 0u; | ||
| 110 | + uint32_t betaIsZero = (betaR == 0.0f && betaI == 0.0f) ? 1u : 0u; | ||
| 111 | + | ||
| 112 | + // Calculate tiling | ||
| 113 | + CgeamTilingData tiling = CalCgeamTilingData( | ||
| 114 | + transa, transb, static_cast<uint32_t>(m), static_cast<uint32_t>(n), alphaR, alphaI, betaR, betaI, | ||
| 115 | + static_cast<uint32_t>(lda), static_cast<uint32_t>(ldb), static_cast<uint32_t>(ldc), alphaIsZero, betaIsZero, | ||
| 116 | + aivCoreNum); | ||
| 117 | + | ||
| 118 | + // Calculate total blocks | ||
| 119 | + uint32_t numBlocks = tiling.colBlocks * tiling.mBlocks; | ||
| 120 | + if (numBlocks == 0) { | ||
| 121 | + numBlocks = 1; | ||
| 122 | + } | ||
| 123 | + | ||
| 124 | + OP_LOGD( | ||
| 125 | + "aclblasCgeam", "Tiling: tileM=%u, colBlocks=%u, mBlocks=%u, totalBlocks=%u", tiling.tileM, tiling.colBlocks, | ||
| 126 | + tiling.mBlocks, numBlocks); | ||
| 127 | + | ||
| 128 | + // Convert pointers to GM_ADDR (complex data is stored as interleaved float pairs). | ||
| 129 | + // When alpha/beta value is 0, A/B may be nullptr; use C as a dummy valid GM address | ||
| 130 | + // to avoid hardware issues with null descriptors. | ||
| 131 | + // (Kernel skips reading A/B based on alphaIsZero/betaIsZero tiling flags.) | ||
| 132 | + GM_ADDR gmC = reinterpret_cast<GM_ADDR>(C); | ||
| 133 | + GM_ADDR gmA = (A != nullptr) ? reinterpret_cast<GM_ADDR>(const_cast<aclblasComplex*>(A)) : gmC; | ||
J 同 sgeam_host.cpp，为将 const aclblasComplex* 转为 GM_ADDR 做 const_cast。   | |||
| 134 | + GM_ADDR gmB = (B != nullptr) ? reinterpret_cast<GM_ADDR>(const_cast<aclblasComplex*>(B)) : gmC; | ||
| 135 | + | ||
| 136 | + // Launch kernel | ||
| 137 | + cgeam_kernel_do(gmA, gmB, gmC, tiling, numBlocks, handle->stream); | ||
| 138 | + | ||
| 139 | + return ACLBLAS_STATUS_SUCCESS; | ||
| 140 | +} | ||
| 141 | + | ||
| 142 | +} // namespace | ||
| 143 | + | ||
| 144 | +extern "C" aclblasStatus_t aclblasCgeam( | ||
| 145 | + aclblasHandle_t handle, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, | ||
| 146 | + const aclblasComplex* alpha, const aclblasComplex* A, int lda, const aclblasComplex* beta, const aclblasComplex* B, | ||
| 147 | + int ldb, aclblasComplex* C, int ldc) | ||
| 148 | +{ | ||
| 149 | + if (handle == nullptr) { | ||
| 150 | + OP_LOGE("aclblasCgeam", "handle is nullptr"); | ||
| 151 | + return ACLBLAS_STATUS_HANDLE_IS_NULLPTR; | ||
| 152 | + } | ||
| 153 | + | ||
| 154 | + aclblasStatus_t status = | ||
| 155 | + ValidateGeamParams("aclblasCgeam", transa, transb, m, n, alpha, A, lda, beta, B, ldb, C, ldc); | ||
| 156 | + if (status != ACLBLAS_STATUS_SUCCESS) { | ||
| 157 | + return status; | ||
| 158 | + } | ||
| 159 | + | ||
| 160 | + if (m == 0 || n == 0) { | ||
| 161 | + return ACLBLAS_STATUS_SUCCESS; | ||
| 162 | + } | ||
| 163 | + | ||
| 164 | + return LaunchCgeamKernel(handle, transa, transb, m, n, alpha, A, lda, beta, B, ldb, C, ldc); | ||
| 165 | +} | ||
| @@ -0,0 +1,453 @@ | |||
| 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 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 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 8 | + */ | ||
| 9 | + | ||
| 10 | +/** | ||
| 11 | + * @file cgeam_kernel.cpp | ||
| 12 | + * @brief Cgeam kernel implementation (SIMD/membase, arch35) | ||
| 13 | + * Computes: C = alpha * op(A) + beta * op(B) for complex matrices | ||
| 14 | + * | ||
| 15 | + * Architecture: TPipe + TQue<VECIN>/TQue<VECOUT> + TBuf<VECCALC> | ||
| 16 | + * EnQue/DeQue provides automatic MTE <-> Vector synchronization. | ||
| 17 | + * | ||
| 18 | + * Complex representation: | ||
| 19 | + * - Interleaved storage: r0,i0,r1,i1,...,rN,iN | ||
| 20 | + * - Each complex matrix element = 2 adjacent floats (real, imag) | ||
| 21 | + * - Separate real/imag buffers managed through individual TQue/TBuf | ||
| 22 | + * | ||
| 23 | + * Buffer planning (each buffer holds tileM floats): | ||
| 24 | + * - inQueueAR_, inQueueAI_ (VECIN): A tile real/imag, MTE2→V sync via TQue | ||
| 25 | + * - calcBufBR_, calcBufBI_ (VECCALC): B tile real/imag when beta!=0; scratch when beta==0 | ||
| 26 | + * - outQueueCR_, outQueueCI_ (VECOUT): C tile real/imag, V→MTE3 sync via TQue | ||
| 27 | + * - calcBufR_, calcBufI_ (VECCALC): Scratch for complex multiply intermediates | ||
| 28 | + * | ||
| 29 | + * Conjugate handling: | ||
| 30 | + * - isConjA: negate A imaginary after DeQue (PipeBarrier<PIPE_V> before compute) | ||
| 31 | + * - isConjB: negate B imaginary after load (PipeBarrier<PIPE_V> before compute) | ||
| 32 | + * | ||
| 33 | + * Single path: one column per iteration (complex interleaved prevents multi-col batching) | ||
| 34 | + * Multi-core decomposition: 2D grid (colBlocks x mBlocks) | ||
| 35 | + */ | ||
| 36 | + | ||
| 37 | + | ||
| 38 | + | ||
| 39 | + | ||
| 40 | + | ||
| 41 | +namespace { | ||
| 42 | + | ||
| 43 | +using namespace AscendC; | ||
| 44 | + | ||
| 45 | +// ============================================================================ | ||
| 46 | +// CgeamAIV: operator class for Cgeam complex vector kernel | ||
| 47 | +// ============================================================================ | ||
| 48 | +class CgeamAIV { | ||
| 49 | +public: | ||
| 50 | + __aicore__ inline CgeamAIV() {} | ||
| 51 | + __aicore__ inline void Init(GM_ADDR A, GM_ADDR B, GM_ADDR C, const CgeamTilingData& tiling, TPipe* pipe); | ||
| 52 | + __aicore__ inline void Process(); | ||
| 53 | + | ||
| 54 | +private: | ||
| 55 | + __aicore__ inline void ProcessTileLoop(); | ||
| 56 | + __aicore__ inline void ComputeCoreRanges(const CgeamTilingData& tiling); | ||
| 57 | + __aicore__ inline void CopyInCplxA( | ||
| 58 | + uint32_t col, uint32_t rowStart, uint32_t curM, int64_t ntCplxStride, int64_t tCplxStrideA); | ||
| 59 | + __aicore__ inline void CopyInCplxB( | ||
| 60 | + uint32_t col, uint32_t rowStart, uint32_t curM, int64_t ntCplxStride, int64_t tCplxStrideB); | ||
| 61 | + __aicore__ inline void ComputeCplx(uint32_t curM); | ||
| 62 | + __aicore__ inline void CopyOutCplx(uint32_t col, uint32_t rowStart, uint32_t curM, int64_t storeStride); | ||
| 63 | + | ||
| 64 | + // GM addresses | ||
| 65 | + GlobalTensor<float> aGm_; | ||
| 66 | + GlobalTensor<float> bGm_; | ||
| 67 | + GlobalTensor<float> cGm_; | ||
| 68 | + | ||
| 69 | + // A input queues (separate real/imag for clean TQue sync) | ||
| 70 | + TQue<TPosition::VECIN, 2> inQueueAR_; // A real tile (EnQue signals MTE2 done) | ||
| 71 | + TQue<TPosition::VECIN, 2> inQueueAI_; // A imag tile (EnQue signals MTE2 done) | ||
| 72 | + | ||
| 73 | + // B input: loaded via TBuf (same pattern as sgeam) | ||
| 74 | + TBuf<TPosition::VECCALC> calcBufBR_; // B real tile (also serves as scratch when beta==0) | ||
| 75 | + TBuf<TPosition::VECCALC> calcBufBI_; // B imag tile (also serves as scratch when beta==0) | ||
| 76 | + | ||
| 77 | + // C output queues | ||
| 78 | + TQue<TPosition::VECOUT, 2> outQueueCR_; // C real tile (EnQue signals Vector done) | ||
| 79 | + TQue<TPosition::VECOUT, 2> outQueueCI_; // C imag tile (EnQue signals Vector done) | ||
| 80 | + | ||
| 81 | + // Scratch for complex multiply | ||
| 82 | + TBuf<TPosition::VECCALC> calcBufR_; // Temp for alpha_i*A or beta_i*B real component | ||
| 83 | + TBuf<TPosition::VECCALC> calcBufI_; // Temp for alpha_i*A or beta_i*B imag component | ||
| 84 | + | ||
| 85 | + // Tiling fields cached from CgeamTilingData | ||
| 86 | + uint32_t m_; | ||
| 87 | + uint32_t n_; | ||
| 88 | + | ||
| 89 | + uint32_t alphaIsZero_; | ||
| 90 | + float alphaR_; | ||
| 91 | + float alphaI_; | ||
| 92 | + aclblasOperation_t opA_; | ||
| 93 | + uint32_t lda_; | ||
| 94 | + bool isConjA_; | ||
| 95 | + | ||
| 96 | + uint32_t betaIsZero_; | ||
| 97 | + float betaR_; | ||
| 98 | + float betaI_; | ||
| 99 | + aclblasOperation_t opB_; | ||
| 100 | + uint32_t ldb_; | ||
| 101 | + bool isConjB_; | ||
| 102 | + | ||
| 103 | + uint32_t ldc_; | ||
| 104 | + uint32_t tileM_; | ||
| 105 | + | ||
| 106 | + // Per-core tile ranges | ||
| 107 | + uint32_t colStart_; | ||
| 108 | + uint32_t colEnd_; | ||
| 109 | + uint32_t mTileStart_; | ||
| 110 | + uint32_t mTileEnd_; | ||
| 111 | +}; | ||
| 112 | + | ||
| 113 | +// ============================================================================ | ||
| 114 | +// ComputeCoreRanges: decode 2D block index into per-core column and m-tile ranges. | ||
| 115 | +// On invalid range, sets colEnd_ or mTileEnd_ to 0 to signal early exit in Process(). | ||
| 116 | +// ============================================================================ | ||
| 117 | +__aicore__ inline void CgeamAIV::ComputeCoreRanges(const CgeamTilingData& tiling) | ||
J ComputeCoreRanges 的逻辑与 sgeam_kernel.cpp:103 几乎完全相同（仅 TilingData 类型不同），可以考虑模板化或提取到公共头文件以减少重复。sgeam 和 cgeam 的 tiling 计算也有类似重复。 | |||
| 118 | +{ | ||
| 119 | + uint32_t blockIdx = GetBlockIdx(); | ||
| 120 | + uint32_t mBlocks = tiling.mBlocks; | ||
| 121 | + uint32_t colBlock = blockIdx / mBlocks; | ||
| 122 | + uint32_t mBlock = blockIdx % mBlocks; | ||
| 123 | + | ||
| 124 | + // Calculate column range | ||
| 125 | + if (colBlock < tiling.remainder) { | ||
| 126 | + colStart_ = colBlock * (tiling.perCoreN + 1); | ||
| 127 | + colEnd_ = colStart_ + tiling.perCoreN + 1; | ||
| 128 | + } else { | ||
| 129 | + colStart_ = colBlock * tiling.perCoreN + tiling.remainder; | ||
| 130 | + colEnd_ = colStart_ + tiling.perCoreN; | ||
| 131 | + } | ||
| 132 | + | ||
| 133 | + // Calculate m-tile range | ||
| 134 | + if (mBlock < tiling.mTileRemainder) { | ||
| 135 | + mTileStart_ = mBlock * (tiling.perCoreMTile + 1); | ||
| 136 | + mTileEnd_ = mTileStart_ + tiling.perCoreMTile + 1; | ||
| 137 | + } else { | ||
| 138 | + mTileStart_ = mBlock * tiling.perCoreMTile + tiling.mTileRemainder; | ||
| 139 | + mTileEnd_ = mTileStart_ + tiling.perCoreMTile; | ||
| 140 | + } | ||
| 141 | + | ||
| 142 | + // Validate column range | ||
| 143 | + if (colStart_ >= n_ || colStart_ >= colEnd_) { | ||
| 144 | + colEnd_ = 0; // Signal Process() to return early | ||
| 145 | + return; | ||
| 146 | + } | ||
| 147 | + if (colEnd_ > n_) { | ||
| 148 | + colEnd_ = n_; | ||
| 149 | + } | ||
| 150 | + | ||
| 151 | + // Validate m-tile range | ||
| 152 | + uint32_t totalMTiles = (m_ + tileM_ - 1) / tileM_; | ||
| 153 | + if (mTileStart_ >= totalMTiles || mTileStart_ >= mTileEnd_) { | ||
| 154 | + mTileEnd_ = 0; // Signal Process() to return early | ||
| 155 | + return; | ||
| 156 | + } | ||
| 157 | + if (mTileEnd_ > totalMTiles) { | ||
| 158 | + mTileEnd_ = totalMTiles; | ||
| 159 | + } | ||
| 160 | +} | ||
| 161 | + | ||
| 162 | +// ============================================================================ | ||
| 163 | +// Init: setup global tensors, compute per-core ranges, allocate UB buffers. | ||
| 164 | +// TPipe pointer passed from kernel entry (R3: TPipe must not be a member). | ||
| 165 | +// ============================================================================ | ||
| 166 | +__aicore__ inline void CgeamAIV::Init(GM_ADDR A, GM_ADDR B, GM_ADDR C, const CgeamTilingData& tiling, TPipe* pipe) | ||
| 167 | +{ | ||
| 168 | + // Cache tiling fields to member variables | ||
| 169 | + m_ = tiling.m; | ||
| 170 | + n_ = tiling.n; | ||
| 171 | + lda_ = tiling.lda; | ||
| 172 | + ldb_ = tiling.ldb; | ||
| 173 | + ldc_ = tiling.ldc; | ||
| 174 | + tileM_ = tiling.tileM; | ||
| 175 | + opA_ = tiling.opA; | ||
| 176 | + opB_ = tiling.opB; | ||
| 177 | + alphaIsZero_ = tiling.alphaIsZero; | ||
| 178 | + betaIsZero_ = tiling.betaIsZero; | ||
| 179 | + alphaR_ = tiling.alphaR; | ||
| 180 | + alphaI_ = tiling.alphaI; | ||
| 181 | + betaR_ = tiling.betaR; | ||
| 182 | + betaI_ = tiling.betaI; | ||
| 183 | + isConjA_ = (opA_ == ACLBLAS_OP_C); | ||
| 184 | + isConjB_ = (opB_ == ACLBLAS_OP_C); | ||
| 185 | + | ||
| 186 | + // Set GM buffers | ||
| 187 | + aGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(A)); | ||
| 188 | + bGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(B)); | ||
| 189 | + cGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(C)); | ||
| 190 | + | ||
| 191 | + // Compute per-core column and m-tile ranges | ||
| 192 | + ComputeCoreRanges(tiling); | ||
J sgeam Init 在 ComputeCoreRanges 后有 | |||
| 193 | + if (colEnd_ == 0 || mTileEnd_ == 0) { | ||
| 194 | + return; | ||
| 195 | + } | ||
| 196 | + | ||
| 197 | + // Buffer size: tileM floats per buffer | ||
| 198 | + uint32_t bufSize = tileM_ * sizeof(float); | ||
| 199 | + | ||
| 200 | + // A input queues (double-buffer, num=2 for MTE2<->V pipeline overlap) | ||
| 201 | + pipe->InitBuffer(inQueueAR_, 2, bufSize); | ||
| 202 | + pipe->InitBuffer(inQueueAI_, 2, bufSize); | ||
| 203 | + | ||
| 204 | + // C output queues (double-buffer, num=2 for V<->MTE3 pipeline overlap) | ||
| 205 | + pipe->InitBuffer(outQueueCR_, 2, bufSize); | ||
| 206 | + pipe->InitBuffer(outQueueCI_, 2, bufSize); | ||
| 207 | + | ||
| 208 | + // B input (TBuf, allocated when beta != 0) | ||
| 209 | + if (!betaIsZero_) { | ||
| 210 | + pipe->InitBuffer(calcBufBR_, bufSize); | ||
| 211 | + pipe->InitBuffer(calcBufBI_, bufSize); | ||
| 212 | + } | ||
| 213 | + | ||
| 214 | + // Complex multiply scratch (only needed when alpha != 0) | ||
| 215 | + if (!alphaIsZero_) { | ||
| 216 | + pipe->InitBuffer(calcBufR_, bufSize); | ||
| 217 | + pipe->InitBuffer(calcBufI_, bufSize); | ||
| 218 | + } | ||
| 219 | +} | ||
| 220 | + | ||
| 221 | +// ============================================================================ | ||
| 222 | +// Process: validate ranges, dispatch to tile loop | ||
| 223 | +// ============================================================================ | ||
| 224 | +__aicore__ inline void CgeamAIV::Process() | ||
| 225 | +{ | ||
| 226 | + if (colStart_ >= colEnd_ || mTileStart_ >= mTileEnd_) { | ||
| 227 | + return; | ||
| 228 | + } | ||
| 229 | + ProcessTileLoop(); | ||
| 230 | +} | ||
| 231 | + | ||
| 232 | +// ============================================================================ | ||
| 233 | +// CopyInCplxA: load A real and imag tiles through their respective VECIN queues. | ||
| 234 | +// ============================================================================ | ||
| 235 | +__aicore__ inline void CgeamAIV::CopyInCplxA( | ||
| 236 | + uint32_t col, uint32_t rowStart, uint32_t curM, int64_t ntCplxStride, int64_t tCplxStrideA) | ||
| 237 | +{ | ||
| 238 | + // CopyIn A real (MTE2) | ||
| 239 | + LocalTensor<float> aR = inQueueAR_.AllocTensor<float>(); | ||
| 240 | + if (!alphaIsZero_) { | ||
| 241 | + uint64_t offsetBase; | ||
| 242 | + int64_t stride; | ||
| 243 | + if (opA_ == ACLBLAS_OP_N) { | ||
| 244 | + offsetBase = static_cast<uint64_t>(col) * static_cast<uint64_t>(lda_) + rowStart; | ||
| 245 | + stride = ntCplxStride; | ||
| 246 | + } else { | ||
| 247 | + offsetBase = static_cast<uint64_t>(col) + static_cast<uint64_t>(rowStart) * static_cast<uint64_t>(lda_); | ||
| 248 | + stride = tCplxStrideA; | ||
| 249 | + } | ||
| 250 | + DataCopyExtParams cp{static_cast<uint16_t>(curM), static_cast<uint32_t>(sizeof(float)), stride, 0, 0}; | ||
| 251 | + DataCopyPadExtParams<float> np; | ||
| 252 | + DataCopyPad<float, PaddingMode::Compact>(aR, aGm_[2 * offsetBase], cp, np); | ||
| 253 | + } | ||
| 254 | + inQueueAR_.EnQue(aR); // Signal MTE2 done for A real | ||
| 255 | + | ||
| 256 | + // CopyIn A imag (MTE2) | ||
| 257 | + LocalTensor<float> aI = inQueueAI_.AllocTensor<float>(); | ||
| 258 | + if (!alphaIsZero_) { | ||
| 259 | + uint64_t offsetBase; | ||
| 260 | + int64_t stride; | ||
| 261 | + if (opA_ == ACLBLAS_OP_N) { | ||
| 262 | + offsetBase = static_cast<uint64_t>(col) * static_cast<uint64_t>(lda_) + rowStart; | ||
| 263 | + stride = ntCplxStride; | ||
| 264 | + } else { | ||
| 265 | + offsetBase = static_cast<uint64_t>(col) + static_cast<uint64_t>(rowStart) * static_cast<uint64_t>(lda_); | ||
| 266 | + stride = tCplxStrideA; | ||
| 267 | + } | ||
| 268 | + DataCopyExtParams cp{static_cast<uint16_t>(curM), static_cast<uint32_t>(sizeof(float)), stride, 0, 0}; | ||
| 269 | + DataCopyPadExtParams<float> np; | ||
| 270 | + DataCopyPad<float, PaddingMode::Compact>(aI, aGm_[2 * offsetBase + 1], cp, np); | ||
| 271 | + } | ||
| 272 | + inQueueAI_.EnQue(aI); // Signal MTE2 done for A imag | ||
| 273 | +} | ||
| 274 | + | ||
| 275 | +// ============================================================================ | ||
| 276 | +// CopyInCplxB: load B real and imag tiles, and apply conjugate if needed. | ||
| 277 | +// ============================================================================ | ||
| 278 | +__aicore__ inline void CgeamAIV::CopyInCplxB( | ||
| 279 | + uint32_t col, uint32_t rowStart, uint32_t curM, int64_t ntCplxStride, int64_t tCplxStrideB) | ||
| 280 | +{ | ||
| 281 | + if (!betaIsZero_) { | ||
| 282 | + LocalTensor<float> bR = calcBufBR_.Get<float>(); | ||
| 283 | + LocalTensor<float> bI = calcBufBI_.Get<float>(); | ||
| 284 | + | ||
| 285 | + uint64_t offsetBase; | ||
| 286 | + int64_t stride; | ||
| 287 | + if (opB_ == ACLBLAS_OP_N) { | ||
| 288 | + offsetBase = static_cast<uint64_t>(col) * static_cast<uint64_t>(ldb_) + rowStart; | ||
| 289 | + stride = ntCplxStride; | ||
| 290 | + } else { | ||
| 291 | + offsetBase = static_cast<uint64_t>(col) + static_cast<uint64_t>(rowStart) * static_cast<uint64_t>(ldb_); | ||
| 292 | + stride = tCplxStrideB; | ||
| 293 | + } | ||
| 294 | + | ||
| 295 | + DataCopyExtParams cp{static_cast<uint16_t>(curM), static_cast<uint32_t>(sizeof(float)), stride, 0, 0}; | ||
| 296 | + DataCopyPadExtParams<float> np; | ||
| 297 | + DataCopyPad<float, PaddingMode::Compact>(bR, bGm_[2 * offsetBase], cp, np); | ||
| 298 | + DataCopyPad<float, PaddingMode::Compact>(bI, bGm_[2 * offsetBase + 1], cp, np); | ||
B 实/虚部经 DataCopyPad 写入 calcBufBR_/calcBufBI_ 后由 ComputeCplx 直接读取，无显式 MTE2→V 同步，且第 297 行共轭取反 Muls 紧跟在异步搬运之后，依赖隐式时序保障（实机大矩阵用例未复现错误），后续调整时易暴露竞争 | |||
| 299 | + event_t eMte2V_ = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V)); | ||
| 300 | + SetFlag<HardEvent::MTE2_V>(eMte2V_); | ||
| 301 | + WaitFlag<HardEvent::MTE2_V>(eMte2V_); | ||
| 302 | + | ||
| 303 | + if (isConjB_) { | ||
| 304 | + Muls<float>(bI, bI, -1.0f, curM); | ||
| 305 | + PipeBarrier<PIPE_V>(); | ||
| 306 | + } | ||
| 307 | + } | ||
| 308 | +} | ||
| 309 | + | ||
| 310 | +// ============================================================================ | ||
| 311 | +// ComputeCplx: Phase 1 (C = alpha * A) + Phase 2 (C += beta * B) complex multiply. | ||
| 312 | +// ============================================================================ | ||
| 313 | +__aicore__ inline void CgeamAIV::ComputeCplx(uint32_t curM) | ||
| 314 | +{ | ||
| 315 | + // DeQue A waits for MTE2 -> V sync | ||
| 316 | + LocalTensor<float> aRv = inQueueAR_.DeQue<float>(); | ||
| 317 | + LocalTensor<float> aIv = inQueueAI_.DeQue<float>(); | ||
| 318 | + | ||
| 319 | + // Conjugate A: negate imaginary part after DeQue | ||
| 320 | + if (isConjA_) { | ||
| 321 | + Muls<float>(aIv, aIv, -1.0f, curM); | ||
| 322 | + PipeBarrier<PIPE_V>(); | ||
| 323 | + } | ||
| 324 | + | ||
| 325 | + LocalTensor<float> cR = outQueueCR_.AllocTensor<float>(); | ||
| 326 | + LocalTensor<float> cI = outQueueCI_.AllocTensor<float>(); | ||
| 327 | + | ||
| 328 | + if (alphaIsZero_) { | ||
| 329 | + // Zero-initialize C (alpha == 0: A not referenced) | ||
| 330 | + Duplicate<float>(cR, 0.0f, curM); | ||
| 331 | + Duplicate<float>(cI, 0.0f, curM); | ||
| 332 | + } else { | ||
| 333 | + // Scratch buffers for complex multiply | ||
| 334 | + LocalTensor<float> tmpR = calcBufR_.Get<float>(); | ||
| 335 | + LocalTensor<float> tmpI = calcBufI_.Get<float>(); | ||
| 336 | + | ||
| 337 | + // Phase 1: C = alpha * A (complex multiplication) | ||
| 338 | + // C_r = alphaR * A_r - alphaI * A_i | ||
| 339 | + // C_i = alphaR * A_i + alphaI * A_r | ||
| 340 | + Muls<float>(tmpR, aRv, alphaI_, curM); // tmpR = alphaI * A_r | ||
| 341 | + Muls<float>(tmpI, aIv, alphaI_, curM); // tmpI = alphaI * A_i | ||
| 342 | + Muls<float>(cR, aRv, alphaR_, curM); // cR = alphaR * A_r | ||
| 343 | + Muls<float>(cI, aIv, alphaR_, curM); // cI = alphaR * A_i | ||
| 344 | + Sub<float>(cR, cR, tmpI, curM); // cR -= alphaI * A_i | ||
| 345 | + Add<float>(cI, cI, tmpR, curM); // cI += alphaI * A_r | ||
| 346 | + } | ||
| 347 | + | ||
| 348 | + if (!betaIsZero_) { | ||
| 349 | + // Phase 2: C += beta * B (complex multiplication) | ||
| 350 | + LocalTensor<float> bR = calcBufBR_.Get<float>(); | ||
| 351 | + LocalTensor<float> bI = calcBufBI_.Get<float>(); | ||
| 352 | + // When alpha != 0, use calcBufR_/I_ as scratch; when alpha == 0, | ||
| 353 | + // reuse A input buffers (aRv/aIv) as scratch since Phase 1 was skipped. | ||
| 354 | + LocalTensor<float> tmpR = alphaIsZero_ ? aRv : calcBufR_.Get<float>(); | ||
| 355 | + LocalTensor<float> tmpI = alphaIsZero_ ? aIv : calcBufI_.Get<float>(); | ||
| 356 | + // C_r += betaR * B_r - betaI * B_i | ||
| 357 | + // C_i += betaR * B_i + betaI * B_r | ||
| 358 | + Muls<float>(tmpR, bR, betaR_, curM); // tmpR = betaR * B_r | ||
| 359 | + Add<float>(cR, cR, tmpR, curM); // cR += betaR * B_r | ||
| 360 | + Muls<float>(tmpI, bI, betaR_, curM); // tmpI = betaR * B_i | ||
| 361 | + Add<float>(cI, cI, tmpI, curM); // cI += betaR * B_i | ||
| 362 | + Muls<float>(tmpR, bR, betaI_, curM); // tmpR = betaI * B_r | ||
| 363 | + Add<float>(cI, cI, tmpR, curM); // cI += betaI * B_r | ||
| 364 | + Muls<float>(tmpI, bI, betaI_, curM); // tmpI = betaI * B_i | ||
| 365 | + Sub<float>(cR, cR, tmpI, curM); // cR -= betaI * B_i | ||
| 366 | + } | ||
| 367 | + | ||
| 368 | + outQueueCR_.EnQue(cR); // Signal Vector done for C real | ||
| 369 | + outQueueCI_.EnQue(cI); // Signal Vector done for C imag | ||
| 370 | + inQueueAR_.FreeTensor(aRv); | ||
| 371 | + inQueueAI_.FreeTensor(aIv); | ||
| 372 | +} | ||
| 373 | + | ||
| 374 | +// ============================================================================ | ||
| 375 | +// CopyOutCplx: store C real and imag tiles to global memory (interleaved). | ||
| 376 | +// ============================================================================ | ||
| 377 | +__aicore__ inline void CgeamAIV::CopyOutCplx(uint32_t col, uint32_t rowStart, uint32_t curM, int64_t storeStride) | ||
| 378 | +{ | ||
| 379 | + // CopyOut C real (MTE3): DeQue waits for V -> MTE3 sync | ||
| 380 | + LocalTensor<float> cRo = outQueueCR_.DeQue<float>(); | ||
| 381 | + { | ||
| 382 | + const uint64_t offsetBase = static_cast<uint64_t>(col) * static_cast<uint64_t>(ldc_) + rowStart; | ||
| 383 | + DataCopyExtParams cp{static_cast<uint16_t>(curM), static_cast<uint32_t>(sizeof(float)), 0, storeStride, 0}; | ||
| 384 | + DataCopyPad<float, PaddingMode::Compact>(cGm_[2 * offsetBase], cRo, cp); | ||
| 385 | + } | ||
| 386 | + outQueueCR_.FreeTensor(cRo); | ||
| 387 | + | ||
| 388 | + // CopyOut C imag (MTE3) | ||
| 389 | + LocalTensor<float> cIo = outQueueCI_.DeQue<float>(); | ||
| 390 | + { | ||
| 391 | + const uint64_t offsetBase = static_cast<uint64_t>(col) * static_cast<uint64_t>(ldc_) + rowStart; | ||
| 392 | + DataCopyExtParams cp{static_cast<uint16_t>(curM), static_cast<uint32_t>(sizeof(float)), 0, storeStride, 0}; | ||
| 393 | + DataCopyPad<float, PaddingMode::Compact>(cGm_[2 * offsetBase + 1], cIo, cp); | ||
| 394 | + } | ||
| 395 | + outQueueCI_.FreeTensor(cIo); | ||
| 396 | +} | ||
| 397 | + | ||
| 398 | +// ============================================================================ | ||
| 399 | +// ProcessTileLoop: complex geam tile processing | ||
| 400 | +// Each tile: CopyIn A -> EnQue -> CopyIn B -> Compute -> EnQue C -> CopyOut C | ||
| 401 | +// ============================================================================ | ||
| 402 | +__aicore__ inline void CgeamAIV::ProcessTileLoop() | ||
| 403 | +{ | ||
| 404 | + // Pre-compute strides for complex strided access | ||
| 405 | + // NoTrans complex: adjacent r (or i) elements are 2 floats apart → srcStride = sizeof(float) | ||
| 406 | + const int64_t ntCplxStride = static_cast<int64_t>(sizeof(float)); | ||
| 407 | + // Trans/ConjTrans complex: each row = ld complex = 2*ld floats → gap = (2*ld-1)*sizeof(float) | ||
| 408 | + const int64_t tCplxStrideA = static_cast<int64_t>(2 * lda_ - 1) * static_cast<int64_t>(sizeof(float)); | ||
| 409 | + const int64_t tCplxStrideB = static_cast<int64_t>(2 * ldb_ - 1) * static_cast<int64_t>(sizeof(float)); | ||
| 410 | + // Store: interleaved complex write, dstStride gap = sizeof(float) | ||
| 411 | + const int64_t storeStride = static_cast<int64_t>(sizeof(float)); | ||
| 412 | + | ||
| 413 | + for (uint32_t col = colStart_; col < colEnd_; ++col) { | ||
| 414 | + for (uint32_t mTile = mTileStart_; mTile < mTileEnd_; ++mTile) { | ||
| 415 | + const uint32_t rowStart = mTile * tileM_; | ||
| 416 | + uint32_t rowEnd = rowStart + tileM_; | ||
| 417 | + if (rowEnd > m_) { | ||
| 418 | + rowEnd = m_; | ||
| 419 | + } | ||
| 420 | + const uint32_t curM = rowEnd - rowStart; | ||
| 421 | + | ||
| 422 | + CopyInCplxA(col, rowStart, curM, ntCplxStride, tCplxStrideA); | ||
| 423 | + CopyInCplxB(col, rowStart, curM, ntCplxStride, tCplxStrideB); | ||
| 424 | + ComputeCplx(curM); | ||
J calcBufBR_/calcBufBI_ 是 TBuf，CopyInCplxB 通过 DataCopyPad（MTE2）加载 B 的实部/虚部，但 ComputeCplx 中 Vector 读取 bR/bI（第343-344行）时缺少 MTE2→V 同步。当 isConjB_ 为真时，第297行 Muls 读取 bI 也没有 MTE2→V 同步，PipeBarrier<PIPE_V> 只保证 V→V 顺序，不覆盖 MTE2→V。建议在 CopyInCplxB 后、ComputeCplx 前插入 PipeBarrier<PIPE_MTE2_V>。 | |||
| 425 | + CopyOutCplx(col, rowStart, curM, storeStride); | ||
| 426 | + } | ||
| 427 | + } | ||
| 428 | +} | ||
| 429 | + | ||
| 430 | +} // namespace | ||
| 431 | + | ||
| 432 | +// ============================================================================ | ||
| 433 | +// Kernel entry point | ||
| 434 | +// ============================================================================ | ||
| 435 | +extern "C" __global__ __aicore__ void cgeam_kernel(GM_ADDR A, GM_ADDR B, GM_ADDR C, CgeamTilingData tiling) | ||
| 436 | +{ | ||
| 437 | + KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); | ||
| 438 | + | ||
| 439 | + if (tiling.m == 0 || tiling.n == 0) { | ||
| 440 | + return; | ||
| 441 | + } | ||
| 442 | + | ||
| 443 | + TPipe pipe; | ||
| 444 | + CgeamAIV op; | ||
| 445 | + op.Init(A, B, C, tiling, &pipe); | ||
| 446 | + op.Process(); | ||
| 447 | +} | ||
| 448 | + | ||
| 449 | +void cgeam_kernel_do( | ||
| 450 | + GM_ADDR A, GM_ADDR B, GM_ADDR C, const CgeamTilingData& tiling, uint32_t numBlocks, aclrtStream stream) | ||
| 451 | +{ | ||
| 452 | + cgeam_kernel<<<numBlocks, nullptr, stream>>>(A, B, C, tiling); | ||
| 453 | +} | ||
| @@ -0,0 +1,37 @@ | |||
| 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 cgeam_kernel.h | ||
| 13 | + * @brief Kernel launcher declaration for Cgeam operator (complex) | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | +/** | ||
| 27 | + * @brief Launch Cgeam kernel asynchronously | ||
| 28 | + * | ||
| 29 | + * @param A GM address of matrix A (complex) | ||
| 30 | + * @param B GM address of matrix B (complex) | ||
| 31 | + * @param C GM address of matrix C (complex, output) | ||
| 32 | + * @param tiling Host-computed tiling data (copied to device by value) | ||
| 33 | + * @param numBlocks Block count for the <<<>>> launch | ||
| 34 | + * @param stream aclrtStream handle | ||
| 35 | + */ | ||
| 36 | +void cgeam_kernel_do( | ||
| 37 | + GM_ADDR A, GM_ADDR B, GM_ADDR C, const CgeamTilingData& tiling, uint32_t numBlocks, aclrtStream stream); | ||
| @@ -0,0 +1,44 @@ | |||
| 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 cgeam_tiling_data.h | ||
| 13 | + * @brief Tiling data structure for Cgeam operator (host/device shared) | ||
| 14 | + * C = alpha * op(A) + beta * op(B) (complex) | ||
| 15 | + */ | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | +struct CgeamTilingData { | ||
| 23 | + uint32_t m; // Number of rows in output matrix C | ||
| 24 | + uint32_t n; // Number of columns in output matrix C | ||
| 25 | + float alphaR; // Real part of scalar alpha | ||
| 26 | + float alphaI; // Imaginary part of scalar alpha | ||
| 27 | + uint32_t alphaIsZero; // Optimized flag indicating alpha == 0 | ||
| 28 | + aclblasOperation_t opA; // Transpose type for matrix A | ||
| 29 | + uint32_t lda; // Leading dimension of matrix A | ||
| 30 | + float betaR; // Real part of scalar beta | ||
| 31 | + float betaI; // Imaginary part of scalar beta | ||
| 32 | + uint32_t betaIsZero; // Optimized flag indicating beta == 0 | ||
| 33 | + aclblasOperation_t opB; // Transpose type for matrix B | ||
| 34 | + uint32_t ldb; // Leading dimension of matrix B | ||
| 35 | + uint32_t ldc; // Leading dimension of matrix C | ||
| 36 | + // 2D multi-core tiling: colBlocks x mBlocks | ||
| 37 | + uint32_t colBlocks; // Number of column blocks | ||
| 38 | + uint32_t perCoreN; // Number of columns processed per core | ||
| 39 | + uint32_t remainder; // Column remainder | ||
| 40 | + uint32_t mBlocks; // Number of row blocks | ||
| 41 | + uint32_t perCoreMTile; // Number of tiles processed per core in m dimension | ||
| 42 | + uint32_t mTileRemainder; // Tile count remainder in m dimension | ||
| 43 | + uint32_t tileM; // Tile size in m dimension (number of elements) | ||
| 44 | +}; | ||
| @@ -0,0 +1,150 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | +inline bool IsGeamScalarZero(const float* p) { return p == nullptr || *p == 0.0f; } | ||
IsGeamScalarZero 将 alpha/beta 空指针视为 0，与仓内 gemm 对空标量指针返回 INVALID_VALUE 的约定不一致，跨算子错误码语义不统一 | |||
| 18 | + | ||
| 19 | +inline bool IsGeamScalarZero(const aclblasComplex* p) { return p == nullptr || (p->real == 0.0f && p->imag == 0.0f); } | ||
| 20 | + | ||
| 21 | +inline aclblasStatus_t ValidateGeamTranspose( | ||
| 22 | + aclblasOperation_t transa, aclblasOperation_t transb, const std::string& op_name) | ||
| 23 | +{ | ||
| 24 | + if (transa != ACLBLAS_OP_N && transa != ACLBLAS_OP_T && transa != ACLBLAS_OP_C) { | ||
| 25 | + OP_LOGE(op_name, "Invalid transa=%d (must be N/T/C)", transa); | ||
| 26 | + return ACLBLAS_STATUS_INVALID_ENUM; | ||
| 27 | + } | ||
| 28 | + if (transb != ACLBLAS_OP_N && transb != ACLBLAS_OP_T && transb != ACLBLAS_OP_C) { | ||
| 29 | + OP_LOGE(op_name, "Invalid transb=%d (must be N/T/C)", transb); | ||
| 30 | + return ACLBLAS_STATUS_INVALID_ENUM; | ||
| 31 | + } | ||
| 32 | + return ACLBLAS_STATUS_SUCCESS; | ||
| 33 | +} | ||
| 34 | + | ||
| 35 | +inline aclblasStatus_t ValidateGeamLd( | ||
| 36 | + int m, int n, int lda, int ldb, int ldc, aclblasOperation_t transa, aclblasOperation_t transb, | ||
| 37 | + const std::string& op_name) | ||
| 38 | +{ | ||
| 39 | + if (transa == ACLBLAS_OP_N) { | ||
| 40 | + if (lda < std::max(1, m)) { | ||
| 41 | + OP_LOGE(op_name, "lda=%d must be >= max(1, m=%d) when transa=N", lda, m); | ||
| 42 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 43 | + } | ||
| 44 | + } else { | ||
| 45 | + if (lda < std::max(1, n)) { | ||
| 46 | + OP_LOGE(op_name, "lda=%d must be >= max(1, n=%d) when transa=T/C", lda, n); | ||
| 47 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 48 | + } | ||
| 49 | + } | ||
| 50 | + if (transb == ACLBLAS_OP_N) { | ||
| 51 | + if (ldb < std::max(1, m)) { | ||
| 52 | + OP_LOGE(op_name, "ldb=%d must be >= max(1, m=%d) when transb=N", ldb, m); | ||
| 53 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 54 | + } | ||
| 55 | + } else { | ||
| 56 | + if (ldb < std::max(1, n)) { | ||
| 57 | + OP_LOGE(op_name, "ldb=%d must be >= max(1, n=%d) when transb=T/C", ldb, n); | ||
| 58 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 59 | + } | ||
| 60 | + } | ||
| 61 | + if (ldc < std::max(1, m)) { | ||
| 62 | + OP_LOGE(op_name, "ldc=%d must be >= max(1, m=%d)", ldc, m); | ||
| 63 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 64 | + } | ||
| 65 | + return ACLBLAS_STATUS_SUCCESS; | ||
| 66 | +} | ||
| 67 | + | ||
| 68 | +template <class T> | ||
| 69 | +inline aclblasStatus_t ValidateGeamPointers( | ||
| 70 | + const T* alpha, const T* A, const T* beta, const T* B, T* C, const std::string& op_name) | ||
| 71 | +{ | ||
| 72 | + if (alpha == nullptr) { | ||
| 73 | + OP_LOGE(op_name, "alpha pointer is nullptr"); | ||
| 74 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 75 | + } | ||
| 76 | + if (beta == nullptr) { | ||
| 77 | + OP_LOGE(op_name, "beta pointer is nullptr"); | ||
| 78 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 79 | + } | ||
| 80 | + if (C == nullptr) { | ||
| 81 | + OP_LOGE(op_name, "C pointer is nullptr"); | ||
| 82 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 83 | + } | ||
| 84 | + if (!IsGeamScalarZero(alpha) && A == nullptr) { | ||
| 85 | + OP_LOGE(op_name, "A pointer is nullptr (alpha != 0)"); | ||
| 86 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 87 | + } | ||
| 88 | + if (!IsGeamScalarZero(beta) && B == nullptr) { | ||
| 89 | + OP_LOGE(op_name, "B pointer is nullptr (beta != 0)"); | ||
| 90 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 91 | + } | ||
| 92 | + return ACLBLAS_STATUS_SUCCESS; | ||
| 93 | +} | ||
| 94 | + | ||
| 95 | +template <class T> | ||
| 96 | +inline aclblasStatus_t ValidateGeamInplace( | ||
| 97 | + const T* A, const T* B, T* C, int lda, int ldb, int ldc, aclblasOperation_t transa, aclblasOperation_t transb, | ||
| 98 | + const std::string& op_name) | ||
| 99 | +{ | ||
| 100 | + if (A != nullptr && C == A) { | ||
| 101 | + if (transa != ACLBLAS_OP_N) { | ||
| 102 | + OP_LOGE(op_name, "In-place C==A requires transa==N (got %d)", static_cast<int>(transa)); | ||
| 103 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 104 | + } | ||
| 105 | + if (lda != ldc) { | ||
| 106 | + OP_LOGE(op_name, "In-place C==A requires lda==ldc (lda=%d, ldc=%d)", lda, ldc); | ||
| 107 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 108 | + } | ||
| 109 | + } | ||
| 110 | + if (B != nullptr && C == B) { | ||
| 111 | + if (transb != ACLBLAS_OP_N) { | ||
| 112 | + OP_LOGE(op_name, "In-place C==B requires transb==N (got %d)", static_cast<int>(transb)); | ||
| 113 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 114 | + } | ||
| 115 | + if (ldb != ldc) { | ||
| 116 | + OP_LOGE(op_name, "In-place C==B requires ldb==ldc (ldb=%d, ldc=%d)", ldb, ldc); | ||
| 117 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 118 | + } | ||
| 119 | + } | ||
| 120 | + return ACLBLAS_STATUS_SUCCESS; | ||
| 121 | +} | ||
| 122 | + | ||
| 123 | +template <class T> | ||
| 124 | +inline aclblasStatus_t ValidateGeamParams( | ||
| 125 | + const std::string& op_name, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, const T* alpha, | ||
| 126 | + const T* A, int lda, const T* beta, const T* B, int ldb, T* C, int ldc) | ||
| 127 | +{ | ||
| 128 | + aclblasStatus_t status; | ||
| 129 | + status = ValidateGeamTranspose(transa, transb, op_name); | ||
| 130 | + if (status != ACLBLAS_STATUS_SUCCESS) { | ||
| 131 | + return status; | ||
| 132 | + } | ||
| 133 | + if (m < 0) { | ||
| 134 | + OP_LOGE(op_name, "m=%d must be >= 0", m); | ||
| 135 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 136 | + } | ||
| 137 | + if (n < 0) { | ||
| 138 | + OP_LOGE(op_name, "n=%d must be >= 0", n); | ||
| 139 | + return ACLBLAS_STATUS_INVALID_VALUE; | ||
| 140 | + } | ||
| 141 | + status = ValidateGeamPointers<T>(alpha, A, beta, B, C, op_name); | ||
| 142 | + if (status != ACLBLAS_STATUS_SUCCESS) { | ||
| 143 | + return status; | ||
| 144 | + } | ||
| 145 | + status = ValidateGeamLd(m, n, lda, ldb, ldc, transa, transb, op_name); | ||
| 146 | + if (status != ACLBLAS_STATUS_SUCCESS) { | ||
| 147 | + return status; | ||
| 148 | + } | ||
| 149 | + return ValidateGeamInplace<T>(A, B, C, lda, ldb, ldc, transa, transb, op_name); | ||
| 150 | +} | ||
| @@ -0,0 +1,158 @@ | |||
| 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 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 8 | + */ | ||
| 9 | + | ||
| 10 | +/** | ||
| 11 | + * @file sgeam_host.cpp | ||
| 12 | + * @brief Host-side implementation for Sgeam operator (float) | ||
| 13 | + * Handles parameter validation, tiling calculation, and kernel launch | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | +namespace { | ||
| 27 | + | ||
| 28 | +constexpr uint32_t UB_SIZE = 248 * 1024; // 248 KB | ||
| 29 | +constexpr uint32_t UB_RESERVE = 256; | ||
| 30 | +constexpr uint32_t ALIGN_UNIT = 8; // 32 bytes / sizeof(float) | ||
| 31 | +constexpr uint32_t SGEAM_BUFFERS = 5; // 2*bufA(TQue num=2) + 2*bufC(TQue num=2) + bufB(TBuf) | ||
| 32 | +// Hardware limit for DataCopyPadExt blockCount parameter. | ||
| 33 | +// When curM (rows processed per core) exceeds this, DataCopyPad blockCount | ||
| 34 | +// exceeds hardware limit and behavior is undefined. | ||
| 35 | +constexpr uint32_t MAX_BLOCK_COUNT = 4095; | ||
| 36 | + | ||
| 37 | +SgeamTilingData CalSgeamTilingData( | ||
| 38 | + aclblasOperation_t transa, aclblasOperation_t transb, uint32_t m, uint32_t n, float alpha, float beta, uint32_t lda, | ||
| 39 | + uint32_t ldb, uint32_t ldc, uint32_t aivCoreNum) | ||
| 40 | +{ | ||
| 41 | + SgeamTilingData tiling{}; | ||
| 42 | + tiling.m = m; | ||
| 43 | + tiling.n = n; | ||
| 44 | + tiling.lda = lda; | ||
| 45 | + tiling.ldb = ldb; | ||
| 46 | + tiling.ldc = ldc; | ||
| 47 | + tiling.alpha = alpha; | ||
| 48 | + tiling.beta = beta; | ||
| 49 | + tiling.alphaIsZero = (alpha == 0.0f) ? 1u : 0u; | ||
| 50 | + tiling.betaIsZero = (beta == 0.0f) ? 1u : 0u; | ||
| 51 | + tiling.opA = transa; | ||
| 52 | + tiling.opB = transb; | ||
| 53 | + | ||
| 54 | + // Calculate tileM based on UB capacity, capped to MAX_BLOCK_COUNT | ||
| 55 | + // to avoid triggering the kernel batch-loop error path when curM > 4095. | ||
| 56 | + uint32_t maxTileM = (UB_SIZE - UB_RESERVE) / (SGEAM_BUFFERS * sizeof(float)); | ||
| 57 | + uint32_t tileM = (maxTileM / ALIGN_UNIT) * ALIGN_UNIT; | ||
| 58 | + tileM = std::min(tileM, MAX_BLOCK_COUNT); | ||
| 59 | + tiling.tileM = tileM; | ||
| 60 | + | ||
| 61 | + // NN multi-column: max columns per inner iteration that fit in UB | ||
| 62 | + // UB budget: SGEAM_BUFFERS × colsIter × tileM × sizeof(float) ≤ UB_SIZE - UB_RESERVE | ||
| 63 | + uint32_t colsIter = (UB_SIZE - UB_RESERVE) / (SGEAM_BUFFERS * static_cast<uint32_t>(sizeof(float)) * tileM); | ||
| 64 | + if (colsIter < 1u) { | ||
| 65 | + colsIter = 1u; | ||
| 66 | + } | ||
| 67 | + tiling.colsIter = colsIter; | ||
| 68 | + | ||
| 69 | + // 2D block decomposition: colBlocks x mBlocks | ||
| 70 | + uint32_t colBlocks = std::min(n, aivCoreNum); | ||
| 71 | + if (colBlocks == 0) { | ||
| 72 | + colBlocks = 1; | ||
| 73 | + } | ||
| 74 | + tiling.colBlocks = colBlocks; | ||
| 75 | + tiling.perCoreN = n / colBlocks; | ||
| 76 | + tiling.remainder = n % colBlocks; | ||
| 77 | + | ||
| 78 | + // Calculate m-blocks | ||
| 79 | + uint32_t totalMTiles = (tiling.tileM > 0) ? ((m + tiling.tileM - 1) / tiling.tileM) : 1; | ||
| 80 | + uint32_t mBlocks = 1; | ||
| 81 | + if (colBlocks < aivCoreNum && totalMTiles > 1) { | ||
| 82 | + uint32_t maxMBlocks = aivCoreNum / colBlocks; | ||
| 83 | + mBlocks = std::min(maxMBlocks, totalMTiles); | ||
| 84 | + } | ||
| 85 | + if (mBlocks == 0) // to suppress 'maybe divide by zero' warning | ||
| 86 | + __builtin_unreachable(); | ||
| 87 | + tiling.mBlocks = mBlocks; | ||
| 88 | + tiling.perCoreMTile = totalMTiles / mBlocks; | ||
| 89 | + tiling.mTileRemainder = totalMTiles % mBlocks; | ||
| 90 | + | ||
| 91 | + return tiling; | ||
| 92 | +} | ||
| 93 | + | ||
| 94 | +aclblasStatus_t LaunchSgeamKernel( | ||
| 95 | + aclblasHandle_t handle, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, const float* alpha, | ||
| 96 | + const float* A, int lda, const float* beta, const float* B, int ldb, float* C, int ldc) | ||
| 97 | +{ | ||
| 98 | + // Get core count | ||
| 99 | + uint32_t aivCoreNum = GetAivCoreCount(); | ||
| 100 | + if (aivCoreNum == 0) { | ||
| 101 | + OP_LOGE("aclblasSgeam", "Failed to get AIV core count"); | ||
| 102 | + return ACLBLAS_STATUS_INTERNAL_ERROR; | ||
| 103 | + } | ||
| 104 | + | ||
| 105 | + float alphaVal = *alpha; | ||
| 106 | + float betaVal = *beta; | ||
| 107 | + | ||
| 108 | + // Calculate tiling | ||
| 109 | + SgeamTilingData tiling = CalSgeamTilingData( | ||
| 110 | + transa, transb, static_cast<uint32_t>(m), static_cast<uint32_t>(n), alphaVal, betaVal, static_cast<uint32_t>(lda), | ||
| 111 | + static_cast<uint32_t>(ldb), static_cast<uint32_t>(ldc), aivCoreNum); | ||
| 112 | + | ||
| 113 | + // Calculate total blocks | ||
| 114 | + uint32_t numBlocks = tiling.colBlocks * tiling.mBlocks; | ||
| 115 | + if (numBlocks == 0) { | ||
| 116 | + numBlocks = 1; | ||
| 117 | + } | ||
| 118 | + | ||
| 119 | + OP_LOGD( | ||
| 120 | + "aclblasSgeam", "Tiling: tileM=%u, colBlocks=%u, mBlocks=%u, totalBlocks=%u", tiling.tileM, tiling.colBlocks, | ||
| 121 | + tiling.mBlocks, numBlocks); | ||
| 122 | + | ||
| 123 | + // Convert pointers to GM_ADDR. When alpha/beta value is 0, A/B may be nullptr; | ||
| 124 | + // use C as a dummy valid GM address to avoid hardware issues with null descriptors. | ||
| 125 | + // (Kernel skips reading A/B based on alphaIsZero/betaIsZero tiling flags.) | ||
| 126 | + GM_ADDR gmC = reinterpret_cast<GM_ADDR>(C); | ||
| 127 | + GM_ADDR gmA = (A != nullptr) ? reinterpret_cast<GM_ADDR>(const_cast<float*>(A)) : gmC; | ||
J 为将 const float* 转为 GM_ADDR 做 const_cast。kernel 只读 A（alpha 非 0 时），运行时安全，但 const_cast 是代码异味。如果 GM_ADDR 能支持 const 限定会更规范。 | |||
| 128 | + GM_ADDR gmB = (B != nullptr) ? reinterpret_cast<GM_ADDR>(const_cast<float*>(B)) : gmC; | ||
| 129 | + | ||
| 130 | + // Launch kernel | ||
| 131 | + sgeam_kernel_do(gmA, gmB, gmC, tiling, numBlocks, handle->stream); | ||
| 132 | + | ||
| 133 | + return ACLBLAS_STATUS_SUCCESS; | ||
| 134 | +} | ||
| 135 | + | ||
| 136 | +} // namespace | ||
| 137 | + | ||
| 138 | +extern "C" aclblasStatus_t aclblasSgeam( | ||
| 139 | + aclblasHandle_t handle, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, const float* alpha, | ||
| 140 | + const float* A, int lda, const float* beta, const float* B, int ldb, float* C, int ldc) | ||
| 141 | +{ | ||
| 142 | + if (handle == nullptr) { | ||
| 143 | + OP_LOGE("aclblasSgeam", "handle is nullptr"); | ||
| 144 | + return ACLBLAS_STATUS_HANDLE_IS_NULLPTR; | ||
| 145 | + } | ||
| 146 | + | ||
| 147 | + aclblasStatus_t status = | ||
| 148 | + ValidateGeamParams("aclblasSgeam", transa, transb, m, n, alpha, A, lda, beta, B, ldb, C, ldc); | ||
| 149 | + if (status != ACLBLAS_STATUS_SUCCESS) { | ||
| 150 | + return status; | ||
| 151 | + } | ||
| 152 | + | ||
| 153 | + if (m == 0 || n == 0) { | ||
| 154 | + return ACLBLAS_STATUS_SUCCESS; | ||
| 155 | + } | ||
| 156 | + | ||
| 157 | + return LaunchSgeamKernel(handle, transa, transb, m, n, alpha, A, lda, beta, B, ldb, C, ldc); | ||
| 158 | +} | ||
| @@ -0,0 +1,445 @@ | |||
| 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 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 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 8 | + */ | ||
| 9 | + | ||
| 10 | +/** | ||
| 11 | + * @file sgeam_kernel.cpp | ||
| 12 | + * @brief Sgeam kernel implementation (SIMD/membase, arch35) | ||
| 13 | + * Computes: C = alpha * op(A) + beta * op(B) | ||
| 14 | + * | ||
| 15 | + * Architecture: TPipe + TQue<VECIN>/TQue<VECOUT> + TBuf<VECCALC> | ||
| 16 | + * EnQue/DeQue provides automatic MTE <-> Vector synchronization. | ||
| 17 | + * | ||
| 18 | + * Buffer planning: | ||
| 19 | + * - inQueue_ (VECIN): A tile data, managed by TQue for MTE2<->V sync | ||
| 20 | + * - outQueue_ (VECOUT): C tile data, managed by TQue for V<->MTE3 sync | ||
| 21 | + * - calcBuf_ (VECCALC): Scratch for B data when beta != 0; unused otherwise | ||
| 22 | + * | ||
| 23 | + * Two dispatch paths: | ||
| 24 | + * PATH A: NN multi-column (colsIter > 1), strided DataCopyPad for NC columns | ||
| 25 | + * PATH B: Single-column / Trans / ConjTrans | ||
| 26 | + * | ||
| 27 | + * Multi-core decomposition: 2D grid (colBlocks x mBlocks) | ||
| 28 | + */ | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | +namespace { | ||
| 35 | + | ||
| 36 | +using namespace AscendC; | ||
| 37 | + | ||
| 38 | +// ============================================================================ | ||
| 39 | +// SgeamAIV: operator class for Sgeam vector kernel | ||
| 40 | +// ============================================================================ | ||
| 41 | +class SgeamAIV { | ||
| 42 | +public: | ||
| 43 | + __aicore__ inline SgeamAIV() {} | ||
| 44 | + __aicore__ inline void Init(GM_ADDR A, GM_ADDR B, GM_ADDR C, const SgeamTilingData& tiling, TPipe* pipe); | ||
| 45 | + __aicore__ inline void Process(); | ||
| 46 | + | ||
| 47 | +private: | ||
| 48 | + __aicore__ inline void ProcessPathA(); | ||
| 49 | + __aicore__ inline void ProcessPathB(); | ||
| 50 | + __aicore__ inline void ComputeCoreRanges(const SgeamTilingData& tiling); | ||
| 51 | + | ||
| 52 | + // PATH A helpers (NN multi-column) | ||
| 53 | + __aicore__ inline void CopyInA_PathA(uint32_t col, uint32_t rowStart, uint32_t curM, uint32_t actualNC); | ||
| 54 | + __aicore__ inline void CopyInB_PathA(uint32_t col, uint32_t rowStart, uint32_t curM, uint32_t actualNC); | ||
| 55 | + __aicore__ inline void Compute_PathA(uint32_t curM, uint32_t actualNC); | ||
| 56 | + __aicore__ inline void CopyOutC_PathA(uint32_t col, uint32_t rowStart, uint32_t curM, uint32_t actualNC); | ||
| 57 | + | ||
| 58 | + // PATH B helpers (single-column / Trans / ConjTrans) | ||
| 59 | + __aicore__ inline void CopyInA_PathB(uint32_t col, uint32_t rowStart, uint32_t curM, int64_t srcStrideA); | ||
| 60 | + __aicore__ inline void CopyInB_PathB(uint32_t col, uint32_t rowStart, uint32_t curM, int64_t srcStrideB); | ||
| 61 | + __aicore__ inline void Compute_PathB(uint32_t curM); | ||
| 62 | + __aicore__ inline void CopyOutC_PathB(uint32_t col, uint32_t rowStart, uint32_t curM); | ||
| 63 | + | ||
| 64 | + // GM addresses | ||
| 65 | + GlobalTensor<float> aGm_; | ||
| 66 | + GlobalTensor<float> bGm_; | ||
| 67 | + GlobalTensor<float> cGm_; | ||
| 68 | + | ||
| 69 | + // Queues & buffer | ||
| 70 | + TQue<TPosition::VECIN, 2> inQueue_; // MTE2 input for A (EnQue signals MTE2 done) | ||
| 71 | + TQue<TPosition::VECOUT, 2> outQueue_; // Vector output for C (EnQue signals Vector done) | ||
| 72 | + TBuf<TPosition::VECCALC> calcBuf_; // Scratch for B data (no MTE sync needed) | ||
| 73 | + | ||
| 74 | + // Tiling fields cached from SgeamTilingData | ||
| 75 | + uint32_t m_; | ||
| 76 | + uint32_t n_; | ||
| 77 | + | ||
| 78 | + uint32_t alphaIsZero_; | ||
| 79 | + float alpha_; | ||
| 80 | + aclblasOperation_t opA_; | ||
| 81 | + uint32_t lda_; | ||
| 82 | + | ||
| 83 | + uint32_t betaIsZero_; | ||
| 84 | + float beta_; | ||
| 85 | + aclblasOperation_t opB_; | ||
| 86 | + uint32_t ldb_; | ||
| 87 | + | ||
| 88 | + uint32_t ldc_; | ||
| 89 | + uint32_t tileM_; | ||
| 90 | + uint32_t colsIter_; | ||
| 91 | + | ||
| 92 | + // Per-core tile ranges | ||
| 93 | + uint32_t colStart_; | ||
| 94 | + uint32_t colEnd_; | ||
| 95 | + uint32_t mTileStart_; | ||
| 96 | + uint32_t mTileEnd_; | ||
| 97 | +}; | ||
| 98 | + | ||
| 99 | +// ============================================================================ | ||
| 100 | +// ComputeCoreRanges: decode 2D block index into per-core column and m-tile ranges. | ||
| 101 | +// On invalid range, sets colEnd_ or mTileEnd_ to 0 to signal early exit in Process(). | ||
| 102 | +// ============================================================================ | ||
| 103 | +__aicore__ inline void SgeamAIV::ComputeCoreRanges(const SgeamTilingData& tiling) | ||
| 104 | +{ | ||
| 105 | + uint32_t blockIdx = GetBlockIdx(); | ||
| 106 | + uint32_t mBlocks = tiling.mBlocks; | ||
| 107 | + uint32_t colBlock = blockIdx / mBlocks; | ||
| 108 | + uint32_t mBlock = blockIdx % mBlocks; | ||
| 109 | + | ||
| 110 | + // Calculate column range | ||
| 111 | + if (colBlock < tiling.remainder) { | ||
| 112 | + colStart_ = colBlock * (tiling.perCoreN + 1); | ||
| 113 | + colEnd_ = colStart_ + tiling.perCoreN + 1; | ||
| 114 | + } else { | ||
| 115 | + colStart_ = colBlock * tiling.perCoreN + tiling.remainder; | ||
| 116 | + colEnd_ = colStart_ + tiling.perCoreN; | ||
| 117 | + } | ||
| 118 | + | ||
| 119 | + // Validate column range | ||
| 120 | + if (colStart_ >= n_ || colStart_ >= colEnd_) { | ||
| 121 | + colEnd_ = 0; // Signal Process() to return early | ||
| 122 | + return; | ||
| 123 | + } | ||
| 124 | + if (colEnd_ > n_) { | ||
| 125 | + colEnd_ = n_; | ||
| 126 | + } | ||
| 127 | + | ||
| 128 | + // Calculate m-tile range | ||
| 129 | + if (mBlock < tiling.mTileRemainder) { | ||
| 130 | + mTileStart_ = mBlock * (tiling.perCoreMTile + 1); | ||
| 131 | + mTileEnd_ = mTileStart_ + tiling.perCoreMTile + 1; | ||
| 132 | + } else { | ||
| 133 | + mTileStart_ = mBlock * tiling.perCoreMTile + tiling.mTileRemainder; | ||
| 134 | + mTileEnd_ = mTileStart_ + tiling.perCoreMTile; | ||
| 135 | + } | ||
| 136 | + | ||
| 137 | + // Validate m-tile range | ||
| 138 | + uint32_t totalMTiles = (m_ + tileM_ - 1) / tileM_; | ||
| 139 | + if (mTileStart_ >= totalMTiles || mTileStart_ >= mTileEnd_) { | ||
| 140 | + mTileEnd_ = 0; // Signal Process() to return early | ||
| 141 | + return; | ||
| 142 | + } | ||
| 143 | + if (mTileEnd_ > totalMTiles) { | ||
| 144 | + mTileEnd_ = totalMTiles; | ||
| 145 | + } | ||
| 146 | +} | ||
| 147 | + | ||
| 148 | +// ============================================================================ | ||
| 149 | +// Init: setup global tensors, compute per-core ranges, allocate UB buffers. | ||
| 150 | +// TPipe pointer passed from kernel entry (R3: TPipe must not be a member). | ||
| 151 | +// ============================================================================ | ||
| 152 | +__aicore__ inline void SgeamAIV::Init(GM_ADDR A, GM_ADDR B, GM_ADDR C, const SgeamTilingData& tiling, TPipe* pipe) | ||
| 153 | +{ | ||
| 154 | + // Cache tiling fields to member variables | ||
| 155 | + m_ = tiling.m; | ||
| 156 | + n_ = tiling.n; | ||
| 157 | + | ||
| 158 | + alphaIsZero_ = tiling.alphaIsZero; | ||
| 159 | + opA_ = tiling.opA; | ||
| 160 | + lda_ = tiling.lda; | ||
| 161 | + alpha_ = tiling.alpha; | ||
| 162 | + | ||
| 163 | + betaIsZero_ = tiling.betaIsZero; | ||
| 164 | + ldb_ = tiling.ldb; | ||
| 165 | + opB_ = tiling.opB; | ||
| 166 | + beta_ = tiling.beta; | ||
| 167 | + | ||
| 168 | + ldc_ = tiling.ldc; | ||
| 169 | + tileM_ = tiling.tileM; | ||
| 170 | + colsIter_ = tiling.colsIter; | ||
| 171 | + | ||
| 172 | + // Set GM buffers | ||
| 173 | + aGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(A)); | ||
| 174 | + bGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(B)); | ||
| 175 | + cGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(C)); | ||
| 176 | + | ||
| 177 | + // Compute per-core column and m-tile ranges | ||
| 178 | + ComputeCoreRanges(tiling); | ||
| 179 | + if (colEnd_ == 0 || mTileEnd_ == 0) { | ||
| 180 | + return; | ||
| 181 | + } | ||
| 182 | + | ||
| 183 | + // Buffer size: PATH A needs colsIter * tileM elements, PATH B needs tileM elements. | ||
| 184 | + // Allocate the maximum (colsIter >= 1, so colsIter * tileM >= tileM). | ||
| 185 | + uint32_t bufElems = colsIter_ * tileM_; | ||
J 所有 buffer 按 colsIter*tileM 分配（PathA 需要），但 PathB 只处理单列，用 tileM 即可。PathB 执行时 buffer 大部分空闲。如果 PathB 是常见路径（非 NN 场景），可考虑按路径动态分配。 | |||
| 186 | + uint32_t bufSize = bufElems * sizeof(float); | ||
| 187 | + | ||
| 188 | + // Double-buffer queues (num=2) enable MTE2<->Vector pipeline overlap | ||
| 189 | + pipe->InitBuffer(inQueue_, 2, bufSize); | ||
| 190 | + pipe->InitBuffer(outQueue_, 2, bufSize); | ||
| 191 | + pipe->InitBuffer(calcBuf_, bufSize); | ||
| 192 | +} | ||
| 193 | + | ||
| 194 | +// ============================================================================ | ||
| 195 | +// Process: dispatch to PATH A or PATH B based on op and colsIter | ||
| 196 | +// ============================================================================ | ||
| 197 | +__aicore__ inline void SgeamAIV::Process() | ||
| 198 | +{ | ||
| 199 | + if (colStart_ >= colEnd_ || mTileStart_ >= mTileEnd_) { | ||
| 200 | + return; | ||
| 201 | + } | ||
| 202 | + | ||
| 203 | + if (opA_ == ACLBLAS_OP_N && opB_ == ACLBLAS_OP_N && colsIter_ > 1) { | ||
| 204 | + ProcessPathA(); | ||
| 205 | + } else { | ||
| 206 | + ProcessPathB(); | ||
| 207 | + } | ||
| 208 | +} | ||
| 209 | + | ||
| 210 | +// ============================================================================ | ||
| 211 | +// CopyInA_PathA: strided DataCopyPad loads NC columns of A at once. | ||
| 212 | +// ============================================================================ | ||
| 213 | +__aicore__ inline void SgeamAIV::CopyInA_PathA(uint32_t col, uint32_t rowStart, uint32_t curM, uint32_t actualNC) | ||
| 214 | +{ | ||
| 215 | + LocalTensor<float> bufA = inQueue_.AllocTensor<float>(); | ||
| 216 | + if (!alphaIsZero_) { | ||
| 217 | + const uint64_t baseOff = static_cast<uint64_t>(col) * static_cast<uint64_t>(lda_) + rowStart; | ||
| 218 | + const int64_t colGap = static_cast<int64_t>(lda_ - curM) * static_cast<int64_t>(sizeof(float)); | ||
| 219 | + DataCopyExtParams cpA{ | ||
| 220 | + static_cast<uint16_t>(actualNC), static_cast<uint32_t>(curM * sizeof(float)), colGap, 0, 0}; | ||
| 221 | + DataCopyPadExtParams<float> npA; | ||
| 222 | + DataCopyPad<float, PaddingMode::Compact>(bufA, aGm_[baseOff], cpA, npA); | ||
| 223 | + } | ||
| 224 | + inQueue_.EnQue(bufA); // Signal MTE2 done for A | ||
| 225 | +} | ||
| 226 | + | ||
| 227 | +// ============================================================================ | ||
| 228 | +// CopyInB_PathA: strided DataCopyPad loads NC columns of B when beta != 0. | ||
| 229 | +// ============================================================================ | ||
| 230 | +__aicore__ inline void SgeamAIV::CopyInB_PathA(uint32_t col, uint32_t rowStart, uint32_t curM, uint32_t actualNC) | ||
| 231 | +{ | ||
| 232 | + if (!betaIsZero_) { | ||
| 233 | + LocalTensor<float> ubB = calcBuf_.Get<float>(); | ||
| 234 | + const uint64_t baseOff = static_cast<uint64_t>(col) * static_cast<uint64_t>(ldb_) + rowStart; | ||
| 235 | + const int64_t colGap = static_cast<int64_t>(ldb_ - curM) * static_cast<int64_t>(sizeof(float)); | ||
| 236 | + DataCopyExtParams cpB{ | ||
| 237 | + static_cast<uint16_t>(actualNC), static_cast<uint32_t>(curM * sizeof(float)), colGap, 0, 0}; | ||
| 238 | + DataCopyPadExtParams<float> npB; | ||
| 239 | + DataCopyPad<float, PaddingMode::Compact>(ubB, bGm_[baseOff], cpB, npB); | ||
B 数据经 DataCopyPad 写入 TBuf calcBuf_ 后未用 EnQue/DeQue 或 SetFlag/WaitFlag 建立 MTE2→V 同步，向量读取 calcBuf_ 依赖隐式时序保障（实机大矩阵用例未复现错误），后续调整流水或平台时序变化时易引入数据竞争 | |||
| 240 | + event_t eMte2V_ = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V)); | ||
| 241 | + SetFlag<HardEvent::MTE2_V>(eMte2V_); | ||
| 242 | + WaitFlag<HardEvent::MTE2_V>(eMte2V_); | ||
| 243 | + } | ||
| 244 | +} | ||
| 245 | + | ||
| 246 | +// ============================================================================ | ||
| 247 | +// Compute_PathA: C = alpha*A + beta*B for the multi-column PATH A case. | ||
| 248 | +// ============================================================================ | ||
| 249 | +__aicore__ inline void SgeamAIV::Compute_PathA(uint32_t curM, uint32_t actualNC) | ||
| 250 | +{ | ||
| 251 | + const uint32_t totalElems = actualNC * curM; | ||
| 252 | + LocalTensor<float> aIn = inQueue_.DeQue<float>(); | ||
| 253 | + LocalTensor<float> cOut = outQueue_.AllocTensor<float>(); | ||
| 254 | + | ||
| 255 | + if (alphaIsZero_ && betaIsZero_) { | ||
| 256 | + Duplicate<float>(cOut, 0.0f, totalElems); | ||
| 257 | + } else if (alphaIsZero_) { | ||
| 258 | + LocalTensor<float> ubB = calcBuf_.Get<float>(); | ||
| 259 | + Muls<float>(cOut, ubB, beta_, totalElems); | ||
| 260 | + } else if (betaIsZero_) { | ||
| 261 | + Muls<float>(cOut, aIn, alpha_, totalElems); | ||
| 262 | + } else { | ||
| 263 | + LocalTensor<float> ubB = calcBuf_.Get<float>(); | ||
| 264 | + Muls<float>(cOut, aIn, alpha_, totalElems); | ||
| 265 | + // Reuse aIn as scratch for beta*B to save a buffer | ||
| 266 | + Muls<float>(aIn, ubB, beta_, totalElems); | ||
J 直接覆写 aIn（从 inQueue_ DeQue 出来的 A buffer）为 beta*B，然后 Add(cOut, cOut, aIn)。虽然逻辑正确且节省一个 scratch buffer，但覆写输入 buffer 的写法容易让读者困惑。建议加一行注释说明此处是故意复用 aIn 作为临时存储。 | |||
| 267 | + Add<float>(cOut, cOut, aIn, totalElems); | ||
| 268 | + } | ||
| 269 | + outQueue_.EnQue(cOut); // Signal Vector done | ||
| 270 | + inQueue_.FreeTensor(aIn); | ||
| 271 | +} | ||
| 272 | + | ||
| 273 | +// ============================================================================ | ||
| 274 | +// CopyOutC_PathA: strided DataCopyPad writes NC columns of C. | ||
| 275 | +// ============================================================================ | ||
| 276 | +__aicore__ inline void SgeamAIV::CopyOutC_PathA(uint32_t col, uint32_t rowStart, uint32_t curM, uint32_t actualNC) | ||
| 277 | +{ | ||
| 278 | + LocalTensor<float> cFinal = outQueue_.DeQue<float>(); | ||
| 279 | + const uint64_t baseOff = static_cast<uint64_t>(col) * static_cast<uint64_t>(ldc_) + rowStart; | ||
| 280 | + const int64_t colGap = static_cast<int64_t>(ldc_ - curM) * static_cast<int64_t>(sizeof(float)); | ||
| 281 | + DataCopyExtParams cpS{static_cast<uint16_t>(actualNC), static_cast<uint32_t>(curM * sizeof(float)), 0, colGap, 0}; | ||
| 282 | + DataCopyPad<float, PaddingMode::Compact>(cGm_[baseOff], cFinal, cpS); | ||
| 283 | + outQueue_.FreeTensor(cFinal); | ||
| 284 | +} | ||
| 285 | + | ||
| 286 | +// ============================================================================ | ||
| 287 | +// PATH A: NN multi-column — strided DataCopyPad loads NC columns at once | ||
| 288 | +// Each tile: CopyIn A -> EnQue -> CopyIn B -> Compute -> EnQue C -> CopyOut C | ||
| 289 | +// ============================================================================ | ||
| 290 | +__aicore__ inline void SgeamAIV::ProcessPathA() | ||
| 291 | +{ | ||
| 292 | + for (uint32_t col = colStart_; col < colEnd_; col += colsIter_) { | ||
| 293 | + const uint32_t actualNC = ((colEnd_ - col) < colsIter_) ? (colEnd_ - col) : colsIter_; | ||
| 294 | + for (uint32_t mTile = mTileStart_; mTile < mTileEnd_; ++mTile) { | ||
| 295 | + const uint32_t rowStart = mTile * tileM_; | ||
| 296 | + uint32_t rowEnd = rowStart + tileM_; | ||
| 297 | + if (rowEnd > m_) { | ||
| 298 | + rowEnd = m_; | ||
| 299 | + } | ||
| 300 | + const uint32_t curM = rowEnd - rowStart; | ||
| 301 | + | ||
| 302 | + CopyInA_PathA(col, rowStart, curM, actualNC); | ||
| 303 | + CopyInB_PathA(col, rowStart, curM, actualNC); | ||
| 304 | + Compute_PathA(curM, actualNC); | ||
J calcBuf_ 是 TBuf，CopyInB_PathA 通过 DataCopyPad（MTE2）将 B 数据搬入 calcBuf_，但 Compute_PathA 中 Vector 直接读取 calcBuf_ 时缺少 MTE2→V 同步。inQueue_ 的 EnQue 只同步了 A 的 MTE2，不覆盖 B 的 DataCopyPad（B 的 MTE2 指令在 EnQue 之后发射）。对照仓库 sdgmm_kernel.cpp:71-73，TBuf+DataCopyPad 后需显式 SetFlag/WaitFlag<MTE2_V> 或 PipeBarrier<PIPE_MTE2_V>。PathB 的 Compute_PathB 存在同样问题。 | |||
| 305 | + CopyOutC_PathA(col, rowStart, curM, actualNC); | ||
| 306 | + } | ||
| 307 | + } | ||
| 308 | +} | ||
| 309 | + | ||
| 310 | +// ============================================================================ | ||
| 311 | +// CopyInA_PathB: single-shot for NN, strided Compact for T/C (one column). | ||
| 312 | +// ============================================================================ | ||
| 313 | +__aicore__ inline void SgeamAIV::CopyInA_PathB(uint32_t col, uint32_t rowStart, uint32_t curM, int64_t srcStrideA) | ||
| 314 | +{ | ||
| 315 | + LocalTensor<float> bufA = inQueue_.AllocTensor<float>(); | ||
| 316 | + if (!alphaIsZero_) { | ||
| 317 | + if (opA_ == ACLBLAS_OP_N) { | ||
| 318 | + const uint64_t offset = static_cast<uint64_t>(col) * static_cast<uint64_t>(lda_) + rowStart; | ||
| 319 | + DataCopyExtParams cp{1, static_cast<uint32_t>(curM * sizeof(float)), 0, 0, 0}; | ||
| 320 | + DataCopyPadExtParams<float> np; | ||
| 321 | + DataCopyPad(bufA, aGm_[offset], cp, np); | ||
| 322 | + } else { | ||
| 323 | + const uint64_t baseOff = | ||
| 324 | + static_cast<uint64_t>(col) + static_cast<uint64_t>(rowStart) * static_cast<uint64_t>(lda_); | ||
| 325 | + DataCopyExtParams cp{static_cast<uint16_t>(curM), static_cast<uint32_t>(sizeof(float)), srcStrideA, 0, 0}; | ||
| 326 | + DataCopyPadExtParams<float> np; | ||
| 327 | + DataCopyPad<float, PaddingMode::Compact>(bufA, aGm_[baseOff], cp, np); | ||
| 328 | + } | ||
| 329 | + } | ||
| 330 | + inQueue_.EnQue(bufA); // Signal MTE2 done for A | ||
| 331 | +} | ||
| 332 | + | ||
| 333 | +// ============================================================================ | ||
| 334 | +// CopyInB_PathB: single-shot for NN, strided Compact for T/C (one column). | ||
| 335 | +// ============================================================================ | ||
| 336 | +__aicore__ inline void SgeamAIV::CopyInB_PathB(uint32_t col, uint32_t rowStart, uint32_t curM, int64_t srcStrideB) | ||
| 337 | +{ | ||
| 338 | + if (!betaIsZero_) { | ||
| 339 | + LocalTensor<float> ubB = calcBuf_.Get<float>(); | ||
| 340 | + if (opB_ == ACLBLAS_OP_N) { | ||
| 341 | + const uint64_t offset = static_cast<uint64_t>(col) * static_cast<uint64_t>(ldb_) + rowStart; | ||
| 342 | + DataCopyExtParams cp{1, static_cast<uint32_t>(curM * sizeof(float)), 0, 0, 0}; | ||
| 343 | + DataCopyPadExtParams<float> np; | ||
| 344 | + DataCopyPad(ubB, bGm_[offset], cp, np); | ||
| 345 | + } else { | ||
| 346 | + const uint64_t baseOff = | ||
| 347 | + static_cast<uint64_t>(col) + static_cast<uint64_t>(rowStart) * static_cast<uint64_t>(ldb_); | ||
| 348 | + DataCopyExtParams cp{static_cast<uint16_t>(curM), static_cast<uint32_t>(sizeof(float)), srcStrideB, 0, 0}; | ||
| 349 | + DataCopyPadExtParams<float> np; | ||
| 350 | + DataCopyPad<float, PaddingMode::Compact>(ubB, bGm_[baseOff], cp, np); | ||
| 351 | + } | ||
| 352 | + event_t eMte2V_ = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V)); | ||
| 353 | + SetFlag<HardEvent::MTE2_V>(eMte2V_); | ||
| 354 | + WaitFlag<HardEvent::MTE2_V>(eMte2V_); | ||
| 355 | + } | ||
| 356 | +} | ||
| 357 | + | ||
| 358 | +// ============================================================================ | ||
| 359 | +// Compute_PathB: C = alpha*A + beta*B for the single-column PATH B case. | ||
| 360 | +// ============================================================================ | ||
| 361 | +__aicore__ inline void SgeamAIV::Compute_PathB(uint32_t curM) | ||
| 362 | +{ | ||
| 363 | + LocalTensor<float> aIn = inQueue_.DeQue<float>(); | ||
| 364 | + LocalTensor<float> cOut = outQueue_.AllocTensor<float>(); | ||
| 365 | + | ||
| 366 | + if (alphaIsZero_ && betaIsZero_) { | ||
| 367 | + Duplicate<float>(cOut, 0.0f, curM); | ||
| 368 | + } else if (alphaIsZero_) { | ||
| 369 | + LocalTensor<float> ubB = calcBuf_.Get<float>(); | ||
| 370 | + Muls<float>(cOut, ubB, beta_, curM); | ||
| 371 | + } else if (betaIsZero_) { | ||
| 372 | + Muls<float>(cOut, aIn, alpha_, curM); | ||
| 373 | + } else { | ||
| 374 | + LocalTensor<float> ubB = calcBuf_.Get<float>(); | ||
| 375 | + Muls<float>(cOut, aIn, alpha_, curM); | ||
| 376 | + // Reuse aIn as scratch for beta*B to save a buffer | ||
| 377 | + Muls<float>(aIn, ubB, beta_, curM); | ||
| 378 | + Add<float>(cOut, cOut, aIn, curM); | ||
| 379 | + } | ||
| 380 | + outQueue_.EnQue(cOut); // Signal Vector done | ||
| 381 | + inQueue_.FreeTensor(aIn); | ||
| 382 | +} | ||
| 383 | + | ||
| 384 | +// ============================================================================ | ||
| 385 | +// CopyOutC_PathB: single-shot DataCopyPad writes one column of C. | ||
| 386 | +// ============================================================================ | ||
| 387 | +__aicore__ inline void SgeamAIV::CopyOutC_PathB(uint32_t col, uint32_t rowStart, uint32_t curM) | ||
| 388 | +{ | ||
| 389 | + LocalTensor<float> cFinal = outQueue_.DeQue<float>(); | ||
| 390 | + const uint64_t offset = static_cast<uint64_t>(col) * static_cast<uint64_t>(ldc_) + rowStart; | ||
| 391 | + DataCopyExtParams cp{1, static_cast<uint32_t>(curM * sizeof(float)), 0, 0, 0}; | ||
| 392 | + DataCopyPad(cGm_[offset], cFinal, cp); | ||
| 393 | + outQueue_.FreeTensor(cFinal); | ||
| 394 | +} | ||
| 395 | + | ||
| 396 | +// ============================================================================ | ||
| 397 | +// PATH B: Single-column / Trans / ConjTrans — one column per iteration | ||
| 398 | +// Each tile: CopyIn A -> EnQue -> CopyIn B -> Compute -> EnQue C -> CopyOut C | ||
| 399 | +// ============================================================================ | ||
| 400 | +__aicore__ inline void SgeamAIV::ProcessPathB() | ||
| 401 | +{ | ||
| 402 | + const int64_t srcStrideA = static_cast<int64_t>(lda_ - 1) * static_cast<int64_t>(sizeof(float)); | ||
| 403 | + const int64_t srcStrideB = static_cast<int64_t>(ldb_ - 1) * static_cast<int64_t>(sizeof(float)); | ||
| 404 | + | ||
| 405 | + for (uint32_t col = colStart_; col < colEnd_; ++col) { | ||
| 406 | + for (uint32_t mTile = mTileStart_; mTile < mTileEnd_; ++mTile) { | ||
| 407 | + const uint32_t rowStart = mTile * tileM_; | ||
| 408 | + uint32_t rowEnd = rowStart + tileM_; | ||
| 409 | + if (rowEnd > m_) { | ||
| 410 | + rowEnd = m_; | ||
| 411 | + } | ||
| 412 | + const uint32_t curM = rowEnd - rowStart; | ||
| 413 | + | ||
| 414 | + CopyInA_PathB(col, rowStart, curM, srcStrideA); | ||
| 415 | + CopyInB_PathB(col, rowStart, curM, srcStrideB); | ||
| 416 | + Compute_PathB(curM); | ||
| 417 | + CopyOutC_PathB(col, rowStart, curM); | ||
| 418 | + } | ||
| 419 | + } | ||
| 420 | +} | ||
| 421 | + | ||
| 422 | +} // namespace | ||
| 423 | + | ||
| 424 | +// ============================================================================ | ||
| 425 | +// Kernel entry point | ||
| 426 | +// ============================================================================ | ||
| 427 | +extern "C" __global__ __aicore__ void sgeam_kernel(GM_ADDR A, GM_ADDR B, GM_ADDR C, SgeamTilingData tiling) | ||
| 428 | +{ | ||
| 429 | + KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); | ||
| 430 | + | ||
| 431 | + if (tiling.m == 0 || tiling.n == 0) { | ||
| 432 | + return; | ||
| 433 | + } | ||
| 434 | + | ||
| 435 | + TPipe pipe; | ||
| 436 | + SgeamAIV op; | ||
| 437 | + op.Init(A, B, C, tiling, &pipe); | ||
| 438 | + op.Process(); | ||
| 439 | +} | ||
| 440 | + | ||
| 441 | +void sgeam_kernel_do( | ||
| 442 | + GM_ADDR A, GM_ADDR B, GM_ADDR C, const SgeamTilingData& tiling, uint32_t numBlocks, aclrtStream stream) | ||
| 443 | +{ | ||
| 444 | + sgeam_kernel<<<numBlocks, nullptr, stream>>>(A, B, C, tiling); | ||
| 445 | +} | ||
| @@ -0,0 +1,37 @@ | |||
| 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 sgeam_kernel.h | ||
| 13 | + * @brief Kernel launcher declaration for Sgeam operator | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | +/** | ||
| 27 | + * @brief Launch Sgeam kernel asynchronously | ||
| 28 | + * | ||
| 29 | + * @param A GM address of matrix A | ||
| 30 | + * @param B GM address of matrix B | ||
| 31 | + * @param C GM address of matrix C (output) | ||
| 32 | + * @param tiling Host-computed tiling data (copied to device by value) | ||
| 33 | + * @param numBlocks Block count for the <<<>>> launch | ||
| 34 | + * @param stream aclrtStream handle | ||
| 35 | + */ | ||
| 36 | +void sgeam_kernel_do( | ||
| 37 | + GM_ADDR A, GM_ADDR B, GM_ADDR C, const SgeamTilingData& tiling, uint32_t numBlocks, aclrtStream stream); | ||
| @@ -0,0 +1,44 @@ | |||
| 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 sgeam_tiling_data.h | ||
| 13 | + * @brief Tiling data structure for Sgeam operator (host/device shared) | ||
| 14 | + * C = alpha * op(A) + beta * op(B) | ||
| 15 | + */ | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | +struct SgeamTilingData { | ||
| 23 | + uint32_t m; // Number of rows in output matrix C | ||
| 24 | + uint32_t n; // Number of columns in output matrix C | ||
| 25 | + float alpha; // Scalar alpha | ||
| 26 | + uint32_t alphaIsZero; // Optimized flag indicating alpha == 0 | ||
| 27 | + aclblasOperation_t opA; // Transpose type for matrix A | ||
| 28 | + uint32_t lda; // Leading dimension of matrix A | ||
| 29 | + float beta; // Scalar beta | ||
| 30 | + uint32_t betaIsZero; // Optimized flag indicating beta == 0 | ||
| 31 | + aclblasOperation_t opB; // Transpose type for matrix B | ||
| 32 | + uint32_t ldb; // Leading dimension of matrix B | ||
| 33 | + uint32_t ldc; // Leading dimension of matrix C | ||
| 34 | + // 2D multi-core tiling: colBlocks x mBlocks | ||
| 35 | + uint32_t colBlocks; // Number of column blocks | ||
| 36 | + uint32_t perCoreN; // Number of columns processed per core | ||
| 37 | + uint32_t remainder; // Column remainder | ||
| 38 | + uint32_t mBlocks; // Number of row blocks | ||
| 39 | + uint32_t perCoreMTile; // Number of tiles processed per core in m dimension | ||
| 40 | + uint32_t mTileRemainder; // Tile count remainder in m dimension | ||
| 41 | + uint32_t tileM; // Tile size in m dimension (number of elements) | ||
| 42 | + uint32_t colsIter; // NN path: columns processed per inner iteration (multi-column) | ||
| 43 | + uint32_t reserved; // Reserved for 8-byte alignment | ||
| 44 | +}; | ||
J 结构体共 19 个 4 字节字段，总大小 76 字节，不是 8 的倍数。cgeam_tiling_data.h 是 80 字节（8 对齐）。建议在末尾补一个 uint32_t 保留字段使大小对齐到 80 字节，与 cgeam 保持一致。 | |||
| @@ -186,9 +186,9 @@ function(_ops_blas_register_test_target target link_lib) | |||
| 186 | ${CMAKE_SOURCE_DIR}/test/utils | 186 | ${CMAKE_SOURCE_DIR}/test/utils |
| 187 | ${CMAKE_CURRENT_SOURCE_DIR} | 187 | ${CMAKE_CURRENT_SOURCE_DIR} |
| 188 | ${CMAKE_SOURCE_DIR}/blas/common/helper | 188 | ${CMAKE_SOURCE_DIR}/blas/common/helper |
| 189 | - ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/op_common/ | 189 | + ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/op_common/ |
| 190 | - ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/base/ | 190 | + ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/base/ |
| 191 | - ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/ | 191 | + ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/ |
| 192 | $ENV{LINUX_INCLUDE_PATH} | 192 | $ENV{LINUX_INCLUDE_PATH} |
| 193 | ${REFBLAS_INCLUDE_DIR} | 193 | ${REFBLAS_INCLUDE_DIR} |
| 194 | ) | 194 | ) |
| @@ -892,6 +892,8 @@ BLAS-like Extension 提供标准 BLAS Level 3 之外的扩展 GEMM 接口（以 | |||
| 892 | 892 | ||
| 893 | | 接口名 | 说明 | | 893 | | 接口名 | 说明 | |
| 894 | |---|---| | 894 | |---|---| |
| 895 | +| [aclblasSgeam](../../blas/geam/README.md) | 单精度矩阵加法：C = alpha * op(A) + beta * op(B) | | ||
| 896 | +| [aclblasCgeam](../../blas/geam/README.md) | 复数矩阵加法：C = alpha * op(A) + beta * op(B) | | ||
| 895 | | [aclblasGemmEx](../../blas/gemm_ex/README.md) | 通用矩阵乘法扩展接口，支持 A/B/C 独立数据类型 | | 897 | | [aclblasGemmEx](../../blas/gemm_ex/README.md) | 通用矩阵乘法扩展接口，支持 A/B/C 独立数据类型 | |
| 896 | | [aclblasGemmBatchedEx](../../blas/gemm_batched_ex/README.md) | 通用矩阵乘法批量扩展接口 | | 898 | | [aclblasGemmBatchedEx](../../blas/gemm_batched_ex/README.md) | 通用矩阵乘法批量扩展接口 | |
| 897 | | [aclblasGemmGroupedBatchedEx](../../blas/gemm_grouped_batched_ex/README.md) | 通用矩阵乘法分组批量扩展接口 | | 899 | | [aclblasGemmGroupedBatchedEx](../../blas/gemm_grouped_batched_ex/README.md) | 通用矩阵乘法分组批量扩展接口 | |
| @@ -478,6 +478,22 @@ aclblasStatus_t aclblasCdgmm( | |||
| 478 | const aclblasComplex* A, int lda, const aclblasComplex* x, int incx, | 478 | const aclblasComplex* A, int lda, const aclblasComplex* x, int incx, |
| 479 | aclblasComplex* C, int ldc); | 479 | aclblasComplex* C, int ldc); |
| 480 | 480 | ||
| 481 | +aclblasStatus_t aclblasSgeam( | ||
| 482 | + aclblasHandle_t handle, | ||
| 483 | + aclblasOperation_t transa, aclblasOperation_t transb, | ||
| 484 | + int m, int n, | ||
| 485 | + const float* alpha, const float* A, int lda, | ||
| 486 | + const float* beta, const float* B, int ldb, | ||
| 487 | + float* C, int ldc); | ||
| 488 | + | ||
| 489 | +aclblasStatus_t aclblasCgeam( | ||
| 490 | + aclblasHandle_t handle, | ||
| 491 | + aclblasOperation_t transa, aclblasOperation_t transb, | ||
| 492 | + int m, int n, | ||
| 493 | + const aclblasComplex* alpha, const aclblasComplex* A, int lda, | ||
| 494 | + const aclblasComplex* beta, const aclblasComplex* B, int ldb, | ||
| 495 | + aclblasComplex* C, int ldc); | ||
| 496 | + | ||
| 481 | aclblasStatus_t aclblasSgemm3m( | 497 | aclblasStatus_t aclblasSgemm3m( |
| 482 | aclblasHandle handle, aclblasOperation_t transA, aclblasOperation_t transB, int m, int n, int k, | 498 | aclblasHandle handle, aclblasOperation_t transA, aclblasOperation_t transB, int m, int n, int k, |
| 483 | const float* alpha, const float* A, int lda, | 499 | const float* alpha, const float* A, int lda, |
| @@ -0,0 +1,11 @@ | |||
| 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 FILE 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 | +ops_blas_add_gtest_tests(${OPS_BLAS} cgeam_test) | ||
| @@ -0,0 +1,115 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | +inline size_t CgeamBufferBytes(aclblasOperation_t trans, int m, int n, int ld) | ||
| 24 | +{ | ||
| 25 | + size_t cols = (trans == ACLBLAS_OP_N) ? static_cast<size_t>(n) : static_cast<size_t>(m); | ||
| 26 | + return static_cast<size_t>(ld) * cols * sizeof(aclblasComplex); | ||
| 27 | +} | ||
| 28 | + | ||
| 29 | +inline std::unique_ptr<DeviceBuffer> CgeamAllocAndCopy(const aclblasComplex* hostPtr, size_t bytes) | ||
| 30 | +{ | ||
| 31 | + if (hostPtr == nullptr) { | ||
| 32 | + return nullptr; | ||
| 33 | + } | ||
| 34 | + auto buf = std::make_unique<DeviceBuffer>(bytes); | ||
| 35 | + buf->copyFromHost(hostPtr, bytes); | ||
| 36 | + return buf; | ||
| 37 | +} | ||
| 38 | + | ||
| 39 | +inline std::unique_ptr<DeviceBuffer> CgeamAllocOutput(size_t cBytes) | ||
| 40 | +{ | ||
| 41 | + auto buf = std::make_unique<DeviceBuffer>(cBytes); | ||
| 42 | + std::vector<float> sentinel(cBytes / sizeof(float), kBlasSentinel); | ||
| 43 | + buf->copyFromHost(sentinel.data(), cBytes); | ||
| 44 | + return buf; | ||
| 45 | +} | ||
| 46 | + | ||
| 47 | +inline aclblasComplex* CgeamGetCPtr(bool inplaceAC, bool inplaceBC, DeviceBuffer* dA, DeviceBuffer* dB, DeviceBuffer* dC) | ||
| 48 | +{ | ||
| 49 | + if (inplaceAC) { | ||
| 50 | + return static_cast<aclblasComplex*>(dA->ptr()); | ||
| 51 | + } | ||
| 52 | + if (inplaceBC) { | ||
| 53 | + return static_cast<aclblasComplex*>(dB->ptr()); | ||
| 54 | + } | ||
| 55 | + return static_cast<aclblasComplex*>(dC->ptr()); | ||
| 56 | +} | ||
| 57 | + | ||
| 58 | +inline void CgeamCopyBack(bool inplaceAC, bool inplaceBC, DeviceBuffer* dA, DeviceBuffer* dB, DeviceBuffer* dC, | ||
| 59 | + aclblasComplex* C, size_t aBytes, size_t bBytes, size_t cBytes) | ||
| 60 | +{ | ||
| 61 | + if (inplaceAC && dA) { | ||
| 62 | + dA->copyToHost(C, aBytes); | ||
| 63 | + } else if (inplaceBC && dB) { | ||
| 64 | + dB->copyToHost(C, bBytes); | ||
| 65 | + } else if (dC) { | ||
| 66 | + dC->copyToHost(C, cBytes); | ||
| 67 | + } | ||
| 68 | +} | ||
| 69 | + | ||
| 70 | +inline aclblasStatus_t aclblasCgeam_npu( | ||
| 71 | + aclblasHandle_t handle, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, | ||
| 72 | + const aclblasComplex* alpha, const aclblasComplex* A, int lda, | ||
| 73 | + const aclblasComplex* beta, const aclblasComplex* B, int ldb, aclblasComplex* C, int ldc) | ||
| 74 | +{ | ||
| 75 | + if (m <= 0 || n <= 0) { | ||
| 76 | + aclblasComplex dummyA = {0.0f, 0.0f}; | ||
| 77 | + aclblasComplex dummyB = {0.0f, 0.0f}; | ||
| 78 | + aclblasComplex dummyC = {0.0f, 0.0f}; | ||
| 79 | + aclblasComplex dummyAlpha = alpha ? *alpha : aclblasComplex{0.0f, 0.0f}; | ||
| 80 | + aclblasComplex dummyBeta = beta ? *beta : aclblasComplex{0.0f, 0.0f}; | ||
| 81 | + return aclblasCgeam(handle, transa, transb, m, n, | ||
| 82 | + &dummyAlpha, &dummyA, std::max(lda, 1), | ||
| 83 | + &dummyBeta, &dummyB, std::max(ldb, 1), | ||
| 84 | + &dummyC, std::max(ldc, 1)); | ||
| 85 | + } | ||
| 86 | + size_t aBytes = CgeamBufferBytes(transa, m, n, lda); | ||
| 87 | + size_t bBytes = CgeamBufferBytes(transb, m, n, ldb); | ||
| 88 | + size_t cBytes = static_cast<size_t>(ldc) * static_cast<size_t>(n) * sizeof(aclblasComplex); | ||
| 89 | + | ||
| 90 | + bool inplaceAC = (A != nullptr && C == A); | ||
| 91 | + bool inplaceBC = (B != nullptr && C == B); | ||
| 92 | + | ||
| 93 | + auto dA = CgeamAllocAndCopy(A, aBytes); | ||
| 94 | + auto dB = CgeamAllocAndCopy(B, bBytes); | ||
| 95 | + | ||
| 96 | + std::unique_ptr<DeviceBuffer> dC; | ||
| 97 | + if (!inplaceAC && !inplaceBC) { | ||
| 98 | + dC = CgeamAllocOutput(cBytes); | ||
| 99 | + } | ||
| 100 | + | ||
| 101 | + aclblasComplex* dC_ptr = CgeamGetCPtr(inplaceAC, inplaceBC, dA.get(), dB.get(), dC.get()); | ||
| 102 | + | ||
| 103 | + aclblasStatus_t ret = aclblasCgeam( | ||
| 104 | + handle, transa, transb, m, n, alpha, | ||
| 105 | + dA ? static_cast<const aclblasComplex*>(dA->ptr()) : nullptr, lda, beta, | ||
| 106 | + dB ? static_cast<const aclblasComplex*>(dB->ptr()) : nullptr, ldb, dC_ptr, ldc); | ||
| 107 | + | ||
| 108 | + if (aclrtSynchronizeDevice() != ACL_SUCCESS) { | ||
| 109 | + return ACLBLAS_STATUS_INTERNAL_ERROR; | ||
| 110 | + } | ||
| 111 | + if (ret == ACLBLAS_STATUS_SUCCESS) { | ||
| 112 | + CgeamCopyBack(inplaceAC, inplaceBC, dA.get(), dB.get(), dC.get(), C, aBytes, bBytes, cBytes); | ||
| 113 | + } | ||
| 114 | + return ret; | ||
| 115 | +} | ||
| @@ -0,0 +1,383 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | +// Generate a column-major complex matrix: m rows, n cols, column stride lda. | ||
| 27 | +// Returns interleaved float vector of size lda * n * 2. | ||
| 28 | +static inline std::vector<float> makeComplexMatrixCM(int m, int n, int lda, const BlasFillMode& fill, uint32_t seed) | ||
| 29 | +{ | ||
| 30 | + const size_t storageSize = static_cast<size_t>(lda) * n * 2; | ||
| 31 | + std::vector<float> data(storageSize, 0.0f); | ||
| 32 | + | ||
| 33 | + if (fill.method == BlasFillMode::M_VALUE) { | ||
| 34 | + for (size_t i = 0; i < storageSize; i++) | ||
| 35 | + data[i] = fill.val1; | ||
| 36 | + return data; | ||
| 37 | + } | ||
| 38 | + | ||
| 39 | + std::mt19937 rngReal(seed ? seed : 42); | ||
| 40 | + std::mt19937 rngImag((seed ? seed : 42) + 2000); | ||
| 41 | + auto genReal = createGenerator(fill, rngReal); | ||
| 42 | + auto genImag = createGenerator(fill, rngImag); | ||
| 43 | + | ||
| 44 | + for (int j = 0; j < n; j++) { | ||
| 45 | + for (int i = 0; i < m; i++) { | ||
| 46 | + size_t idx = (static_cast<size_t>(j) * lda + i) * 2; | ||
| 47 | + size_t flatIdx = static_cast<size_t>(j) * m + i; | ||
| 48 | + data[idx] = genReal->at(flatIdx); | ||
| 49 | + data[idx + 1] = genImag->at(flatIdx); | ||
| 50 | + } | ||
| 51 | + } | ||
| 52 | + return data; | ||
| 53 | +} | ||
| 54 | + | ||
| 55 | +// -- Test fixture ------------------------------------------------------------ | ||
| 56 | +class CgeamArch35Test : public BlasTest<CgeamParam> {}; | ||
| 57 | + | ||
| 58 | +// Helper: 4x4 complex matrix as interleaved float vector | ||
| 59 | +static inline std::vector<float> c4x4(float v = 1.0f) { return std::vector<float>(4 * 4 * 2, v); } | ||
| 60 | + | ||
| 61 | +// -- TEST_F: null handle (not in CSV) --------------------------------------- | ||
| 62 | +TEST_F(CgeamArch35Test, NullHandle) | ||
| 63 | +{ | ||
| 64 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 65 | + auto A = c4x4(), B = c4x4(), C = c4x4(0.0f); | ||
| 66 | + EXPECT_EQ( | ||
| 67 | + aclblasCgeam( | ||
| 68 | + nullptr, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, reinterpret_cast<const aclblasComplex*>(A.data()), 4, | ||
| 69 | + &beta, reinterpret_cast<const aclblasComplex*>(B.data()), 4, reinterpret_cast<aclblasComplex*>(C.data()), | ||
| 70 | + 4), | ||
| 71 | + ACLBLAS_STATUS_HANDLE_IS_NULLPTR); | ||
| 72 | +} | ||
| 73 | + | ||
| 74 | +// -- TEST_F: A=nullptr (host validates and returns error) -------------------- | ||
| 75 | +TEST_F(CgeamArch35Test, NullA) | ||
| 76 | +{ | ||
| 77 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 78 | + auto B = c4x4(), C = c4x4(0.0f); | ||
| 79 | + EXPECT_EQ( | ||
| 80 | + aclblasCgeam( | ||
| 81 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, nullptr, 4, &beta, | ||
| 82 | + reinterpret_cast<const aclblasComplex*>(B.data()), 4, reinterpret_cast<aclblasComplex*>(C.data()), 4), | ||
| 83 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 84 | +} | ||
| 85 | + | ||
| 86 | +// -- TEST_F: B=nullptr (host validates and returns error) -------------------- | ||
| 87 | +TEST_F(CgeamArch35Test, NullB) | ||
| 88 | +{ | ||
| 89 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 90 | + auto A = c4x4(), C = c4x4(0.0f); | ||
| 91 | + EXPECT_EQ( | ||
| 92 | + aclblasCgeam( | ||
| 93 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, | ||
| 94 | + reinterpret_cast<const aclblasComplex*>(A.data()), 4, &beta, nullptr, 4, | ||
| 95 | + reinterpret_cast<aclblasComplex*>(C.data()), 4), | ||
| 96 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 97 | +} | ||
| 98 | + | ||
| 99 | +// -- TEST_F: C=nullptr (host validates and returns error) -------------------- | ||
| 100 | +TEST_F(CgeamArch35Test, NullC) | ||
| 101 | +{ | ||
| 102 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 103 | + auto A = c4x4(), B = c4x4(); | ||
| 104 | + EXPECT_EQ( | ||
| 105 | + aclblasCgeam( | ||
| 106 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, | ||
| 107 | + reinterpret_cast<const aclblasComplex*>(A.data()), 4, &beta, | ||
| 108 | + reinterpret_cast<const aclblasComplex*>(B.data()), 4, nullptr, 4), | ||
| 109 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 110 | +} | ||
| 111 | + | ||
| 112 | +// -- TEST_F: alpha=nullptr returns INVALID_VALUE (host validates) ------------ | ||
| 113 | +TEST_F(CgeamArch35Test, NullAlpha) | ||
| 114 | +{ | ||
| 115 | + aclblasComplex beta = {1.0f, 0.0f}; | ||
| 116 | + auto A = c4x4(2.0f), B = c4x4(3.0f), C = c4x4(kBlasSentinel); | ||
| 117 | + EXPECT_EQ( | ||
| 118 | + aclblasCgeam( | ||
| 119 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, nullptr, | ||
| 120 | + reinterpret_cast<aclblasComplex*>(A.data()), 4, &beta, reinterpret_cast<aclblasComplex*>(B.data()), 4, | ||
| 121 | + reinterpret_cast<aclblasComplex*>(C.data()), 4), | ||
| 122 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 123 | +} | ||
| 124 | + | ||
| 125 | +// -- TEST_F: beta=nullptr returns INVALID_VALUE (host validates) ------------- | ||
| 126 | +TEST_F(CgeamArch35Test, NullBeta) | ||
| 127 | +{ | ||
| 128 | + aclblasComplex alpha = {1.0f, 0.0f}; | ||
| 129 | + auto A = c4x4(2.0f), B = c4x4(3.0f), C = c4x4(kBlasSentinel); | ||
| 130 | + EXPECT_EQ( | ||
| 131 | + aclblasCgeam( | ||
| 132 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, | ||
| 133 | + reinterpret_cast<aclblasComplex*>(A.data()), 4, nullptr, reinterpret_cast<aclblasComplex*>(B.data()), 4, | ||
| 134 | + reinterpret_cast<aclblasComplex*>(C.data()), 4), | ||
| 135 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 136 | +} | ||
| 137 | + | ||
| 138 | +// -- TEST_F: transa invalid 0xFF (host validates) ---------------------------- | ||
| 139 | +TEST_F(CgeamArch35Test, TransaInvalid) | ||
| 140 | +{ | ||
| 141 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 142 | + auto A = c4x4(), B = c4x4(), C = c4x4(0.0f); | ||
| 143 | + EXPECT_EQ( | ||
| 144 | + aclblasCgeam( | ||
| 145 | + CgeamArch35Test::handle_, static_cast<aclblasOperation_t>(0xFF), ACLBLAS_OP_N, 4, 4, &alpha, | ||
| 146 | + reinterpret_cast<const aclblasComplex*>(A.data()), 4, &beta, | ||
| 147 | + reinterpret_cast<const aclblasComplex*>(B.data()), 4, reinterpret_cast<aclblasComplex*>(C.data()), 4), | ||
| 148 | + ACLBLAS_STATUS_INVALID_ENUM); | ||
| 149 | +} | ||
| 150 | + | ||
| 151 | +// -- TEST_F: transb invalid 0xFF (host validates) ---------------------------- | ||
| 152 | +TEST_F(CgeamArch35Test, TransbInvalid) | ||
| 153 | +{ | ||
| 154 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 155 | + auto A = c4x4(), B = c4x4(), C = c4x4(0.0f); | ||
| 156 | + EXPECT_EQ( | ||
| 157 | + aclblasCgeam( | ||
| 158 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, static_cast<aclblasOperation_t>(0xFF), 4, 4, &alpha, | ||
| 159 | + reinterpret_cast<const aclblasComplex*>(A.data()), 4, &beta, | ||
| 160 | + reinterpret_cast<const aclblasComplex*>(B.data()), 4, reinterpret_cast<aclblasComplex*>(C.data()), 4), | ||
| 161 | + ACLBLAS_STATUS_INVALID_ENUM); | ||
| 162 | +} | ||
| 163 | + | ||
| 164 | +// -- TEST_F: m<0 (host validates and returns error) ------------------------- | ||
| 165 | +TEST_F(CgeamArch35Test, MNegative) | ||
| 166 | +{ | ||
| 167 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 168 | + EXPECT_EQ( | ||
| 169 | + aclblasCgeam( | ||
| 170 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, -1, 4, &alpha, nullptr, 1, &beta, nullptr, 1, nullptr, | ||
| 171 | + 1), | ||
| 172 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 173 | +} | ||
| 174 | + | ||
| 175 | +// -- TEST_F: n<0 (host validates and returns error) ------------------------- | ||
| 176 | +TEST_F(CgeamArch35Test, NNegative) | ||
| 177 | +{ | ||
| 178 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 179 | + EXPECT_EQ( | ||
| 180 | + aclblasCgeam( | ||
| 181 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, -1, &alpha, nullptr, 4, &beta, nullptr, 4, nullptr, | ||
| 182 | + 4), | ||
| 183 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 184 | +} | ||
| 185 | + | ||
| 186 | +// -- TEST_F: lda too small when transa=N (host validates) -------------------- | ||
| 187 | +TEST_F(CgeamArch35Test, LdaTooSmallN) | ||
| 188 | +{ | ||
| 189 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 190 | + auto A = c4x4(), B = c4x4(), C = c4x4(0.0f); | ||
| 191 | + EXPECT_EQ( | ||
| 192 | + aclblasCgeam( | ||
| 193 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, | ||
| 194 | + reinterpret_cast<const aclblasComplex*>(A.data()), 3, &beta, | ||
| 195 | + reinterpret_cast<const aclblasComplex*>(B.data()), 4, reinterpret_cast<aclblasComplex*>(C.data()), 4), | ||
| 196 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 197 | +} | ||
| 198 | + | ||
| 199 | +// -- TEST_F: ldc too small (host validates and returns error) ---------------- | ||
| 200 | +TEST_F(CgeamArch35Test, LdcTooSmall) | ||
| 201 | +{ | ||
| 202 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 203 | + std::vector<float> A(8 * 4 * 2, 1.0f), B(8 * 4 * 2, 1.0f), C(8 * 4 * 2, 0.0f); | ||
| 204 | + EXPECT_EQ( | ||
| 205 | + aclblasCgeam( | ||
| 206 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 8, 4, &alpha, | ||
| 207 | + reinterpret_cast<const aclblasComplex*>(A.data()), 8, &beta, | ||
| 208 | + reinterpret_cast<const aclblasComplex*>(B.data()), 8, reinterpret_cast<aclblasComplex*>(C.data()), 7), | ||
| 209 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 210 | +} | ||
| 211 | + | ||
| 212 | +// -- TEST_F: in-place C==A with transa!=N (host validates) ------------------- | ||
| 213 | +TEST_F(CgeamArch35Test, InplaceATransInvalid) | ||
| 214 | +{ | ||
| 215 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 216 | + std::vector<float> A(8 * 8 * 2, 1.0f), B(8 * 8 * 2, 1.0f); | ||
| 217 | + EXPECT_EQ( | ||
| 218 | + aclblasCgeam( | ||
| 219 | + CgeamArch35Test::handle_, ACLBLAS_OP_T, ACLBLAS_OP_N, 8, 8, &alpha, | ||
| 220 | + reinterpret_cast<const aclblasComplex*>(A.data()), 8, &beta, | ||
| 221 | + reinterpret_cast<const aclblasComplex*>(B.data()), 8, reinterpret_cast<aclblasComplex*>(A.data()), 8), | ||
| 222 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 223 | +} | ||
| 224 | + | ||
| 225 | +// -- TEST_F: in-place C==B with transb!=N (host validates) ------------------- | ||
| 226 | +TEST_F(CgeamArch35Test, InplaceBTransInvalid) | ||
| 227 | +{ | ||
| 228 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 229 | + std::vector<float> A(8 * 8 * 2, 1.0f), B(8 * 8 * 2, 1.0f); | ||
| 230 | + EXPECT_EQ( | ||
| 231 | + aclblasCgeam( | ||
| 232 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_T, 8, 8, &alpha, | ||
| 233 | + reinterpret_cast<const aclblasComplex*>(A.data()), 8, &beta, | ||
| 234 | + reinterpret_cast<const aclblasComplex*>(B.data()), 8, reinterpret_cast<aclblasComplex*>(B.data()), 8), | ||
| 235 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 236 | +} | ||
| 237 | + | ||
| 238 | +// -- TEST_F: in-place C==A with lda!=ldc (host validates) -------------------- | ||
| 239 | +TEST_F(CgeamArch35Test, InplaceALdaNeLdc) | ||
| 240 | +{ | ||
| 241 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 242 | + std::vector<float> A(10 * 8 * 2, 1.0f), B(10 * 8 * 2, 1.0f); | ||
| 243 | + EXPECT_EQ( | ||
| 244 | + aclblasCgeam( | ||
| 245 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 8, 8, &alpha, | ||
| 246 | + reinterpret_cast<const aclblasComplex*>(A.data()), 10, &beta, | ||
| 247 | + reinterpret_cast<const aclblasComplex*>(B.data()), 8, reinterpret_cast<aclblasComplex*>(A.data()), 8), | ||
| 248 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 249 | +} | ||
| 250 | + | ||
| 251 | +// -- TEST_F: in-place C==B with ldb!=ldc (host validates) -------------------- | ||
| 252 | +TEST_F(CgeamArch35Test, InplaceBLdbNeLdc) | ||
| 253 | +{ | ||
| 254 | + aclblasComplex alpha = {1.0f, 0.0f}, beta = {1.0f, 0.0f}; | ||
| 255 | + std::vector<float> A(10 * 8 * 2, 1.0f), B(10 * 8 * 2, 1.0f); | ||
| 256 | + EXPECT_EQ( | ||
| 257 | + aclblasCgeam( | ||
| 258 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 8, 8, &alpha, | ||
| 259 | + reinterpret_cast<const aclblasComplex*>(A.data()), 8, &beta, | ||
| 260 | + reinterpret_cast<const aclblasComplex*>(B.data()), 10, reinterpret_cast<aclblasComplex*>(B.data()), 8), | ||
| 261 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 262 | +} | ||
| 263 | + | ||
| 264 | +// -- TEST_F: alpha=nullptr with A=nullptr returns INVALID_VALUE ------------- | ||
| 265 | +TEST_F(CgeamArch35Test, NullAlphaAllowsNullA) | ||
| 266 | +{ | ||
| 267 | + aclblasComplex beta = {1.0f, 0.0f}; | ||
| 268 | + auto B = c4x4(), C = c4x4(kBlasSentinel); | ||
| 269 | + EXPECT_EQ( | ||
| 270 | + aclblasCgeam( | ||
| 271 | + CgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, nullptr, nullptr, 4, &beta, | ||
| 272 | + reinterpret_cast<aclblasComplex*>(B.data()), 4, reinterpret_cast<aclblasComplex*>(C.data()), 4), | ||
| 273 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 274 | +} | ||
| 275 | + | ||
| 276 | +// -- CSV parameterised test suite ------------------------------------------- | ||
| 277 | +INSTANTIATE_TEST_SUITE_P( | ||
| 278 | + Cgeam, CgeamArch35Test, ::testing::ValuesIn(GetCasesFromCsv<CgeamParam>(ReplaceFileExtension2Csv(__FILE__))), | ||
| 279 | + PrintCaseInfoString<CgeamParam>); | ||
| 280 | + | ||
| 281 | +// -- Helper: generate A or B input complex matrix ----------------------------- | ||
| 282 | +static std::vector<float> CgeamGenerateInput(const CgeamParam& p, bool isA, int extraSeed) | ||
| 283 | +{ | ||
| 284 | + int nullFlag = isA ? p.nullA : p.nullB; | ||
| 285 | + if (nullFlag != 0 || p.m <= 0 || p.n <= 0) { | ||
| 286 | + return {}; | ||
| 287 | + } | ||
| 288 | + aclblasOperation_t trans = isA ? p.transa : p.transb; | ||
| 289 | + int ld = isA ? p.lda : p.ldb; | ||
| 290 | + const BlasFillMode& fill = isA ? p.aFill : p.bFill; | ||
| 291 | + int rows = (trans == ACLBLAS_OP_N) ? p.m : p.n; | ||
| 292 | + int cols = (trans == ACLBLAS_OP_N) ? p.n : p.m; | ||
| 293 | + return makeComplexMatrixCM(rows, cols, ld, fill, p.randomSeed + extraSeed); | ||
| 294 | +} | ||
| 295 | + | ||
| 296 | +// -- Helper: prepare C buffer (in-place aware, complex) ----------------------- | ||
| 297 | +static aclblasComplex* CgeamPrepareC(const CgeamParam& p, const aclblasComplex* aPtr, | ||
| 298 | + const aclblasComplex* bPtr, std::vector<float>& cHost) | ||
| 299 | +{ | ||
| 300 | + if (p.inplace == 1 && aPtr) { | ||
| 301 | + return const_cast<aclblasComplex*>(aPtr); | ||
| 302 | + } | ||
| 303 | + if (p.inplace == 2 && bPtr) { | ||
| 304 | + return const_cast<aclblasComplex*>(bPtr); | ||
| 305 | + } | ||
| 306 | + if (p.nullC == 0 && p.m > 0 && p.n > 0) { | ||
| 307 | + cHost.assign(static_cast<size_t>(p.ldc) * static_cast<size_t>(p.n) * 2, kBlasSentinel); | ||
| 308 | + } | ||
| 309 | + return cHost.empty() ? nullptr : reinterpret_cast<aclblasComplex*>(cHost.data()); | ||
| 310 | +} | ||
| 311 | + | ||
| 312 | +// -- Helper: compute golden reference (complex) -------------------------------- | ||
| 313 | +static std::vector<float> CgeamComputeGolden( | ||
| 314 | + const CgeamParam& p, aclblasComplex alpha, aclblasComplex beta, | ||
| 315 | + const aclblasComplex* aPtr, const aclblasComplex* bPtr, aclblasComplex* cPtr, | ||
| 316 | + const std::vector<float>& aHost, const std::vector<float>& bHost, const std::vector<float>& cHost) | ||
| 317 | +{ | ||
| 318 | + std::vector<float> goldenC; | ||
| 319 | + if (p.expectResult != ACLBLAS_STATUS_SUCCESS || p.m <= 0 || p.n <= 0 || !cPtr) { | ||
| 320 | + return goldenC; | ||
| 321 | + } | ||
| 322 | + size_t goldenSize = cHost.empty() ? (p.inplace == 1 ? aHost.size() : | ||
| 323 | + p.inplace == 2 ? bHost.size() : 0) : | ||
| 324 | + cHost.size(); | ||
| 325 | + goldenC.assign(goldenSize, kBlasSentinel); | ||
| 326 | + aclblasGeam_cpu<std::complex<float>>( | ||
| 327 | + p.transa, p.transb, static_cast<std::size_t>(p.m), static_cast<std::size_t>(p.n), | ||
| 328 | + *reinterpret_cast<const std::complex<float>*>(&alpha), reinterpret_cast<const std::complex<float>*>(aPtr), | ||
| 329 | + static_cast<std::size_t>(p.lda), *reinterpret_cast<const std::complex<float>*>(&beta), | ||
| 330 | + reinterpret_cast<const std::complex<float>*>(bPtr), static_cast<std::size_t>(p.ldb), | ||
| 331 | + reinterpret_cast<std::complex<float>*>(goldenC.data()), static_cast<std::size_t>(p.ldc)); | ||
| 332 | + return goldenC; | ||
| 333 | +} | ||
| 334 | + | ||
| 335 | +// -- Helper: verify output precision (complex) --------------------------------- | ||
| 336 | +static void CgeamVerifyOutput( | ||
| 337 | + const CgeamParam& p, const std::vector<float>& goldenC, | ||
| 338 | + const std::vector<float>& aHost, const std::vector<float>& bHost, const std::vector<float>& cHost) | ||
| 339 | +{ | ||
| 340 | + VerifyConfig cfg; | ||
| 341 | + cfg.mode = PrecisionMode::MERE_MARE; | ||
| 342 | + cfg.mereThreshold = p.mereThreshold; | ||
| 343 | + cfg.mareMultiplier = p.mareMultiplier; | ||
| 344 | + const float* outPtr = nullptr; | ||
| 345 | + size_t outSize = 0; | ||
| 346 | + if (p.inplace == 1) { | ||
| 347 | + outPtr = aHost.data(); | ||
| 348 | + outSize = aHost.size(); | ||
| 349 | + } else if (p.inplace == 2) { | ||
| 350 | + outPtr = bHost.data(); | ||
| 351 | + outSize = bHost.size(); | ||
| 352 | + } else { | ||
| 353 | + outPtr = cHost.data(); | ||
| 354 | + outSize = cHost.size(); | ||
| 355 | + } | ||
| 356 | + if (outSize > 0) { | ||
| 357 | + EXPECT_TRUE(Verifier::verifyVector(outPtr, goldenC.data(), outSize, 1, cfg, p.caseName)); | ||
| 358 | + } | ||
| 359 | +} | ||
| 360 | + | ||
| 361 | +// -- TEST_P: 5-step CSV-driven flow ----------------------------------------- | ||
| 362 | +TEST_P(CgeamArch35Test, CsvDriven) | ||
| 363 | +{ | ||
| 364 | + const auto& p = GetParam(); | ||
| 365 | + aclblasComplex alpha = {p.alphaFill.val1, p.alphaFill.val2}; | ||
| 366 | + aclblasComplex beta = {p.betaFill.val1, p.betaFill.val2}; | ||
| 367 | + | ||
| 368 | + std::vector<float> aHost = CgeamGenerateInput(p, true, 0); | ||
| 369 | + std::vector<float> bHost = CgeamGenerateInput(p, false, 1); | ||
| 370 | + const aclblasComplex* aPtr = aHost.empty() ? nullptr : reinterpret_cast<const aclblasComplex*>(aHost.data()); | ||
| 371 | + const aclblasComplex* bPtr = bHost.empty() ? nullptr : reinterpret_cast<const aclblasComplex*>(bHost.data()); | ||
| 372 | + | ||
| 373 | + std::vector<float> cHost; | ||
| 374 | + aclblasComplex* cPtr = CgeamPrepareC(p, aPtr, bPtr, cHost); | ||
| 375 | + | ||
| 376 | + std::vector<float> goldenC = CgeamComputeGolden(p, alpha, beta, aPtr, bPtr, cPtr, aHost, bHost, cHost); | ||
| 377 | + | ||
| 378 | + aclblasStatus_t ret = aclblasCgeam_npu( | ||
| 379 | + CgeamArch35Test::handle_, p.transa, p.transb, p.m, p.n, &alpha, aPtr, p.lda, &beta, bPtr, p.ldb, cPtr, p.ldc); | ||
| 380 | + EXPECT_EQ(ret, ACLBLAS_STATUS_SUCCESS); | ||
| 381 | + | ||
| 382 | + CgeamVerifyOutput(p, goldenC, aHost, bHost, cHost); | ||
| 383 | +} | ||
| @@ -0,0 +1,40 @@ | |||
| 1 | +case_name,description,transa,transb,m,n,lda,ldb,ldc,alpha_fill,beta_fill,a_fill,b_fill,nullA,nullB,nullC,mere_threshold,mare_multiplier,random_seed,inplace | ||
| 2 | +TC_CL0_01,NN 2x2 basic complex,N,N,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 3 | +TC_CL0_02,NT 2x2 basic complex,N,T,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 4 | +TC_CL0_03,NC 2x2 basic complex,N,C,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 5 | +TC_CL0_04,TN 2x2 basic complex,T,N,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 6 | +TC_CL0_05,TT 2x2 basic complex,T,T,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 7 | +TC_CL0_06,TC 2x2 basic complex,T,C,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 8 | +TC_CL0_07,CN 2x2 basic complex,C,N,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 9 | +TC_CL0_08,CT 2x2 basic complex,C,T,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 10 | +TC_CL0_09,CC 2x2 basic complex,C,C,2,2,2,2,2,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 11 | +TC_CL0_10,TC 4x4 conj vs trans distinction,T,C,4,4,4,4,4,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 12 | +TC_CL0_11,CT 4x4 conj vs trans distinction,C,T,4,4,4,4,4,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 13 | +TC_CL0_12,CC 4x4 dual conjugate,C,C,4,4,4,4,4,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 14 | +TC_CL0_13,alpha=(0.5 0.5) complex scaling,N,N,4,4,4,4,4,VALUE_NORM_0.5_0.5,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 15 | +TC_CL0_14,beta=(0 1) pure imaginary,N,N,4,4,4,4,4,VALUE_NORM_1_0,VALUE_NORM_0_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 16 | +TC_CL0_15,alpha=(0 0) beta=(1 1),N,N,4,4,4,4,4,VALUE_NORM_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 17 | +TC_CL0_16,NN 4x4 alpha=0 beta=1 short-circuit,N,N,4,4,4,4,4,VALUE_NORM_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 18 | +TC_CL0_17,NN 4x4 alpha=1 beta=0 short-circuit,N,N,4,4,4,4,4,VALUE_NORM_1_0,VALUE_NORM_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 19 | +TC_CL0_21,m=0 early return complex,N,N,0,4,1,1,1,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0,0,42,0 | ||
| 20 | +TC_CL0_22,n=0 early return complex,N,N,4,0,4,4,4,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0,0,42,0 | ||
| 21 | +TC_CL1_01,m=1 n=8 single-row NN complex,N,N,1,8,1,1,1,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 22 | +TC_CL1_02,m=8 n=1 single-col NN complex,N,N,8,1,8,8,8,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 23 | +TC_CL1_03,m=1024 n=1024 large data NN complex,N,N,1024,1024,1024,1024,1024,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_5_5,RANDOM_5_5,0,0,0,0.0001220703125,10.0,42,0 | ||
| 24 | +TC_CL1_04,ld padding +1 all matrices NN complex,N,N,8,8,9,9,9,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 25 | +TC_CL1_05,alpha=(0 1) pure-imag scaling CN complex,C,N,4,4,4,4,4,VALUE_NORM_0_1,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 26 | +TC_CL1_06,beta=(0.5 0.5) complex scaling NN,N,N,8,8,8,8,8,VALUE_NORM_1_0,VALUE_NORM_0.5_0.5,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 27 | +TC_CL1_07,conj correctness CN beta=0 isolate A,C,N,4,4,4,4,4,VALUE_NORM_1_0,VALUE_NORM_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 28 | +TC_CL1_08,dual conjugate CC complex m=n=8,C,C,8,8,8,8,8,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 29 | +TC_CL1_09,alpha=INF real part complex NN,N,N,4,4,4,4,4,VALUE_NORM_INF,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 30 | +TC_CL1_11,m=127 n=63 non-power2 TT complex,T,T,127,63,63,63,127,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 31 | +TC_CL1_12,m=8 n=16 non-square NN complex,N,N,8,16,8,8,8,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 32 | +TC_CL1_13,m=5000 n=1 NN complex batch loop,N,N,5000,1,5000,5000,5000,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 33 | +TC_CL1_14,in-place C==A transa=N lda=ldc complex,N,N,8,8,8,8,8,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,1 | ||
| 34 | +TC_CL1_15,in-place C==B transb=N ldb=ldc complex,N,N,8,8,8,8,8,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,2 | ||
| 35 | +TC_CL1_18,in-place C==A large N=64 complex,N,N,64,64,64,64,64,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,1 | ||
| 36 | +TC_CL1_19,in-place C==A alpha=0 beta=0 complex,N,N,8,8,8,8,8,VALUE_NORM_0,VALUE_NORM_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,1 | ||
| 37 | +TC_CL1_20,in-place C==A large N=128 complex,N,N,128,128,128,128,128,VALUE_NORM_1_0,VALUE_NORM_1_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,1 | ||
| 38 | +TC_CL1_21,alpha=0 A=nullptr NN complex,N,N,8,8,8,8,8,VALUE_NORM_0,VALUE_NORM_1_0,NULLPTR,RANDOM_1_1,1,0,0,0.0001220703125,10.0,42,0 | ||
| 39 | +TC_CL1_22,alpha=0 beta=0 both nullptr NN complex,N,N,8,8,8,8,8,VALUE_NORM_0,VALUE_NORM_0,NULLPTR,NULLPTR,1,1,0,0.0001220703125,10.0,42,0 | ||
| 40 | +TC_CL1_23,beta=0 B=nullptr alpha=1 NN complex,N,N,8,8,8,8,8,VALUE_NORM_1_0,VALUE_NORM_0,RANDOM_1_1,NULLPTR,0,1,0,0.0001220703125,10.0,42,0 | ||
| @@ -0,0 +1,54 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | +struct CgeamParam : public BlasTestParamBase { | ||
| 19 | + aclblasOperation_t transa = ACLBLAS_OP_N; | ||
| 20 | + aclblasOperation_t transb = ACLBLAS_OP_N; | ||
| 21 | + int m = 0; | ||
| 22 | + int n = 0; | ||
| 23 | + int lda = 0; | ||
| 24 | + int ldb = 0; | ||
| 25 | + int ldc = 0; | ||
| 26 | + // alpha/beta encoded as BlasFillMode: val1=real, val2=imag | ||
| 27 | + BlasFillMode alphaFill = BlasFillMode("VALUE_NORM_1"); | ||
| 28 | + BlasFillMode betaFill = BlasFillMode("VALUE_NORM_1"); | ||
| 29 | + BlasFillMode aFill = BlasFillMode("RANDOM_1_1"); | ||
| 30 | + BlasFillMode bFill = BlasFillMode("RANDOM_1_1"); | ||
| 31 | + int nullA = 0; | ||
| 32 | + int nullB = 0; | ||
| 33 | + int nullC = 0; | ||
| 34 | + int inplace = 0; // 0=out-of-place, 1=C==A, 2=C==B | ||
| 35 | + | ||
| 36 | + CgeamParam(const csv_map& csv) : BlasTestParamBase(csv) | ||
| 37 | + { | ||
| 38 | + transa = parseOpTrans(ReadMap(csv, "transa", "N")); | ||
| 39 | + transb = parseOpTrans(ReadMap(csv, "transb", "N")); | ||
| 40 | + m = parseInt(ReadMap(csv, "m", "0")); | ||
| 41 | + n = parseInt(ReadMap(csv, "n", "0")); | ||
| 42 | + lda = parseInt(ReadMap(csv, "lda", std::to_string((transa == ACLBLAS_OP_N) ? std::max(1, m) : std::max(1, n)))); | ||
| 43 | + ldb = parseInt(ReadMap(csv, "ldb", std::to_string((transb == ACLBLAS_OP_N) ? std::max(1, m) : std::max(1, n)))); | ||
| 44 | + ldc = parseInt(ReadMap(csv, "ldc", std::to_string(std::max(1, m)))); | ||
| 45 | + alphaFill = BlasFillMode(ReadMap(csv, "alpha_fill", "VALUE_NORM_1")); | ||
| 46 | + betaFill = BlasFillMode(ReadMap(csv, "beta_fill", "VALUE_NORM_1")); | ||
| 47 | + aFill = BlasFillMode(ReadMap(csv, "a_fill", "RANDOM_1_1")); | ||
| 48 | + bFill = BlasFillMode(ReadMap(csv, "b_fill", "RANDOM_1_1")); | ||
| 49 | + nullA = parseInt(ReadMap(csv, "nullA", "0")); | ||
| 50 | + nullB = parseInt(ReadMap(csv, "nullB", "0")); | ||
| 51 | + nullC = parseInt(ReadMap(csv, "nullC", "0")); | ||
| 52 | + inplace = parseInt(ReadMap(csv, "inplace", "0")); | ||
| 53 | + } | ||
| 54 | +}; | ||
| @@ -0,0 +1,74 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | +// todo: 优化 cmake 让 family/dtype/arch/test.cpp 可以直接找到 family/golden.h | ||
| 18 | + | ||
| 19 | +// std::conj(float) 返回 complex<float> ， complex 不能自动转换成 float | ||
| 20 | +// 这里定义一个直接返回 float 的 conj | ||
| 21 | +template <class T> | ||
| 22 | +constexpr T conjugate(T x) | ||
| 23 | +{ | ||
| 24 | + return x; | ||
| 25 | +} | ||
| 26 | +template <class T> | ||
| 27 | +constexpr std::complex<T> conjugate(std::complex<T> x) | ||
| 28 | +{ | ||
| 29 | + return std::conj(x); | ||
| 30 | +} | ||
| 31 | + | ||
| 32 | +static_assert(conjugate(1.0f) == 1.0f); | ||
| 33 | +// std::conj is constexpr since c++20 | ||
| 34 | +// using namespace std::complex_literals; | ||
| 35 | +// add this code if upgrade to c++20 `static_assert(conjugate(1.0if) == -1.0if);` | ||
| 36 | + | ||
| 37 | +template <class T> | ||
| 38 | +inline void aclblasGeam_cpu( | ||
| 39 | + aclblasOperation_t transa, aclblasOperation_t transb, std::size_t m, std::size_t n, T alpha, const T* A, | ||
| 40 | + std::size_t lda, T beta, const T* B, std::size_t ldb, T* C, std::size_t ldc) | ||
| 41 | +{ | ||
| 42 | + for (std::size_t j = 0; j < n; j++) { | ||
| 43 | + for (std::size_t i = 0; i < m; i++) { | ||
| 44 | + // maybe A == C or B == C | ||
| 45 | + T cij = 0; | ||
| 46 | + if (alpha != 0.0f) { | ||
| 47 | + // ignore nan in A, when alpha == 0 | ||
| 48 | + T aij; | ||
| 49 | + if (transa == ACLBLAS_OP_N) { | ||
| 50 | + aij = A[i + j * lda]; | ||
| 51 | + } else { | ||
| 52 | + aij = A[j + i * lda]; | ||
| 53 | + } | ||
| 54 | + if (transa == ACLBLAS_OP_C) { | ||
| 55 | + aij = conjugate(aij); | ||
| 56 | + } | ||
| 57 | + cij += alpha * aij; | ||
| 58 | + } | ||
| 59 | + if (beta != 0.0f) { | ||
| 60 | + T bij; | ||
| 61 | + if (transb == ACLBLAS_OP_N) { | ||
| 62 | + bij = B[i + j * ldb]; | ||
| 63 | + } else { | ||
| 64 | + bij = B[j + i * ldb]; | ||
| 65 | + } | ||
| 66 | + if (transb == ACLBLAS_OP_C) { | ||
| 67 | + bij = conjugate(bij); | ||
| 68 | + } | ||
| 69 | + cij += beta * bij; | ||
| 70 | + } | ||
| 71 | + C[i + j * ldc] = cij; | ||
| 72 | + } | ||
| 73 | + } | ||
| 74 | +} | ||
| @@ -0,0 +1,11 @@ | |||
| 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 FILE 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 | +ops_blas_add_gtest_tests(${OPS_BLAS} sgeam_test) | ||
| @@ -0,0 +1,113 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | +inline size_t SgeamBufferBytes(aclblasOperation_t trans, int m, int n, int ld) | ||
| 24 | +{ | ||
| 25 | + size_t cols = (trans == ACLBLAS_OP_N) ? static_cast<size_t>(n) : static_cast<size_t>(m); | ||
| 26 | + return static_cast<size_t>(ld) * cols * sizeof(float); | ||
| 27 | +} | ||
| 28 | + | ||
| 29 | +inline std::unique_ptr<DeviceBuffer> SgeamAllocAndCopy(const float* hostPtr, size_t bytes) | ||
| 30 | +{ | ||
| 31 | + if (hostPtr == nullptr) { | ||
| 32 | + return nullptr; | ||
| 33 | + } | ||
| 34 | + auto buf = std::make_unique<DeviceBuffer>(bytes); | ||
| 35 | + buf->copyFromHost(hostPtr, bytes); | ||
| 36 | + return buf; | ||
| 37 | +} | ||
| 38 | + | ||
| 39 | +inline std::unique_ptr<DeviceBuffer> SgeamAllocOutput(size_t cBytes) | ||
| 40 | +{ | ||
| 41 | + auto buf = std::make_unique<DeviceBuffer>(cBytes); | ||
| 42 | + std::vector<float> sentinel(cBytes / sizeof(float), kBlasSentinel); | ||
| 43 | + buf->copyFromHost(sentinel.data(), cBytes); | ||
| 44 | + return buf; | ||
| 45 | +} | ||
| 46 | + | ||
| 47 | +inline float* SgeamGetCPtr(bool inplaceAC, bool inplaceBC, DeviceBuffer* dA, DeviceBuffer* dB, DeviceBuffer* dC) | ||
| 48 | +{ | ||
| 49 | + if (inplaceAC) { | ||
| 50 | + return static_cast<float*>(dA->ptr()); | ||
| 51 | + } | ||
| 52 | + if (inplaceBC) { | ||
| 53 | + return static_cast<float*>(dB->ptr()); | ||
| 54 | + } | ||
| 55 | + return static_cast<float*>(dC->ptr()); | ||
| 56 | +} | ||
| 57 | + | ||
| 58 | +inline void SgeamCopyBack(bool inplaceAC, bool inplaceBC, DeviceBuffer* dA, DeviceBuffer* dB, DeviceBuffer* dC, | ||
| 59 | + float* C, size_t aBytes, size_t bBytes, size_t cBytes) | ||
| 60 | +{ | ||
| 61 | + if (inplaceAC && dA) { | ||
| 62 | + dA->copyToHost(C, aBytes); | ||
| 63 | + } else if (inplaceBC && dB) { | ||
| 64 | + dB->copyToHost(C, bBytes); | ||
| 65 | + } else if (dC) { | ||
| 66 | + dC->copyToHost(C, cBytes); | ||
| 67 | + } | ||
| 68 | +} | ||
| 69 | + | ||
| 70 | +inline aclblasStatus_t aclblasSgeam_npu( | ||
| 71 | + aclblasHandle_t handle, aclblasOperation_t transa, aclblasOperation_t transb, int m, int n, | ||
| 72 | + const float* alpha, const float* A, int lda, | ||
| 73 | + const float* beta, const float* B, int ldb, float* C, int ldc) | ||
| 74 | +{ | ||
| 75 | + if (m <= 0 || n <= 0) { | ||
| 76 | + float dummyA = 0.0f, dummyB = 0.0f, dummyC = 0.0f; | ||
| 77 | + float dummyAlpha = alpha ? *alpha : 0.0f; | ||
| 78 | + float dummyBeta = beta ? *beta : 0.0f; | ||
| 79 | + return aclblasSgeam(handle, transa, transb, m, n, | ||
| 80 | + &dummyAlpha, &dummyA, std::max(lda, 1), | ||
| 81 | + &dummyBeta, &dummyB, std::max(ldb, 1), | ||
| 82 | + &dummyC, std::max(ldc, 1)); | ||
| 83 | + } | ||
| 84 | + size_t aBytes = SgeamBufferBytes(transa, m, n, lda); | ||
| 85 | + size_t bBytes = SgeamBufferBytes(transb, m, n, ldb); | ||
| 86 | + size_t cBytes = static_cast<size_t>(ldc) * static_cast<size_t>(n) * sizeof(float); | ||
| 87 | + | ||
| 88 | + bool inplaceAC = (A != nullptr && C == A); | ||
| 89 | + bool inplaceBC = (B != nullptr && C == B); | ||
| 90 | + | ||
| 91 | + auto dA = SgeamAllocAndCopy(A, aBytes); | ||
| 92 | + auto dB = SgeamAllocAndCopy(B, bBytes); | ||
| 93 | + | ||
| 94 | + std::unique_ptr<DeviceBuffer> dC; | ||
| 95 | + if (!inplaceAC && !inplaceBC) { | ||
| 96 | + dC = SgeamAllocOutput(cBytes); | ||
| 97 | + } | ||
| 98 | + | ||
| 99 | + float* dC_ptr = SgeamGetCPtr(inplaceAC, inplaceBC, dA.get(), dB.get(), dC.get()); | ||
| 100 | + | ||
| 101 | + aclblasStatus_t ret = aclblasSgeam( | ||
| 102 | + handle, transa, transb, m, n, alpha, | ||
| 103 | + dA ? static_cast<const float*>(dA->ptr()) : nullptr, lda, beta, | ||
| 104 | + dB ? static_cast<const float*>(dB->ptr()) : nullptr, ldb, dC_ptr, ldc); | ||
| 105 | + | ||
| 106 | + if (aclrtSynchronizeDevice() != ACL_SUCCESS) { | ||
| 107 | + return ACLBLAS_STATUS_INTERNAL_ERROR; | ||
| 108 | + } | ||
| 109 | + if (ret == ACLBLAS_STATUS_SUCCESS) { | ||
| 110 | + SgeamCopyBack(inplaceAC, inplaceBC, dA.get(), dB.get(), dC.get(), C, aBytes, bBytes, cBytes); | ||
| 111 | + } | ||
| 112 | + return ret; | ||
| 113 | +} | ||
| @@ -0,0 +1,366 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | +// -- Test fixture ------------------------------------------------------------ | ||
| 25 | +class SgeamArch35Test : public BlasTest<SgeamParam> {}; | ||
| 26 | + | ||
| 27 | +// -- TEST_F: null handle (not in CSV) --------------------------------------- | ||
| 28 | +TEST_F(SgeamArch35Test, NullHandle) | ||
| 29 | +{ | ||
| 30 | + float alpha = 1.0f, beta = 1.0f; | ||
| 31 | + std::vector<float> A(4 * 4, 1.0f); | ||
| 32 | + std::vector<float> B(4 * 4, 1.0f); | ||
| 33 | + std::vector<float> C(4 * 4, 0.0f); | ||
| 34 | + EXPECT_EQ( | ||
| 35 | + aclblasSgeam(nullptr, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, A.data(), 4, &beta, B.data(), 4, C.data(), 4), | ||
| 36 | + ACLBLAS_STATUS_HANDLE_IS_NULLPTR); | ||
| 37 | +} | ||
| 38 | + | ||
| 39 | +// -- TEST_F: A=nullptr (host validates and returns error) -------------------- | ||
| 40 | +TEST_F(SgeamArch35Test, NullA) | ||
| 41 | +{ | ||
| 42 | + float alpha = 1.0f, beta = 1.0f; | ||
| 43 | + std::vector<float> B(4 * 4, 1.0f); | ||
| 44 | + std::vector<float> C(4 * 4, 0.0f); | ||
| 45 | + EXPECT_EQ( | ||
| 46 | + aclblasSgeam( | ||
| 47 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, nullptr, 4, &beta, B.data(), 4, | ||
| 48 | + C.data(), 4), | ||
| 49 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 50 | +} | ||
| 51 | + | ||
| 52 | +// -- TEST_F: B=nullptr (host validates and returns error) -------------------- | ||
| 53 | +TEST_F(SgeamArch35Test, NullB) | ||
| 54 | +{ | ||
| 55 | + float alpha = 1.0f, beta = 1.0f; | ||
| 56 | + std::vector<float> A(4 * 4, 1.0f); | ||
| 57 | + std::vector<float> C(4 * 4, 0.0f); | ||
| 58 | + EXPECT_EQ( | ||
| 59 | + aclblasSgeam( | ||
| 60 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, A.data(), 4, &beta, nullptr, 4, | ||
| 61 | + C.data(), 4), | ||
| 62 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 63 | +} | ||
| 64 | + | ||
| 65 | +// -- TEST_F: C=nullptr (host validates and returns error) -------------------- | ||
| 66 | +TEST_F(SgeamArch35Test, NullC) | ||
| 67 | +{ | ||
| 68 | + float alpha = 1.0f, beta = 1.0f; | ||
| 69 | + std::vector<float> A(4 * 4, 1.0f); | ||
| 70 | + std::vector<float> B(4 * 4, 1.0f); | ||
| 71 | + EXPECT_EQ( | ||
| 72 | + aclblasSgeam( | ||
| 73 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, A.data(), 4, &beta, B.data(), 4, | ||
| 74 | + nullptr, 4), | ||
| 75 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 76 | +} | ||
| 77 | + | ||
| 78 | +// -- TEST_F: alpha=nullptr returns INVALID_VALUE (host validates) ------------ | ||
| 79 | +TEST_F(SgeamArch35Test, NullAlpha) | ||
| 80 | +{ | ||
| 81 | + float beta = 1.0f; | ||
| 82 | + std::vector<float> A(4 * 4, 2.0f); | ||
| 83 | + std::vector<float> B(4 * 4, 3.0f); | ||
| 84 | + std::vector<float> C(4 * 4, kBlasSentinel); | ||
| 85 | + EXPECT_EQ( | ||
| 86 | + aclblasSgeam( | ||
| 87 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, nullptr, A.data(), 4, &beta, B.data(), 4, | ||
| 88 | + C.data(), 4), | ||
| 89 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 90 | +} | ||
| 91 | + | ||
| 92 | +// -- TEST_F: beta=nullptr returns INVALID_VALUE (host validates) ------------- | ||
| 93 | +TEST_F(SgeamArch35Test, NullBeta) | ||
| 94 | +{ | ||
| 95 | + float alpha = 1.0f; | ||
| 96 | + std::vector<float> A(4 * 4, 2.0f); | ||
| 97 | + std::vector<float> B(4 * 4, 3.0f); | ||
| 98 | + std::vector<float> C(4 * 4, kBlasSentinel); | ||
| 99 | + EXPECT_EQ( | ||
| 100 | + aclblasSgeam( | ||
| 101 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, A.data(), 4, nullptr, B.data(), 4, | ||
| 102 | + C.data(), 4), | ||
| 103 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 104 | +} | ||
| 105 | + | ||
| 106 | +// -- TEST_F: transa invalid 0xFF (host validates) ---------------------------- | ||
| 107 | +TEST_F(SgeamArch35Test, TransaInvalid) | ||
| 108 | +{ | ||
| 109 | + float alpha = 1.0f, beta = 1.0f; | ||
| 110 | + std::vector<float> A(4 * 4, 1.0f); | ||
| 111 | + std::vector<float> B(4 * 4, 1.0f); | ||
| 112 | + std::vector<float> C(4 * 4, 0.0f); | ||
| 113 | + EXPECT_EQ( | ||
| 114 | + aclblasSgeam( | ||
| 115 | + SgeamArch35Test::handle_, static_cast<aclblasOperation_t>(0xFF), ACLBLAS_OP_N, 4, 4, &alpha, A.data(), 4, | ||
| 116 | + &beta, B.data(), 4, C.data(), 4), | ||
| 117 | + ACLBLAS_STATUS_INVALID_ENUM); | ||
| 118 | +} | ||
| 119 | + | ||
| 120 | +// -- TEST_F: transb invalid 0xFF (host validates) ---------------------------- | ||
| 121 | +TEST_F(SgeamArch35Test, TransbInvalid) | ||
| 122 | +{ | ||
| 123 | + float alpha = 1.0f, beta = 1.0f; | ||
| 124 | + std::vector<float> A(4 * 4, 1.0f); | ||
| 125 | + std::vector<float> B(4 * 4, 1.0f); | ||
| 126 | + std::vector<float> C(4 * 4, 0.0f); | ||
| 127 | + EXPECT_EQ( | ||
| 128 | + aclblasSgeam( | ||
| 129 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, static_cast<aclblasOperation_t>(0xFF), 4, 4, &alpha, A.data(), 4, | ||
| 130 | + &beta, B.data(), 4, C.data(), 4), | ||
| 131 | + ACLBLAS_STATUS_INVALID_ENUM); | ||
| 132 | +} | ||
| 133 | + | ||
| 134 | +// -- TEST_F: m<0 (host validates and returns error) ------------------------- | ||
| 135 | +TEST_F(SgeamArch35Test, MNegative) | ||
| 136 | +{ | ||
| 137 | + float alpha = 1.0f, beta = 1.0f; | ||
| 138 | + EXPECT_EQ( | ||
| 139 | + aclblasSgeam( | ||
| 140 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, -1, 4, &alpha, nullptr, 1, &beta, nullptr, 1, nullptr, | ||
| 141 | + 1), | ||
| 142 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 143 | +} | ||
| 144 | + | ||
| 145 | +// -- TEST_F: n<0 (host validates and returns error) ------------------------- | ||
| 146 | +TEST_F(SgeamArch35Test, NNegative) | ||
| 147 | +{ | ||
| 148 | + float alpha = 1.0f, beta = 1.0f; | ||
| 149 | + EXPECT_EQ( | ||
| 150 | + aclblasSgeam( | ||
| 151 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, -1, &alpha, nullptr, 4, &beta, nullptr, 4, nullptr, | ||
| 152 | + 4), | ||
| 153 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 154 | +} | ||
| 155 | + | ||
| 156 | +// -- TEST_F: lda too small when transa=N (host validates) -------------------- | ||
| 157 | +TEST_F(SgeamArch35Test, LdaTooSmallN) | ||
| 158 | +{ | ||
| 159 | + float alpha = 1.0f, beta = 1.0f; | ||
| 160 | + std::vector<float> A(4 * 4, 1.0f); | ||
| 161 | + std::vector<float> B(4 * 4, 1.0f); | ||
| 162 | + std::vector<float> C(4 * 4, 0.0f); | ||
| 163 | + EXPECT_EQ( | ||
| 164 | + aclblasSgeam( | ||
| 165 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, &alpha, A.data(), 3, &beta, B.data(), 4, | ||
| 166 | + C.data(), 4), | ||
| 167 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 168 | +} | ||
| 169 | + | ||
| 170 | +// -- TEST_F: lda too small when transa=T (lda<n required) ------------------- | ||
| 171 | +TEST_F(SgeamArch35Test, LdaTooSmallT) | ||
| 172 | +{ | ||
| 173 | + float alpha = 1.0f, beta = 1.0f; | ||
| 174 | + std::vector<float> A(4 * 8, 1.0f); | ||
| 175 | + std::vector<float> B(4 * 8, 1.0f); | ||
| 176 | + std::vector<float> C(4 * 8, 0.0f); | ||
| 177 | + EXPECT_EQ( | ||
| 178 | + aclblasSgeam( | ||
| 179 | + SgeamArch35Test::handle_, ACLBLAS_OP_T, ACLBLAS_OP_N, 4, 8, &alpha, A.data(), 7, &beta, B.data(), 8, | ||
| 180 | + C.data(), 4), | ||
| 181 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 182 | +} | ||
| 183 | + | ||
| 184 | +// -- TEST_F: ldc too small (ldc < m) (host validates) ----------------------- | ||
| 185 | +TEST_F(SgeamArch35Test, LdcTooSmall) | ||
| 186 | +{ | ||
| 187 | + float alpha = 1.0f, beta = 1.0f; | ||
| 188 | + std::vector<float> A(8 * 4, 1.0f); | ||
| 189 | + std::vector<float> B(8 * 4, 1.0f); | ||
| 190 | + std::vector<float> C(8 * 4, 0.0f); | ||
| 191 | + EXPECT_EQ( | ||
| 192 | + aclblasSgeam( | ||
| 193 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 8, 4, &alpha, A.data(), 8, &beta, B.data(), 8, | ||
| 194 | + C.data(), 7), | ||
| 195 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 196 | +} | ||
| 197 | + | ||
| 198 | +// -- TEST_F: in-place C==A with transa!=N (host validates) ------------------- | ||
| 199 | +TEST_F(SgeamArch35Test, InplaceATransInvalid) | ||
| 200 | +{ | ||
| 201 | + float alpha = 1.0f, beta = 1.0f; | ||
| 202 | + std::vector<float> A(8 * 8, 1.0f); | ||
| 203 | + std::vector<float> B(8 * 8, 1.0f); | ||
| 204 | + EXPECT_EQ( | ||
| 205 | + aclblasSgeam( | ||
| 206 | + SgeamArch35Test::handle_, ACLBLAS_OP_T, ACLBLAS_OP_N, 8, 8, &alpha, A.data(), 8, &beta, B.data(), 8, | ||
| 207 | + A.data(), 8), | ||
| 208 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 209 | +} | ||
| 210 | + | ||
| 211 | +// -- TEST_F: in-place C==B with transb!=N (host validates) ------------------- | ||
| 212 | +TEST_F(SgeamArch35Test, InplaceBTransInvalid) | ||
| 213 | +{ | ||
| 214 | + float alpha = 1.0f, beta = 1.0f; | ||
| 215 | + std::vector<float> A(8 * 8, 1.0f); | ||
| 216 | + std::vector<float> B(8 * 8, 1.0f); | ||
| 217 | + EXPECT_EQ( | ||
| 218 | + aclblasSgeam( | ||
| 219 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_T, 8, 8, &alpha, A.data(), 8, &beta, B.data(), 8, | ||
| 220 | + B.data(), 8), | ||
| 221 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 222 | +} | ||
| 223 | + | ||
| 224 | +// -- TEST_F: in-place C==A with lda!=ldc (host validates) -------------------- | ||
| 225 | +TEST_F(SgeamArch35Test, InplaceALdaNeLdc) | ||
| 226 | +{ | ||
| 227 | + float alpha = 1.0f, beta = 1.0f; | ||
| 228 | + std::vector<float> A(10 * 8, 1.0f); | ||
| 229 | + std::vector<float> B(10 * 8, 1.0f); | ||
| 230 | + EXPECT_EQ( | ||
| 231 | + aclblasSgeam( | ||
| 232 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 8, 8, &alpha, A.data(), 10, &beta, B.data(), 8, | ||
| 233 | + A.data(), 8), | ||
| 234 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 235 | +} | ||
| 236 | + | ||
| 237 | +// -- TEST_F: in-place C==B with ldb!=ldc (host validates) -------------------- | ||
| 238 | +TEST_F(SgeamArch35Test, InplaceBLdbNeLdc) | ||
| 239 | +{ | ||
| 240 | + float alpha = 1.0f, beta = 1.0f; | ||
| 241 | + std::vector<float> A(10 * 8, 1.0f); | ||
| 242 | + std::vector<float> B(10 * 8, 1.0f); | ||
| 243 | + EXPECT_EQ( | ||
| 244 | + aclblasSgeam( | ||
| 245 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 8, 8, &alpha, A.data(), 8, &beta, B.data(), 10, | ||
| 246 | + B.data(), 8), | ||
| 247 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 248 | +} | ||
| 249 | + | ||
| 250 | +// -- TEST_F: alpha=nullptr with A=nullptr returns INVALID_VALUE ------------- | ||
| 251 | +TEST_F(SgeamArch35Test, NullAlphaAllowsNullA) | ||
| 252 | +{ | ||
| 253 | + float beta = 1.0f; | ||
| 254 | + std::vector<float> B(4 * 4, 1.0f); | ||
| 255 | + std::vector<float> C(4 * 4, kBlasSentinel); | ||
| 256 | + EXPECT_EQ( | ||
| 257 | + aclblasSgeam( | ||
| 258 | + SgeamArch35Test::handle_, ACLBLAS_OP_N, ACLBLAS_OP_N, 4, 4, nullptr, nullptr, 4, &beta, B.data(), 4, | ||
| 259 | + C.data(), 4), | ||
| 260 | + ACLBLAS_STATUS_INVALID_VALUE); | ||
| 261 | +} | ||
| 262 | + | ||
| 263 | +// -- CSV parameterised test suite ------------------------------------------- | ||
| 264 | +INSTANTIATE_TEST_SUITE_P( | ||
| 265 | + Sgeam, SgeamArch35Test, ::testing::ValuesIn(GetCasesFromCsv<SgeamParam>(ReplaceFileExtension2Csv(__FILE__))), | ||
| 266 | + PrintCaseInfoString<SgeamParam>); | ||
| 267 | + | ||
| 268 | +// -- Helper: generate A or B input matrix --------------------------------------- | ||
| 269 | +static std::vector<float> SgeamGenerateInput(const SgeamParam& p, bool isA, float extraSeed) | ||
| 270 | +{ | ||
| 271 | + int nullFlag = isA ? p.nullA : p.nullB; | ||
| 272 | + if (nullFlag != 0 || p.m <= 0 || p.n <= 0) { | ||
| 273 | + return {}; | ||
| 274 | + } | ||
| 275 | + aclblasOperation_t trans = isA ? p.transa : p.transb; | ||
| 276 | + int ld = isA ? p.lda : p.ldb; | ||
| 277 | + const BlasFillMode& fill = isA ? p.aFill : p.bFill; | ||
| 278 | + int rows = (trans == ACLBLAS_OP_N) ? p.m : p.n; | ||
| 279 | + int cols = (trans == ACLBLAS_OP_N) ? p.n : p.m; | ||
| 280 | + return makeBlasMatrix(rows, cols, ld, fill, p.randomSeed + static_cast<uint32_t>(extraSeed)); | ||
| 281 | +} | ||
| 282 | + | ||
| 283 | +// -- Helper: prepare C buffer (in-place aware) -------------------------------- | ||
| 284 | +static float* SgeamPrepareC(const SgeamParam& p, const float* aPtr, const float* bPtr, std::vector<float>& cHost) | ||
| 285 | +{ | ||
| 286 | + if (p.inplace == 1 && aPtr) { | ||
| 287 | + return const_cast<float*>(aPtr); | ||
| 288 | + } | ||
| 289 | + if (p.inplace == 2 && bPtr) { | ||
| 290 | + return const_cast<float*>(bPtr); | ||
| 291 | + } | ||
| 292 | + if (p.nullC == 0 && p.m > 0 && p.n > 0) { | ||
| 293 | + cHost.assign(static_cast<size_t>(p.ldc) * static_cast<size_t>(p.n), kBlasSentinel); | ||
| 294 | + } | ||
| 295 | + return cHost.empty() ? nullptr : cHost.data(); | ||
| 296 | +} | ||
| 297 | + | ||
| 298 | +// -- Helper: compute golden reference ------------------------------------------ | ||
| 299 | +static std::vector<float> SgeamComputeGolden( | ||
| 300 | + const SgeamParam& p, float alpha, float beta, const float* aPtr, const float* bPtr, float* cPtr, | ||
| 301 | + const std::vector<float>& aHost, const std::vector<float>& bHost, const std::vector<float>& cHost) | ||
| 302 | +{ | ||
| 303 | + std::vector<float> goldenC; | ||
| 304 | + if (p.expectResult != ACLBLAS_STATUS_SUCCESS || p.m <= 0 || p.n <= 0 || !cPtr) { | ||
| 305 | + return goldenC; | ||
| 306 | + } | ||
| 307 | + size_t goldenSize = cHost.empty() ? (p.inplace == 1 ? aHost.size() : | ||
| 308 | + p.inplace == 2 ? bHost.size() : 0) : | ||
| 309 | + cHost.size(); | ||
| 310 | + goldenC.assign(goldenSize, kBlasSentinel); | ||
| 311 | + aclblasGeam_cpu<float>( | ||
| 312 | + p.transa, p.transb, static_cast<std::size_t>(p.m), static_cast<std::size_t>(p.n), alpha, aPtr, | ||
| 313 | + static_cast<std::size_t>(p.lda), beta, bPtr, static_cast<std::size_t>(p.ldb), goldenC.data(), | ||
| 314 | + static_cast<std::size_t>(p.ldc)); | ||
| 315 | + return goldenC; | ||
| 316 | +} | ||
| 317 | + | ||
| 318 | +// -- Helper: verify output precision ------------------------------------------ | ||
| 319 | +static void SgeamVerifyOutput( | ||
| 320 | + const SgeamParam& p, const std::vector<float>& goldenC, | ||
| 321 | + const std::vector<float>& aHost, const std::vector<float>& bHost, const std::vector<float>& cHost) | ||
| 322 | +{ | ||
| 323 | + VerifyConfig cfg; | ||
| 324 | + cfg.mode = PrecisionMode::MERE_MARE; | ||
| 325 | + cfg.mereThreshold = p.mereThreshold; | ||
| 326 | + cfg.mareMultiplier = p.mareMultiplier; | ||
| 327 | + const float* outPtr = nullptr; | ||
| 328 | + size_t outSize = 0; | ||
| 329 | + if (p.inplace == 1) { | ||
| 330 | + outPtr = aHost.data(); | ||
| 331 | + outSize = aHost.size(); | ||
| 332 | + } else if (p.inplace == 2) { | ||
| 333 | + outPtr = bHost.data(); | ||
| 334 | + outSize = bHost.size(); | ||
| 335 | + } else { | ||
| 336 | + outPtr = cHost.data(); | ||
| 337 | + outSize = cHost.size(); | ||
| 338 | + } | ||
| 339 | + if (outSize > 0) { | ||
| 340 | + EXPECT_TRUE(Verifier::verifyVector(outPtr, goldenC.data(), outSize, 1, cfg, p.caseName)); | ||
| 341 | + } | ||
| 342 | +} | ||
| 343 | + | ||
| 344 | +// -- TEST_P: 5-step CSV-driven flow ----------------------------------------- | ||
| 345 | +TEST_P(SgeamArch35Test, CsvDriven) | ||
| 346 | +{ | ||
| 347 | + const auto& p = GetParam(); | ||
| 348 | + float alpha = p.alphaFill.val1; | ||
| 349 | + float beta = p.betaFill.val1; | ||
| 350 | + | ||
| 351 | + std::vector<float> aHost = SgeamGenerateInput(p, true, 0); | ||
| 352 | + std::vector<float> bHost = SgeamGenerateInput(p, false, 1); | ||
| 353 | + const float* aPtr = aHost.empty() ? nullptr : aHost.data(); | ||
| 354 | + const float* bPtr = bHost.empty() ? nullptr : bHost.data(); | ||
| 355 | + | ||
| 356 | + std::vector<float> cHost; | ||
| 357 | + float* cPtr = SgeamPrepareC(p, aPtr, bPtr, cHost); | ||
| 358 | + | ||
| 359 | + std::vector<float> goldenC = SgeamComputeGolden(p, alpha, beta, aPtr, bPtr, cPtr, aHost, bHost, cHost); | ||
| 360 | + | ||
| 361 | + aclblasStatus_t ret = aclblasSgeam_npu( | ||
| 362 | + SgeamArch35Test::handle_, p.transa, p.transb, p.m, p.n, &alpha, aPtr, p.lda, &beta, bPtr, p.ldb, cPtr, p.ldc); | ||
| 363 | + EXPECT_EQ(static_cast<int>(ret), ACLBLAS_STATUS_SUCCESS); | ||
| 364 | + | ||
| 365 | + SgeamVerifyOutput(p, goldenC, aHost, bHost, cHost); | ||
| 366 | +} | ||
| @@ -0,0 +1,71 @@ | |||
| 1 | +case_name,description,transa,transb,m,n,lda,ldb,ldc,alpha_fill,beta_fill,a_fill,b_fill,nullA,nullB,nullC,mere_threshold,mare_multiplier,random_seed,inplace | ||
| 2 | +TC_L0_01,NN 2x2 basic,N,N,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 3 | +TC_L0_02,NT 2x2 basic,N,T,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 4 | +TC_L0_03,NC 2x2 basic,N,C,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 5 | +TC_L0_04,TN 2x2 basic,T,N,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 6 | +TC_L0_05,TT 2x2 basic,T,T,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 7 | +TC_L0_06,TC 2x2 basic,T,C,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 8 | +TC_L0_07,CN 2x2 basic,C,N,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 9 | +TC_L0_08,CT 2x2 basic,C,T,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 10 | +TC_L0_09,CC 2x2 basic,C,C,2,2,2,2,2,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 11 | +TC_L0_10,NN 4x4 basic,N,N,4,4,4,4,4,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 12 | +TC_L0_11,NN 1x1 minimal,N,N,1,1,1,1,1,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 13 | +TC_L0_12,m=0 early return,N,N,0,4,1,1,1,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0,0,42,0 | ||
| 14 | +TC_L0_13,n=0 early return,N,N,4,0,4,4,4,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0,0,42,0 | ||
| 15 | +TC_L0_14,alpha=0 beta=1,N,N,4,4,4,4,4,VALUE_NORM_0,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 16 | +TC_L0_15,alpha=1 beta=0 (short-circuit),N,N,4,4,4,4,4,VALUE_NORM_1,VALUE_NORM_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 17 | +TC_L0_16,alpha=-1 beta=-1,N,N,4,4,4,4,4,VALUE_NORM_N1,VALUE_NORM_N1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 18 | +TC_L0_25,T minimal ld (lda>=n),T,T,4,8,8,8,4,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 19 | +TC_L0_28,alpha=0 beta=0 both short-circuit,N,N,4,4,4,4,4,VALUE_NORM_0,VALUE_NORM_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 20 | +TC_L0_29,alpha=0.5 beta=2,N,N,4,4,4,4,4,VALUE_NORM_0.5,VALUE_NORM_2,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 21 | +TC_L0_30,NN m>n tall matrix,N,N,16,4,16,16,16,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 22 | +TC_L0_31,TT m>n tall matrix,T,T,16,4,4,4,16,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 23 | +TC_L0_32,NN m<n wide matrix,N,N,4,16,4,4,4,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 24 | +TC_L1_01,m=1 n=8 single-row NN,N,N,1,8,1,1,1,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 25 | +TC_L1_02,m=1 n=8 single-row NT,N,T,1,8,1,8,1,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 26 | +TC_L1_03,m=1 n=8 single-row NC,N,C,1,8,1,8,1,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 27 | +TC_L1_04,m=1 n=8 single-row TN,T,N,1,8,8,1,1,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 28 | +TC_L1_05,m=1 n=8 single-row TT,T,T,1,8,8,8,1,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 29 | +TC_L1_06,m=1 n=8 single-row CN,C,N,1,8,8,1,1,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 30 | +TC_L1_07,m=8 n=1 single-col NN,N,N,8,1,8,8,8,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 31 | +TC_L1_08,m=8 n=1 single-col TN,T,N,8,1,1,8,8,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 32 | +TC_L1_09,m=8 n=1 single-col CN,C,N,8,1,1,8,8,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 33 | +TC_L1_10,m=8 n=16 non-square NN,N,N,8,16,8,8,8,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 34 | +TC_L1_11,m=8 n=16 non-square TT,T,T,8,16,16,16,8,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 35 | +TC_L1_12,m=16 n=8 non-square NN,N,N,16,8,16,16,16,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 36 | +TC_L1_13,m=16 n=8 non-square TT,T,T,16,8,8,8,16,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 37 | +TC_L1_14,m=16 n=8 non-square CC,C,C,16,8,8,8,16,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 38 | +TC_L1_15,m=64 n=64 shape gradient NN,N,N,64,64,64,64,64,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 39 | +TC_L1_16,m=64 n=64 shape gradient TT,T,T,64,64,64,64,64,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 40 | +TC_L1_17,m=128 n=128 tail-m-tile NN,N,N,128,128,128,128,128,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 41 | +TC_L1_18,m=128 n=128 tail-m-tile CC,C,C,128,128,128,128,128,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 42 | +TC_L1_19,m=256 n=256 shape gradient NN,N,N,256,256,256,256,256,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 43 | +TC_L1_20,m=1024 n=1024 large data NN,N,N,1024,1024,1024,1024,1024,VALUE_NORM_1,VALUE_NORM_1,RANDOM_5_5,RANDOM_5_5,0,0,0,0.0001220703125,10.0,42,0 | ||
| 44 | +TC_L1_21,m=1024 n=1024 large data TT,T,T,1024,1024,1024,1024,1024,VALUE_NORM_1,VALUE_NORM_1,RANDOM_5_5,RANDOM_5_5,0,0,0,0.0001220703125,10.0,42,0 | ||
| 45 | +TC_L1_22,m=1024 n=1024 large data CC,C,C,1024,1024,1024,1024,1024,VALUE_NORM_1,VALUE_NORM_1,RANDOM_5_5,RANDOM_5_5,0,0,0,0.0001220703125,10.0,42,0 | ||
| 46 | +TC_L1_23,ld padding +1 all matrices NN,N,N,8,8,9,9,9,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 47 | +TC_L1_24,ld padding +32 all matrices NN,N,N,8,8,40,40,40,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 48 | +TC_L1_25,ld padding T mode lda=n*2,T,T,4,8,16,8,4,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 49 | +TC_L1_26,large ld padding 32 all NN,N,N,4,4,32,32,32,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 50 | +TC_L1_27,alpha=+INF beta=1,N,N,4,4,4,4,4,VALUE_NORM_INF,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 51 | +TC_L1_28,alpha=-INF beta=1,N,N,4,4,4,4,4,VALUE_NORM_NINF,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 52 | +TC_L1_29,alpha=NAN beta=1 NAN propagation,N,N,4,4,4,4,4,VALUE_NORM_NAN,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 53 | +TC_L1_30,beta=+INF alpha=1 INF propagation,N,N,4,4,4,4,4,VALUE_NORM_1,VALUE_NORM_INF,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 54 | +TC_L1_31,beta=NAN alpha=1 NAN propagation,N,N,4,4,4,4,4,VALUE_NORM_1,VALUE_NORM_NAN,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 55 | +TC_L1_32,m=127 n=63 non-power2 TT,T,T,127,63,63,63,127,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 56 | +TC_L1_33,m=33 n=255 non-power2 NN,N,N,33,255,33,33,33,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 57 | +TC_L1_34,alpha=0 beta=0.5 beta-only output TT,T,T,4,4,4,4,4,VALUE_NORM_0,VALUE_NORM_0.5,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 58 | +TC_L1_35,alpha=2 beta=-1 differential NN,N,N,8,8,8,8,8,VALUE_NORM_2,VALUE_NORM_N1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 59 | +TC_L1_36,A all zero B random NN,N,N,8,8,8,8,8,VALUE_NORM_1,VALUE_NORM_1,VALUE_NORM_0,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 60 | +TC_L1_37,B all zero A random NN,N,N,8,8,8,8,8,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,VALUE_NORM_0,0,0,0,0.0001220703125,10.0,42,0 | ||
| 61 | +TC_L1_38,INDEX fill equidistant NN,N,N,8,8,8,8,8,VALUE_NORM_1,VALUE_NORM_1,INDEX,INDEX,0,0,0,0.0001220703125,10.0,42,0 | ||
| 62 | +TC_L1_39,INDEX_ALTER fill NN,N,N,8,8,8,8,8,VALUE_NORM_1,VALUE_NORM_1,INDEX_ALTER,INDEX_ALTER,0,0,0,0.0001220703125,10.0,42,0 | ||
| 63 | +TC_L1_40,m=5000 n=1 TT batch loop curM>4095,T,T,5000,1,1,1,5000,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,0 | ||
| 64 | +TC_L1_41,in-place C==A transa=N lda=ldc,N,N,8,8,8,8,8,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,1 | ||
| 65 | +TC_L1_42,in-place C==B transb=N ldb=ldc,N,N,8,8,8,8,8,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,2 | ||
| 66 | +TC_L1_47,in-place C==A alpha=0 beta=0,N,N,8,8,8,8,8,VALUE_NORM_0,VALUE_NORM_0,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,1 | ||
| 67 | +TC_L1_48,in-place C==A large N=128,N,N,128,128,128,128,128,VALUE_NORM_1,VALUE_NORM_1,RANDOM_1_1,RANDOM_1_1,0,0,0,0.0001220703125,10.0,42,1 | ||
| 68 | +TC_L1_49,alpha=0 A=nullptr NN,N,N,8,8,8,8,8,VALUE_NORM_0,VALUE_NORM_1,NULLPTR,RANDOM_1_1,1,0,0,0.0001220703125,10.0,42,0 | ||
| 69 | +TC_L1_50,alpha=0 A=nullptr TT,T,T,16,4,4,4,16,VALUE_NORM_0,VALUE_NORM_1,NULLPTR,RANDOM_1_1,1,0,0,0.0001220703125,10.0,42,0 | ||
| 70 | +TC_L1_51,alpha=0 beta=0 both nullptr NN,N,N,8,8,8,8,8,VALUE_NORM_0,VALUE_NORM_0,NULLPTR,NULLPTR,1,1,0,0.0001220703125,10.0,42,0 | ||
| 71 | +TC_L1_53,beta=0 B=nullptr alpha=1 NN,N,N,8,8,8,8,8,VALUE_NORM_1,VALUE_NORM_0,RANDOM_1_1,NULLPTR,0,1,0,0.0001220703125,10.0,42,0 | ||
| @@ -0,0 +1,53 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | +struct SgeamParam : public BlasTestParamBase { | ||
| 19 | + aclblasOperation_t transa = ACLBLAS_OP_N; | ||
| 20 | + aclblasOperation_t transb = ACLBLAS_OP_N; | ||
| 21 | + int m = 0; | ||
| 22 | + int n = 0; | ||
| 23 | + int lda = 0; | ||
| 24 | + int ldb = 0; | ||
| 25 | + int ldc = 0; | ||
| 26 | + BlasFillMode alphaFill = BlasFillMode("VALUE_NORM_1"); | ||
| 27 | + BlasFillMode betaFill = BlasFillMode("VALUE_NORM_1"); | ||
| 28 | + BlasFillMode aFill = BlasFillMode("RANDOM_1_1"); | ||
| 29 | + BlasFillMode bFill = BlasFillMode("RANDOM_1_1"); | ||
| 30 | + int nullA = 0; | ||
| 31 | + int nullB = 0; | ||
| 32 | + int nullC = 0; | ||
| 33 | + int inplace = 0; // 0=out-of-place, 1=C==A, 2=C==B | ||
| 34 | + | ||
| 35 | + SgeamParam(const csv_map& csv) : BlasTestParamBase(csv) | ||
| 36 | + { | ||
| 37 | + transa = parseOpTrans(ReadMap(csv, "transa", "N")); | ||
| 38 | + transb = parseOpTrans(ReadMap(csv, "transb", "N")); | ||
| 39 | + m = parseInt(ReadMap(csv, "m", "0")); | ||
| 40 | + n = parseInt(ReadMap(csv, "n", "0")); | ||
| 41 | + lda = parseInt(ReadMap(csv, "lda", std::to_string((transa == ACLBLAS_OP_N) ? std::max(1, m) : std::max(1, n)))); | ||
| 42 | + ldb = parseInt(ReadMap(csv, "ldb", std::to_string((transb == ACLBLAS_OP_N) ? std::max(1, m) : std::max(1, n)))); | ||
| 43 | + ldc = parseInt(ReadMap(csv, "ldc", std::to_string(std::max(1, m)))); | ||
| 44 | + alphaFill = BlasFillMode(ReadMap(csv, "alpha_fill", "VALUE_NORM_1")); | ||
| 45 | + betaFill = BlasFillMode(ReadMap(csv, "beta_fill", "VALUE_NORM_1")); | ||
| 46 | + aFill = BlasFillMode(ReadMap(csv, "a_fill", "RANDOM_1_1")); | ||
| 47 | + bFill = BlasFillMode(ReadMap(csv, "b_fill", "RANDOM_1_1")); | ||
| 48 | + nullA = parseInt(ReadMap(csv, "nullA", "0")); | ||
| 49 | + nullB = parseInt(ReadMap(csv, "nullB", "0")); | ||
| 50 | + nullC = parseInt(ReadMap(csv, "nullC", "0")); | ||
| 51 | + inplace = parseInt(ReadMap(csv, "inplace", "0")); | ||
| 52 | + } | ||
| 53 | +}; | ||
主机按 beta!=0 时 6 个缓冲区（CGEAM_BUFFERS）计算 tileM 容量上限，而 kernel 实际无条件分配 8 个缓冲区（4 个队列 + calcBufBR_/BI_ + 2 个乘法 scratch），容量预算与 kernel 实现不符，当前仅靠 MAX_BLOCK_COUNT=4095 上限恰好不超 UB，上限一旦调整 tileM 将静默超出 UB 容量