已合并
将matmul_compress_dequant算子从math仓迁到nn仓 #2728
wmg1创建于 5月14日
将matmul_compress_dequant算子从math仓迁到nn仓 #2728
已合并
共 37 个文件变更+0-4456
| @@ -710,8 +710,6 @@ conversion@ops-math: | |||
| 710 | - ops/ops-math/conversion/diag_flat/op_host/diag_flat_def.cpp | 710 | - ops/ops-math/conversion/diag_flat/op_host/diag_flat_def.cpp |
| 711 | - ops/ops-math/conversion/diag_flat/op_host/diag_flat_tiling.cpp | 711 | - ops/ops-math/conversion/diag_flat/op_host/diag_flat_tiling.cpp |
| 712 | - ops/ops-math/conversion/diag_flat/op_graph/ | 712 | - ops/ops-math/conversion/diag_flat/op_graph/ |
| 713 | - - ops/ops-math/conversion/matmul_v2_compress_dequant/ | ||
| 714 | - - ops/ops-math/conversion/matmul_v2_compress_dequant/op_host/op_api/aclnn_matmul_compress_dequant.cpp | ||
| 715 | - ops/ops-math/math/tile/op_api/ | 713 | - ops/ops-math/math/tile/op_api/ |
| 716 | - ops/ops-math/math/tile/docs/ | 714 | - ops/ops-math/math/tile/docs/ |
| 717 | - ops/ops-math/math/tile/op_graph/ | 715 | - ops/ops-math/math/tile/op_graph/ |
| @@ -1,20 +0,0 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | ||
| 2 | -# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | -# This file is a part of the CANN Open Software. | ||
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | -# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | -# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | -# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | -# ---------------------------------------------------------------------------- | ||
| 11 | - | ||
| 12 | -file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 13 | -if(NOT ENABLE_TEST AND NOT BENCHMARK) | ||
| 14 | - list(REMOVE_ITEM CURRENT_DIRS tests) | ||
| 15 | -endif() | ||
| 16 | -foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 17 | - if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 18 | - add_subdirectory(${SUB_DIR}) | ||
| 19 | - endif() | ||
| 20 | -endforeach() | ||
| @@ -1,114 +0,0 @@ | |||
| 1 | -# MatMulV2CompressDequant | ||
| 2 | - | ||
| 3 | -## 产品支持情况 | ||
| 4 | - | ||
| 5 | -| 产品 | 是否支持 | | ||
| 6 | -| :----------------------------------------------------------- | :------: | | ||
| 7 | -| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | × | | ||
| 8 | -| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | × | | ||
| 9 | - | ||
| 10 | -## 功能说明 | ||
| 11 | - | ||
| 12 | -- **算子功能**：进行矩阵乘计算时，可先通过msModelSlim工具对右矩阵进行无损压缩，减少内存占用，然后通过本接口完成无损解压缩、矩阵乘和反量化计算。 | ||
| 13 | -- **计算公式**： | ||
| 14 | - | ||
| 15 | - ```text | ||
| 16 | - x2_unzip = unzip(x2, compressIndex, compressInfo) | ||
| 17 | - result = (x1 @ x2_unzip + bias) * deqScale | ||
| 18 | - ``` | ||
| 19 | - | ||
| 20 | - 其中x2表示右矩阵经过msModelSlim工具压缩后的一维数据，x2_unzip是接口内部进行无损解压缩后的数据（与原始右矩阵数据一致）。 | ||
| 21 | - | ||
| 22 | -## 参数说明约束说明 | ||
| 23 | - | ||
| 24 | -<table style="undefined;table-layout: fixed; width: 869px"><colgroup> | ||
| 25 | -<col style="width: 144px"> | ||
| 26 | -<col style="width: 166px"> | ||
| 27 | -<col style="width: 343px"> | ||
| 28 | -<col style="width: 114px"> | ||
| 29 | -<col style="width: 102px"> | ||
| 30 | -</colgroup> | ||
| 31 | -<thead> | ||
| 32 | - <tr> | ||
| 33 | - <th>参数名</th> | ||
| 34 | - <th>输入/输出/属性</th> | ||
| 35 | - <th>描述</th> | ||
| 36 | - <th>数据类型</th> | ||
| 37 | - <th>数据格式</th> | ||
| 38 | - </tr></thead> | ||
| 39 | -<tbody> | ||
| 40 | - <tr> | ||
| 41 | - <td>x1</td> | ||
| 42 | - <td>输入张量</td> | ||
| 43 | - <td>矩阵乘的左输入，2维张量。</td> | ||
| 44 | - <td>INT8</td> | ||
| 45 | - <td>ND</td> | ||
| 46 | - </tr> | ||
| 47 | - <tr> | ||
| 48 | - <td>x2</td> | ||
| 49 | - <td>输入张量</td> | ||
| 50 | - <td>压缩后的矩阵乘右输入，1维张量。</td> | ||
| 51 | - <td>INT8</td> | ||
| 52 | - <td>ND</td> | ||
| 53 | - </tr> | ||
| 54 | - <tr> | ||
| 55 | - <td>compressIndex</td> | ||
| 56 | - <td>输入张量</td> | ||
| 57 | - <td>矩阵乘右输入的压缩索引表，1维张量。</td> | ||
| 58 | - <td>INT8</td> | ||
| 59 | - <td>ND</td> | ||
| 60 | - </tr> | ||
| 61 | - <tr> | ||
| 62 | - <td>bias</td> | ||
| 63 | - <td>输入张量</td> | ||
| 64 | - <td>偏置项，支持空指针传入。</td> | ||
| 65 | - <td>INT32</td> | ||
| 66 | - <td>ND</td> | ||
| 67 | - </tr> | ||
| 68 | - <tr> | ||
| 69 | - <td>deqScale</td> | ||
| 70 | - <td>输入张量</td> | ||
| 71 | - <td>反量化参数，数据类型为UINT64。</td> | ||
| 72 | - <td>UINT64</td> | ||
| 73 | - <td>ND</td> | ||
| 74 | - </tr> | ||
| 75 | - <tr> | ||
| 76 | - <td>offsetW</td> | ||
| 77 | - <td>输入张量</td> | ||
| 78 | - <td>矩阵乘右输入的偏移量，当前仅支持空指针传入。</td> | ||
| 79 | - <td>INT8</td> | ||
| 80 | - <td>ND</td> | ||
| 81 | - </tr> | ||
| 82 | - <tr> | ||
| 83 | - <td>offsetX</td> | ||
| 84 | - <td>输入属性</td> | ||
| 85 | - <td>矩阵乘左输入的偏移量，当前仅支持0。</td> | ||
| 86 | - <td>INT32</td> | ||
| 87 | - <td>-</td> | ||
| 88 | - </tr> | ||
| 89 | - <tr> | ||
| 90 | - <td>compressInfo</td> | ||
| 91 | - <td>输入数组</td> | ||
| 92 | - <td>压缩数据相关信息，包括压缩块信息和原始shape等。</td> | ||
| 93 | - <td>INT64</td> | ||
| 94 | - <td>-</td> | ||
| 95 | - </tr> | ||
| 96 | - <tr> | ||
| 97 | - <td>out</td> | ||
| 98 | - <td>输出张量</td> | ||
| 99 | - <td>计算结果输出。</td> | ||
| 100 | - <td>FLOAT16</td> | ||
| 101 | - <td>ND</td> | ||
| 102 | - </tr> | ||
| 103 | -</tbody></table> | ||
| 104 | - | ||
| 105 | -- x1和x2_unzip的Reduce维度大小必须相等。 | ||
| 106 | -- 所有输入张量不支持非连续的Tensor。 | ||
| 107 | -- deqScale需要将原始float类型参数转换为UINT64数据格式。 | ||
| 108 | -- 当前offsetW仅支持空指针，offsetX仅支持0。 | ||
| 109 | - | ||
| 110 | -## 调用说明 | ||
| 111 | - | ||
| 112 | -| 调用方式 | 样例代码 | 说明 | | ||
| 113 | -| --------- | ------------------------------------------------------------ | ------------------------------------------------------------ | | ||
| 114 | -| aclnn接口 | [test_aclnn_matmul_compress_dequant](./tests/ut/op_host/op_api/test_aclnn_matmul_compress_dequant.cpp) | 通过[aclnnMatmulCompressDequant](docs/aclnnMatmulCompressDequant.md)接口方式调用MatmulCompressDequant算子。 | | ||
| @@ -1,610 +0,0 @@ | |||
| 1 | -# aclnnMatmulCompressDequant | ||
| 2 | - | ||
| 3 | -[📄 查看源码](https://gitcode.com/cann/ops-math/tree/master/conversion/matmul_v2_compress_dequant) | ||
| 4 | - | ||
| 5 | -## 产品支持情况 | ||
| 6 | - | ||
| 7 | -| 产品 | 是否支持 | | ||
| 8 | -| :----------------------------------------------------------- | :------: | | ||
| 9 | -| <term>Ascend 950PR/Ascend 950DT</term> | × | | ||
| 10 | -| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | × | | ||
| 11 | -| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | × | | ||
| 12 | -| <term>Atlas 200I/500 A2 推理产品</term> | × | | ||
| 13 | -| <term>Atlas 推理系列产品 </term> | √ | | ||
| 14 | -| <term>Atlas 训练系列产品</term> | × | | ||
| 15 | -| <term>Atlas 200/300/500 推理产品</term> | × | | ||
| 16 | - | ||
| 17 | -## 功能说明 | ||
| 18 | - | ||
| 19 | -- 接口功能：进行l@r矩阵乘计算时，可先通过msModelSlim工具对r矩阵进行无损压缩，减少r矩阵的内存占用大小，然后通过本接口完成无损解压缩、矩阵乘、反量化计算。 | ||
| 20 | -- 计算公式： | ||
| 21 | - | ||
| 22 | - $$ | ||
| 23 | - x2\_unzip = unzip(x2, compressIndex, compressInfo)\\ | ||
| 24 | - result=(x1 @ x2\_unzip + bias)*deqScale | ||
| 25 | - $$ | ||
| 26 | - | ||
| 27 | - 其中x2表示r矩阵经过msModelSlim工具进行压缩后的一维数据，compressIndex以及compressInfo表示压缩算法相关的信息，$x2\_unzip$是本接口内部进行无损解压缩后的数据（与原始r矩阵数据一致），压缩和调用本接口的详细使用样例参考[调用示例](#调用示例)。 | ||
| 28 | - | ||
| 29 | -## 函数原型 | ||
| 30 | - | ||
| 31 | -每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMatmulCompressDequantGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMatmulCompressDequant”接口执行计算。 | ||
| 32 | - | ||
| 33 | -```cpp | ||
| 34 | -aclnnStatus aclnnMatmulCompressDequantGetWorkspaceSize( | ||
| 35 | - const aclTensor* x1, | ||
| 36 | - const aclTensor* x2, | ||
| 37 | - const aclTensor* compressIndex, | ||
| 38 | - const aclTensor* bias, | ||
| 39 | - const aclTensor* deqScale, | ||
| 40 | - const aclTensor* offsetW, | ||
| 41 | - int offsetX, | ||
| 42 | - const aclIntArray* compressInfo, | ||
| 43 | - aclTensor* out, | ||
| 44 | - uint64_t* workspaceSize, | ||
| 45 | - aclOpExecutor** executor) | ||
| 46 | -``` | ||
| 47 | - | ||
| 48 | -```cpp | ||
| 49 | -aclnnStatus aclnnMatmulCompressDequant( | ||
| 50 | - void* workspace, | ||
| 51 | - uint64_t workspaceSize, | ||
| 52 | - aclOpExecutor* executor, | ||
| 53 | - aclrtStream stream) | ||
| 54 | -``` | ||
| 55 | - | ||
| 56 | -## aclnnMatmulCompressDequantGetWorkspaceSize | ||
| 57 | - | ||
| 58 | -- **参数说明** | ||
| 59 | - <table style="undefined;table-layout: fixed; width: 1475px"><colgroup> | ||
| 60 | - <col style="width: 167px"> | ||
| 61 | - <col style="width: 123px"> | ||
| 62 | - <col style="width: 325px"> | ||
| 63 | - <col style="width: 230px"> | ||
| 64 | - <col style="width: 128px"> | ||
| 65 | - <col style="width: 118px"> | ||
| 66 | - <col style="width: 239px"> | ||
| 67 | - <col style="width: 145px"> | ||
| 68 | - </colgroup> | ||
| 69 | - <thead> | ||
| 70 | - <tr> | ||
| 71 | - <th>参数名</th> | ||
| 72 | - <th>输入/输出</th> | ||
| 73 | - <th>描述</th> | ||
| 74 | - <th>使用说明</th> | ||
| 75 | - <th>数据类型</th> | ||
| 76 | - <th>数据格式</th> | ||
| 77 | - <th>维度(shape)</th> | ||
| 78 | - <th>非连续tensor</th> | ||
| 79 | - </tr></thead> | ||
| 80 | - <tbody> | ||
| 81 | - <tr> | ||
| 82 | - <td>x1</td> | ||
| 83 | - <td>输入</td> | ||
| 84 | - <td>表示矩阵乘的左输入。</td> | ||
| 85 | - <td>-</td> | ||
| 86 | - <td>INT8</td> | ||
| 87 | - <td>ND</td> | ||
| 88 | - <td>2</td> | ||
| 89 | - <td>-</td> | ||
| 90 | - </tr> | ||
| 91 | - <tr> | ||
| 92 | - <td>x2</td> | ||
| 93 | - <td>输入</td> | ||
| 94 | - <td>表示压缩后的矩阵乘的右输入，为通过msModelSlim工具中weight_compression模块压缩后的输入。</td> | ||
| 95 | - <td>-</td> | ||
| 96 | - <td>INT8</td> | ||
| 97 | - <td>ND</td> | ||
| 98 | - <td>1</td> | ||
| 99 | - <td>-</td> | ||
| 100 | - </tr> | ||
| 101 | - <tr> | ||
| 102 | - <td>compressIndex</td> | ||
| 103 | - <td>输入</td> | ||
| 104 | - <td>表示矩阵乘右输入的压缩索引表。</td> | ||
| 105 | - <td>通过示例中的msModelSlim工具中获取</td> | ||
| 106 | - <td>INT8</td> | ||
| 107 | - <td>ND</td> | ||
| 108 | - <td>1</td> | ||
| 109 | - <td>-</td> | ||
| 110 | - </tr> | ||
| 111 | - <tr> | ||
| 112 | - <td>bias</td> | ||
| 113 | - <td>输入</td> | ||
| 114 | - <td>参与矩阵乘计算的偏置项。</td> | ||
| 115 | - <td>支持空指针传入。</td> | ||
| 116 | - <td>INT8</td> | ||
| 117 | - <td>ND</td> | ||
| 118 | - <td>2维，shape仅支持(1, n)或者(n)，其中n为输出shape(m, n)的n</td> | ||
| 119 | - <td>-</td> | ||
| 120 | - </tr> | ||
| 121 | - <tr> | ||
| 122 | - <td>deqScale</td> | ||
| 123 | - <td>输入</td> | ||
| 124 | - <td>表示反量化参数。</td> | ||
| 125 | - <td>tensor中的值为float通过下述示例中转换后的UINT64的数据。</td> | ||
| 126 | - <td>UINT64</td> | ||
| 127 | - <td>ND</td> | ||
| 128 | - <td>2维，shape支持(1, n)或者(1, 1), 其中n为输出shape(m, n)中的n。</td> | ||
| 129 | - <td>-</td> | ||
| 130 | - </tr> | ||
| 131 | - <tr> | ||
| 132 | - <td>offsetW</td> | ||
| 133 | - <td>输入</td> | ||
| 134 | - <td>标量，表示矩阵乘右输入的偏移量。</td> | ||
| 135 | - <td>当前仅支持空指针传入。</td> | ||
| 136 | - <td>INT8</td> | ||
| 137 | - <td>-</td> | ||
| 138 | - <td>与x2_unzip一致。</td> | ||
| 139 | - <td>-</td> | ||
| 140 | - </tr> | ||
| 141 | - <tr> | ||
| 142 | - <td>offsetX</td> | ||
| 143 | - <td>输入</td> | ||
| 144 | - <td>标量，表示矩阵乘左输入的偏移量。</td> | ||
| 145 | - <td>当前仅支持0。</td> | ||
| 146 | - <td>INT32</td> | ||
| 147 | - <td>-</td> | ||
| 148 | - <td>-</td> | ||
| 149 | - <td>-</td> | ||
| 150 | - </tr> | ||
| 151 | - <tr> | ||
| 152 | - <td>compressInfo</td> | ||
| 153 | - <td>输入</td> | ||
| 154 | - <td>整型数据列表，数据类型为INT64。其中包括压缩块信息tilingN、tilingK（通过msModelSlim工具中weight_compression模块压缩后获取，分别表示压缩前shape(n, k)在n方向和k方向上一个基本压缩块的大小），压缩前x2矩阵原始shape（shape为2维，用(n, k)表示），以及压缩块遍历方向的标识。</td> | ||
| 155 | - <td>-</td> | ||
| 156 | - <td>INT64</td> | ||
| 157 | - <td>-</td> | ||
| 158 | - <td>-</td> | ||
| 159 | - <td>-</td> | ||
| 160 | - </tr> | ||
| 161 | - <tr> | ||
| 162 | - <td>out</td> | ||
| 163 | - <td>输出</td> | ||
| 164 | - <td>计算输出。</td> | ||
| 165 | - <td>-</td> | ||
| 166 | - <td>FLOAT16</td> | ||
| 167 | - <td>ND</td> | ||
| 168 | - <td>2</td> | ||
| 169 | - <td>-</td> | ||
| 170 | - </tr> | ||
| 171 | - <tr> | ||
| 172 | - <td>workspaceSize</td> | ||
| 173 | - <td>出参</td> | ||
| 174 | - <td>返回需要在Device侧申请的workspace大小。</td> | ||
| 175 | - <td>-</td> | ||
| 176 | - <td>-</td> | ||
| 177 | - <td>-</td> | ||
| 178 | - <td>-</td> | ||
| 179 | - <td>-</td> | ||
| 180 | - </tr> | ||
| 181 | - <tr> | ||
| 182 | - <td>executor</td> | ||
| 183 | - <td>出参</td> | ||
| 184 | - <td>返回op执行器，包含了算子计算流程。</td> | ||
| 185 | - <td>-</td> | ||
| 186 | - <td>-</td> | ||
| 187 | - <td>-</td> | ||
| 188 | - <td>-</td> | ||
| 189 | - <td>-</td> | ||
| 190 | - </tr> | ||
| 191 | - </tbody> | ||
| 192 | - </table> | ||
| 193 | - | ||
| 194 | -- **返回值** | ||
| 195 | - | ||
| 196 | - aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。 | ||
| 197 | - | ||
| 198 | - 第一段接口完成入参校验，出现以下场景时报错： | ||
| 199 | - | ||
| 200 | - <table style="undefined;table-layout: fixed; width: 1475px"><colgroup> | ||
| 201 | - <col style="width: 300px"> | ||
| 202 | - <col style="width: 200px"> | ||
| 203 | - <col style="width: 975px"> | ||
| 204 | - </colgroup> | ||
| 205 | - <thead> | ||
| 206 | - <tr> | ||
| 207 | - <th>返回值</th> | ||
| 208 | - <th>错误码</th> | ||
| 209 | - <th>描述</th> | ||
| 210 | - </tr></thead> | ||
| 211 | - <tbody> | ||
| 212 | - <tr> | ||
| 213 | - <td>ACLNN_ERR_PARAM_NULLPTR</td> | ||
| 214 | - <td>161001</td> | ||
| 215 | - <td>传入的x1、x2或out是空指针。</td> | ||
| 216 | - </tr> | ||
| 217 | - <tr> | ||
| 218 | - <td rowspan="3">ACLNN_ERR_PARAM_INVALID</td> | ||
| 219 | - <td rowspan="3">161002</td> | ||
| 220 | - <td>x1或x2的数据类型和数据格式不在支持的范围之内。</td> | ||
| 221 | - </tr> | ||
| 222 | - <tr> | ||
| 223 | - <td>x1或x2无法做数据类型推导。</td> | ||
| 224 | - </tr> | ||
| 225 | - <tr> | ||
| 226 | - <td>推导出的数据类型无法转换为指定输出out的类型。</td> | ||
| 227 | - </tr> | ||
| 228 | - </tbody> | ||
| 229 | - </table> | ||
| 230 | - | ||
| 231 | -## aclnnMatmulCompressDequant | ||
| 232 | - | ||
| 233 | -- **参数说明** | ||
| 234 | - <table style="undefined;table-layout: fixed; width: 1475px"><colgroup> | ||
| 235 | - <col style="width: 300px"> | ||
| 236 | - <col style="width: 200px"> | ||
| 237 | - <col style="width: 975px"> | ||
| 238 | - </colgroup> | ||
| 239 | - <thead> | ||
| 240 | - <tr> | ||
| 241 | - <th>参数名</th> | ||
| 242 | - <th>输入/输出</th> | ||
| 243 | - <th>描述</th> | ||
| 244 | - </tr></thead> | ||
| 245 | - <tbody> | ||
| 246 | - <tr> | ||
| 247 | - <td>workspace</td> | ||
| 248 | - <td>输入</td> | ||
| 249 | - <td>在Device侧申请的workspace内存地址。</td> | ||
| 250 | - </tr> | ||
| 251 | - <tr> | ||
| 252 | - <td>workspaceSize</td> | ||
| 253 | - <td>输入</td> | ||
| 254 | - <td>在Device侧申请的workspace大小，由第一段接口aclnnMatmulCompressDequantGetWorkspaceSize获取。</td> | ||
| 255 | - </tr> | ||
| 256 | - <tr> | ||
| 257 | - <td>executor</td> | ||
| 258 | - <td>输入</td> | ||
| 259 | - <td>op执行器，包含了算子计算流程。</td> | ||
| 260 | - </tr> | ||
| 261 | - <tr> | ||
| 262 | - <td>stream</td> | ||
| 263 | - <td>输入</td> | ||
| 264 | - <td>指定执行任务的Stream。</td> | ||
| 265 | - </tr> | ||
| 266 | - </tbody> | ||
| 267 | - </table> | ||
| 268 | - | ||
| 269 | -- **返回值** | ||
| 270 | - | ||
| 271 | - aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。 | ||
| 272 | - | ||
| 273 | -## 约束说明 | ||
| 274 | - | ||
| 275 | -- 确定性计算： | ||
| 276 | - - aclnnMatmulCompressDequant默认确定性实现。 | ||
| 277 | - | ||
| 278 | -## 调用示例 | ||
| 279 | - | ||
| 280 | -- **准备压缩前的数据** | ||
| 281 | - | ||
| 282 | - 假设通过脚本gen_data.py生成输入数据，示例如下，仅供参考： | ||
| 283 | - | ||
| 284 | - ```python | ||
| 285 | - import numpy as np | ||
| 286 | - import os | ||
| 287 | - import sys | ||
| 288 | - from numpy import random | ||
| 289 | - | ||
| 290 | - def write2file(data, path): | ||
| 291 | - with open(path, 'wb') as f: | ||
| 292 | - data.tofile(f) | ||
| 293 | - | ||
| 294 | - if not os.path.exists("./data"): | ||
| 295 | - os.mkdir("./data") | ||
| 296 | - | ||
| 297 | - if len(sys.argv) != 4: | ||
| 298 | - print("Usage: python gen_data.py m k n") | ||
| 299 | - sys.exit(1) | ||
| 300 | - | ||
| 301 | - m = int(sys.argv[1]) | ||
| 302 | - k = int(sys.argv[2]) | ||
| 303 | - n = int(sys.argv[3]) | ||
| 304 | - | ||
| 305 | - if m <= 0 or k <= 0 or n <= 0: | ||
| 306 | - print("Error: m, k and n must be positive integers.") | ||
| 307 | - sys.exit(1) | ||
| 308 | - | ||
| 309 | - # 随机生成矩阵mat1，shape为(m,k ) | ||
| 310 | - mat1 = random.randn(m, k).astype(np.int8) | ||
| 311 | - write2file(mat1, "./data/mat1.bin") | ||
| 312 | - | ||
| 313 | - # 随机生成矩阵mat2，shape为(n, k) | ||
| 314 | - mat2 = random.randint(0, 100, size=(n, k)).astype(np.int8) | ||
| 315 | - np.save("./data/weight.npy", {'weight': mat2}) | ||
| 316 | - os.chmod("./data/weight.npy", 0o0640) | ||
| 317 | - | ||
| 318 | - # 生成output | ||
| 319 | - output = np.random.randn(m, n).astype(np.float16) | ||
| 320 | - write2file(output, "./data/output.bin") | ||
| 321 | - | ||
| 322 | - # 生成bias | ||
| 323 | - bias = random.randn(n).astype(np.float32) | ||
| 324 | - write2file(bias, "./data/bias.bin") | ||
| 325 | - | ||
| 326 | - # 生成deq_scale | ||
| 327 | - deq_scale = random.randn(n).astype(np.float32) | ||
| 328 | - write2file(deq_scale, "./data/deqScale_ori.bin") | ||
| 329 | - deq_scale_int64 = np.fromfile("./data/deqScale_ori.bin", dtype=np.int32).astype(np.int64) | ||
| 330 | - deq_scale_int64.tofile("./data/deqScale.bin") | ||
| 331 | - ``` | ||
| 332 | - | ||
| 333 | - 执行gen_data.py，假设mat1和mat2的shape入参为m=512、k=1024、n=1024。 | ||
| 334 | - | ||
| 335 | - ```shell | ||
| 336 | - python3 gen_data.py 512 1024 1024 | ||
| 337 | - ``` | ||
| 338 | - | ||
| 339 | -- **对数据进行预处理** | ||
| 340 | - | ||
| 341 | - - **原始权重通过msModelSlim压缩工具生成压缩后的x2、compressIndex以及compressInfo** | ||
| 342 | - | ||
| 343 | - 使用以下接口时，需对CANN包中msModelSlim压缩工具进行编译，具体操作参考[Gitee msit仓](https://gitee.com/ascend/msit/tree/master/msmodelslim)中msmodelslim/pytorch/weight_compression目录下的README.md。 | ||
| 344 | - | ||
| 345 | - ```python | ||
| 346 | - from msmodelslim.pytorch.weight_compression import CompressConfig, Compressor | ||
| 347 | - | ||
| 348 | - compress_config = CompressConfig(do_pseudo_sparse=False, sparse_ratio=1) | ||
| 349 | - compressor = Compressor(compress_config, weight_path=weight_path) | ||
| 350 | - | ||
| 351 | - compress_weight, compress_index, compress_info = compressor.run() | ||
| 352 | - # 压缩后的权重，对应aclnnMatmulCompressDequantGetWorkspaceSize接口的x2 | ||
| 353 | - compressor.export(compress_weight, './data/weight') | ||
| 354 | - # 压缩权重的索引，对应aclnnMatmulCompressDequantGetWorkspaceSize接口的compressIndex | ||
| 355 | - compressor.export(compress_index, './data/index') | ||
| 356 | - # 压缩数据的相关信息，对应aclnnMatmulCompressDequantGetWorkspaceSize接口的compressInfo | ||
| 357 | - compressor.export(compress_info, './data/compress_info') | ||
| 358 | - ``` | ||
| 359 | - | ||
| 360 | - - **将原始float类型的反量化参数deqscale进行转换， 得到aclnn接口需要的uint64数据** | ||
| 361 | - | ||
| 362 | - deqScale原始为float类型，以int32读取并转换为int64 | ||
| 363 | - | ||
| 364 | - ```python | ||
| 365 | - import numpy as np | ||
| 366 | - data = np.fromfile('./deqScale_original.bin', dtype=np.int32).astype(np.int64) | ||
| 367 | - data.tofile('./deqScale.bin') | ||
| 368 | - ``` | ||
| 369 | - | ||
| 370 | -- **调用aclnn接口运算** | ||
| 371 | - | ||
| 372 | - 示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。 | ||
| 373 | - | ||
| 374 | -```Cpp | ||
| 375 | -#include <iostream> | ||
| 376 | -#include <vector> | ||
| 377 | -#include <acl/acl.h> | ||
| 378 | -#include <aclnnop/aclnn_matmul_compress_dequant.h> | ||
| 379 | -#include <fstream> | ||
| 380 | -#include <unistd.h> | ||
| 381 | -#include <sys/stat.h> | ||
| 382 | -#include <stdio.h> | ||
| 383 | -#include <cstdlib> | ||
| 384 | -#include <string> | ||
| 385 | - | ||
| 386 | -#define CHECK_RET(cond, return_expr) \ | ||
| 387 | - do { \ | ||
| 388 | - if (!(cond)) { \ | ||
| 389 | - return_expr; \ | ||
| 390 | - } \ | ||
| 391 | - } while (0) | ||
| 392 | - | ||
| 393 | -#define LOG_PRINT(message, ...) \ | ||
| 394 | - do { \ | ||
| 395 | - printf(message, ##__VA_ARGS__); \ | ||
| 396 | - } while (0) | ||
| 397 | - | ||
| 398 | -int64_t GetShapeSize(const std::vector<int64_t>& shape) { | ||
| 399 | - int64_t shapeSize = 1; | ||
| 400 | - for (auto i : shape) { | ||
| 401 | - shapeSize *= i; | ||
| 402 | - } | ||
| 403 | - return shapeSize; | ||
| 404 | -} | ||
| 405 | - | ||
| 406 | -int Init(int32_t deviceId, aclrtStream* stream) { | ||
| 407 | - // 固定写法，资源初始化 | ||
| 408 | - auto ret = aclInit(nullptr); | ||
| 409 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 410 | - ret = aclrtSetDevice(deviceId); | ||
| 411 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 412 | - ret = aclrtCreateStream(stream); | ||
| 413 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 414 | - return 0; | ||
| 415 | -} | ||
| 416 | - | ||
| 417 | -int ReadBinFileNNop(std::string filePath, void* buffer, size_t bufferSize) | ||
| 418 | -{ | ||
| 419 | - struct stat sBuf; | ||
| 420 | - int fileStatus = stat(filePath.data(), &sBuf); | ||
| 421 | - CHECK_RET(fileStatus == ACL_SUCCESS, LOG_PRINT("Failed to get file %s\n", filePath); return -1); | ||
| 422 | - | ||
| 423 | - std::ifstream file; | ||
| 424 | - file.open(filePath, std::ios::binary); | ||
| 425 | - CHECK_RET(file.is_open(), LOG_PRINT("Open file failed.\n"); return -1); | ||
| 426 | - | ||
| 427 | - file.seekg(0, file.end); | ||
| 428 | - uint64_t binFileBufferLen = file.tellg(); | ||
| 429 | - CHECK_RET(binFileBufferLen > 0, | ||
| 430 | - std::cout<<"File size is 0.\n"; | ||
| 431 | - file.close(); | ||
| 432 | - return -1); | ||
| 433 | - | ||
| 434 | - file.seekg(0, file.beg); | ||
| 435 | - file.read(static_cast<char *>(buffer), binFileBufferLen); | ||
| 436 | - file.close(); | ||
| 437 | - return ACL_SUCCESS; | ||
| 438 | -} | ||
| 439 | - | ||
| 440 | -int CreateAclTensor(std::string filePath, const std::vector<int64_t>& shape, int typeSize, | ||
| 441 | - void** deviceAddr, aclDataType dataType, aclTensor** tensor) { | ||
| 442 | - auto size = GetShapeSize(shape) * typeSize; | ||
| 443 | - // 调用aclrtMalloc申请device侧内存 | ||
| 444 | - auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 445 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 446 | - | ||
| 447 | - // 调用aclrtMallocHost申请host侧内存 | ||
| 448 | - void* binBufferHost = nullptr; | ||
| 449 | - ret = aclrtMallocHost(&binBufferHost, size); | ||
| 450 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMallocHost failed. ERROR: %d\n", ret); return ret); | ||
| 451 | - | ||
| 452 | - // 读取文件 | ||
| 453 | - ret = ReadBinFileNNop(filePath, binBufferHost, size); | ||
| 454 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("ReadBinFileNNop failed. ERROR: %d\n", ret); return ret); | ||
| 455 | - | ||
| 456 | - // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 | ||
| 457 | - ret = aclrtMemcpy(*deviceAddr, size, binBufferHost, size, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 458 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); | ||
| 459 | - | ||
| 460 | - // 计算连续tensor的strides | ||
| 461 | - std::vector<int64_t> strides(shape.size(), 1); | ||
| 462 | - for (int64_t i = shape.size() - 2; i >= 0; i--) { | ||
| 463 | - strides[i] = shape[i + 1] * strides[i + 1]; | ||
| 464 | - } | ||
| 465 | - | ||
| 466 | - // 调用aclCreateTensor接口创建aclTensor | ||
| 467 | - *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, | ||
| 468 | - shape.data(), shape.size(), *deviceAddr); | ||
| 469 | - return 0; | ||
| 470 | -} | ||
| 471 | - | ||
| 472 | -int main(int argc, char* argv[]) { | ||
| 473 | - // 1. （固定写法）device/stream初始化，参考acl API手册 | ||
| 474 | - // 根据自己的实际device填写deviceId | ||
| 475 | - int32_t deviceId = 0; | ||
| 476 | - aclrtStream stream; | ||
| 477 | - auto ret = Init(deviceId, &stream); | ||
| 478 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); | ||
| 479 | - | ||
| 480 | - if (argc != 6) { | ||
| 481 | - std::cerr << "Error: Invalid number of arguments. Usage: <program> m k n wCompressedSize indexSize" << std::endl; | ||
| 482 | - return -1; | ||
| 483 | - } | ||
| 484 | - | ||
| 485 | - // 2. 构造输入与输出，需要根据API的接口自定义构造 | ||
| 486 | - int m = atoi(argv[1]); | ||
| 487 | - int k = atoi(argv[2]); | ||
| 488 | - int n = atoi(argv[3]); | ||
| 489 | - // wShape是右矩阵压缩后数据的大小 | ||
| 490 | - int wCompressedSize = atoi(argv[4]); | ||
| 491 | - // indexShape是压缩索引数据的大小 | ||
| 492 | - int indexSize = atoi(argv[5]); | ||
| 493 | - | ||
| 494 | - if (m <= 0 || k <= 0 || n <= 0 || wCompressedSize <= 0 || indexSize <= 0) { | ||
| 495 | - std::cerr << "Error: m, k, n, wCompressedSize and indexSize must be positive integers." << std::endl; | ||
| 496 | - return -1; | ||
| 497 | - } | ||
| 498 | - | ||
| 499 | - std::vector<int64_t> mat1Shape = {m, k}; | ||
| 500 | - std::vector<int64_t> mat2CompressedShape = {wCompressedSize}; | ||
| 501 | - std::vector<int64_t> indexShape = {indexSize}; | ||
| 502 | - std::vector<int64_t> biasShape = {n}; | ||
| 503 | - std::vector<int64_t> deqScaleShape = {n}; | ||
| 504 | - std::vector<int64_t> outputShape = {m, n}; | ||
| 505 | - | ||
| 506 | - std::vector<int64_t> compressInfoHostData = {8, 8, k, n, 1}; | ||
| 507 | - | ||
| 508 | - void* mat1DeviceAddr = nullptr; | ||
| 509 | - void* mat2CompressedDeviceAddr = nullptr; | ||
| 510 | - void* indexDeviceAddr = nullptr; | ||
| 511 | - void* biasDeviceAddr = nullptr; | ||
| 512 | - void* deqScaleDeviceAddr = nullptr; | ||
| 513 | - void* outputDeviceAddr = nullptr; | ||
| 514 | - | ||
| 515 | - aclTensor* mat1 = nullptr; | ||
| 516 | - aclTensor* mat2Compressed = nullptr; | ||
| 517 | - aclTensor* index = nullptr; | ||
| 518 | - aclTensor* bias = nullptr; | ||
| 519 | - aclTensor* deqScale = nullptr; | ||
| 520 | - aclTensor* output = nullptr; | ||
| 521 | - aclIntArray* compressInfo = nullptr; | ||
| 522 | - | ||
| 523 | - std::string rootPath = "./data/"; | ||
| 524 | - | ||
| 525 | - // 创建mat1 aclTensor | ||
| 526 | - std::string mat1FilePath = rootPath + "mat1.bin"; | ||
| 527 | - ret = CreateAclTensor(mat1FilePath, mat1Shape, sizeof(int8_t), &mat1DeviceAddr, aclDataType::ACL_INT8, &mat1); | ||
| 528 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Create mat1 tensor failed. ERROR: %d\n", ret); return ret); | ||
| 529 | - // 创建mat2Compressed aclTensor | ||
| 530 | - std::string mat2FilePath = rootPath + "weight/weight.dat"; | ||
| 531 | - ret = CreateAclTensor(mat2FilePath, mat2CompressedShape, sizeof(int8_t), &mat2CompressedDeviceAddr, | ||
| 532 | - aclDataType::ACL_INT8, &mat2Compressed); | ||
| 533 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Create mat2 tensor failed. ERROR: %d\n", ret); return ret); | ||
| 534 | - // 创建index aclTensor | ||
| 535 | - std::string indexFilePath = rootPath + "index/weight.dat"; | ||
| 536 | - ret = CreateAclTensor(indexFilePath, indexShape, sizeof(int8_t), &indexDeviceAddr, aclDataType::ACL_INT8, &index); | ||
| 537 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Create index tensor failed. ERROR: %d\n", ret); return ret); | ||
| 538 | - // 创建bias aclTensor | ||
| 539 | - std::string biasFilePath = rootPath + "bias.bin"; | ||
| 540 | - ret = CreateAclTensor(biasFilePath, biasShape, sizeof(int32_t), &biasDeviceAddr, aclDataType::ACL_INT32, &bias); | ||
| 541 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Create bias tensor failed. ERROR: %d\n", ret); return ret); | ||
| 542 | - // 创建deqScale aclTensor | ||
| 543 | - std::string deqScaleFilePath = rootPath + "deqScale.bin"; | ||
| 544 | - ret = CreateAclTensor(deqScaleFilePath, deqScaleShape, sizeof(int32_t), &deqScaleDeviceAddr, aclDataType::ACL_UINT64, | ||
| 545 | - &deqScale); | ||
| 546 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Create deqScale tensor failed. ERROR: %d\n", ret); return ret); | ||
| 547 | - // 创建compressInfo | ||
| 548 | - compressInfo = aclCreateIntArray(compressInfoHostData.data(), aclDataType::ACL_INT64); | ||
| 549 | - // 创建out aclTensor | ||
| 550 | - std::string outputFilePath = rootPath + "output.bin"; | ||
| 551 | - ret = CreateAclTensor(outputFilePath, outputShape, 2, &outputDeviceAddr, aclDataType::ACL_FLOAT16, &output); | ||
| 552 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Create output tensor failed. ERROR: %d\n", ret); return ret); | ||
| 553 | - | ||
| 554 | - int32_t offsetX = 0; | ||
| 555 | - | ||
| 556 | - // 3. 调用CANN算子库API，需要修改为具体的Api名称 | ||
| 557 | - uint64_t workspaceSize = 0; | ||
| 558 | - aclOpExecutor* executor; | ||
| 559 | - // 调用aclnnMm第一段接口 | ||
| 560 | - ret = aclnnMatmulCompressDequantGetWorkspaceSize(mat1, mat2Compressed, index, bias, deqScale, nullptr, offsetX, compressInfo, output, &workspaceSize, &executor); | ||
| 561 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMatmulCompressDequantGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 562 | - // 根据第一段接口计算出的workspaceSize申请device内存 | ||
| 563 | - void* workspaceAddr = nullptr; | ||
| 564 | - if (workspaceSize > 0) { | ||
| 565 | - ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 566 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret); | ||
| 567 | - } | ||
| 568 | - // 调用aclnnMm第二段接口 | ||
| 569 | - ret = aclnnMatmulCompressDequant(workspaceAddr, workspaceSize, executor, stream); | ||
| 570 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMatmulCompressDequant failed. ERROR: %d\n", ret); return ret); | ||
| 571 | - | ||
| 572 | - // 4. （固定写法）同步等待任务执行结束 | ||
| 573 | - ret = aclrtSynchronizeStream(stream); | ||
| 574 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 575 | - | ||
| 576 | - // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改 | ||
| 577 | - auto size = GetShapeSize(outputShape); | ||
| 578 | - std::vector<float> resultData(size, 0); | ||
| 579 | - ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outputDeviceAddr, | ||
| 580 | - size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 581 | - CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); | ||
| 582 | - for (int64_t i = 0; i < size; i++) { | ||
| 583 | - LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]); | ||
| 584 | - } | ||
| 585 | - | ||
| 586 | - // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改 | ||
| 587 | - aclDestroyTensor(mat1); | ||
| 588 | - aclDestroyTensor(mat2Compressed); | ||
| 589 | - aclDestroyTensor(index); | ||
| 590 | - aclDestroyTensor(bias); | ||
| 591 | - aclDestroyTensor(deqScale); | ||
| 592 | - aclDestroyTensor(output); | ||
| 593 | - aclDestroyIntArray(compressInfo); | ||
| 594 | - | ||
| 595 | - // 7.释放硬件资源，需要根据具体API的接口定义修改 | ||
| 596 | - aclrtFree(mat1DeviceAddr); | ||
| 597 | - aclrtFree(mat2CompressedDeviceAddr); | ||
| 598 | - aclrtFree(indexDeviceAddr); | ||
| 599 | - aclrtFree(biasDeviceAddr); | ||
| 600 | - aclrtFree(deqScaleDeviceAddr); | ||
| 601 | - aclrtFree(outputDeviceAddr); | ||
| 602 | - if (workspaceSize > 0) { | ||
| 603 | - aclrtFree(workspaceAddr); | ||
| 604 | - } | ||
| 605 | - aclrtDestroyStream(stream); | ||
| 606 | - aclrtResetDevice(deviceId); | ||
| 607 | - aclFinalize(); | ||
| 608 | - return 0; | ||
| 609 | -} | ||
| 610 | -``` | ||
| @@ -1,13 +0,0 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | ||
| 2 | -# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | -# This file is a part of the CANN Open Software. | ||
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | -# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | -# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | -# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | -# ---------------------------------------------------------------------------- | ||
| 11 | - | ||
| 12 | -add_modules_sources(OPTYPE matmul_v2_compress_dequant ACLNNTYPE aclnn_exclude) | ||
| 13 | -add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE matmul_v2_compress_dequant) | ||
| @@ -1,95 +0,0 @@ | |||
| 1 | -{ | ||
| 2 | - "op_type": "MatMulV2CompressDequant", | ||
| 3 | - "op_list": [ | ||
| 4 | - { | ||
| 5 | - "bin_filename": "MatmulV2CompressDequant_Int8", | ||
| 6 | - "inputs": [ | ||
| 7 | - { | ||
| 8 | - "name": "x1", | ||
| 9 | - "index": 0, | ||
| 10 | - "dtype": "int8", | ||
| 11 | - "format": "ND", | ||
| 12 | - "paramType": "required", | ||
| 13 | - "shape": [-2] | ||
| 14 | - }, | ||
| 15 | - { | ||
| 16 | - "name": "x2", | ||
| 17 | - "index": 1, | ||
| 18 | - "dtype": "int8", | ||
| 19 | - "format": "ND", | ||
| 20 | - "paramType": "required", | ||
| 21 | - "shape": [-2] | ||
| 22 | - }, | ||
| 23 | - { | ||
| 24 | - "name": "compress_index", | ||
| 25 | - "index": 2, | ||
| 26 | - "dtype": "int8", | ||
| 27 | - "format": "ND", | ||
| 28 | - "paramType": "required", | ||
| 29 | - "shape": [-2] | ||
| 30 | - }, | ||
| 31 | - { | ||
| 32 | - "name": "deq_scale", | ||
| 33 | - "index": 3, | ||
| 34 | - "dtype": "uint64", | ||
| 35 | - "format": "ND", | ||
| 36 | - "paramType": "required", | ||
| 37 | - "shape": [-2] | ||
| 38 | - }, | ||
| 39 | - { | ||
| 40 | - "name": "bias", | ||
| 41 | - "index": 4, | ||
| 42 | - "dtype": "int32", | ||
| 43 | - "format": "ND", | ||
| 44 | - "paramType": "required", | ||
| 45 | - "shape": [-2] | ||
| 46 | - }, | ||
| 47 | - { | ||
| 48 | - "name": "offset_w", | ||
| 49 | - "index": 5, | ||
| 50 | - "dtype": "int8", | ||
| 51 | - "format": "ND", | ||
| 52 | - "paramType": "optional", | ||
| 53 | - "shape": [-2] | ||
| 54 | - } | ||
| 55 | - ], | ||
| 56 | - "outputs": [ | ||
| 57 | - { | ||
| 58 | - "name": "out", | ||
| 59 | - "index": 0, | ||
| 60 | - "dtype": "float16", | ||
| 61 | - "format": "ND", | ||
| 62 | - "paramType": "required", | ||
| 63 | - "shape": [-2] | ||
| 64 | - } | ||
| 65 | - ], | ||
| 66 | - "attrs": [ | ||
| 67 | - { | ||
| 68 | - "name": "transpose_x1", | ||
| 69 | - "dtype": "bool", | ||
| 70 | - "value": null | ||
| 71 | - }, | ||
| 72 | - { | ||
| 73 | - "name": "transpose_x2", | ||
| 74 | - "dtype": "bool", | ||
| 75 | - "value": null | ||
| 76 | - }, | ||
| 77 | - { | ||
| 78 | - "name": "compress_info", | ||
| 79 | - "dtype": "list_int", | ||
| 80 | - "value": null | ||
| 81 | - }, | ||
| 82 | - { | ||
| 83 | - "name": "offset_x", | ||
| 84 | - "dtype": "int", | ||
| 85 | - "value": null | ||
| 86 | - }, | ||
| 87 | - { | ||
| 88 | - "name": "alg", | ||
| 89 | - "dtype": "string", | ||
| 90 | - "value": null | ||
| 91 | - } | ||
| 92 | - ] | ||
| 93 | - } | ||
| 94 | - ] | ||
| 95 | -} | ||
| @@ -1,74 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | -/*! | ||
| 12 | - * \file matmul_v2_compress_dequant_def.cpp | ||
| 13 | - * \brief MatMulV2CompressDequant op definition for ascend310p | ||
| 14 | - */ | ||
| 15 | - | ||
| 16 | - | ||
| 17 | -namespace ops { | ||
| 18 | -class MatMulV2CompressDequant : public OpDef { | ||
| 19 | -public: | ||
| 20 | - explicit MatMulV2CompressDequant(const char* name) : OpDef(name) | ||
| 21 | - { | ||
| 22 | - this->Input("x1") | ||
| 23 | - .ParamType(REQUIRED) | ||
| 24 | - .DataType({ge::DT_INT8}) | ||
| 25 | - .Format({ge::FORMAT_ND}) | ||
| 26 | - .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 27 | - this->Input("x2") | ||
| 28 | - .ParamType(REQUIRED) | ||
| 29 | - .DataType({ge::DT_INT8}) | ||
| 30 | - .Format({ge::FORMAT_ND}) | ||
| 31 | - .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 32 | - this->Input("compress_index") | ||
| 33 | - .ParamType(REQUIRED) | ||
| 34 | - .DataType({ge::DT_INT8}) | ||
| 35 | - .Format({ge::FORMAT_ND}) | ||
| 36 | - .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 37 | - this->Input("deq_scale") | ||
| 38 | - .ParamType(REQUIRED) | ||
| 39 | - .DataType({ge::DT_UINT64}) | ||
| 40 | - .Format({ge::FORMAT_ND}) | ||
| 41 | - .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 42 | - this->Input("bias") | ||
| 43 | - .ParamType(REQUIRED) | ||
| 44 | - .DataType({ge::DT_INT32}) | ||
| 45 | - .Format({ge::FORMAT_ND}) | ||
| 46 | - .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 47 | - this->Input("offset_w") | ||
| 48 | - .ParamType(OPTIONAL) | ||
| 49 | - .DataType({ge::DT_INT8}) | ||
| 50 | - .Format({ge::FORMAT_ND}) | ||
| 51 | - .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 52 | - this->Attr("transpose_x1").AttrType(OPTIONAL).Bool(false); | ||
| 53 | - this->Attr("transpose_x2").AttrType(OPTIONAL).Bool(false); | ||
| 54 | - this->Attr("compress_info").AttrType(OPTIONAL).ListInt({1, 1, 1, 1, 1}); | ||
| 55 | - this->Attr("offset_x").AttrType(OPTIONAL).Int(0); | ||
| 56 | - this->Attr("alg").AttrType(OPTIONAL).String("weight_unzip"); | ||
| 57 | - this->Output("out") | ||
| 58 | - .ParamType(REQUIRED) | ||
| 59 | - .DataType({ge::DT_FLOAT16}) | ||
| 60 | - .Format({ge::FORMAT_ND}) | ||
| 61 | - .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 62 | - | ||
| 63 | - OpAICoreConfig config_310p; | ||
| 64 | - config_310p.DynamicCompileStaticFlag(true) | ||
| 65 | - .DynamicRankSupportFlag(true) | ||
| 66 | - .DynamicShapeSupportFlag(true) | ||
| 67 | - .NeedCheckSupportFlag(false) | ||
| 68 | - .ExtendCfgInfo("opFile.value", "matmul_v2_compress_dequant"); | ||
| 69 | - this->AICore().AddConfig("ascend310p", config_310p); | ||
| 70 | - } | ||
| 71 | -}; | ||
| 72 | - | ||
| 73 | -OP_ADD(MatMulV2CompressDequant); | ||
| 74 | -} // namespace ops | ||
| @@ -1,299 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | -/*! | ||
| 12 | - * \file matmul_v2_compress_dequant_tiling.cpp | ||
| 13 | - * \brief Tiling implementation for matmul_v2_compress_dequant | ||
| 14 | - * Core algorithm ported from pp_matmul_i8_nz_tiling.cpp | ||
| 15 | - */ | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| 24 | -namespace optiling { | ||
| 25 | -namespace { | ||
| 26 | - | ||
| 27 | -// ============================================================ | ||
| 28 | -// Constants (aligned with pp_matmul common_tiling.h) | ||
| 29 | -// ============================================================ | ||
| 30 | -constexpr uint32_t BLOCK_SIZE_16 = 16; | ||
| 31 | -constexpr uint32_t BLOCK_SIZE_INT8_K = 32; | ||
| 32 | -constexpr uint32_t BASE_BLOCK_STEP = 2; | ||
| 33 | -constexpr uint32_t FP16_BYTE_SIZE = 2; | ||
| 34 | -constexpr uint32_t FP32_BYTE_SIZE = 4; | ||
| 35 | -constexpr float INT8_DTYPE_SIZE = 1.0f; | ||
| 36 | -constexpr uint32_t AXES_ALIGN_SIZE = 512; | ||
| 37 | -constexpr uint32_t AXES_ALIGN_SIZE_INT8 = 256; | ||
| 38 | -constexpr uint32_t L0AB_PP_BUF_LEN_FP16 = 131072; // 128 KB | ||
| 39 | -constexpr uint32_t L1AB_PP_BUF_LEN_INT8_SPARSE = 163840; // 160 KB | ||
| 40 | -constexpr uint32_t ALIGNMENT_16 = 16; | ||
| 41 | -constexpr uint32_t CONST_3 = 3; | ||
| 42 | -constexpr uint32_t CONST_4 = 4; | ||
| 43 | -constexpr uint32_t CONST_256 = 256; | ||
| 44 | -constexpr uint32_t TRANS_B_MASK = 0b001000; | ||
| 45 | -constexpr float L2_BW_RATIO = 5.0f; | ||
| 46 | -constexpr uint64_t DEFAULT_L2_SIZE = 48UL * 1024 * 1024; // 48 MB | ||
| 47 | -constexpr uint64_t DEFAULT_L0C_SIZE = 256UL * 1024; // 256 KB | ||
| 48 | - | ||
| 49 | -constexpr size_t TRANSPOSE_X1_ATTR_IDX = 0; | ||
| 50 | -constexpr size_t TRANSPOSE_X2_ATTR_IDX = 1; | ||
| 51 | -constexpr size_t COMPRESS_INFO_ATTR_IDX = 2; | ||
| 52 | - | ||
| 53 | -constexpr size_t CI_IDX_TILING_K = 0; | ||
| 54 | -constexpr size_t CI_IDX_TILING_N = 1; | ||
| 55 | -constexpr size_t CI_IDX_K = 2; | ||
| 56 | -constexpr size_t CI_IDX_N = 3; | ||
| 57 | -constexpr size_t COMPRESS_INFO_MIN_ELEMS = 4; | ||
| 58 | - | ||
| 59 | -constexpr size_t INPUT_IDX_X1 = 0; | ||
| 60 | -constexpr size_t INPUT_IDX_BIAS = 4; | ||
| 61 | - | ||
| 62 | -// ============================================================ | ||
| 63 | -// Utility functions | ||
| 64 | -// ============================================================ | ||
| 65 | - | ||
| 66 | - | ||
| 67 | -inline uint32_t CeilDivU32(uint32_t dividend, uint32_t divisor) | ||
| 68 | -{ | ||
| 69 | - if (divisor == 0 || dividend + divisor - 1 < dividend) { | ||
| 70 | - return dividend; | ||
| 71 | - } | ||
| 72 | - return (dividend + divisor - 1) / divisor; | ||
| 73 | -} | ||
| 74 | - | ||
| 75 | - | ||
| 76 | -// // ============================================================ | ||
| 77 | -// // Compute k0 and compressOverlapN | ||
| 78 | -// // Ported from PpTilingData310P::End | ||
| 79 | -// // ============================================================ | ||
| 80 | -uint32_t ComputeK0(uint32_t n0, uint32_t n, | ||
| 81 | - bool isCompress, uint32_t tilingNVal, | ||
| 82 | - uint32_t &compressOverlapN) | ||
| 83 | -{ | ||
| 84 | - compressOverlapN = 0; | ||
| 85 | - | ||
| 86 | - if (isCompress) { | ||
| 87 | - if (n0 == 0) { | ||
| 88 | - return 1; | ||
| 89 | - } | ||
| 90 | - uint32_t nTail = n % n0; | ||
| 91 | - uint32_t compressNTile = (tilingNVal > 0) ? (CeilDivU32(nTail, ALIGNMENT_16) % tilingNVal) : 0; | ||
| 92 | - compressOverlapN = (compressNTile == 0) ? 0 : tilingNVal - compressNTile; | ||
| 93 | - } | ||
| 94 | - return 0; | ||
| 95 | -} | ||
| 96 | - | ||
| 97 | -// ============================================================ | ||
| 98 | -// Extract M, K, N, batchSize from TilingContext | ||
| 99 | -// Handles both ND shapes (2D: [M, K]) and NZ shapes (4D) | ||
| 100 | -// ============================================================ | ||
| 101 | -ge::graphStatus ExtractMatmulDims(const gert::TilingContext *context, uint32_t &batchSize, uint32_t &m, uint32_t &k, uint32_t &n, | ||
| 102 | - const int64_t *compressInfoData, size_t compressInfoCount) | ||
| 103 | -{ | ||
| 104 | - const auto *x1ShapePtr = context->GetInputShape(INPUT_IDX_X1); | ||
| 105 | - if (x1ShapePtr == nullptr) { | ||
| 106 | - OP_LOGE("MatMulV2CompressDequant", "x1 shape is nullptr."); | ||
| 107 | - return ge::GRAPH_FAILED; | ||
| 108 | - } | ||
| 109 | - const auto &x1Shape = x1ShapePtr->GetOriginShape(); | ||
| 110 | - | ||
| 111 | - m = static_cast<uint32_t>(x1Shape.GetDim(0)); | ||
| 112 | - k = static_cast<uint32_t>(x1Shape.GetDim(1)); | ||
| 113 | - // 输入是个2维 | ||
| 114 | - batchSize = 1; | ||
| 115 | - | ||
| 116 | - n = 0; | ||
| 117 | - if (compressInfoCount > CI_IDX_N && compressInfoData != nullptr) { | ||
| 118 | - n = static_cast<uint32_t>(compressInfoData[CI_IDX_N]); | ||
| 119 | - } | ||
| 120 | - if (n == 0) { | ||
| 121 | - const auto *outShapePtr = context->GetOutputShape(0); | ||
| 122 | - if (outShapePtr != nullptr) { | ||
| 123 | - const auto &outShape = outShapePtr->GetStorageShape(); | ||
| 124 | - size_t outNdim = outShape.GetDimNum(); | ||
| 125 | - if (outNdim >= 4) { | ||
| 126 | - n = static_cast<uint32_t>(outShape.GetDim(outNdim - 4) * | ||
| 127 | - outShape.GetDim(outNdim - 1)); | ||
| 128 | - } else if (outNdim >= 2) { | ||
| 129 | - n = static_cast<uint32_t>(outShape.GetDim(outNdim - 1)); | ||
| 130 | - } | ||
| 131 | - } | ||
| 132 | - } | ||
| 133 | - | ||
| 134 | - if (m == 0 || k == 0 || n == 0) { | ||
| 135 | - OP_LOGE("MatMulV2CompressDequant", "Invalid dims: m=%u, k=%u, n=%u.", m, k, n); | ||
| 136 | - return ge::GRAPH_FAILED; | ||
| 137 | - } | ||
| 138 | - return ge::GRAPH_SUCCESS; | ||
| 139 | -} | ||
| 140 | - | ||
| 141 | -} // namespace | ||
| 142 | - | ||
| 143 | - | ||
| 144 | -static ge::graphStatus TbmmEinsumTilingFunc(gert::TilingContext* context, | ||
| 145 | - uint32_t &batchSize, uint32_t &m, uint32_t &k, uint32_t &n, | ||
| 146 | - uint32_t& tilingK, uint32_t& tilingN) | ||
| 147 | -{ | ||
| 148 | - if (context == nullptr) { | ||
| 149 | - OP_LOGE("MatMulV2CompressDequant", "context is null."); | ||
| 150 | - return ge::GRAPH_FAILED; | ||
| 151 | - } | ||
| 152 | - auto compileInfo = reinterpret_cast<const MatmulV2CompressDequantCompileInfo *>(context->GetCompileInfo()); | ||
| 153 | - uint32_t coreNum = static_cast<uint32_t>(compileInfo->aivNum); | ||
| 154 | - | ||
| 155 | - size_t sysWorkspaceSize = static_cast<size_t>(24 * 1024 * 1024); // 24M same as ppmatmul tiling | ||
| 156 | - size_t* currentWorkSpace = context->GetWorkspaceSizes(1); | ||
| 157 | - currentWorkSpace[0] = sysWorkspaceSize; | ||
| 158 | - | ||
| 159 | - optiling::pp_matmul::PpMatMulDefault tbmmEinsumTiling(context); | ||
| 160 | - auto inputDType = context->GetInputDesc(0)->GetDataType(); | ||
| 161 | - | ||
| 162 | - tbmmEinsumTiling.matMulInfo_.isInt8 = (inputDType == ge::DT_INT8); | ||
| 163 | - tbmmEinsumTiling.matMulInfo_.inDtype = ge::GetSizeByDataType(inputDType); | ||
| 164 | - (void)tbmmEinsumTiling.GetHardwareInfo(); | ||
| 165 | - | ||
| 166 | - // 设置一个m, k, n | ||
| 167 | - tbmmEinsumTiling.matMulInfo_.m = m; | ||
| 168 | - tbmmEinsumTiling.matMulInfo_.k = k; | ||
| 169 | - tbmmEinsumTiling.matMulInfo_.n = n; | ||
| 170 | - tbmmEinsumTiling.matMulInfo_.batchSize = batchSize; | ||
| 171 | - (void)tbmmEinsumTiling.GetMatMulTilingData(); | ||
| 172 | - (void)tbmmEinsumTiling.PrintTiling(); | ||
| 173 | - | ||
| 174 | - uint32_t compressOverlapN = 0; | ||
| 175 | - (void)ComputeK0(tbmmEinsumTiling.ppMatmulDefaultTilingData_.opShape.n0, | ||
| 176 | - tbmmEinsumTiling.ppMatmulDefaultTilingData_.opShape.n, true, tilingN, compressOverlapN); | ||
| 177 | - // PostTiling | ||
| 178 | - uint32_t blockDim = std::min(tbmmEinsumTiling.ppMatmulDefaultTilingData_.coreLoop, (uint64_t)coreNum); | ||
| 179 | - context->SetBlockDim(blockDim); | ||
| 180 | - MatmulV2CompressDequantTilingData tiling; | ||
| 181 | - tiling.set_batchSize(batchSize); | ||
| 182 | - tiling.set_m(tbmmEinsumTiling.ppMatmulDefaultTilingData_.opShape.m); | ||
| 183 | - tiling.set_k(tbmmEinsumTiling.ppMatmulDefaultTilingData_.opShape.k); | ||
| 184 | - tiling.set_n(tbmmEinsumTiling.ppMatmulDefaultTilingData_.opShape.n); | ||
| 185 | - tiling.set_m0(tbmmEinsumTiling.ppMatmulDefaultTilingData_.opShape.m0); | ||
| 186 | - tiling.set_k0(tbmmEinsumTiling.ppMatmulDefaultTilingData_.opShape.k0); | ||
| 187 | - tiling.set_n0(tbmmEinsumTiling.ppMatmulDefaultTilingData_.opShape.n0); | ||
| 188 | - tiling.set_mLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.mLoop); | ||
| 189 | - tiling.set_kLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.kLoop); | ||
| 190 | - tiling.set_nLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.nLoop); | ||
| 191 | - tiling.set_coreLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.coreLoop); | ||
| 192 | - tiling.set_swizzlCount(tbmmEinsumTiling.ppMatmulDefaultTilingData_.swizzlCount); | ||
| 193 | - tiling.set_tilingK(tilingK); | ||
| 194 | - tiling.set_tilingN(tilingN); | ||
| 195 | - tiling.set_compressOverlapN(compressOverlapN); | ||
| 196 | - tiling.set_tilingKey(0); | ||
| 197 | - tiling.set_blockDimVal(blockDim); | ||
| 198 | - tiling.set_splitK(0); | ||
| 199 | - tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); | ||
| 200 | - context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); | ||
| 201 | - | ||
| 202 | - return ge::GRAPH_SUCCESS; | ||
| 203 | -} | ||
| 204 | - | ||
| 205 | - | ||
| 206 | -// ============================================================ | ||
| 207 | -// Main tiling function | ||
| 208 | -// ============================================================ | ||
| 209 | -ge::graphStatus TilingForMatmulV2CompressDequant(gert::TilingContext *context) | ||
| 210 | -{ | ||
| 211 | - OP_LOGI(context->GetNodeName(), "TbmmEinsum Tiling start."); | ||
| 212 | - | ||
| 213 | - uint32_t batchSize = 1; | ||
| 214 | - uint32_t m = 0, k = 0, n = 0; | ||
| 215 | - const int64_t *compressInfoData = nullptr; | ||
| 216 | - size_t compressInfoCount = 0; | ||
| 217 | - auto *attrs = context->GetAttrs(); | ||
| 218 | - const auto *compressInfoVec = attrs->GetAttrPointer<gert::ContinuousVector>(COMPRESS_INFO_ATTR_IDX); | ||
| 219 | - compressInfoCount = compressInfoVec->GetSize(); | ||
| 220 | - compressInfoData = reinterpret_cast<const int64_t *>(compressInfoVec->GetData()); | ||
| 221 | - | ||
| 222 | - (void)ExtractMatmulDims(context, batchSize, m, k, n, | ||
| 223 | - compressInfoData, compressInfoCount); | ||
| 224 | - | ||
| 225 | - uint32_t tilingKVal = 0; | ||
| 226 | - uint32_t tilingNVal = 0; | ||
| 227 | - if (compressInfoData != nullptr && compressInfoCount >= 2) { | ||
| 228 | - tilingKVal = static_cast<uint32_t>(compressInfoData[CI_IDX_TILING_K]); | ||
| 229 | - tilingNVal = static_cast<uint32_t>(compressInfoData[CI_IDX_TILING_N]); | ||
| 230 | - } | ||
| 231 | - TbmmEinsumTilingFunc(context, batchSize, m, k, n, tilingKVal, tilingNVal); | ||
| 232 | - | ||
| 233 | - auto compileInfo = reinterpret_cast<const MatmulV2CompressDequantCompileInfo *>(context->GetCompileInfo()); | ||
| 234 | - | ||
| 235 | - // ---- Workspace ---- | ||
| 236 | - size_t *currentWorkspace = context->GetWorkspaceSizes(1); | ||
| 237 | - currentWorkspace[0] = compileInfo->workSpaceSize; | ||
| 238 | - return ge::GRAPH_SUCCESS; | ||
| 239 | -} | ||
| 240 | - | ||
| 241 | -// ============================================================ | ||
| 242 | -// TilingPrepare – collect compile-time hardware info | ||
| 243 | -// ============================================================ | ||
| 244 | -ge::graphStatus TilingPrepareForMatmulV2CompressDequant(gert::TilingParseContext *context) | ||
| 245 | -{ | ||
| 246 | - if (context == nullptr) { | ||
| 247 | - OP_LOGE("MatMulV2CompressDequant", "TilingParse context is nullptr."); | ||
| 248 | - return ge::GRAPH_FAILED; | ||
| 249 | - } | ||
| 250 | - OP_LOGD(context, "TilingPrepareForMatmulV2CompressDequant start."); | ||
| 251 | - | ||
| 252 | - fe::PlatFormInfos* platformInfo = context->GetPlatformInfo(); | ||
| 253 | - if (platformInfo == nullptr) { | ||
| 254 | - OP_LOGE("MatMulV2CompressDequant", "platformInfoPtr is null"); | ||
| 255 | - return ge::GRAPH_FAILED; | ||
| 256 | - } | ||
| 257 | - | ||
| 258 | - auto compileInfoPtr = context->GetCompiledInfo<MatmulV2CompressDequantCompileInfo>(); | ||
| 259 | - if (compileInfoPtr == nullptr) { | ||
| 260 | - OP_LOGE("MatMulV2CompressDequant", "compileInfoPtr is null"); | ||
| 261 | - return ge::GRAPH_FAILED; | ||
| 262 | - } | ||
| 263 | - auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo); | ||
| 264 | - platformInfo->GetPlatformRes("version", "SoC_version", compileInfoPtr->socVersionStr); | ||
| 265 | - std::string val; | ||
| 266 | - std::string dataMoveL12Bt; | ||
| 267 | - platformInfo->GetPlatformRes("AICoreintrinsicDtypeMap", "Intrinsic_fix_pipe_l0c2out", val); | ||
| 268 | - platformInfo->GetPlatformRes("AICoreintrinsicDtypeMap", "Intrinsic_data_move_l12bt", dataMoveL12Bt); | ||
| 269 | - compileInfoPtr->supportL0c2out = !val.empty(); | ||
| 270 | - compileInfoPtr->supportL12BtBf16 = (dataMoveL12Bt.find("bf16") != std::string::npos); | ||
| 271 | - compileInfoPtr->workSpaceSize = ascendcPlatform.GetLibApiWorkSpaceSize(); | ||
| 272 | - compileInfoPtr->aivNum = ascendcPlatform.GetCoreNumAiv(); | ||
| 273 | - compileInfoPtr->aicNum = ascendcPlatform.GetCoreNumAic(); | ||
| 274 | - compileInfoPtr->socVersion = ascendcPlatform.GetSocVersion(); | ||
| 275 | - compileInfoPtr->npuArch = ascendcPlatform.GetCurNpuArch(); | ||
| 276 | - compileInfoPtr->btSize = compileInfoPtr->supportL0c2out ? 1024UL : 0UL; // 1024 is btSize | ||
| 277 | - compileInfoPtr->btSize = compileInfoPtr->supportL12BtBf16 ? 4096UL : compileInfoPtr->btSize; // 4096 is btSize | ||
| 278 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfoPtr->ubSize); | ||
| 279 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L1, compileInfoPtr->l1Size); | ||
| 280 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_A, compileInfoPtr->l0ASize); | ||
| 281 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_B, compileInfoPtr->l0BSize); | ||
| 282 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_C, compileInfoPtr->l0CSize); | ||
| 283 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L2, compileInfoPtr->l2Size); | ||
| 284 | - | ||
| 285 | - compileInfoPtr->isRegbase = Ops::Math::OpTiling::IsRegbaseSocVersion(context); | ||
| 286 | - | ||
| 287 | - OP_LOGD(context, "TilingPrepare end: aivNum=%lu ubSize=%lu l0CSize=%lu l2Size=%lu", | ||
| 288 | - compileInfoPtr->aivNum, compileInfoPtr->ubSize, compileInfoPtr->l0CSize, compileInfoPtr->l2Size); | ||
| 289 | - return ge::GRAPH_SUCCESS; | ||
| 290 | -} | ||
| 291 | - | ||
| 292 | -// ============================================================ | ||
| 293 | -// Registration | ||
| 294 | -// ============================================================ | ||
| 295 | -IMPL_OP_OPTILING(MatMulV2CompressDequant) | ||
| 296 | - .Tiling(TilingForMatmulV2CompressDequant) | ||
| 297 | - .TilingParse<MatmulV2CompressDequantCompileInfo>(TilingPrepareForMatmulV2CompressDequant); | ||
| 298 | - | ||
| 299 | -} // namespace optiling | ||
| @@ -1,74 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | -/*! | ||
| 12 | - * \file matmul_v2_compress_dequant_tiling.h | ||
| 13 | - * \brief Tiling data and registration for matmul_v2_compress_dequant | ||
| 14 | - */ | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| 24 | - | ||
| 25 | -namespace optiling { | ||
| 26 | - | ||
| 27 | -BEGIN_TILING_DATA_DEF(MatmulV2CompressDequantTilingData) | ||
| 28 | - TILING_DATA_FIELD_DEF(uint32_t, batchSize); | ||
| 29 | - TILING_DATA_FIELD_DEF(uint32_t, m); | ||
| 30 | - TILING_DATA_FIELD_DEF(uint32_t, k); | ||
| 31 | - TILING_DATA_FIELD_DEF(uint32_t, n); | ||
| 32 | - TILING_DATA_FIELD_DEF(uint32_t, m0); | ||
| 33 | - TILING_DATA_FIELD_DEF(uint32_t, k0); | ||
| 34 | - TILING_DATA_FIELD_DEF(uint32_t, n0); | ||
| 35 | - TILING_DATA_FIELD_DEF(uint32_t, mLoop); | ||
| 36 | - TILING_DATA_FIELD_DEF(uint32_t, kLoop); | ||
| 37 | - TILING_DATA_FIELD_DEF(uint32_t, nLoop); | ||
| 38 | - TILING_DATA_FIELD_DEF(uint32_t, coreLoop); | ||
| 39 | - TILING_DATA_FIELD_DEF(uint32_t, swizzlCount); | ||
| 40 | - TILING_DATA_FIELD_DEF(uint32_t, tilingK); | ||
| 41 | - TILING_DATA_FIELD_DEF(uint32_t, tilingN); | ||
| 42 | - TILING_DATA_FIELD_DEF(uint32_t, compressOverlapN); | ||
| 43 | - TILING_DATA_FIELD_DEF(uint32_t, tilingKey); | ||
| 44 | - TILING_DATA_FIELD_DEF(uint32_t, blockDimVal); | ||
| 45 | - TILING_DATA_FIELD_DEF(uint32_t, splitK); | ||
| 46 | -END_TILING_DATA_DEF; | ||
| 47 | - | ||
| 48 | -REGISTER_TILING_DATA_CLASS(MatMulV2CompressDequant, MatmulV2CompressDequantTilingData) | ||
| 49 | - | ||
| 50 | -struct MatmulV2CompressDequantCompileInfo { | ||
| 51 | - uint64_t aicNum{0UL}; | ||
| 52 | - uint64_t aivNum{0UL}; | ||
| 53 | - uint64_t ubSize = 0; | ||
| 54 | - uint64_t l1Size = 0; | ||
| 55 | - uint64_t l2Size = 0; | ||
| 56 | - uint64_t l0CSize = 0; | ||
| 57 | - uint64_t l0ASize{0UL}; | ||
| 58 | - uint64_t l0BSize{0UL}; | ||
| 59 | - uint64_t btSize{0UL}; | ||
| 60 | - float cubeFreq{0}; | ||
| 61 | - NpuArch npuArch; | ||
| 62 | - platform_ascendc::SocVersion socVersion; | ||
| 63 | - std::string socVersionStr = ""; | ||
| 64 | - bool supportL0c2out = false; | ||
| 65 | - bool supportL12BtBf16 = false; | ||
| 66 | - uint64_t workSpaceSize = 0; | ||
| 67 | - bool isRegbase = false; | ||
| 68 | -}; | ||
| 69 | - | ||
| 70 | -ge::graphStatus TilingForMatmulV2CompressDequant(gert::TilingContext *context); | ||
| 71 | - | ||
| 72 | -} // namespace optiling | ||
| 73 | - | ||
| 74 | - | ||
| @@ -1,306 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2025 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 aclnn_matmul_compress_dequant.cpp | ||
| 13 | - * \brief | ||
| 14 | - */ | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| 24 | - | ||
| 25 | - | ||
| 26 | - | ||
| 27 | - | ||
| 28 | - | ||
| 29 | - | ||
| 30 | - | ||
| 31 | - | ||
| 32 | - | ||
| 33 | -using namespace op; | ||
| 34 | - | ||
| 35 | -extern "C" { | ||
| 36 | - | ||
| 37 | - | ||
| 38 | - | ||
| 39 | - | ||
| 40 | - | ||
| 41 | -static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {DataType::DT_INT8}; | ||
| 42 | -static const std::initializer_list<op::DataType> BIAS_DTYPE_SUPPORT_LIST = {DataType::DT_INT32}; | ||
| 43 | -static const std::initializer_list<op::DataType> DEQ_SCALE_DTYPE_SUPPORT_LIST = {DataType::DT_UINT64}; | ||
| 44 | -static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST = {DataType::DT_FLOAT16}; | ||
| 45 | - | ||
| 46 | -static const std::string ALG_WEIGHT_UNZIP = "weight_unzip"; | ||
| 47 | -static const std::string ALG_UNKNOWN = "unknown"; | ||
| 48 | - | ||
| 49 | -static op::FVector<int64_t> GetShape(const aclTensor *tensor) { | ||
| 50 | - op::FVector<int64_t> shape; | ||
| 51 | - if (tensor == nullptr) { | ||
| 52 | - shape.push_back(1); | ||
| 53 | - OP_LOGW("The input tensor of Func GetShape is nullptr"); | ||
| 54 | - return shape; | ||
| 55 | - } | ||
| 56 | - if (tensor->GetViewShape().GetDimNum() == 0U) { | ||
| 57 | - shape.push_back(1); | ||
| 58 | - } else { | ||
| 59 | - size_t dimNum = tensor->GetViewShape().GetDimNum(); | ||
| 60 | - for (size_t idx = 0U; idx < dimNum; idx++) { | ||
| 61 | - int64_t tmpVal = tensor->GetViewShape().GetDim(idx); | ||
| 62 | - shape.push_back(tmpVal); | ||
| 63 | - } | ||
| 64 | - } | ||
| 65 | - return shape; | ||
| 66 | -} | ||
| 67 | - | ||
| 68 | -enum UnzipMode | ||
| 69 | -{ | ||
| 70 | - UNKNOWN = 0, | ||
| 71 | - WEIGHT_UNZIP = 1 | ||
| 72 | -}; | ||
| 73 | - | ||
| 74 | -inline static const std::string &GetAlgStr(int algMode) { | ||
| 75 | - if (algMode == UnzipMode::WEIGHT_UNZIP) { | ||
| 76 | - return ALG_WEIGHT_UNZIP; | ||
| 77 | - } | ||
| 78 | - return ALG_UNKNOWN; | ||
| 79 | -} | ||
| 80 | - | ||
| 81 | -struct MatmulUnzipInput | ||
| 82 | -{ | ||
| 83 | - const aclTensor *x1; | ||
| 84 | - const aclTensor *x2; | ||
| 85 | - const aclTensor *compressIndex; | ||
| 86 | - const aclTensor *bias; | ||
| 87 | - const aclTensor *deqScale; | ||
| 88 | - const aclTensor *offsetW; | ||
| 89 | -}; | ||
| 90 | - | ||
| 91 | - | ||
| 92 | -inline static bool CheckNotNull(MatmulUnzipInput matmulUnzipInput, const aclIntArray *compressInfo, | ||
| 93 | - const aclTensor* out) | ||
| 94 | -{ | ||
| 95 | - OP_CHECK_NULL(matmulUnzipInput.x1, return false); | ||
| 96 | - OP_CHECK_NULL(matmulUnzipInput.x2, return false); | ||
| 97 | - OP_CHECK_NULL(matmulUnzipInput.compressIndex, return false); | ||
| 98 | - OP_CHECK_NULL(matmulUnzipInput.bias, return false); | ||
| 99 | - OP_CHECK_NULL(matmulUnzipInput.deqScale, return false); | ||
| 100 | - OP_CHECK_NULL(compressInfo, return false); | ||
| 101 | - OP_CHECK_NULL(out, return false); | ||
| 102 | - return true; | ||
| 103 | -} | ||
| 104 | - | ||
| 105 | -inline static bool CheckDtypeValid(MatmulUnzipInput matmulUnzipInput, const aclTensor *out) { | ||
| 106 | - OP_CHECK_DTYPE_NOT_SUPPORT(matmulUnzipInput.x1, DTYPE_SUPPORT_LIST, return false); | ||
| 107 | - OP_CHECK_DTYPE_NOT_SUPPORT(matmulUnzipInput.x2, DTYPE_SUPPORT_LIST, return false); | ||
| 108 | - OP_CHECK_DTYPE_NOT_SUPPORT(matmulUnzipInput.compressIndex, DTYPE_SUPPORT_LIST, return false); | ||
| 109 | - OP_CHECK_DTYPE_NOT_SUPPORT(matmulUnzipInput.bias, BIAS_DTYPE_SUPPORT_LIST, return false); | ||
| 110 | - OP_CHECK_DTYPE_NOT_SUPPORT(matmulUnzipInput.deqScale, DEQ_SCALE_DTYPE_SUPPORT_LIST, return false); | ||
| 111 | - OP_CHECK_DTYPE_NOT_SUPPORT(out, OUT_DTYPE_SUPPORT_LIST, return false); | ||
| 112 | - return true; | ||
| 113 | -} | ||
| 114 | - | ||
| 115 | -static bool CheckShapeValid(const aclTensor* x1, const aclTensor* x2, const aclIntArray *compressInfo) | ||
| 116 | -{ | ||
| 117 | - op::Shape x1Shape = x1->GetViewShape(); | ||
| 118 | - op::Shape x2Shape = x2->GetViewShape(); | ||
| 119 | - auto dimTensor1 = x1Shape.GetDimNum(); | ||
| 120 | - auto dimTensor2 = x2Shape.GetDimNum(); | ||
| 121 | - int64_t x1KDim = 0; | ||
| 122 | - int64_t x2KDim = 0; | ||
| 123 | - | ||
| 124 | - if (dimTensor1 != 2 || dimTensor2 != 1) { // ND format dims > 2 for x1 | ||
| 125 | - OP_LOGE(ACLNN_ERR_PARAM_INVALID, "MatmulUnzip not support x1 shape %s, x2 shape %s", | ||
| 126 | - op::ToString(x1Shape).GetString(), | ||
| 127 | - op::ToString(x2Shape).GetString()); | ||
| 128 | - return false; | ||
| 129 | - } else { | ||
| 130 | - x1KDim = x1Shape.GetDim(dimTensor1 -1); | ||
| 131 | - x2KDim = (*compressInfo)[K_DIMENSION_INDEX]; | ||
| 132 | - if (x1KDim != x2KDim) { | ||
| 133 | - OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The k-axis of the two inputs are different %s, %s", | ||
| 134 | - op::ToString(x1Shape).GetString(), op::ToString(x2Shape).GetString()); | ||
| 135 | - return false; | ||
| 136 | - } | ||
| 137 | - } | ||
| 138 | - | ||
| 139 | - return true; | ||
| 140 | -} | ||
| 141 | - | ||
| 142 | -inline static aclnnStatus CheckParam(MatmulUnzipInput matmulUnzipInput, const aclIntArray *compressInfo, | ||
| 143 | - const aclTensor *out) { | ||
| 144 | - // 1. 检查参数是否为空指针 | ||
| 145 | - CHECK_RET(CheckNotNull(matmulUnzipInput, compressInfo, out), ACLNN_ERR_PARAM_NULLPTR); | ||
| 146 | - // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验 | ||
| 147 | - CHECK_RET(CheckDtypeValid(matmulUnzipInput, out), ACLNN_ERR_PARAM_INVALID); | ||
| 148 | - // 3. 检查Shape是否支持 | ||
| 149 | - CHECK_RET(CheckShapeValid(matmulUnzipInput.x1, matmulUnzipInput.x2, compressInfo), | ||
| 150 | - ACLNN_ERR_PARAM_INVALID); | ||
| 151 | - | ||
| 152 | - return ACLNN_SUCCESS; | ||
| 153 | -} | ||
| 154 | - | ||
| 155 | -static const aclTensor* ProcessEmptyTensor(const aclTensor* x1, const aclTensor* out, | ||
| 156 | - aclOpExecutor* executor) | ||
| 157 | -{ | ||
| 158 | - // 获取shape信息 | ||
| 159 | - op::Shape outShape = out->GetViewShape(); | ||
| 160 | - auto output = executor->AllocTensor(outShape, x1->GetDataType()); | ||
| 161 | - if (output->IsEmpty()) { | ||
| 162 | - OP_LOGI("Returning an empty tensor without actually doing calculation"); | ||
| 163 | - return output; | ||
| 164 | - } | ||
| 165 | - FVector<int64_t> fillShape = GetShape(output); | ||
| 166 | - const aclTensor* dims = executor->ConvertToTensor(fillShape.data(), fillShape.size(), op::DataType::DT_INT64); | ||
| 167 | - aclIntArray* shapeArray = executor->AllocIntArray(fillShape.data(), fillShape.size()); | ||
| 168 | - const aclScalar* valueScalar = executor->AllocScalar(0); | ||
| 169 | - const aclTensor* valueTensor = executor->ConvertToTensor(valueScalar, out->GetDataType()); | ||
| 170 | - auto fillTensor = l0op::Fill(dims, valueTensor, shapeArray, executor); | ||
| 171 | - return fillTensor; | ||
| 172 | -} | ||
| 173 | - | ||
| 174 | -inline static const aclTensor *TensorReformat(const aclTensor *x, const op::Format& format, aclOpExecutor *executor) { | ||
| 175 | - auto formatTensor = executor == nullptr ? const_cast<aclTensor *>(x) | ||
| 176 | - : executor->CreateView(x, x->GetViewShape(), x->GetViewOffset()); | ||
| 177 | - formatTensor->SetViewFormat(format); | ||
| 178 | - formatTensor->SetOriginalFormat(format); | ||
| 179 | - formatTensor->SetStorageFormat(format); | ||
| 180 | - return formatTensor; | ||
| 181 | -} | ||
| 182 | - | ||
| 183 | -static aclnnStatus InputsContiguousAndTransFormat(const aclTensor *tensor, const aclTensor *&reformatedTensor, | ||
| 184 | - const std::string &tensorName, aclOpExecutor *executor) | ||
| 185 | -{ | ||
| 186 | - if (tensor == nullptr) { | ||
| 187 | - return ACLNN_SUCCESS; | ||
| 188 | - } | ||
| 189 | - op::Format tensorFormat = tensor->GetStorageFormat(); | ||
| 190 | - if (tensorFormat != Format::FORMAT_FRACTAL_NZ) { | ||
| 191 | - reformatedTensor = l0op::Contiguous(tensor, executor); | ||
| 192 | - CHECK_COND(reformatedTensor != nullptr, ACLNN_ERR_INNER_NULLPTR, "%s Contiguous failed.", tensorName.c_str()); | ||
| 193 | - | ||
| 194 | - bool is310P = GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND310P; | ||
| 195 | - // only in 310P, tensor should be transdata to NZ format | ||
| 196 | - if (!is310P) { | ||
| 197 | - return ACLNN_SUCCESS; | ||
| 198 | - } | ||
| 199 | - | ||
| 200 | - reformatedTensor = l0op::TransData(reformatedTensor, Format::FORMAT_FRACTAL_NZ, 1, executor); | ||
| 201 | - CHECK_COND(reformatedTensor != nullptr, ACLNN_ERR_INNER_NULLPTR, "%s TransData failed.", tensorName.c_str()); | ||
| 202 | - return ACLNN_SUCCESS; | ||
| 203 | - } | ||
| 204 | - reformatedTensor = tensor; | ||
| 205 | - return ACLNN_SUCCESS; | ||
| 206 | -} | ||
| 207 | - | ||
| 208 | -static const aclTensor *BuildMatMulUnzipGraph(MatmulUnzipInput matmulUnzipInput, const int offsetX, | ||
| 209 | - const aclIntArray *compressInfo, aclTensor *out, | ||
| 210 | - aclOpExecutor *executor) { | ||
| 211 | - /* | ||
| 212 | - * x1 x2 | ||
| 213 | - * | | | ||
| 214 | - * x1FractalNZ x2ReformatFractalZ | ||
| 215 | - * | | | ||
| 216 | - * \ / | ||
| 217 | - * matmulv2_compress_unzip -- compressIndex, bias, deqScale(format?), offsetW, offsetX, compressInfo | ||
| 218 | - * | | ||
| 219 | - * out | ||
| 220 | - * | | ||
| 221 | - * outND | ||
| 222 | - */ | ||
| 223 | - | ||
| 224 | - // 空tensor 处理 | ||
| 225 | - if (matmulUnzipInput.x1->IsEmpty() || matmulUnzipInput.x2->IsEmpty()) { | ||
| 226 | - auto emptyOut = ProcessEmptyTensor(matmulUnzipInput.x1, out, executor); | ||
| 227 | - CHECK_RET(emptyOut != nullptr, nullptr); | ||
| 228 | - return emptyOut; | ||
| 229 | - } | ||
| 230 | - const aclTensor *reformatedX = nullptr; | ||
| 231 | - (void)InputsContiguousAndTransFormat(matmulUnzipInput.x1, reformatedX, "x1", executor); | ||
| 232 | - matmulUnzipInput.bias = l0op::Contiguous(matmulUnzipInput.bias, executor); | ||
| 233 | - CHECK_RET(matmulUnzipInput.bias != nullptr, nullptr); | ||
| 234 | - const aclTensor *deqScale5HD = matmulUnzipInput.deqScale; | ||
| 235 | - if (matmulUnzipInput.deqScale->Numel() % DEQUANT_SCALE_ALIGN_SIZE == 0) { | ||
| 236 | - deqScale5HD = TensorReformat(matmulUnzipInput.deqScale, op::Format::FORMAT_NC1HWC0, executor); | ||
| 237 | - } else { | ||
| 238 | - OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Dequant Scale is invalid Data."); | ||
| 239 | - return ProcessEmptyTensor(matmulUnzipInput.x1, out, executor); | ||
| 240 | - } | ||
| 241 | - const aclTensor *x2ReFormatFractalZ = TensorReformat(matmulUnzipInput.x2, op::Format::FORMAT_FRACTAL_Z, executor); | ||
| 242 | - const aclTensor *matmulOut = l0op::MatMulCompressDequant(reformatedX, x2ReFormatFractalZ, | ||
| 243 | - matmulUnzipInput.compressIndex, deqScale5HD, | ||
| 244 | - matmulUnzipInput.bias, nullptr, false, false, compressInfo, | ||
| 245 | - offsetX, GetAlgStr(UnzipMode::WEIGHT_UNZIP), executor); | ||
| 246 | - CHECK_RET(matmulOut != nullptr, nullptr); | ||
| 247 | - // TransData out format from NZ to ND | ||
| 248 | - auto matmulOutND = l0op::TransData(matmulOut, op::Format::FORMAT_ND, 1, executor); | ||
| 249 | - CHECK_RET(matmulOutND != nullptr, nullptr); | ||
| 250 | - | ||
| 251 | - return matmulOutND; | ||
| 252 | -} | ||
| 253 | - | ||
| 254 | -aclnnStatus aclnnMatmulCompressDequantGetWorkspaceSize(const aclTensor *x1, const aclTensor *x2, | ||
| 255 | - const aclTensor *compressIndex, const aclTensor *bias, | ||
| 256 | - const aclTensor *deqScale, const aclTensor *offsetW, | ||
| 257 | - int offsetX, const aclIntArray *compressInfo, | ||
| 258 | - aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor) { | ||
| 259 | - L2_DFX_PHASE_1(aclnnMatmulCompressDequant, | ||
| 260 | - DFX_IN(x1, x2, compressIndex, bias, deqScale, offsetX, compressInfo), | ||
| 261 | - DFX_OUT(out)); | ||
| 262 | - // 固定写法，创建OpExecutor | ||
| 263 | - auto uniqueExecutor = CREATE_EXECUTOR(); | ||
| 264 | - CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR); | ||
| 265 | - // 入参初始化 | ||
| 266 | - MatmulUnzipInput matmulUnzipInput; | ||
| 267 | - matmulUnzipInput.x1 = x1; | ||
| 268 | - matmulUnzipInput.x2 = x2; | ||
| 269 | - matmulUnzipInput.compressIndex = compressIndex; | ||
| 270 | - matmulUnzipInput.bias = bias; | ||
| 271 | - matmulUnzipInput.deqScale = deqScale; | ||
| 272 | - matmulUnzipInput.offsetW = offsetW; | ||
| 273 | - | ||
| 274 | - // 入参检查 | ||
| 275 | - auto ret = CheckParam(matmulUnzipInput, compressInfo, out); | ||
| 276 | - CHECK_RET(ret == ACLNN_SUCCESS, ret); | ||
| 277 | - | ||
| 278 | - // 构建matmul_unzip计算图 | ||
| 279 | - auto matmulOut = BuildMatMulUnzipGraph(matmulUnzipInput, offsetX, compressInfo, out, uniqueExecutor.get()); | ||
| 280 | - CHECK_RET(matmulOut != nullptr, ACLNN_ERR_INNER_NULLPTR); | ||
| 281 | - if (matmulOut->IsEmpty()) { | ||
| 282 | - // 当输出为空tensor的场景，空tensor处理 | ||
| 283 | - *workspaceSize = 0; | ||
| 284 | - uniqueExecutor.ReleaseTo(executor); | ||
| 285 | - return ACLNN_SUCCESS; | ||
| 286 | - } | ||
| 287 | - | ||
| 288 | - auto viewCopyResult = l0op::ViewCopy(matmulOut, out, uniqueExecutor.get()); | ||
| 289 | - CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR); | ||
| 290 | - | ||
| 291 | - // 获取workspace | ||
| 292 | - *workspaceSize = uniqueExecutor->GetWorkspaceSize(); | ||
| 293 | - uniqueExecutor.ReleaseTo(executor); | ||
| 294 | - return ACLNN_SUCCESS; | ||
| 295 | -} | ||
| 296 | - | ||
| 297 | -aclnnStatus aclnnMatmulCompressDequant(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, | ||
| 298 | - aclrtStream stream) { | ||
| 299 | - L2_DFX_PHASE_2(aclnnMatmulCompressDequant); | ||
| 300 | - // 固定写法，调用框架能力，完成计算 | ||
| 301 | - return CommonOpExecutorRun(workspace, workspaceSize, executor, stream); | ||
| 302 | -} | ||
| 303 | - | ||
| 304 | - | ||
| 305 | -} | ||
| 306 | - | ||
| @@ -1,41 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2025 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 | -extern "C" { | ||
| 18 | - | ||
| 19 | - | ||
| 20 | -/** | ||
| 21 | - * @brief aclnnMatmulCompressDequant的第一段接口，根据具体的计算流程，计算workspace大小。 | ||
| 22 | - * @domain aclnn_ops_infer | ||
| 23 | - */ | ||
| 24 | -ACLNN_API aclnnStatus aclnnMatmulCompressDequantGetWorkspaceSize(const aclTensor* x1, const aclTensor* x2, | ||
| 25 | - const aclTensor* compressIndex, const aclTensor* bias, | ||
| 26 | - const aclTensor* deqScale, const aclTensor* offsetW, | ||
| 27 | - int offsetX, const aclIntArray* compressInfo, | ||
| 28 | - aclTensor* out, uint64_t* workspaceSize, | ||
| 29 | - aclOpExecutor** executor); | ||
| 30 | - | ||
| 31 | -/** | ||
| 32 | - * @brief aclnnMatmulCompressDequant的第二段接口，用于执行计算。 | ||
| 33 | - */ | ||
| 34 | -ACLNN_API aclnnStatus aclnnMatmulCompressDequant(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, | ||
| 35 | - aclrtStream stream); | ||
| 36 | - | ||
| 37 | - | ||
| 38 | -} | ||
| 39 | - | ||
| 40 | - | ||
| 41 | - | ||
| @@ -1,44 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2025 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 | -using namespace op; | ||
| 20 | - | ||
| 21 | -namespace l0op { | ||
| 22 | - | ||
| 23 | -OP_TYPE_REGISTER(MatMulV2CompressDequant); | ||
| 24 | -// 用大小写区分是否可以区分两个不同的kernel | ||
| 25 | -const aclTensor *MatMulCompressDequant(const aclTensor *x1, const aclTensor *x2, const aclTensor *compressIndex, | ||
| 26 | - const aclTensor *deqScale, const aclTensor *bias, const aclTensor *offsetW, | ||
| 27 | - const bool transposeX1, const bool transposeX2, const aclIntArray *compressInfo, | ||
| 28 | - const int offsetX, const std::string &algStr, aclOpExecutor *executor) { | ||
| 29 | - L0_DFX(MatMulCompressDequant, x1, x2, compressIndex, deqScale, bias, offsetW, transposeX1, transposeX2, | ||
| 30 | - compressInfo, offsetX, algStr); | ||
| 31 | - auto mmCompressDequantOut = executor->AllocTensor(DataType::DT_FLOAT16, Format::FORMAT_FRACTAL_NZ, Format::FORMAT_ND); | ||
| 32 | - // 是否可以复用tbe的infershape | ||
| 33 | - auto ret = INFER_SHAPE(MatMulV2CompressDequant, OP_INPUT(x1, x2, compressIndex, deqScale, bias, offsetW), | ||
| 34 | - OP_OUTPUT(mmCompressDequantOut), OP_ATTR(transposeX1, transposeX2, compressInfo, offsetX, algStr)); | ||
| 35 | - OP_CHECK_INFERSHAPE(ret != ACLNN_SUCCESS, return nullptr, "MatMulCompressDequant InferShape failed."); | ||
| 36 | - // 使用新的MatmulV2CompressDequant内核 | ||
| 37 | - ret = ADD_TO_LAUNCHER_LIST_AICORE(MatMulV2CompressDequant, OP_INPUT(x1, x2, compressIndex, deqScale, bias, offsetW), | ||
| 38 | - OP_OUTPUT(mmCompressDequantOut), | ||
| 39 | - OP_ATTR(transposeX1, transposeX2, compressInfo, offsetX, algStr)); | ||
| 40 | - OP_CHECK_ADD_TO_LAUNCHER_LIST_AICORE(ret != ACLNN_SUCCESS, return nullptr, | ||
| 41 | - "MatMulCompressDequant ADD_TO_LAUNCHER_LIST_AICORE failed."); | ||
| 42 | - return mmCompressDequantOut; | ||
| 43 | -}; | ||
| 44 | -} // namespace l0op | ||
| @@ -1,25 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2025 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 | -namespace l0op { | ||
| 16 | - | ||
| 17 | -const aclTensor *MatMulCompressDequant(const aclTensor *x1, const aclTensor *x2, const aclTensor *compressIndex, | ||
| 18 | - const aclTensor *deqScale, const aclTensor *bias, const aclTensor *offsetW, | ||
| 19 | - const bool transposeX1, const bool transposeX2, const aclIntArray *compressInfo, | ||
| 20 | - const int offsetX, const std::string &algStr, aclOpExecutor *executor); | ||
| 21 | - | ||
| 22 | -} // namespace l0op | ||
| 23 | - | ||
| 24 | - | ||
| 25 | - | ||
| @@ -1,243 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2025 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 pp_matmul_common_tiling.h | ||
| 13 | - * \brief | ||
| 14 | - */ | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| 24 | -namespace optiling { | ||
| 25 | -namespace pp_matmul { | ||
| 26 | - | ||
| 27 | -constexpr uint64_t FP16_SIZE = 2; | ||
| 28 | -constexpr uint64_t FP32_SIZE = 4; | ||
| 29 | -constexpr uint64_t BLOCK_SIZE = 16; | ||
| 30 | -constexpr uint64_t BLOCK_SIZE_INT8_K = 32; | ||
| 31 | -constexpr uint64_t BASE_BLOCK_STEP = 2; | ||
| 32 | -constexpr uint64_t AXES_ALIGN_SIZE = 512; | ||
| 33 | -constexpr uint64_t AXES_ALIGN_SIZE_INT8 = 256; | ||
| 34 | -constexpr uint64_t ND_SHAPE_SIZE = 2; | ||
| 35 | -constexpr uint64_t NZ_SHAPE_SIZE = 4; | ||
| 36 | -constexpr uint64_t CUBE_BLOCK_SIZE = 256; | ||
| 37 | -constexpr uint64_t CUBE_BLOCK_SIZE_INT8 = 512; | ||
| 38 | -constexpr uint64_t L1AB_PINGPONG_BUFFER_SIZE = 262144; | ||
| 39 | -constexpr uint64_t L0AB_PINGPONG_BUFFER_SIZE_INT8 = 262144; // 131072 * 2 = 256 KB | ||
| 40 | -constexpr uint64_t L0AB_PINGPONG_BUFFER_SIZE_FP16 = 131072; // 128 KB | ||
| 41 | -constexpr uint64_t L1AB_PINGPONG_BUFFER_SIZE_INT8_SPARSE = 163840; // 160 * 1024 | ||
| 42 | -constexpr uint64_t UB_LIMIT_SIZE_910A = 131072; // 128 * 1024 | ||
| 43 | - | ||
| 44 | -template <uint64_t DIV> | ||
| 45 | -inline uint64_t CeilDiv(uint64_t num) | ||
| 46 | -{ | ||
| 47 | - if (DIV == 0UL || num + DIV - 1UL < num) { | ||
| 48 | - return num; | ||
| 49 | - } | ||
| 50 | - return (num + DIV - 1UL) / DIV; | ||
| 51 | -} | ||
| 52 | - | ||
| 53 | -inline uint64_t CeilDiv(uint64_t dividend, uint64_t divisor) | ||
| 54 | -{ | ||
| 55 | - if (divisor == 0UL || dividend + divisor - 1UL < dividend) { | ||
| 56 | - return dividend; | ||
| 57 | - } | ||
| 58 | - return (dividend + divisor - 1UL) / divisor; | ||
| 59 | -} | ||
| 60 | - | ||
| 61 | -template <uint64_t RND> | ||
| 62 | -inline uint64_t Round(uint64_t num) | ||
| 63 | -{ | ||
| 64 | - if (RND == 0UL || num + RND - 1UL < num) { | ||
| 65 | - return num; | ||
| 66 | - } | ||
| 67 | - return (num + RND - 1UL) / RND * RND; | ||
| 68 | -} | ||
| 69 | - | ||
| 70 | -inline uint64_t RoundUp(uint64_t num, uint64_t rnd) | ||
| 71 | -{ | ||
| 72 | - if (rnd == 0UL || num + rnd - 1UL < num) { | ||
| 73 | - return num; | ||
| 74 | - } | ||
| 75 | - return (num + rnd - 1UL) / rnd * rnd; | ||
| 76 | -} | ||
| 77 | - | ||
| 78 | -inline int64_t RoundUp(int64_t num, int64_t rnd) | ||
| 79 | -{ | ||
| 80 | - if (rnd == 0UL || static_cast<int64_t>(num + rnd - 1) < num) { | ||
| 81 | - return num; | ||
| 82 | - } | ||
| 83 | - return static_cast<int64_t>((num + rnd - 1) / rnd * rnd); | ||
| 84 | -} | ||
| 85 | - | ||
| 86 | -inline uint64_t RoundDown(uint64_t num, uint64_t rnd) | ||
| 87 | -{ | ||
| 88 | - if (rnd == 0UL) { | ||
| 89 | - return 0UL; | ||
| 90 | - } | ||
| 91 | - return num / rnd * rnd; | ||
| 92 | -} | ||
| 93 | - | ||
| 94 | -inline uint64_t GetN0TilingLimit(bool compressFlag, uint64_t tilingN, const platform_ascendc::SocVersion &platformType) | ||
| 95 | -{ | ||
| 96 | - if (compressFlag) { | ||
| 97 | - return std::min(tilingN * BLOCK_SIZE, AXES_ALIGN_SIZE_INT8); | ||
| 98 | - } else { | ||
| 99 | - return (platformType == platform_ascendc::SocVersion::ASCEND310P || platformType == platform_ascendc::SocVersion::ASCEND910) | ||
| 100 | - ? AXES_ALIGN_SIZE | ||
| 101 | - : AXES_ALIGN_SIZE_INT8; | ||
| 102 | - } | ||
| 103 | -} | ||
| 104 | - | ||
| 105 | -template <typename OpShareType> | ||
| 106 | -inline uint64_t GetN0TilingInit(const OpShareType &opShape, bool compressFlag, | ||
| 107 | - uint64_t tilingN) | ||
| 108 | -{ | ||
| 109 | - const uint64_t RND = 16UL; | ||
| 110 | - if (compressFlag) { | ||
| 111 | - if (tilingN * BLOCK_SIZE > opShape.n) { | ||
| 112 | - return Round<RND>(opShape.n); | ||
| 113 | - } else { | ||
| 114 | - return tilingN * BLOCK_SIZE; | ||
| 115 | - } | ||
| 116 | - } else { | ||
| 117 | - return BLOCK_SIZE; | ||
| 118 | - } | ||
| 119 | -} | ||
| 120 | - | ||
| 121 | -template <bool PRI_FLAG> | ||
| 122 | -inline bool IsExceedTilingLimit(uint64_t axes0, uint64_t priAxes0, | ||
| 123 | - uint64_t n0TilingLimit, platform_ascendc::SocVersion platformType, | ||
| 124 | - uint64_t basicBlockSize) | ||
| 125 | -{ | ||
| 126 | - return (PRI_FLAG && axes0 > n0TilingLimit) || (!PRI_FLAG && priAxes0 > n0TilingLimit) || | ||
| 127 | - (platformType == platform_ascendc::SocVersion::ASCEND910 && basicBlockSize > UB_LIMIT_SIZE_910A); | ||
| 128 | -} | ||
| 129 | - | ||
| 130 | -template <bool PRI_FLAG, typename OpShareType> | ||
| 131 | -inline void SetOpShapeAxesInfo(OpShareType &opShape, uint64_t priAxes0, uint64_t axes0) | ||
| 132 | -{ | ||
| 133 | - opShape.m0 = PRI_FLAG ? priAxes0 : axes0; | ||
| 134 | - opShape.n0 = PRI_FLAG ? axes0 : priAxes0; | ||
| 135 | -} | ||
| 136 | - | ||
| 137 | -template <typename HardwareType, typename OpShapeType, typename MatMulInfoType> | ||
| 138 | -inline float CostFunc(const HardwareType &hwInfor, OpShapeType &shape, const MatMulInfoType &mmInfo) | ||
| 139 | -{ | ||
| 140 | - float aCoef = 1; | ||
| 141 | - float bCoef = 1; | ||
| 142 | - float bwCoef = static_cast<float>(hwInfor.l2BandWidth) / static_cast<float>(hwInfor.hbmBandWidth); | ||
| 143 | - uint64_t mLoop = CeilDiv(shape.m, shape.m0); | ||
| 144 | - uint64_t nLoop = CeilDiv(shape.n, shape.n0); | ||
| 145 | - if (mLoop == 0UL || nLoop == 0UL) { | ||
| 146 | - return 1; | ||
| 147 | - } | ||
| 148 | - uint64_t coreNeed = shape.batchSize * mLoop * nLoop; | ||
| 149 | - uint64_t blockDim = std::min(coreNeed, hwInfor.coreNum); | ||
| 150 | - uint64_t mOnce = blockDim < nLoop ? shape.m0 : blockDim / nLoop * shape.m0; | ||
| 151 | - uint64_t nOnce = blockDim < nLoop ? hwInfor.coreNum * shape.n0 : shape.n; | ||
| 152 | - (void)mmInfo; | ||
| 153 | - if (mOnce * shape.k * 2 > hwInfor.l2Size) { | ||
| 154 | - aCoef = bwCoef; | ||
| 155 | - } | ||
| 156 | - if (nOnce * shape.k * 2 > hwInfor.l2Size) { | ||
| 157 | - bCoef = bwCoef; | ||
| 158 | - } | ||
| 159 | - return 1 / (aCoef * static_cast<float>(shape.n0)) + 1 / (bCoef * static_cast<float>(shape.m0)); | ||
| 160 | -} | ||
| 161 | - | ||
| 162 | -// OpShareType is OpShape, TilingType is PpTilingData, HardwareType is HardwareType, MatMulInfoType is MatMulInfo | ||
| 163 | -template <bool PRI_FLAG, typename OpShareType, typename TilingType, typename HardwareType, typename MatMulInfoType> | ||
| 164 | -void TilingFunc(OpShareType &opShape, TilingType &tilingParam, const HardwareType &hwInfor, | ||
| 165 | - const MatMulInfoType &mmInfo, bool compressFlag = false, const uint64_t tilingN = 1) | ||
| 166 | -{ | ||
| 167 | - float costMin = 1; | ||
| 168 | - const float CONST_2 = 2.0; | ||
| 169 | - const uint64_t CONST_16 = 16UL; | ||
| 170 | - uint64_t roundBase = | ||
| 171 | - static_cast<uint64_t>(pow(2, ceil(log(CeilDiv(PRI_FLAG ? opShape.n : opShape.m, CONST_16)))) * CONST_16); | ||
| 172 | - uint64_t priAxes = RoundUp(PRI_FLAG ? opShape.m : opShape.n, CONST_16); | ||
| 173 | - uint64_t axes = RoundUp(PRI_FLAG ? opShape.n : opShape.m, roundBase); | ||
| 174 | - float axes0Max = static_cast<float>(AXES_ALIGN_SIZE) / mmInfo.inDtype; | ||
| 175 | - auto platformType = hwInfor.socVersion; | ||
| 176 | - if (mmInfo.isInt8 && (platformType == platform_ascendc::SocVersion::ASCEND310P || platformType == platform_ascendc::SocVersion::ASCEND910)) { | ||
| 177 | - axes0Max /= CONST_2; | ||
| 178 | - } | ||
| 179 | - uint64_t n0TilingInit = GetN0TilingInit(opShape, compressFlag, tilingN); | ||
| 180 | - uint64_t n0TilingLimit = GetN0TilingLimit(compressFlag, tilingN, platformType); | ||
| 181 | - uint64_t priAxes0Init = PRI_FLAG ? BLOCK_SIZE : n0TilingInit; | ||
| 182 | - uint64_t axes0Init = PRI_FLAG ? n0TilingInit : BLOCK_SIZE; | ||
| 183 | - bool isAscend310P = platformType == platform_ascendc::SocVersion::ASCEND310P; | ||
| 184 | - for (uint64_t priAxes0 = priAxes0Init; priAxes0 <= priAxes && priAxes0 <= axes0Max; priAxes0 *= BASE_BLOCK_STEP) { | ||
| 185 | - for (uint64_t axes0 = axes0Init; axes0 <= axes && axes0 <= axes0Max; axes0 *= BASE_BLOCK_STEP) { | ||
| 186 | - uint64_t basicBlockSize = priAxes0 * axes0 * FP32_SIZE; | ||
| 187 | - if (basicBlockSize > hwInfor.l0cSize) { | ||
| 188 | - continue; | ||
| 189 | - } | ||
| 190 | - if (mmInfo.isInt8 && | ||
| 191 | - IsExceedTilingLimit<PRI_FLAG>(axes0, priAxes0, n0TilingLimit, platformType, basicBlockSize)) { | ||
| 192 | - continue; | ||
| 193 | - } | ||
| 194 | - SetOpShapeAxesInfo<PRI_FLAG>(opShape, priAxes0, axes0); | ||
| 195 | - float cost = CostFunc<HardwareType, OpShareType, MatMulInfoType>(hwInfor, opShape, mmInfo); | ||
| 196 | - if (cost >= costMin) { | ||
| 197 | - continue; | ||
| 198 | - } | ||
| 199 | - costMin = cost; | ||
| 200 | - tilingParam.SetBaseOp(hwInfor.coreNum, hwInfor.l0cSize, opShape.m0, opShape.n0, mmInfo, isAscend310P); | ||
| 201 | - } | ||
| 202 | - } | ||
| 203 | -} | ||
| 204 | - | ||
| 205 | -template <typename PpTilingDataType> | ||
| 206 | -uint64_t Swizzl(PpTilingDataType &tilingData) | ||
| 207 | -{ | ||
| 208 | - uint64_t swizzlDirect = 0UL; | ||
| 209 | - uint64_t swizzlCount = 1UL; | ||
| 210 | - float m0 = tilingData.opShape.m0; | ||
| 211 | - float n0 = tilingData.opShape.n0; | ||
| 212 | - float m = tilingData.opShape.m; | ||
| 213 | - float k = tilingData.opShape.k; | ||
| 214 | - float n = tilingData.opShape.n; | ||
| 215 | - float mincost = m * k + k * n; | ||
| 216 | - | ||
| 217 | - for (uint32_t i = 1; i <= tilingData.blockDim; ++i) { | ||
| 218 | - int c = static_cast<int32_t>((tilingData.blockDim + i - 1) / i); | ||
| 219 | - float cost; | ||
| 220 | - // B0 + A < A0 + B | ||
| 221 | - if (i * n0 + m < m0 * c + n) { | ||
| 222 | - swizzlDirect = 1UL; // Nz | ||
| 223 | - cost = n0 * i + m0 * c; | ||
| 224 | - if (cost <= mincost) { | ||
| 225 | - mincost = cost; | ||
| 226 | - swizzlCount = i; | ||
| 227 | - } | ||
| 228 | - } else { | ||
| 229 | - swizzlDirect = 0UL; // Zn | ||
| 230 | - cost = m0 * i + n0 * c; | ||
| 231 | - if (cost < mincost) { | ||
| 232 | - mincost = cost; | ||
| 233 | - swizzlCount = i; | ||
| 234 | - } | ||
| 235 | - } | ||
| 236 | - } | ||
| 237 | - tilingData.swizzlDirect = swizzlDirect; | ||
| 238 | - tilingData.swizzlCount = swizzlCount; | ||
| 239 | - return swizzlDirect; | ||
| 240 | -} | ||
| 241 | -} // namespace pp_matmul | ||
| 242 | -} // namespace optiling | ||
| 243 | - | ||
| @@ -1,184 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2025 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 pp_matmul_default.h | ||
| 13 | - * \brief | ||
| 14 | - */ | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| 24 | - | ||
| 25 | - | ||
| 26 | - | ||
| 27 | -namespace { | ||
| 28 | -constexpr uint64_t L1_DESCALE_BUFFER_SIZE_MAX = 6144; | ||
| 29 | -constexpr uint64_t CONST_3 = 3; | ||
| 30 | -constexpr uint64_t CONST_4 = 4; | ||
| 31 | -constexpr uint64_t CONST_16 = 16; | ||
| 32 | -constexpr uint64_t CONST_32 = 32; | ||
| 33 | -constexpr uint64_t CONST_256 = 256; | ||
| 34 | -constexpr uint64_t CONST_512 = 512; | ||
| 35 | - | ||
| 36 | -constexpr size_t DIM_2 = 2; | ||
| 37 | -constexpr size_t DIM_3 = 3; | ||
| 38 | -constexpr size_t DIM_4 = 4; | ||
| 39 | -constexpr size_t COMPUTE_TYPE_IDX_CANN = 3; | ||
| 40 | -constexpr size_t COMPUTE_TYPE_IDX_ATB = 10; | ||
| 41 | -constexpr size_t EINSUM_MODE = 4; | ||
| 42 | - | ||
| 43 | -constexpr uint32_t DTYPE_BIT_COUNT = 2; | ||
| 44 | -constexpr uint32_t FORMAT_BIT_COUNT = 1; | ||
| 45 | -constexpr uint32_t COMPUTE_TYPE_BIT_COUNT = 3; | ||
| 46 | -constexpr uint32_t MAX_ATTRS_NUM = 4; | ||
| 47 | -constexpr uint32_t EN_SHUFFFLE_IDX_ATB = 7; | ||
| 48 | -} | ||
| 49 | - | ||
| 50 | -namespace optiling { | ||
| 51 | -namespace pp_matmul { | ||
| 52 | - | ||
| 53 | -class PpMatMulDefault{ | ||
| 54 | -public: | ||
| 55 | - explicit PpMatMulDefault(gert::TilingContext* context) : context_(context) {} | ||
| 56 | - virtual ~PpMatMulDefault() = default; | ||
| 57 | - | ||
| 58 | - void GetHardwareInfo(); | ||
| 59 | - bool GetMatMulTilingData(); | ||
| 60 | - void PrintTiling(); | ||
| 61 | - gert::TilingContext *context_ = nullptr; | ||
| 62 | - MatMulInfo matMulInfo_; | ||
| 63 | - PpMatmulDefaultTilingData ppMatmulDefaultTilingData_{}; | ||
| 64 | - HardwareInfo hardwareInfo_; | ||
| 65 | - uint64_t kernelKey_; | ||
| 66 | -}; | ||
| 67 | - | ||
| 68 | - | ||
| 69 | -void PpMatmulDefaultTilingData::SetBaseShape(uint64_t batchSize, uint64_t m, uint64_t k, uint64_t n) { | ||
| 70 | - opShape.batchSize = batchSize; | ||
| 71 | - opShape.m = m; | ||
| 72 | - opShape.k = k; | ||
| 73 | - opShape.n = n; | ||
| 74 | -} | ||
| 75 | - | ||
| 76 | -void PpMatmulDefaultTilingData::SetBaseOp(uint64_t coreNum, uint64_t l0cSize, uint64_t mBase, uint64_t nBase, const MatMulInfo &mmInfo, bool isAscend310P) { | ||
| 77 | - opShape.m0 = mBase; | ||
| 78 | - opShape.n0 = nBase; | ||
| 79 | - mLoop = CeilDiv(opShape.m, opShape.m0); | ||
| 80 | - nLoop = CeilDiv(opShape.n, opShape.n0); | ||
| 81 | - coreLoop = opShape.batchSize * mLoop * nLoop; | ||
| 82 | - if (!isAscend310P && mLoop == 1UL && mmInfo.transB && static_cast<uint64_t>(coreLoop % coreNum) < | ||
| 83 | - static_cast<uint64_t>(coreNum / CONST_4) * CONST_3) { | ||
| 84 | - mBase = RoundUp(opShape.m, CONST_16); | ||
| 85 | - opShape.m0 = mBase; | ||
| 86 | - uint64_t maxN0 = l0cSize / (mBase * sizeof(float)); | ||
| 87 | - uint64_t x = CeilDiv(opShape.n, coreNum); | ||
| 88 | - uint64_t y = CeilDiv(x, maxN0); | ||
| 89 | - nBase = RoundUp(CeilDiv(x, y), CONST_16); | ||
| 90 | - uint64_t rqdL0CSize = mBase * nBase * sizeof(float); | ||
| 91 | - if (rqdL0CSize < l0cSize && | ||
| 92 | - (mBase + nBase) * CONST_256 * sizeof(uint16_t) < L1AB_PINGPONG_BUFFER_SIZE) { | ||
| 93 | - opShape.n0 = nBase; | ||
| 94 | - nLoop = CeilDiv(opShape.n, opShape.n0); | ||
| 95 | - coreLoop = opShape.batchSize * nLoop; | ||
| 96 | - } | ||
| 97 | - } | ||
| 98 | - blockDim = std::min(coreLoop, coreNum); | ||
| 99 | -} | ||
| 100 | - | ||
| 101 | -void PpMatmulDefaultTilingData::End(const MatMulInfo &mmInfo, bool isAscend310P) { | ||
| 102 | - uint64_t shapeSum = opShape.m0 + opShape.n0; | ||
| 103 | - if (!isAscend310P) { | ||
| 104 | - uint64_t k0Max = shapeSum == 0UL | ||
| 105 | - ? L1AB_PINGPONG_BUFFER_SIZE | ||
| 106 | - : static_cast<uint64_t>(static_cast<float>(L1AB_PINGPONG_BUFFER_SIZE) | ||
| 107 | - / (shapeSum * mmInfo.inDtype)); | ||
| 108 | - opShape.k0 = k0Max < CUBE_BLOCK_SIZE ? RoundDown(k0Max, BLOCK_SIZE) : RoundDown(k0Max, CUBE_BLOCK_SIZE); | ||
| 109 | - if (opShape.k0 > CONST_512) { | ||
| 110 | - opShape.k0 = RoundDown(opShape.k0, CONST_512); | ||
| 111 | - } | ||
| 112 | - } else { | ||
| 113 | - uint32_t k0Max = (shapeSum == 0UL) ? UB_LIMIT_SIZE_910A : (UB_LIMIT_SIZE_910A / shapeSum); | ||
| 114 | - opShape.k0 = k0Max < CUBE_BLOCK_SIZE ? k0Max / BLOCK_SIZE * BLOCK_SIZE : \ | ||
| 115 | - k0Max / CUBE_BLOCK_SIZE * CUBE_BLOCK_SIZE; // k0Max less than 256, matrix 16 | ||
| 116 | - } | ||
| 117 | - // 删除 | ||
| 118 | - opShape.k0 = 512; | ||
| 119 | - kLoop = CeilDiv(opShape.k, opShape.k0); | ||
| 120 | -} | ||
| 121 | - | ||
| 122 | - | ||
| 123 | -void PpMatMulDefault::GetHardwareInfo() | ||
| 124 | -{ | ||
| 125 | - auto platformInfo = context_->GetPlatformInfo(); | ||
| 126 | - if (platformInfo == nullptr) { | ||
| 127 | - OP_LOGE("[PpMatMul]", "platformInfo is nullptr"); | ||
| 128 | - return; | ||
| 129 | - } | ||
| 130 | - HardwareInfo hardwareInfo; | ||
| 131 | - | ||
| 132 | - auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo); | ||
| 133 | - platformInfo->GetPlatformRes("version", "SoC_version", hardwareInfo.socVersionStr); | ||
| 134 | - | ||
| 135 | - hardwareInfo.coreNum = static_cast<uint64_t>(ascendcPlatform.GetCoreNumAic()); | ||
| 136 | - hardwareInfo.socVersion = ascendcPlatform.GetSocVersion(); | ||
| 137 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L2, hardwareInfo.l2Size); | ||
| 138 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L1, hardwareInfo.l1Size); | ||
| 139 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_A, hardwareInfo.l0aSize); | ||
| 140 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_B, hardwareInfo.l0bSize); | ||
| 141 | - ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_C, hardwareInfo.l0cSize); | ||
| 142 | - hardwareInfo_ = hardwareInfo; | ||
| 143 | -} | ||
| 144 | - | ||
| 145 | - | ||
| 146 | -bool PpMatMulDefault::GetMatMulTilingData() | ||
| 147 | -{ | ||
| 148 | - ppMatmulDefaultTilingData_.SetBaseShape(matMulInfo_.batchSize, matMulInfo_.m, matMulInfo_.k, matMulInfo_.n); | ||
| 149 | - OpShape opShape = ppMatmulDefaultTilingData_.opShape; | ||
| 150 | - if (opShape.m < opShape.n) { | ||
| 151 | - TilingFunc<false, OpShape, PpMatmulDefaultTilingData, HardwareInfo, MatMulInfo>(opShape, ppMatmulDefaultTilingData_, hardwareInfo_, matMulInfo_, true, 8); | ||
| 152 | - } else { | ||
| 153 | - TilingFunc<true, OpShape, PpMatmulDefaultTilingData, HardwareInfo, MatMulInfo>(opShape, ppMatmulDefaultTilingData_, hardwareInfo_, matMulInfo_, true, 8); | ||
| 154 | - } | ||
| 155 | - Swizzl<PpMatmulDefaultTilingData>(ppMatmulDefaultTilingData_); | ||
| 156 | - ppMatmulDefaultTilingData_.End(matMulInfo_, hardwareInfo_.socVersion == platform_ascendc::SocVersion::ASCEND310P); | ||
| 157 | - return true; | ||
| 158 | -} | ||
| 159 | - | ||
| 160 | - | ||
| 161 | -void PpMatMulDefault::PrintTiling() { | ||
| 162 | - OP_LOGD(context_->GetNodeName(), "PpMatMul batchSize: %ld.", ppMatmulDefaultTilingData_.opShape.batchSize); | ||
| 163 | - OP_LOGD(context_->GetNodeName(), "PpMatMul m: %ld.", ppMatmulDefaultTilingData_.opShape.m); | ||
| 164 | - OP_LOGD(context_->GetNodeName(), "PpMatMul k: %ld.", ppMatmulDefaultTilingData_.opShape.k); | ||
| 165 | - OP_LOGD(context_->GetNodeName(), "PpMatMul n: %ld.", ppMatmulDefaultTilingData_.opShape.n); | ||
| 166 | - OP_LOGD(context_->GetNodeName(), "PpMatMul m0: %ld.", ppMatmulDefaultTilingData_.opShape.m0); | ||
| 167 | - OP_LOGD(context_->GetNodeName(), "PpMatMul k0: %ld.", ppMatmulDefaultTilingData_.opShape.k0); | ||
| 168 | - OP_LOGD(context_->GetNodeName(), "PpMatMul n0: %ld.", ppMatmulDefaultTilingData_.opShape.n0); | ||
| 169 | - OP_LOGD(context_->GetNodeName(), "PpMatMul mLoop: %ld.", ppMatmulDefaultTilingData_.mLoop); | ||
| 170 | - OP_LOGD(context_->GetNodeName(), "PpMatMul kLoop: %ld.", ppMatmulDefaultTilingData_.kLoop); | ||
| 171 | - OP_LOGD(context_->GetNodeName(), "PpMatMul nLoop: %ld.", ppMatmulDefaultTilingData_.nLoop); | ||
| 172 | - OP_LOGD(context_->GetNodeName(), "PpMatMul coreLoop: %ld.", ppMatmulDefaultTilingData_.coreLoop); | ||
| 173 | - OP_LOGD(context_->GetNodeName(), "PpMatMul swizzlCount: %ld.", ppMatmulDefaultTilingData_.swizzlCount); | ||
| 174 | - OP_LOGD(context_->GetNodeName(), "PpMatMul tilingKey: %d.", ppMatmulDefaultTilingData_.tilingKey); | ||
| 175 | - OP_LOGD(context_->GetNodeName(), "PpMatMul blockDim: %ld.", ppMatmulDefaultTilingData_.blockDim); | ||
| 176 | - OP_LOGD(context_->GetNodeName(), "PpMatMul swizzlDirect: %ld.", ppMatmulDefaultTilingData_.swizzlDirect); | ||
| 177 | - OP_LOGD(context_->GetNodeName(), "PpMatMul splitk: %ld.", ppMatmulDefaultTilingData_.splitk); | ||
| 178 | - OP_LOGD(context_->GetNodeName(), "PpMatMul enShuffleK: %ld.", ppMatmulDefaultTilingData_.enShuffleK); | ||
| 179 | -} | ||
| 180 | - | ||
| 181 | - | ||
| 182 | -} | ||
| 183 | -} | ||
| 184 | - | ||
| @@ -1,88 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2025 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 pp_matmul_info.h | ||
| 13 | - * \brief | ||
| 14 | - */ | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| 24 | -namespace optiling { | ||
| 25 | -namespace pp_matmul { | ||
| 26 | -struct MatMulInfo { | ||
| 27 | - uint64_t batchSize{0}; | ||
| 28 | - uint64_t m{0}; // 实际输入的 m | ||
| 29 | - uint64_t n{0}; // 实际输入的 n | ||
| 30 | - uint64_t k{0}; // 实际输入的 k | ||
| 31 | - ge::DataType dtypeA = ge::DT_FLOAT16; | ||
| 32 | - ge::DataType dtypeB = ge::DT_FLOAT16; | ||
| 33 | - ge::DataType dtypeC = ge::DT_FLOAT16; | ||
| 34 | - ge::Format formatA = ge::FORMAT_ND; | ||
| 35 | - ge::Format formatB = ge::FORMAT_ND; | ||
| 36 | - ge::Format formatC = ge::FORMAT_ND; | ||
| 37 | - uint64_t transA{0}; | ||
| 38 | - uint64_t transB{0}; | ||
| 39 | - bool biasFlag{0}; // false: 0, true: 1 | ||
| 40 | - bool isInt8{0}; // 是否shi int8融合 | ||
| 41 | - float inDtype{0}; | ||
| 42 | - float outDtype{0}; | ||
| 43 | -}; | ||
| 44 | - | ||
| 45 | -struct HardwareInfo { | ||
| 46 | - uint64_t coreNum{0}; | ||
| 47 | - uint64_t l2Size{0}; | ||
| 48 | - uint64_t l1Size{0}; | ||
| 49 | - uint64_t l0aSize{0}; | ||
| 50 | - uint64_t l0bSize{0}; | ||
| 51 | - uint64_t l0cSize{0}; | ||
| 52 | - uint64_t hbmBandWidth{1}; | ||
| 53 | - uint64_t l2BandWidth{5};// 5x faster than hbm. | ||
| 54 | - platform_ascendc::SocVersion socVersion = platform_ascendc::SocVersion::ASCEND910B; | ||
| 55 | - std::string socVersionStr = ""; | ||
| 56 | -}; | ||
| 57 | - | ||
| 58 | -struct OpShape { | ||
| 59 | - uint64_t batchSize{0}; | ||
| 60 | - uint64_t m{0}; | ||
| 61 | - uint64_t k{0}; | ||
| 62 | - uint64_t n{0}; | ||
| 63 | - uint64_t m0{0}; | ||
| 64 | - uint64_t k0{0}; | ||
| 65 | - uint64_t n0{0}; | ||
| 66 | -}; | ||
| 67 | - | ||
| 68 | -struct PpMatmulDefaultTilingData { | ||
| 69 | - OpShape opShape{}; | ||
| 70 | - uint64_t mLoop{1}; | ||
| 71 | - uint64_t kLoop{1}; | ||
| 72 | - uint64_t nLoop{1}; | ||
| 73 | - uint64_t coreLoop{1}; | ||
| 74 | - uint64_t swizzlCount{1}; | ||
| 75 | - uint32_t tilingKey{0}; | ||
| 76 | - uint64_t blockDim{1}; | ||
| 77 | - uint64_t swizzlDirect{0}; | ||
| 78 | - uint64_t splitk{0}; | ||
| 79 | - uint64_t enShuffleK{0}; | ||
| 80 | - | ||
| 81 | - void SetBaseShape(uint64_t batchSize, uint64_t m, uint64_t k, uint64_t n); | ||
| 82 | - void SetBaseOp(uint64_t coreNum, uint64_t l0cSize, uint64_t mBase, uint64_t nBase, const MatMulInfo &mmInfo, bool isAscend310P); | ||
| 83 | - void End(const MatMulInfo &mmInfo, bool isAscend310P); | ||
| 84 | -}; | ||
| 85 | - | ||
| 86 | -} // namespace pp_matmul | ||
| 87 | -} // namespace optiling | ||
| 88 | - | ||
| @@ -1,23 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | - | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| @@ -1,112 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | - | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| 24 | -template <uint32_t ALIGN, typename T = uint32_t> | ||
| 25 | -inline __aicore__ T RoundUp(const T val) | ||
| 26 | -{ | ||
| 27 | - static_assert(ALIGN != 0, "align must not be zero"); | ||
| 28 | - static_assert(std::is_arithmetic<T>::value, "T must be an arithmetic type"); | ||
| 29 | - T align = ALIGN; | ||
| 30 | - if (val + align - 1 < val) { | ||
| 31 | - return val; | ||
| 32 | - } | ||
| 33 | - return (val + align - 1) / align * align; | ||
| 34 | -} | ||
| 35 | - | ||
| 36 | -template <typename T> | ||
| 37 | -inline __aicore__ T RoundUp(const T val, const T align) | ||
| 38 | -{ | ||
| 39 | - static_assert(std::is_arithmetic<T>::value, "T must be an arithmetic type"); | ||
| 40 | - if (align == 0 || val + align - 1 < val) { | ||
| 41 | - return val; | ||
| 42 | - } | ||
| 43 | - return (val + align - 1) / align * align; | ||
| 44 | -} | ||
| 45 | - | ||
| 46 | -template <uint32_t DIVISOR, typename T = uint32_t> | ||
| 47 | -inline __aicore__ T CeilDiv(const T dividend) | ||
| 48 | -{ | ||
| 49 | - static_assert(DIVISOR != 0, "align must not be zero"); | ||
| 50 | - static_assert(std::is_arithmetic<T>::value, "T must be an arithmetic type"); | ||
| 51 | - T divisor = DIVISOR; | ||
| 52 | - if (dividend + divisor - 1 < dividend) { | ||
| 53 | - return dividend; | ||
| 54 | - } | ||
| 55 | - return (dividend + divisor - 1) / divisor; | ||
| 56 | -} | ||
| 57 | - | ||
| 58 | -template <typename T> | ||
| 59 | -constexpr T T_MAX = std::numeric_limits<T>::max(); | ||
| 60 | - | ||
| 61 | -template <typename T> | ||
| 62 | -inline __aicore__ T CeilDiv(const T dividend, const T divisor) | ||
| 63 | -{ | ||
| 64 | - static_assert(std::is_arithmetic<T>::value, "T must be an arithmetic type"); | ||
| 65 | - if (divisor == 0 || dividend + divisor - 1 < dividend) { | ||
| 66 | - return T_MAX<T>; | ||
| 67 | - } | ||
| 68 | - return (dividend + divisor - 1) / divisor; | ||
| 69 | -} | ||
| 70 | - | ||
| 71 | -template <typename T> | ||
| 72 | -__aicore__ inline T Min(const T lhs, const T rhs) | ||
| 73 | -{ | ||
| 74 | - return lhs < rhs ? lhs : rhs; | ||
| 75 | -} | ||
| 76 | - | ||
| 77 | -template <typename Dtype> __aicore__ __attribute__((always_inline)) inline uint32_t BlockSize() | ||
| 78 | -{ | ||
| 79 | - return 32 / sizeof(Dtype); | ||
| 80 | -} | ||
| 81 | - | ||
| 82 | -template <typename Dtype> __aicore__ __attribute__((always_inline)) inline uint32_t MatrixSize() | ||
| 83 | -{ | ||
| 84 | - return 512 / sizeof(Dtype); | ||
| 85 | -} | ||
| 86 | - | ||
| 87 | -template <typename Dtype> __aicore__ __attribute__((always_inline)) inline uint64_t BlockSizeRoundUp(uint64_t num) | ||
| 88 | -{ | ||
| 89 | - return (num + BlockSize<Dtype>() - 1) / BlockSize<Dtype>() * BlockSize<Dtype>(); | ||
| 90 | -} | ||
| 91 | - | ||
| 92 | -template <typename Dtype> __aicore__ __attribute__((always_inline)) inline uint64_t NumBlocksRoundUp(uint64_t num) | ||
| 93 | -{ | ||
| 94 | - return (num + BlockSize<Dtype>() - 1) / BlockSize<Dtype>(); | ||
| 95 | -} | ||
| 96 | - | ||
| 97 | -template <typename Dtype> __aicore__ __attribute__((always_inline)) inline uint64_t MatrixSizeRoundUp(uint64_t num) | ||
| 98 | -{ | ||
| 99 | - return (num + MatrixSize<Dtype>() - 1) / MatrixSize<Dtype>() * MatrixSize<Dtype>(); | ||
| 100 | -} | ||
| 101 | - | ||
| 102 | -template <typename Dtype> __aicore__ __attribute__((always_inline)) inline uint64_t NumMatrixsRoundUp(uint64_t num) | ||
| 103 | -{ | ||
| 104 | - return (num + MatrixSize<Dtype>() - 1) / MatrixSize<Dtype>(); | ||
| 105 | -} | ||
| 106 | - | ||
| 107 | -template <typename Dtype> __aicore__ __attribute__((always_inline)) inline uint64_t L0HalfSize() | ||
| 108 | -{ | ||
| 109 | - return 32 * 1024 / sizeof(Dtype); | ||
| 110 | -} | ||
| 111 | - | ||
| 112 | - | ||
| @@ -1,34 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | -enum class ArchType { ASCEND_V220, ASCEND_V200, ASCEND_M200 }; | ||
| 14 | - | ||
| 15 | -template <ArchType ArchTag> | ||
| 16 | -struct HardwareInfo { | ||
| 17 | - static uint32_t const l2BW = 5; | ||
| 18 | - static uint32_t const hbmBW = 1; | ||
| 19 | - static uint32_t const supportMix = 0; | ||
| 20 | - static uint32_t const l1Size = 512 * 1024; | ||
| 21 | - static uint32_t const l0ASize = 64 * 1024; | ||
| 22 | - static uint32_t const l0BSize = 64 * 1024; | ||
| 23 | - static uint32_t const l0CSize = 128 * 1024; | ||
| 24 | - static uint32_t const l2Size = 192 * 1024 * 1024; | ||
| 25 | - static uint32_t const biasSize = 1024; | ||
| 26 | - static uint32_t const fixBufSize = 7 * 1024; | ||
| 27 | - static uint32_t const ubSize = 192 * 1024; | ||
| 28 | - static uint32_t const fractalSize = 512; | ||
| 29 | - static uint32_t const l1l0BlockSize = 32; | ||
| 30 | - static uint32_t const btBlockSize = 64; | ||
| 31 | - static uint32_t const fbBlockSize = 128; | ||
| 32 | -}; | ||
| 33 | - | ||
| 34 | - | ||
| @@ -1,117 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | - | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | -///////////////////////////////////////////////////// | ||
| 20 | -// gm_to_l1 | ||
| 21 | -///////////////////////////////////////////////////// | ||
| 22 | -template <ArchType ArchTag, typename DataType, DataFormat FormatInGM, DataFormat FormatInL1> | ||
| 23 | -struct gm_to_l1 { | ||
| 24 | - __aicore__ gm_to_l1(AscendC::LocalTensor<DataType> l1Tensor, | ||
| 25 | - AscendC::GlobalTensor<DataType> gmTensor, | ||
| 26 | - uint32_t nTileActual, | ||
| 27 | - uint32_t nTileCeil, | ||
| 28 | - uint32_t nVal, | ||
| 29 | - uint32_t dTileActual, | ||
| 30 | - uint32_t dTileCeil, | ||
| 31 | - uint32_t dVal) {}; | ||
| 32 | -}; | ||
| 33 | - | ||
| 34 | -///////////////////////////////////////////////////// | ||
| 35 | -// l1_to_l0_a | ||
| 36 | -///////////////////////////////////////////////////// | ||
| 37 | -template <ArchType ArchTag, typename DataType, bool IsTransPose, DataFormat DFmtIn, DataFormat DFmtOut> | ||
| 38 | -struct l1_to_l0_a { | ||
| 39 | - __aicore__ l1_to_l0_a(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 40 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 41 | - uint32_t mTileCeil, | ||
| 42 | - uint32_t kPartCeil, | ||
| 43 | - uint32_t mSrcStride, | ||
| 44 | - uint32_t kSrcStride, | ||
| 45 | - uint32_t mDstStride, | ||
| 46 | - uint32_t kDstStride) {}; | ||
| 47 | -}; | ||
| 48 | - | ||
| 49 | -///////////////////////////////////////////////////// | ||
| 50 | -// l1_to_l0_b | ||
| 51 | -///////////////////////////////////////////////////// | ||
| 52 | -template <ArchType ArchTag, typename DataType, bool IsTransPose, DataFormat DFmtIn, DataFormat DFmtOut> | ||
| 53 | -struct l1_to_l0_b { | ||
| 54 | - __aicore__ l1_to_l0_b(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 55 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 56 | - uint32_t nTileCeil, | ||
| 57 | - uint32_t kPartCeil, | ||
| 58 | - uint32_t nSrcStride, | ||
| 59 | - uint32_t kSrcStride, | ||
| 60 | - uint32_t nDstStride, | ||
| 61 | - uint32_t kDstStride) {}; | ||
| 62 | -}; | ||
| 63 | - | ||
| 64 | -///////////////////////////////////////////////////// | ||
| 65 | -// l0c_to_gm | ||
| 66 | -///////////////////////////////////////////////////// | ||
| 67 | -template <ArchType ArchTag, DataFormat OutFormatType, typename OutDataType, typename L0CDataType> | ||
| 68 | -struct l0c_to_gm { | ||
| 69 | - __aicore__ l0c_to_gm(AscendC::GlobalTensor<OutDataType> gmTensor, | ||
| 70 | - AscendC::LocalTensor<L0CDataType> l0cTensor, | ||
| 71 | - uint32_t mTileActual, | ||
| 72 | - uint32_t nTileActual, | ||
| 73 | - uint32_t mTileCeil, | ||
| 74 | - uint32_t nActual, | ||
| 75 | - uint8_t unitFlag = 0) {}; | ||
| 76 | -}; | ||
| 77 | - | ||
| 78 | -///////////////////////////////////////////////////// | ||
| 79 | -// l0c_to_l1 | ||
| 80 | -///////////////////////////////////////////////////// | ||
| 81 | -template <ArchType ArchTag, DataFormat LayoutOut, typename ElementOut, typename ElementIn> | ||
| 82 | -struct l0c_to_l1 { | ||
| 83 | - __aicore__ l0c_to_l1(AscendC::LocalTensor<ElementOut> l1Tensor, | ||
| 84 | - AscendC::LocalTensor<ElementIn> l0cTensor, | ||
| 85 | - AscendC::LocalTensor<uint64_t> deqTensor, | ||
| 86 | - uint32_t mTileActual, | ||
| 87 | - uint32_t nTileActual, | ||
| 88 | - uint32_t mTileCeil, | ||
| 89 | - uint32_t nActual) {}; | ||
| 90 | -}; | ||
| 91 | - | ||
| 92 | -template <ArchType ArchTag, typename DataType> | ||
| 93 | -struct l1_to_bt { | ||
| 94 | - __aicore__ l1_to_bt(uint64_t dst, | ||
| 95 | - const AscendC::LocalTensor<DataType> &src, | ||
| 96 | - uint16_t convControl, | ||
| 97 | - uint16_t nBurst, | ||
| 98 | - uint16_t lenBurst, | ||
| 99 | - uint16_t srcGap, | ||
| 100 | - uint16_t dstGap) {}; | ||
| 101 | -}; | ||
| 102 | - | ||
| 103 | -template <ArchType ArchTag, typename DataType> | ||
| 104 | -struct l1_to_fb { | ||
| 105 | - __aicore__ l1_to_fb(AscendC::LocalTensor<DataType> &dst, | ||
| 106 | - AscendC::LocalTensor<DataType> &src, | ||
| 107 | - uint16_t burstNum, | ||
| 108 | - uint16_t burstLen, | ||
| 109 | - uint16_t srcGap, | ||
| 110 | - uint16_t dstGap) {}; | ||
| 111 | -}; | ||
| 112 | - | ||
| 113 | - | ||
| 114 | - | ||
| 115 | - | ||
| 116 | - | ||
| 117 | - | ||
Dconversion/matmul_v2_compress_dequant/op_kernel/kernels/utils/kernel/iterators/gm_to_l1_iterator.inc+0-161
| @@ -1,161 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | -#include "../iterator.h" | ||
| 11 | - | ||
| 12 | -// Partial specialization for V220, ND_in, ND_out | ||
| 13 | -template <ArchType ArchTag, typename DataType> | ||
| 14 | -struct gm_to_l1<ArchTag, DataType, DataFormat::ND, DataFormat::ND> { | ||
| 15 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 16 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 17 | - | ||
| 18 | - __aicore__ gm_to_l1(AscendC::LocalTensor<DataType> l1Tensor, | ||
| 19 | - AscendC::GlobalTensor<DataType> gmTensor, | ||
| 20 | - uint32_t nTileActual, | ||
| 21 | - uint32_t nTileCeil, | ||
| 22 | - uint32_t nVal, | ||
| 23 | - uint32_t dTileActual, | ||
| 24 | - uint32_t dTileCeil, | ||
| 25 | - uint32_t dVal) | ||
| 26 | - { | ||
| 27 | - AscendC::DataCopy(l1Tensor, // dst | ||
| 28 | - gmTensor, // src | ||
| 29 | - AscendC::DataCopyParams(1, // nBurst | ||
| 30 | - CeilDiv<BLOCK_SIZE>(nTileActual * dTileActual), // lenBurst | ||
| 31 | - 0, // srcGap | ||
| 32 | - 0)); // dstGap | ||
| 33 | - }; | ||
| 34 | -}; | ||
| 35 | - | ||
| 36 | -// Partial specialization for NZ_in, NZ_out | ||
| 37 | -template <ArchType ArchTag, typename DataType> | ||
| 38 | -struct gm_to_l1<ArchTag, DataType, DataFormat::NZ, DataFormat::NZ> { | ||
| 39 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 40 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 41 | - static constexpr uint32_t STRIDE_LIMIT = 65536; | ||
| 42 | - | ||
| 43 | - __aicore__ gm_to_l1(AscendC::LocalTensor<DataType> l1Tensor, | ||
| 44 | - AscendC::GlobalTensor<DataType> gmTensor, | ||
| 45 | - uint32_t nTileActual, | ||
| 46 | - uint32_t nTileCeil, | ||
| 47 | - uint32_t nVal, | ||
| 48 | - uint32_t dTileActual, | ||
| 49 | - uint32_t dTileCeil, | ||
| 50 | - uint32_t dVal) | ||
| 51 | - { | ||
| 52 | - uint64_t srcStride = nVal - nTileCeil; | ||
| 53 | - if (srcStride < STRIDE_LIMIT) { | ||
| 54 | - AscendC::DataCopy(l1Tensor, // dst | ||
| 55 | - gmTensor, // src | ||
| 56 | - AscendC::DataCopyParams(dTileCeil / BLOCK_SIZE, // nBurst | ||
| 57 | - nTileCeil, // lenBurst | ||
| 58 | - srcStride, // srcGap | ||
| 59 | - 0)); // dstGap | ||
| 60 | - } else { | ||
| 61 | - for (uint64_t i = 0; i < dTileCeil / BLOCK_SIZE; i++) { | ||
| 62 | - uint64_t dstOffset = i * nTileCeil * BLOCK_SIZE; | ||
| 63 | - uint64_t srcOffset = i * nVal * BLOCK_SIZE; | ||
| 64 | - AscendC::DataCopy(l1Tensor[dstOffset], // dst | ||
| 65 | - gmTensor[srcOffset], // src | ||
| 66 | - AscendC::DataCopyParams(1, // nBurst | ||
| 67 | - nTileCeil, // lenBurst | ||
| 68 | - 0, // srcGap | ||
| 69 | - 0)); // dstGap | ||
| 70 | - } | ||
| 71 | - } | ||
| 72 | - }; | ||
| 73 | -}; | ||
| 74 | - | ||
| 75 | -// Partial specialization for V220, ND_in, ND_out | ||
| 76 | -template <ArchType ArchTag, typename DataType> | ||
| 77 | -struct gm_to_l1<ArchTag, DataType, DataFormat::ND, DataFormat::NZ> { | ||
| 78 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 79 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 80 | - static constexpr uint32_t STRIDE_LIMIT = 65536; | ||
| 81 | - | ||
| 82 | - __aicore__ gm_to_l1(AscendC::LocalTensor<DataType> l1Tensor, | ||
| 83 | - AscendC::GlobalTensor<DataType> gmTensor, | ||
| 84 | - uint32_t nTileActual, | ||
| 85 | - uint32_t nTileCeil, | ||
| 86 | - uint32_t nVal, | ||
| 87 | - uint32_t dTileActual, | ||
| 88 | - uint32_t dTileCeil, | ||
| 89 | - uint32_t dVal) | ||
| 90 | - { | ||
| 91 | - if (dVal < STRIDE_LIMIT) { | ||
| 92 | - AscendC::DataCopy(l1Tensor, | ||
| 93 | - gmTensor, | ||
| 94 | - AscendC::Nd2NzParams(1, // ndNum | ||
| 95 | - nTileActual, // nValue | ||
| 96 | - dTileActual, // dValue | ||
| 97 | - 0, // srcNdMatrixStride, unused | ||
| 98 | - dVal, // srcDValue | ||
| 99 | - nTileCeil, // dstNzC0Stride | ||
| 100 | - 1, // dstNzNStride | ||
| 101 | - 0)); // dstNzMatrixStride, unused | ||
| 102 | - } else { | ||
| 103 | - for (uint32_t i = 0; i < nTileActual; i++) { | ||
| 104 | - AscendC::DataCopy(l1Tensor[i * BLOCK_SIZE], | ||
| 105 | - gmTensor[i * dVal], | ||
| 106 | - AscendC::Nd2NzParams(1, // ndNum | ||
| 107 | - 1, // nValue | ||
| 108 | - dTileActual, // dValue | ||
| 109 | - 0, // srcNdMatrixStride, unused | ||
| 110 | - 0, // srcDValue | ||
| 111 | - nTileCeil, // dstNzC0Stride | ||
| 112 | - 0, // dstNzNStride | ||
| 113 | - 0)); // dstNzMatrixStride, unused | ||
| 114 | - } | ||
| 115 | - } | ||
| 116 | - }; | ||
| 117 | -}; | ||
| 118 | - | ||
| 119 | -// Partial specialization for V220, ND_in, NZ_out | ||
| 120 | -template <ArchType ArchTag, typename DataType> | ||
| 121 | -struct gm_to_l1<ArchTag, DataType, DataFormat::ND, DataFormat::ZN> { | ||
| 122 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 123 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 124 | - static constexpr uint32_t STRIDE_LIMIT = 65536; | ||
| 125 | - | ||
| 126 | - __aicore__ gm_to_l1(AscendC::LocalTensor<DataType> l1Tensor, | ||
| 127 | - AscendC::GlobalTensor<DataType> gmTensor, | ||
| 128 | - uint32_t nTileActual, | ||
| 129 | - uint32_t nTileCeil, | ||
| 130 | - uint32_t nVal, | ||
| 131 | - uint32_t dTileActual, | ||
| 132 | - uint32_t dTileCeil, | ||
| 133 | - uint32_t dVal) | ||
| 134 | - { | ||
| 135 | - if (dVal < STRIDE_LIMIT) { | ||
| 136 | - AscendC::DataCopy(l1Tensor, | ||
| 137 | - gmTensor, | ||
| 138 | - AscendC::Nd2NzParams(1, // ndNum | ||
| 139 | - nTileActual, // nValue | ||
| 140 | - dTileActual, // dValue | ||
| 141 | - 0, // srcNdMatrixStride, unused | ||
| 142 | - dVal, // srcDValue | ||
| 143 | - nTileCeil, // dstNzC0Stride | ||
| 144 | - 1, // dstNzNStride | ||
| 145 | - 0)); // dstNzMatrixStride, unused | ||
| 146 | - } else { | ||
| 147 | - for (uint32_t i = 0; i < nTileActual; ++i) { | ||
| 148 | - AscendC::DataCopy(l1Tensor, | ||
| 149 | - gmTensor, | ||
| 150 | - AscendC::Nd2NzParams(1, // ndNum | ||
| 151 | - 1, // nValue | ||
| 152 | - dTileActual, // dValue | ||
| 153 | - 0, // srcNdMatrixStride, unused | ||
| 154 | - 0, // srcDValue | ||
| 155 | - nTileCeil, // dstNzC0Stride | ||
| 156 | - 0, // dstNzNStride | ||
| 157 | - 0)); // dstNzMatrixStride, unused | ||
| 158 | - } | ||
| 159 | - } | ||
| 160 | - }; | ||
| 161 | -}; | ||
Dconversion/matmul_v2_compress_dequant/op_kernel/kernels/utils/kernel/iterators/gm_to_ub_iterator.inc+0-87
| @@ -1,87 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | -#include "../iterator.h" | ||
| 11 | - | ||
| 12 | -template <ArchType ArchTag, typename DType> struct gm_to_ub { | ||
| 13 | - __aicore__ inline gm_to_ub(AscendC::LocalTensor<DType> dstTensor, AscendC::GlobalTensor<DType> srcTensor, | ||
| 14 | - uint8_t sid, uint16_t nBurst, uint16_t lenBurst, uint16_t srcStride, uint16_t dstStride) | ||
| 15 | - { | ||
| 16 | - AscendC::DataCopy(dstTensor, srcTensor, AscendC::DataCopyParams(nBurst, lenBurst, srcStride, dstStride)); | ||
| 17 | - }; | ||
| 18 | -}; | ||
| 19 | - | ||
| 20 | -template <ArchType ArchTag, typename DType> struct gm_to_ub_align { | ||
| 21 | - __aicore__ inline gm_to_ub_align(AscendC::LocalTensor<DType> dstTensor, AscendC::GlobalTensor<DType> srcTensor, | ||
| 22 | - uint8_t sid, uint16_t nBurst, uint32_t lenBurst, uint8_t leftPaddingNum, | ||
| 23 | - uint8_t rightPaddingNum, uint32_t srcGap, uint32_t dstGap) | ||
| 24 | - { | ||
| 25 | - AscendC::DataCopyPad(dstTensor, srcTensor, AscendC::DataCopyExtParams(nBurst, lenBurst, srcGap, dstGap, 0), | ||
| 26 | - AscendC::DataCopyPadExtParams<DType>(false, leftPaddingNum, rightPaddingNum, 0)); | ||
| 27 | - }; | ||
| 28 | -}; | ||
| 29 | - | ||
| 30 | -template <ArchType ArchTag, typename DType> struct ub_to_ub { | ||
| 31 | - __aicore__ inline ub_to_ub(AscendC::LocalTensor<DType> dstTensor, AscendC::LocalTensor<DType> srcTensor, | ||
| 32 | - uint8_t sid, uint16_t nBurst, uint16_t lenBurst, uint16_t srcStride, uint16_t dstStride) | ||
| 33 | - { | ||
| 34 | - AscendC::DataCopy(dstTensor, srcTensor, AscendC::DataCopyParams(nBurst, lenBurst, srcStride, dstStride)); | ||
| 35 | - }; | ||
| 36 | -}; | ||
| 37 | - | ||
| 38 | -template <ArchType ArchTag, typename DataType, DataFormat InDataFormat = DataFormat::ND, | ||
| 39 | - DataFormat OutDataFormat = DataFormat::ND> | ||
| 40 | -struct ub_to_gm { | ||
| 41 | - __aicore__ inline ub_to_gm(AscendC::GlobalTensor<DataType> dstTensor, AscendC::LocalTensor<DataType> srcTensor, | ||
| 42 | - uint8_t sid, uint16_t nBurst, uint16_t lenBurst, uint16_t srcStride, uint16_t dstStride) | ||
| 43 | - { | ||
| 44 | - AscendC::DataCopy(dstTensor, srcTensor, AscendC::DataCopyParams(nBurst, lenBurst, srcStride, dstStride)); | ||
| 45 | - }; | ||
| 46 | -}; | ||
| 47 | - | ||
| 48 | -template <ArchType ArchTag, typename DataType> struct ub_to_gm<ArchTag, DataType, DataFormat::NZ, DataFormat::NZ> { | ||
| 49 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 50 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 51 | - | ||
| 52 | - __aicore__ ub_to_gm(AscendC::GlobalTensor<DataType> gmTensor, AscendC::LocalTensor<DataType> ubTensor, | ||
| 53 | - uint32_t nTileActual, uint32_t nTileCeil, uint32_t nVal, uint32_t dTileActual, | ||
| 54 | - uint32_t dTileCeil, uint32_t dVal) | ||
| 55 | - { | ||
| 56 | - constexpr uint32_t STRIDE_LIMIT = 65536; | ||
| 57 | - uint64_t dstStride = nVal - nTileCeil; | ||
| 58 | - if (dstStride < STRIDE_LIMIT) { | ||
| 59 | - AscendC::DataCopy(gmTensor, // dst | ||
| 60 | - ubTensor, // src | ||
| 61 | - AscendC::DataCopyParams(dTileCeil / BLOCK_SIZE, // nBurst | ||
| 62 | - nTileCeil, // lenBurst | ||
| 63 | - 0, // srcGap | ||
| 64 | - dstStride)); // dstGap | ||
| 65 | - } else { | ||
| 66 | - for (uint64_t i = 0; i < dTileCeil / BLOCK_SIZE; ++i) { | ||
| 67 | - uint64_t dstOffset = i * nVal * BLOCK_SIZE; | ||
| 68 | - uint64_t srcOffset = i * nTileCeil * BLOCK_SIZE; | ||
| 69 | - AscendC::DataCopy(gmTensor[dstOffset], // dst | ||
| 70 | - ubTensor[srcOffset], // src | ||
| 71 | - AscendC::DataCopyParams(1, // nBurst | ||
| 72 | - nTileCeil, // lenBurst | ||
| 73 | - 0, // srcGap | ||
| 74 | - 0)); // dstGap | ||
| 75 | - } | ||
| 76 | - } | ||
| 77 | - }; | ||
| 78 | -}; | ||
| 79 | - | ||
| 80 | -template <ArchType ArchTag, typename DType> struct ub_to_gm_align { | ||
| 81 | - __aicore__ inline ub_to_gm_align(AscendC::GlobalTensor<DType> dstTensor, AscendC::LocalTensor<DType> srcTensor, | ||
| 82 | - uint8_t sid, uint16_t nBurst, uint32_t lenBurst, uint8_t leftPaddingNum, | ||
| 83 | - uint8_t rightPaddingNum, uint32_t srcGap, uint32_t dstGap) | ||
| 84 | - { | ||
| 85 | - AscendC::DataCopyPad(dstTensor, srcTensor, AscendC::DataCopyExtParams(nBurst, lenBurst, srcGap, dstGap, 0)); | ||
| 86 | - }; | ||
| 87 | -}; | ||
| @@ -1,69 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | -#include "../iterator.h" | ||
| 11 | - | ||
| 12 | -///////////////////////////////////////////////////// | ||
| 13 | -// l0c_to_ub | ||
| 14 | -///////////////////////////////////////////////////// | ||
| 15 | - | ||
| 16 | -// Partial specialization ZN, half, int32_t | ||
| 17 | -template <ArchType ArchTag, typename ElementIn, typename ElementOut, bool MatrixMode = true> | ||
| 18 | -struct l0c_to_ub { | ||
| 19 | - __aicore__ l0c_to_ub(AscendC::LocalTensor<ElementOut> ubTensor, | ||
| 20 | - AscendC::LocalTensor<ElementIn> l0cTensor, | ||
| 21 | - uint16_t nBurst, | ||
| 22 | - uint16_t lenBurst, | ||
| 23 | - uint16_t srcStride, | ||
| 24 | - uint16_t dstStride) | ||
| 25 | - { | ||
| 26 | - constexpr auto mode = | ||
| 27 | - MatrixMode ? AscendC::BlockMode::BLOCK_MODE_MATRIX : AscendC::BlockMode::BLOCK_MODE_VECTOR; | ||
| 28 | - AscendC::DataCopy(ubTensor, | ||
| 29 | - l0cTensor, | ||
| 30 | - AscendC::DataCopyParams(nBurst, // count | ||
| 31 | - lenBurst, // len | ||
| 32 | - srcStride, // srcStrideIn | ||
| 33 | - dstStride), // dstStrideIn | ||
| 34 | - AscendC::DataCopyEnhancedParams(mode, // blockModeIn | ||
| 35 | - AscendC::DeqScale::DEQ_NONE, // deqScaleIn | ||
| 36 | - 0, // deqValueIn | ||
| 37 | - 0, // sidStoreModeIn | ||
| 38 | - false, // isReluIn | ||
| 39 | - pad_t::PAD_NONE, // padModeIn | ||
| 40 | - 0) // padValueIn | ||
| 41 | - ); | ||
| 42 | - }; | ||
| 43 | -}; | ||
| 44 | - | ||
| 45 | -template <ArchType ArchTag> | ||
| 46 | -struct l0c_to_ub<ArchTag, int32_t, half> { | ||
| 47 | - __aicore__ l0c_to_ub(AscendC::LocalTensor<half> ubTensor, | ||
| 48 | - AscendC::LocalTensor<int32_t> l0cTensor, | ||
| 49 | - uint16_t nBurst, | ||
| 50 | - uint16_t lenBurst, | ||
| 51 | - uint16_t srcStride, | ||
| 52 | - uint16_t dstStride) | ||
| 53 | - { | ||
| 54 | - AscendC::DataCopy(ubTensor, | ||
| 55 | - l0cTensor, | ||
| 56 | - AscendC::DataCopyParams(nBurst, // count | ||
| 57 | - lenBurst, // len | ||
| 58 | - srcStride, // srcStrideIn | ||
| 59 | - dstStride), // dstStrideIn | ||
| 60 | - AscendC::DataCopyEnhancedParams(AscendC::BlockMode::BLOCK_MODE_MATRIX, // blockModeIn | ||
| 61 | - AscendC::DeqScale::VDEQ16, // deqScaleIn | ||
| 62 | - 0, // deqValueIn | ||
| 63 | - 0, // sidStoreModeIn | ||
| 64 | - false, // isReluIn | ||
| 65 | - pad_t::PAD_NONE, // padModeIn | ||
| 66 | - 0) // padValueIn | ||
| 67 | - ); | ||
| 68 | - }; | ||
| 69 | -}; | ||
Dconversion/matmul_v2_compress_dequant/op_kernel/kernels/utils/kernel/iterators/l1_to_l0_iterator.inc+0-308
| @@ -1,308 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | -#include "../iterator.h" | ||
| 11 | - | ||
| 12 | -///////////////////////////////////////////////////// | ||
| 13 | -// l1_to_l0_a | ||
| 14 | -///////////////////////////////////////////////////// | ||
| 15 | - | ||
| 16 | -// Partial specialization for vector | ||
| 17 | -template <ArchType ArchTag, typename DataType, bool IsTransPose> | ||
| 18 | -struct l1_to_l0_a<ArchTag, DataType, IsTransPose, DataFormat::VECTOR, DataFormat::VECTOR> { | ||
| 19 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 20 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(DataType); | ||
| 21 | - | ||
| 22 | - __aicore__ l1_to_l0_a(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 23 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 24 | - uint32_t mTileCeil, | ||
| 25 | - uint32_t kPartCeil, | ||
| 26 | - uint32_t mSrcStride, | ||
| 27 | - uint32_t kSrcStride, | ||
| 28 | - uint32_t mDstStride, | ||
| 29 | - uint32_t kDstStride) | ||
| 30 | - { | ||
| 31 | - AscendC::LoadData(l0Tensor, | ||
| 32 | - l1Tensor, | ||
| 33 | - AscendC::LoadData2dParams(0, // baseIdx | ||
| 34 | - kPartCeil, // repeat | ||
| 35 | - kSrcStride, // srcStride | ||
| 36 | - 0, // sid | ||
| 37 | - kDstStride, // dstStride | ||
| 38 | - IsTransPose, // transpose | ||
| 39 | - 0)); // addrCalMode | ||
| 40 | - }; | ||
| 41 | -}; | ||
| 42 | - | ||
| 43 | -// Partial specialization for no transpose, not vector | ||
| 44 | -template <ArchType ArchTag, typename DataType> | ||
| 45 | -struct l1_to_l0_a<ArchTag, DataType, false, DataFormat::ZN, DataFormat::ZZ> { | ||
| 46 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 47 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 48 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(DataType); | ||
| 49 | - static constexpr uint32_t BLOCK_NUM_PER_FRACTAL = HardwareParams::fractalSize / HardwareParams::l1l0BlockSize; | ||
| 50 | - | ||
| 51 | - __aicore__ l1_to_l0_a(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 52 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 53 | - uint32_t mTileCeil, | ||
| 54 | - uint32_t kPartCeil, | ||
| 55 | - uint32_t mSrcStride, | ||
| 56 | - uint32_t kSrcStride, | ||
| 57 | - uint32_t mDstStride, | ||
| 58 | - uint32_t kDstStride) | ||
| 59 | - { | ||
| 60 | - for (uint32_t i = 0; i < mTileCeil / BLOCK_NUM_PER_FRACTAL; i++) { | ||
| 61 | - AscendC::LoadData(l0Tensor[i * mDstStride * FRACTAL_SIZE], // dst | ||
| 62 | - l1Tensor[i * mSrcStride * FRACTAL_SIZE], // src | ||
| 63 | - AscendC::LoadData2dParams(0, // baseIdx | ||
| 64 | - static_cast<uint16_t>(kPartCeil / BLOCK_SIZE), // repeat | ||
| 65 | - kSrcStride, // srcStride | ||
| 66 | - 0, // sid | ||
| 67 | - kDstStride - 1, // dstStride | ||
| 68 | - false, // transpose | ||
| 69 | - 0)); // addrCalMode | ||
| 70 | - } | ||
| 71 | - }; | ||
| 72 | -}; | ||
| 73 | - | ||
| 74 | -// Partial specialization for transpose, not vector | ||
| 75 | -template <ArchType ArchTag, typename DataType> | ||
| 76 | -struct l1_to_l0_a<ArchTag, DataType, true, DataFormat::ZN, DataFormat::ZZ> { | ||
| 77 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 78 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 79 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(DataType); | ||
| 80 | - static constexpr uint32_t BLOCK_NUM_PER_FRACTAL = HardwareParams::fractalSize / HardwareParams::l1l0BlockSize; | ||
| 81 | - | ||
| 82 | - __aicore__ l1_to_l0_a(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 83 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 84 | - uint32_t mTileCeil, | ||
| 85 | - uint32_t kPartCeil, | ||
| 86 | - uint32_t mSrcStride, | ||
| 87 | - uint32_t kSrcStride, | ||
| 88 | - uint32_t mDstStride, | ||
| 89 | - uint32_t kDstStride) | ||
| 90 | - { | ||
| 91 | - for (uint32_t i = 0; i < mTileCeil / BLOCK_SIZE; i++) { | ||
| 92 | - AscendC::LoadData(l0Tensor[i * mDstStride * FRACTAL_SIZE], | ||
| 93 | - l1Tensor[i * mSrcStride * FRACTAL_SIZE], | ||
| 94 | - AscendC::LoadData2dParams(0, | ||
| 95 | - static_cast<uint16_t>(kPartCeil / BLOCK_NUM_PER_FRACTAL), | ||
| 96 | - kSrcStride, | ||
| 97 | - 0, | ||
| 98 | - kDstStride - 1, | ||
| 99 | - true, | ||
| 100 | - 0)); | ||
| 101 | - } | ||
| 102 | - }; | ||
| 103 | -}; | ||
| 104 | - | ||
| 105 | -template <ArchType ArchTag, typename DataType> | ||
| 106 | -struct l1_to_l0_a<ArchTag, DataType, false, DataFormat::NZ, DataFormat::ZZ> { | ||
| 107 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 108 | - // 16 * 32 | ||
| 109 | - static constexpr uint32_t ROW_BLOCK_SIZE = 16; | ||
| 110 | - static constexpr uint32_t COL_BLOCK_SIZE = 32 / sizeof(DataType); | ||
| 111 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(DataType); | ||
| 112 | - static constexpr uint32_t BLOCK_NUM_PER_FRACTAL = HardwareParams::fractalSize / HardwareParams::l1l0BlockSize; | ||
| 113 | - | ||
| 114 | - __aicore__ l1_to_l0_a(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 115 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 116 | - uint32_t mTileCeil, | ||
| 117 | - uint32_t kPartCeil, | ||
| 118 | - uint32_t mSrcStride, | ||
| 119 | - uint32_t kSrcStride, | ||
| 120 | - uint32_t mDstStride, | ||
| 121 | - uint32_t kDstStride) | ||
| 122 | - { | ||
| 123 | - for (uint32_t i = 0; i < mTileCeil / ROW_BLOCK_SIZE; i++) { | ||
| 124 | - AscendC::LoadData(l0Tensor[i * ROW_BLOCK_SIZE * kPartCeil], | ||
| 125 | - l1Tensor[i * FRACTAL_SIZE], | ||
| 126 | - AscendC::LoadData2dParams(0, | ||
| 127 | - static_cast<uint16_t>(kPartCeil / COL_BLOCK_SIZE), | ||
| 128 | - mTileCeil / ROW_BLOCK_SIZE, | ||
| 129 | - 0, | ||
| 130 | - 0, | ||
| 131 | - false, | ||
| 132 | - 0)); | ||
| 133 | - } | ||
| 134 | - }; | ||
| 135 | -}; | ||
| 136 | - | ||
| 137 | -template <> | ||
| 138 | -struct l1_to_l0_a<ArchType::ASCEND_V220, int8_t, true, DataFormat::ZN, DataFormat::ZZ> { | ||
| 139 | - using HardwareParams = HardwareInfo<ArchType::ASCEND_V220>; | ||
| 140 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(int8_t); // 32 | ||
| 141 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(int8_t); // 512 | ||
| 142 | - static constexpr uint32_t BLOCK_NUM_PER_FRACTAL = HardwareParams::fractalSize / HardwareParams::l1l0BlockSize; // 16 | ||
| 143 | - static constexpr uint32_t NUM_FRACTAL_PER_ITER = 2; | ||
| 144 | - __aicore__ l1_to_l0_a(AscendC::LocalTensor<int8_t> l0Tensor, | ||
| 145 | - AscendC::LocalTensor<int8_t> l1Tensor, | ||
| 146 | - uint32_t mTileCeil, | ||
| 147 | - uint32_t kPartCeil, | ||
| 148 | - uint32_t mSrcStride, | ||
| 149 | - uint32_t kSrcStride, | ||
| 150 | - uint32_t mDstStride, | ||
| 151 | - uint32_t kDstStride) | ||
| 152 | - { | ||
| 153 | - for (uint64_t i = 0; i < mTileCeil / (BLOCK_NUM_PER_FRACTAL * NUM_FRACTAL_PER_ITER); ++i) { | ||
| 154 | - AscendC::LoadDataWithTranspose( | ||
| 155 | - l0Tensor[i * mDstStride * FRACTAL_SIZE * NUM_FRACTAL_PER_ITER], // dstLocalTensor | ||
| 156 | - l1Tensor[i * mSrcStride * FRACTAL_SIZE], // srcLocalTensor | ||
| 157 | - AscendC::LoadData2dTransposeParams(0, // baseIdx | ||
| 158 | - static_cast<uint16_t>(CeilDiv<BLOCK_SIZE>(kPartCeil)), // repeat | ||
| 159 | - kSrcStride, // srcStride | ||
| 160 | - 0, // dstGap | ||
| 161 | - mDstStride - 1)); // dstFracGap | ||
| 162 | - } | ||
| 163 | - } | ||
| 164 | -}; | ||
| 165 | - | ||
| 166 | -///////////////////////////////////////////////////// | ||
| 167 | -// l1_to_l0_b | ||
| 168 | -///////////////////////////////////////////////////// | ||
| 169 | - | ||
| 170 | -// Partial specialization for vector | ||
| 171 | -template <ArchType ArchTag, typename DataType, bool IsTransPose> | ||
| 172 | -struct l1_to_l0_b<ArchTag, DataType, IsTransPose, DataFormat::VECTOR, DataFormat::VECTOR> { | ||
| 173 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 174 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(DataType); | ||
| 175 | - | ||
| 176 | - __aicore__ l1_to_l0_b(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 177 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 178 | - uint32_t nTileCeil, | ||
| 179 | - uint32_t kPartCeil, | ||
| 180 | - uint32_t nSrcStride, | ||
| 181 | - uint32_t kSrcStride, | ||
| 182 | - uint32_t nDstStride, | ||
| 183 | - uint32_t kDstStride) | ||
| 184 | - { | ||
| 185 | - AscendC::LoadData( | ||
| 186 | - l0Tensor, l1Tensor, AscendC::LoadData2dParams(0, kPartCeil, kSrcStride, 0, kDstStride, IsTransPose, 0)); | ||
| 187 | - }; | ||
| 188 | -}; | ||
| 189 | - | ||
| 190 | -template <ArchType ArchTag> | ||
| 191 | -struct l1_to_l0_b<ArchTag, int8_t, true, DataFormat::NZ, DataFormat::ZN> { | ||
| 192 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 193 | - using DataType = int8_t; | ||
| 194 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 195 | - | ||
| 196 | - __aicore__ l1_to_l0_b(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 197 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 198 | - uint32_t nTileCeil, | ||
| 199 | - uint32_t kPartCeil, | ||
| 200 | - uint32_t nSrcStride, | ||
| 201 | - uint32_t kSrcStride, | ||
| 202 | - uint32_t nDstStride, | ||
| 203 | - uint32_t kDstStride) | ||
| 204 | - { | ||
| 205 | - for (uint32_t i = 0; i < nTileCeil / BLOCK_SIZE; i++) { | ||
| 206 | - AscendC::LoadDataWithTranspose(l0Tensor[i * kPartCeil * BLOCK_SIZE], | ||
| 207 | - l1Tensor[i * BLOCK_SIZE * BLOCK_SIZE], | ||
| 208 | - AscendC::LoadData2dTransposeParams(0, // startIndexIn | ||
| 209 | - kPartCeil / BLOCK_SIZE, // repeatTimesIn | ||
| 210 | - nTileCeil / BLOCK_SIZE, // srcStrideIn | ||
| 211 | - 1, // dstGapIn | ||
| 212 | - 0, // dstfracGapIn | ||
| 213 | - 0) // addrModeIn | ||
| 214 | - ); | ||
| 215 | - } | ||
| 216 | - }; | ||
| 217 | -}; | ||
| 218 | - | ||
| 219 | -// Partial specialization for no transpose, not vector | ||
| 220 | -template <ArchType ArchTag, typename DataType> | ||
| 221 | -struct l1_to_l0_b<ArchTag, DataType, false, DataFormat::ZN, DataFormat::NZ> { | ||
| 222 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 223 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 224 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(DataType); | ||
| 225 | - static constexpr uint32_t BLOCK_NUM_PER_FRACTAL = HardwareParams::fractalSize / HardwareParams::l1l0BlockSize; | ||
| 226 | - | ||
| 227 | - __aicore__ l1_to_l0_b(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 228 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 229 | - uint32_t nTileCeil, | ||
| 230 | - uint32_t kPartCeil, | ||
| 231 | - uint32_t nSrcStride, | ||
| 232 | - uint32_t kSrcStride, | ||
| 233 | - uint32_t nDstStride, | ||
| 234 | - uint32_t kDstStride) | ||
| 235 | - { | ||
| 236 | - for (uint32_t i = 0; i < kPartCeil / BLOCK_NUM_PER_FRACTAL; i++) { | ||
| 237 | - AscendC::LoadData(l0Tensor[i * kDstStride * FRACTAL_SIZE], | ||
| 238 | - l1Tensor[i * kSrcStride * FRACTAL_SIZE], | ||
| 239 | - AscendC::LoadData2dParams(0, // baseIdx | ||
| 240 | - static_cast<uint16_t>(nTileCeil / BLOCK_SIZE), // repeat | ||
| 241 | - nSrcStride, // srcStride | ||
| 242 | - 0, // sid | ||
| 243 | - nDstStride - 1, // dstStride | ||
| 244 | - true, // transpose | ||
| 245 | - 0)); // addrCalMode | ||
| 246 | - } | ||
| 247 | - }; | ||
| 248 | -}; | ||
| 249 | - | ||
| 250 | -// Partial specialization for transpose, not vector | ||
| 251 | -template <ArchType ArchTag, typename DataType> | ||
| 252 | -struct l1_to_l0_b<ArchTag, DataType, true, DataFormat::ZN, DataFormat::NZ> { | ||
| 253 | - using HardwareParams = HardwareInfo<ArchTag>; | ||
| 254 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(DataType); | ||
| 255 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(DataType); | ||
| 256 | - static constexpr uint32_t BLOCK_NUM_PER_FRACTAL = HardwareParams::fractalSize / HardwareParams::l1l0BlockSize; | ||
| 257 | - | ||
| 258 | - __aicore__ l1_to_l0_b(AscendC::LocalTensor<DataType> l0Tensor, | ||
| 259 | - AscendC::LocalTensor<DataType> l1Tensor, | ||
| 260 | - uint32_t nTileCeil, | ||
| 261 | - uint32_t kPartCeil, | ||
| 262 | - uint32_t nSrcStride, | ||
| 263 | - uint32_t kSrcStride, | ||
| 264 | - uint32_t nDstStride, | ||
| 265 | - uint32_t kDstStride) | ||
| 266 | - { | ||
| 267 | - AscendC::LoadData( | ||
| 268 | - l0Tensor, | ||
| 269 | - l1Tensor, | ||
| 270 | - AscendC::LoadData2dParams(0, // baseIdx | ||
| 271 | - static_cast<uint16_t>(kPartCeil * nTileCeil / FRACTAL_SIZE), // repeat | ||
| 272 | - 1, // srcStride | ||
| 273 | - 0, // sid | ||
| 274 | - 0, // dstStride | ||
| 275 | - false, // transpose | ||
| 276 | - 0)); // addr_cal_mode_t | ||
| 277 | - }; | ||
| 278 | -}; | ||
| 279 | - | ||
| 280 | -template <> | ||
| 281 | -struct l1_to_l0_b<ArchType::ASCEND_V220, int8_t, false, DataFormat::ZN, DataFormat::NZ> { | ||
| 282 | - using HardwareParams = HardwareInfo<ArchType::ASCEND_V220>; | ||
| 283 | - static constexpr uint32_t BLOCK_SIZE = HardwareParams::l1l0BlockSize / sizeof(int8_t); // 32 | ||
| 284 | - static constexpr uint32_t FRACTAL_SIZE = HardwareParams::fractalSize / sizeof(int8_t); // 16 | ||
| 285 | - static constexpr uint32_t BLOCK_NUM_PER_FRACTAL = HardwareParams::fractalSize / HardwareParams::l1l0BlockSize; | ||
| 286 | - static constexpr uint32_t NUM_FRACTAL_PER_ITER = 2; | ||
| 287 | - | ||
| 288 | - __aicore__ l1_to_l0_b(AscendC::LocalTensor<int8_t> l0Tensor, | ||
| 289 | - AscendC::LocalTensor<int8_t> l1Tensor, | ||
| 290 | - uint32_t nTileCeil, | ||
| 291 | - uint32_t kPartCeil, | ||
| 292 | - uint32_t nSrcStride, | ||
| 293 | - uint32_t kSrcStride, | ||
| 294 | - uint32_t nDstStride, | ||
| 295 | - uint32_t kDstStride) | ||
| 296 | - { | ||
| 297 | - for (uint64_t i = 0; i < kPartCeil / (BLOCK_NUM_PER_FRACTAL * NUM_FRACTAL_PER_ITER); ++i) { | ||
| 298 | - AscendC::LoadDataWithTranspose( | ||
| 299 | - l0Tensor[i * kDstStride * FRACTAL_SIZE], // dstLocalTensor | ||
| 300 | - l1Tensor[i * kSrcStride * FRACTAL_SIZE * NUM_FRACTAL_PER_ITER], // srcLocalTensor | ||
| 301 | - AscendC::LoadData2dTransposeParams(0, // baseIdx | ||
| 302 | - static_cast<uint16_t>(CeilDiv<BLOCK_SIZE>(nTileCeil)), // repeat | ||
| 303 | - nSrcStride / NUM_FRACTAL_PER_ITER, // srcStride | ||
| 304 | - 1, // dstGap | ||
| 305 | - 0)); // dstFracGap | ||
| 306 | - } | ||
| 307 | - }; | ||
| 308 | -}; | ||
| @@ -1,16 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | -enum class DataFormat { ND = 0, NZ, ZN, ZZ, NN, VECTOR }; | ||
| 15 | - | ||
| 16 | - | ||
| @@ -1,82 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | - | ||
| 15 | - | ||
| 16 | - | ||
| 17 | -enum class BufferType { ASCEND_UB, ASCEND_CB, ASCEND_L0A, ASCEND_L0B, ASCEND_L0C, ASCEND_MAX }; | ||
| 18 | - | ||
| 19 | -template <BufferType BufferType_> | ||
| 20 | -__aicore__ constexpr AscendC::TPosition GetPosition() | ||
| 21 | -{ | ||
| 22 | - if constexpr (BufferType_ == BufferType::ASCEND_UB) { | ||
| 23 | - return AscendC::TPosition::VECIN; | ||
| 24 | - } else if constexpr (BufferType_ == BufferType::ASCEND_CB) { | ||
| 25 | - return AscendC::TPosition::A1; | ||
| 26 | - } else if constexpr (BufferType_ == BufferType::ASCEND_L0A) { | ||
| 27 | - return AscendC::TPosition::A2; | ||
| 28 | - } else if constexpr (BufferType_ == BufferType::ASCEND_L0B) { | ||
| 29 | - return AscendC::TPosition::B2; | ||
| 30 | - } else if constexpr (BufferType_ == BufferType::ASCEND_L0C) { | ||
| 31 | - return AscendC::TPosition::CO1; | ||
| 32 | - } | ||
| 33 | - return AscendC::TPosition::GM; | ||
| 34 | -} | ||
| 35 | - | ||
| 36 | -template <ArchType ArchTag> | ||
| 37 | -struct AsdopsBuffer { | ||
| 38 | -public: | ||
| 39 | - __aicore__ AsdopsBuffer() | ||
| 40 | - { | ||
| 41 | - constexpr uint32_t bufferSize[(uint32_t)BufferType::ASCEND_MAX] = {HardwareInfo<ArchTag>::ubSize, | ||
| 42 | - HardwareInfo<ArchTag>::l1Size, | ||
| 43 | - HardwareInfo<ArchTag>::l0ASize, | ||
| 44 | - HardwareInfo<ArchTag>::l0BSize, | ||
| 45 | - HardwareInfo<ArchTag>::l0CSize}; | ||
| 46 | - | ||
| 47 | - tensor[(uint32_t)BufferType::ASCEND_UB].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_UB]); | ||
| 48 | - tensor[(uint32_t)BufferType::ASCEND_UB].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::VECIN); | ||
| 49 | - | ||
| 50 | - tensor[(uint32_t)BufferType::ASCEND_CB].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_CB]); | ||
| 51 | - tensor[(uint32_t)BufferType::ASCEND_CB].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::A1); | ||
| 52 | - tensor[(uint32_t)BufferType::ASCEND_L0A].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_L0A]); | ||
| 53 | - tensor[(uint32_t)BufferType::ASCEND_L0A].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::A2); | ||
| 54 | - tensor[(uint32_t)BufferType::ASCEND_L0B].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_L0B]); | ||
| 55 | - tensor[(uint32_t)BufferType::ASCEND_L0B].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::B2); | ||
| 56 | - tensor[(uint32_t)BufferType::ASCEND_L0C].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_L0C]); | ||
| 57 | - tensor[(uint32_t)BufferType::ASCEND_L0C].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::CO1); | ||
| 58 | - | ||
| 59 | - tensor[(uint32_t)BufferType::ASCEND_UB].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_UB]); | ||
| 60 | - tensor[(uint32_t)BufferType::ASCEND_UB].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::VECIN); | ||
| 61 | - tensor[(uint32_t)BufferType::ASCEND_CB].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_CB]); | ||
| 62 | - tensor[(uint32_t)BufferType::ASCEND_CB].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::A1); | ||
| 63 | - tensor[(uint32_t)BufferType::ASCEND_L0A].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_L0A]); | ||
| 64 | - tensor[(uint32_t)BufferType::ASCEND_L0A].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::A2); | ||
| 65 | - tensor[(uint32_t)BufferType::ASCEND_L0B].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_L0B]); | ||
| 66 | - tensor[(uint32_t)BufferType::ASCEND_L0B].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::B2); | ||
| 67 | - tensor[(uint32_t)BufferType::ASCEND_L0C].InitBuffer(0, bufferSize[(uint32_t)BufferType::ASCEND_L0C]); | ||
| 68 | - tensor[(uint32_t)BufferType::ASCEND_L0C].address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::CO1); | ||
| 69 | - | ||
| 70 | - }; | ||
| 71 | - | ||
| 72 | - template <BufferType BufferType_, typename DstDataType = half> | ||
| 73 | - __aicore__ AscendC::LocalTensor<DstDataType> GetBuffer(const uint32_t offset) const | ||
| 74 | - { | ||
| 75 | - return tensor[(uint32_t)BufferType_][offset].template ReinterpretCast<DstDataType>(); | ||
| 76 | - } | ||
| 77 | - | ||
| 78 | -public: | ||
| 79 | - AscendC::LocalTensor<uint8_t> tensor[(uint32_t)BufferType::ASCEND_MAX]; | ||
| 80 | -}; | ||
| 81 | - | ||
| 82 | - | ||
| @@ -1,87 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | - | ||
| 15 | - | ||
| 16 | -template <ArchType ArchTag, typename ElementA, typename ElementB, typename AccDTypeC, bool IsTransposeA> | ||
| 17 | -struct mmad { | ||
| 18 | - __aicore__ mmad(AscendC::LocalTensor<AccDTypeC> l0cTensor, | ||
| 19 | - AscendC::LocalTensor<ElementA> l0aTensor, | ||
| 20 | - AscendC::LocalTensor<ElementB> l0bTensor, | ||
| 21 | - uint32_t mTileActual, | ||
| 22 | - uint32_t nTileActual, | ||
| 23 | - uint32_t kPartActual, | ||
| 24 | - bool initC, | ||
| 25 | - uint8_t unitFlag = 0) {}; | ||
| 26 | - | ||
| 27 | - __aicore__ mmad(AscendC::LocalTensor<AccDTypeC> l0cTensor, | ||
| 28 | - AscendC::LocalTensor<ElementA> l0aTensor, | ||
| 29 | - AscendC::LocalTensor<ElementB> l0bTensor, | ||
| 30 | - uint64_t biasBt, | ||
| 31 | - uint32_t mTileActual, | ||
| 32 | - uint32_t nTileActual, | ||
| 33 | - uint32_t kPartActual, | ||
| 34 | - bool initC, | ||
| 35 | - uint8_t unitFlag = 0) {}; | ||
| 36 | -}; | ||
| 37 | - | ||
| 38 | -// Partial specialization for V220, int8_t, not_vector_A, not TransposeA | ||
| 39 | -template <ArchType ArchTag, typename AccDTypeC, typename ElementA, typename ElementB> | ||
| 40 | -struct mmad<ArchTag, ElementA, ElementB, AccDTypeC, false> { | ||
| 41 | - __aicore__ mmad(AscendC::LocalTensor<AccDTypeC> l0cTensor, | ||
| 42 | - AscendC::LocalTensor<ElementA> l0aTensor, | ||
| 43 | - AscendC::LocalTensor<ElementB> l0bTensor, | ||
| 44 | - uint32_t mTileActual, | ||
| 45 | - uint32_t nTileActual, | ||
| 46 | - uint32_t kPartActual, | ||
| 47 | - bool initC, | ||
| 48 | - uint8_t unitFlag = 0) | ||
| 49 | - { | ||
| 50 | - AscendC::Mmad(l0cTensor, // C | ||
| 51 | - l0aTensor, // A | ||
| 52 | - l0bTensor, // B | ||
| 53 | - AscendC::MmadParams(mTileActual, // m | ||
| 54 | - nTileActual, // n | ||
| 55 | - kPartActual, // k | ||
| 56 | - unitFlag, // unitFlag | ||
| 57 | - false, // cmatrixSource | ||
| 58 | - initC)); // cmatrixInitVal | ||
| 59 | - }; | ||
| 60 | - | ||
| 61 | - __aicore__ mmad(AscendC::LocalTensor<AccDTypeC> l0cTensor, | ||
| 62 | - AscendC::LocalTensor<ElementA> l0aTensor, | ||
| 63 | - AscendC::LocalTensor<ElementB> l0bTensor, | ||
| 64 | - uint64_t biasBt, | ||
| 65 | - uint32_t mTileActual, | ||
| 66 | - uint32_t nTileActual, | ||
| 67 | - uint32_t kPartActual, | ||
| 68 | - bool initC, | ||
| 69 | - uint8_t unitFlag = 0) | ||
| 70 | - { | ||
| 71 | - AscendC::LocalTensor<AccDTypeC> biasTensor; | ||
| 72 | - biasTensor.InitBuffer(biasBt, nTileActual); | ||
| 73 | - biasTensor.address_.logicPos = static_cast<uint8_t>(AscendC::TPosition::C2); | ||
| 74 | - AscendC::Mmad(l0cTensor, // C | ||
| 75 | - l0aTensor, // A | ||
| 76 | - l0bTensor, // B | ||
| 77 | - biasTensor, // bt | ||
| 78 | - AscendC::MmadParams(mTileActual, // m | ||
| 79 | - nTileActual, // n | ||
| 80 | - kPartActual, // k | ||
| 81 | - unitFlag, // unitFlag | ||
| 82 | - true, // cmatrixSource | ||
| 83 | - false)); // cmatrixInitVal | ||
| 84 | - }; | ||
| 85 | -}; | ||
| 86 | - | ||
| 87 | - | ||
| @@ -1,423 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | - | ||
| 15 | - | ||
| 16 | -///////////////////////////////////////////////////// | ||
| 17 | -// vcgadd | ||
| 18 | -///////////////////////////////////////////////////// | ||
| 19 | -template <ArchType ArchTag, typename DType> | ||
| 20 | -__aicore__ inline void cgadd_v(AscendC::LocalTensor<DType> dst, | ||
| 21 | - AscendC::LocalTensor<DType> src, | ||
| 22 | - const int32_t repeat, | ||
| 23 | - const int32_t dstRepStride, | ||
| 24 | - const int32_t srcBlkStride, | ||
| 25 | - const int32_t srcRepStride) | ||
| 26 | -{ | ||
| 27 | - AscendC::BlockReduceSum<DType, false>(dst, src, repeat, 0, dstRepStride, srcBlkStride, srcRepStride); | ||
| 28 | -} | ||
| 29 | - | ||
| 30 | -///////////////////////////////////////////////////// | ||
| 31 | -// vadd | ||
| 32 | -///////////////////////////////////////////////////// | ||
| 33 | -template <ArchType ArchTag, typename DType> | ||
| 34 | -__aicore__ inline void add_v(AscendC::LocalTensor<DType> dst, | ||
| 35 | - AscendC::LocalTensor<DType> src0, | ||
| 36 | - AscendC::LocalTensor<DType> src1, | ||
| 37 | - uint8_t repeat, | ||
| 38 | - uint8_t dstBlockStride, | ||
| 39 | - uint8_t src0BlockStride, | ||
| 40 | - uint8_t src1BlockStride, | ||
| 41 | - uint8_t dstRepeatStride, | ||
| 42 | - uint8_t src0RepeatStride, | ||
| 43 | - uint8_t src1RepeatStride) | ||
| 44 | -{ | ||
| 45 | - AscendC::Add<DType, false>( | ||
| 46 | - dst, | ||
| 47 | - src0, | ||
| 48 | - src1, | ||
| 49 | - (uint64_t)0, | ||
| 50 | - repeat, | ||
| 51 | - AscendC::BinaryRepeatParams( | ||
| 52 | - dstBlockStride, src0BlockStride, src1BlockStride, dstRepeatStride, src0RepeatStride, src1RepeatStride)); | ||
| 53 | -} | ||
| 54 | - | ||
| 55 | -///////////////////////////////////////////////////// | ||
| 56 | -// vadds | ||
| 57 | -///////////////////////////////////////////////////// | ||
| 58 | -template <ArchType ArchTag, typename DType> | ||
| 59 | -__aicore__ inline void adds_v(AscendC::LocalTensor<DType> dst, | ||
| 60 | - AscendC::LocalTensor<DType> src, | ||
| 61 | - DType scalarValue, | ||
| 62 | - uint8_t repeat, | ||
| 63 | - uint8_t dstBlockStride, | ||
| 64 | - uint8_t srcBlockStride, | ||
| 65 | - uint8_t dstRepeatStride, | ||
| 66 | - uint8_t srcRepeatStride) | ||
| 67 | -{ | ||
| 68 | - AscendC::Adds<DType, false>( | ||
| 69 | - dst, | ||
| 70 | - src, | ||
| 71 | - scalarValue, | ||
| 72 | - (uint64_t)0, | ||
| 73 | - repeat, | ||
| 74 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 75 | -} | ||
| 76 | - | ||
| 77 | -///////////////////////////////////////////////////// | ||
| 78 | -// vcadd | ||
| 79 | -///////////////////////////////////////////////////// | ||
| 80 | -template <ArchType ArchTag, typename DType> | ||
| 81 | -__aicore__ inline void cadd_v(AscendC::LocalTensor<DType> dst, | ||
| 82 | - AscendC::LocalTensor<DType> src, | ||
| 83 | - uint8_t repeat, | ||
| 84 | - uint16_t dstRepeatStride, | ||
| 85 | - uint16_t srcBlockStride, | ||
| 86 | - uint16_t srcRepeatStride) | ||
| 87 | -{ | ||
| 88 | - AscendC::RepeatReduceSum<DType, false>(dst, src, repeat, 0, 0, srcBlockStride, dstRepeatStride, srcRepeatStride); | ||
| 89 | -} | ||
| 90 | -///////////////////////////////////////////////////// | ||
| 91 | -// vbrcb | ||
| 92 | -///////////////////////////////////////////////////// | ||
| 93 | -template <ArchType ArchTag, typename DType> | ||
| 94 | -__aicore__ inline void brcb_v(AscendC::LocalTensor<DType> dst, | ||
| 95 | - AscendC::LocalTensor<DType> src, | ||
| 96 | - uint16_t dstBlockStride, | ||
| 97 | - uint16_t dstRepeatStride, | ||
| 98 | - uint8_t repeat) | ||
| 99 | -{ | ||
| 100 | - AscendC::Brcb(dst, src, repeat, AscendC::BrcbRepeatParams(dstBlockStride, dstRepeatStride)); | ||
| 101 | -} | ||
| 102 | - | ||
| 103 | -///////////////////////////////////////////////////// | ||
| 104 | -// vcmax | ||
| 105 | -///////////////////////////////////////////////////// | ||
| 106 | -template <ArchType ArchTag, typename DType, AscendC::ReduceOrder OrderType> | ||
| 107 | -__aicore__ inline void cmax_v(AscendC::LocalTensor<DType> dst, | ||
| 108 | - AscendC::LocalTensor<DType> src, | ||
| 109 | - uint8_t repeat, | ||
| 110 | - uint16_t dstRepeatStride, | ||
| 111 | - uint16_t srcBlockStride, | ||
| 112 | - uint16_t srcRepeatStride) | ||
| 113 | -{ | ||
| 114 | - | ||
| 115 | - AscendC::WholeReduceMax<DType, false>( | ||
| 116 | - dst, src, (int32_t)0, repeat, dstRepeatStride, srcBlockStride, srcRepeatStride, OrderType); | ||
| 117 | - | ||
| 118 | - AscendC::WholeReduceMax<DType, false>( | ||
| 119 | - dst, src, (int32_t)0, repeat, dstRepeatStride, srcBlockStride, srcRepeatStride); | ||
| 120 | - | ||
| 121 | -} | ||
| 122 | - | ||
| 123 | -///////////////////////////////////////////////////// | ||
| 124 | -// vconv | ||
| 125 | -///////////////////////////////////////////////////// | ||
| 126 | -template <ArchType ArchTag, typename DTypeIn, typename DTypeOut> | ||
| 127 | -__aicore__ inline void conv_v(AscendC::LocalTensor<DTypeOut> dst, | ||
| 128 | - AscendC::LocalTensor<DTypeIn> src, | ||
| 129 | - uint8_t repeat, | ||
| 130 | - uint16_t dstBlockStride, | ||
| 131 | - uint16_t srcBlockStride, | ||
| 132 | - uint16_t dstRepeatStride, | ||
| 133 | - uint16_t srcRepeatStride) | ||
| 134 | -{ | ||
| 135 | - if constexpr (std::is_same<DTypeIn, float>::value && std::is_same<DTypeOut, __bf16>::value) { | ||
| 136 | - AscendC::Cast<DTypeOut, DTypeIn, false>( | ||
| 137 | - dst, | ||
| 138 | - src, | ||
| 139 | - AscendC::RoundMode::CAST_RINT, | ||
| 140 | - (uint64_t)0, | ||
| 141 | - repeat, | ||
| 142 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 143 | - } else { | ||
| 144 | - AscendC::Cast<DTypeOut, DTypeIn, false>( | ||
| 145 | - dst, | ||
| 146 | - src, | ||
| 147 | - AscendC::RoundMode::CAST_NONE, | ||
| 148 | - (uint64_t)0, | ||
| 149 | - repeat, | ||
| 150 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 151 | - } | ||
| 152 | -} | ||
| 153 | - | ||
| 154 | -///////////////////////////////////////////////////// | ||
| 155 | -// vconv_f322bf16r | ||
| 156 | -///////////////////////////////////////////////////// | ||
| 157 | -template <ArchType ArchTag, typename DTypeIn, typename DTypeOut> | ||
| 158 | -__aicore__ inline void convr_v(AscendC::LocalTensor<DTypeOut> dst, | ||
| 159 | - AscendC::LocalTensor<DTypeIn> src, | ||
| 160 | - uint8_t repeat, | ||
| 161 | - uint16_t dstBlockStride, | ||
| 162 | - uint16_t srcBlockStride, | ||
| 163 | - uint16_t dstRepeatStride, | ||
| 164 | - uint16_t srcRepeatStride) | ||
| 165 | -{ | ||
| 166 | - AscendC::Cast<DTypeOut, DTypeIn, false>( | ||
| 167 | - dst, | ||
| 168 | - src, | ||
| 169 | - AscendC::RoundMode::CAST_RINT, | ||
| 170 | - (uint64_t)0, | ||
| 171 | - repeat, | ||
| 172 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 173 | -} | ||
| 174 | - | ||
| 175 | -///////////////////////////////////////////////////// | ||
| 176 | -// vdiv | ||
| 177 | -///////////////////////////////////////////////////// | ||
| 178 | -template <ArchType ArchTag, typename DType> | ||
| 179 | -__aicore__ inline void div_v(AscendC::LocalTensor<DType> dst, | ||
| 180 | - AscendC::LocalTensor<DType> src0, | ||
| 181 | - AscendC::LocalTensor<DType> src1, | ||
| 182 | - uint8_t repeat, | ||
| 183 | - uint8_t dstBlockStride, | ||
| 184 | - uint8_t src0BlockStride, | ||
| 185 | - uint8_t src1BlockStride, | ||
| 186 | - uint8_t dstRepeatStride, | ||
| 187 | - uint8_t src0RepeatStride, | ||
| 188 | - uint8_t src1RepeatStride) | ||
| 189 | -{ | ||
| 190 | - AscendC::Div<DType, false>( | ||
| 191 | - dst, | ||
| 192 | - src0, | ||
| 193 | - src1, | ||
| 194 | - (uint64_t)0, | ||
| 195 | - repeat, | ||
| 196 | - AscendC::BinaryRepeatParams( | ||
| 197 | - dstBlockStride, src0BlockStride, src1BlockStride, dstRepeatStride, src0RepeatStride, src1RepeatStride)); | ||
| 198 | -} | ||
| 199 | - | ||
| 200 | -///////////////////////////////////////////////////// | ||
| 201 | -// vexp | ||
| 202 | -///////////////////////////////////////////////////// | ||
| 203 | -template <ArchType ArchTag, typename DType> | ||
| 204 | -__aicore__ inline void exp_v(AscendC::LocalTensor<DType> dst, | ||
| 205 | - AscendC::LocalTensor<DType> src, | ||
| 206 | - uint8_t repeat, | ||
| 207 | - uint16_t dstBlockStride, | ||
| 208 | - uint16_t srcBlockStride, | ||
| 209 | - uint16_t dstRepeatStride, | ||
| 210 | - uint16_t srcRepeatStride) | ||
| 211 | -{ | ||
| 212 | - AscendC::Exp<DType, false>( | ||
| 213 | - dst, | ||
| 214 | - src, | ||
| 215 | - (uint64_t)0, | ||
| 216 | - repeat, | ||
| 217 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 218 | -} | ||
| 219 | - | ||
| 220 | -///////////////////////////////////////////////////// | ||
| 221 | -// vmax | ||
| 222 | -///////////////////////////////////////////////////// | ||
| 223 | -template <ArchType ArchTag, typename DType> | ||
| 224 | -__aicore__ inline void max_v(AscendC::LocalTensor<DType> dst, | ||
| 225 | - AscendC::LocalTensor<DType> src0, | ||
| 226 | - AscendC::LocalTensor<DType> src1, | ||
| 227 | - uint8_t repeat, | ||
| 228 | - uint8_t dstBlockStride, | ||
| 229 | - uint8_t src0BlockStride, | ||
| 230 | - uint8_t src1BlockStride, | ||
| 231 | - uint8_t dstRepeatStride, | ||
| 232 | - uint8_t src0RepeatStride, | ||
| 233 | - uint8_t src1RepeatStride) | ||
| 234 | -{ | ||
| 235 | - AscendC::Max<DType, false>( | ||
| 236 | - dst, | ||
| 237 | - src0, | ||
| 238 | - src1, | ||
| 239 | - (uint64_t)0, | ||
| 240 | - repeat, | ||
| 241 | - AscendC::BinaryRepeatParams( | ||
| 242 | - dstBlockStride, src0BlockStride, src1BlockStride, dstRepeatStride, src0RepeatStride, src1RepeatStride)); | ||
| 243 | -} | ||
| 244 | - | ||
| 245 | -///////////////////////////////////////////////////// | ||
| 246 | -// vmul | ||
| 247 | -///////////////////////////////////////////////////// | ||
| 248 | -template <ArchType ArchTag, typename DType> | ||
| 249 | -__aicore__ inline void mul_v(AscendC::LocalTensor<DType> dst, | ||
| 250 | - AscendC::LocalTensor<DType> src0, | ||
| 251 | - AscendC::LocalTensor<DType> src1, | ||
| 252 | - uint8_t repeat, | ||
| 253 | - uint8_t dstBlockStride, | ||
| 254 | - uint8_t src0BlockStride, | ||
| 255 | - uint8_t src1BlockStride, | ||
| 256 | - uint8_t dstRepeatStride, | ||
| 257 | - uint8_t src0RepeatStride, | ||
| 258 | - uint8_t src1RepeatStride) | ||
| 259 | -{ | ||
| 260 | - AscendC::Mul<DType, false>( | ||
| 261 | - dst, | ||
| 262 | - src0, | ||
| 263 | - src1, | ||
| 264 | - (uint64_t)0, | ||
| 265 | - repeat, | ||
| 266 | - AscendC::BinaryRepeatParams( | ||
| 267 | - dstBlockStride, src0BlockStride, src1BlockStride, dstRepeatStride, src0RepeatStride, src1RepeatStride)); | ||
| 268 | -} | ||
| 269 | - | ||
| 270 | -///////////////////////////////////////////////////// | ||
| 271 | -// vmuls | ||
| 272 | -///////////////////////////////////////////////////// | ||
| 273 | -template <ArchType ArchTag, typename DType> | ||
| 274 | -__aicore__ inline void muls_v(AscendC::LocalTensor<DType> dst, | ||
| 275 | - AscendC::LocalTensor<DType> src0, | ||
| 276 | - DType src1, | ||
| 277 | - uint8_t repeat, | ||
| 278 | - uint16_t dstBlockStride, | ||
| 279 | - uint16_t srcBlockStride, | ||
| 280 | - uint16_t dstRepeatStride, | ||
| 281 | - uint16_t srcRepeatStride) | ||
| 282 | -{ | ||
| 283 | - AscendC::Muls<DType, false>( | ||
| 284 | - dst, | ||
| 285 | - src0, | ||
| 286 | - src1, | ||
| 287 | - (uint64_t)0, | ||
| 288 | - repeat, | ||
| 289 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 290 | -} | ||
| 291 | - | ||
| 292 | -///////////////////////////////////////////////////// | ||
| 293 | -// vsub | ||
| 294 | -///////////////////////////////////////////////////// | ||
| 295 | -template <ArchType ArchTag, typename DType> | ||
| 296 | -__aicore__ inline void sub_v(AscendC::LocalTensor<DType> dst, | ||
| 297 | - AscendC::LocalTensor<DType> src0, | ||
| 298 | - AscendC::LocalTensor<DType> src1, | ||
| 299 | - uint8_t repeat, | ||
| 300 | - uint8_t dstBlockStride, | ||
| 301 | - uint8_t src0BlockStride, | ||
| 302 | - uint8_t src1BlockStride, | ||
| 303 | - uint8_t dstRepeatStride, | ||
| 304 | - uint8_t src0RepeatStride, | ||
| 305 | - uint8_t src1RepeatStride) | ||
| 306 | -{ | ||
| 307 | - AscendC::Sub<DType, false>( | ||
| 308 | - dst, | ||
| 309 | - src0, | ||
| 310 | - src1, | ||
| 311 | - (uint64_t)0, | ||
| 312 | - repeat, | ||
| 313 | - AscendC::BinaryRepeatParams( | ||
| 314 | - dstBlockStride, src0BlockStride, src1BlockStride, dstRepeatStride, src0RepeatStride, src1RepeatStride)); | ||
| 315 | -} | ||
| 316 | - | ||
| 317 | -///////////////////////////////////////////////////// | ||
| 318 | -// vmaxs | ||
| 319 | -///////////////////////////////////////////////////// | ||
| 320 | -template <ArchType ArchTag, typename DType> | ||
| 321 | -__aicore__ inline void maxs_v(AscendC::LocalTensor<DType> dst, | ||
| 322 | - AscendC::LocalTensor<DType> src0, | ||
| 323 | - DType src1, | ||
| 324 | - uint8_t repeat, | ||
| 325 | - uint16_t dstBlockStride, | ||
| 326 | - uint16_t srcBlockStride, | ||
| 327 | - uint16_t dstRepeatStride, | ||
| 328 | - uint16_t srcRepeatStride) | ||
| 329 | -{ | ||
| 330 | - AscendC::Maxs<DType, false>( | ||
| 331 | - dst, | ||
| 332 | - src0, | ||
| 333 | - src1, | ||
| 334 | - (uint64_t)0, | ||
| 335 | - repeat, | ||
| 336 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 337 | -} | ||
| 338 | - | ||
| 339 | -///////////////////////////////////////////////////// | ||
| 340 | -// vmins | ||
| 341 | -///////////////////////////////////////////////////// | ||
| 342 | -template <ArchType ArchTag, typename DType> | ||
| 343 | -__aicore__ inline void mins_v(AscendC::LocalTensor<DType> dst, | ||
| 344 | - AscendC::LocalTensor<DType> src0, | ||
| 345 | - DType src1, | ||
| 346 | - uint8_t repeat, | ||
| 347 | - uint16_t dstBlockStride, | ||
| 348 | - uint16_t srcBlockStride, | ||
| 349 | - uint16_t dstRepeatStride, | ||
| 350 | - uint16_t srcRepeatStride) | ||
| 351 | -{ | ||
| 352 | - AscendC::Mins<DType, false>( | ||
| 353 | - dst, | ||
| 354 | - src0, | ||
| 355 | - src1, | ||
| 356 | - (uint64_t)0, | ||
| 357 | - repeat, | ||
| 358 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 359 | -} | ||
| 360 | - | ||
| 361 | -///////////////////////////////////////////////////// | ||
| 362 | -// vsqrt | ||
| 363 | -///////////////////////////////////////////////////// | ||
| 364 | -template <ArchType ArchTag, typename DType> | ||
| 365 | -__aicore__ inline void sqrt_v(AscendC::LocalTensor<DType> dst, | ||
| 366 | - AscendC::LocalTensor<DType> src, | ||
| 367 | - uint8_t repeat, | ||
| 368 | - uint16_t dstBlockStride, | ||
| 369 | - uint16_t srcBlockStride, | ||
| 370 | - uint16_t dstRepeatStride, | ||
| 371 | - uint16_t srcRepeatStride) | ||
| 372 | -{ | ||
| 373 | - AscendC::Sqrt<DType, false>( | ||
| 374 | - dst, | ||
| 375 | - src, | ||
| 376 | - (uint64_t)0, | ||
| 377 | - repeat, | ||
| 378 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 379 | -} | ||
| 380 | - | ||
| 381 | -///////////////////////////////////////////////////// | ||
| 382 | -// vln | ||
| 383 | -///////////////////////////////////////////////////// | ||
| 384 | -template <ArchType ArchTag, typename DType> | ||
| 385 | -__aicore__ inline void ln_v(AscendC::LocalTensor<DType> dst, | ||
| 386 | - AscendC::LocalTensor<DType> src, | ||
| 387 | - uint8_t repeat, | ||
| 388 | - uint16_t dstBlockStride, | ||
| 389 | - uint16_t srcBlockStride, | ||
| 390 | - uint16_t dstRepeatStride, | ||
| 391 | - uint16_t srcRepeatStride) | ||
| 392 | -{ | ||
| 393 | - AscendC::Ln<DType, false>( | ||
| 394 | - dst, | ||
| 395 | - src, | ||
| 396 | - (uint64_t)0, | ||
| 397 | - repeat, | ||
| 398 | - AscendC::UnaryRepeatParams(dstBlockStride, srcBlockStride, dstRepeatStride, srcRepeatStride)); | ||
| 399 | -} | ||
| 400 | - | ||
| 401 | -///////////////////////////////////////////////////// | ||
| 402 | -// vtranspose | ||
| 403 | -///////////////////////////////////////////////////// | ||
| 404 | -template <ArchType ArchTag, typename DType> | ||
| 405 | -__aicore__ inline void tranpose_v(AscendC::LocalTensor<DType> dst, AscendC::LocalTensor<DType> src) | ||
| 406 | -{ | ||
| 407 | - AscendC::Transpose(dst, src); | ||
| 408 | -} | ||
| 409 | - | ||
| 410 | -///////////////////////////////////////////////////// | ||
| 411 | -// vcgmax | ||
| 412 | -///////////////////////////////////////////////////// | ||
| 413 | -template <ArchType ArchTag, typename DType> | ||
| 414 | -__aicore__ inline void cgmax_v(AscendC::LocalTensor<DType> dst, | ||
| 415 | - AscendC::LocalTensor<DType> src, | ||
| 416 | - const int32_t repeat, | ||
| 417 | - const int32_t dstRepStride, | ||
| 418 | - const int32_t srcBlkStride, | ||
| 419 | - const int32_t srcRepStride) | ||
| 420 | -{ | ||
| 421 | - AscendC::BlockReduceMax<DType, false>(dst, src, repeat, 0, dstRepStride, srcBlkStride, srcRepStride); | ||
| 422 | -} | ||
| 423 | - | ||
| @@ -1,70 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | -template <typename IN_DTYPE> | ||
| 14 | -__aicore__ inline void CreateCaMatrix(const AscendC::LocalTensor<IN_DTYPE> &dst, | ||
| 15 | - const uint16_t repeats, | ||
| 16 | - const uint16_t blockNum, | ||
| 17 | - const uint16_t dstGap, | ||
| 18 | - const IN_DTYPE initValue) | ||
| 19 | -{ | ||
| 20 | - AscendC::InitConstValue<IN_DTYPE>(dst, | ||
| 21 | - AscendC::InitConstValueParams<IN_DTYPE>(repeats, blockNum, dstGap, initValue)); | ||
| 22 | -} | ||
| 23 | -__aicore__ inline void SetFftsBaseAddr(uint64_t config) | ||
| 24 | -{ | ||
| 25 | - AscendC::SetSyncBaseAddr(config); | ||
| 26 | -} | ||
| 27 | -template <typename IN_DTYPE> | ||
| 28 | -__aicore__ inline void SetPadding(IN_DTYPE padValue) | ||
| 29 | -{ | ||
| 30 | - AscendC::SetLoadDataPaddingValue<IN_DTYPE>(padValue); | ||
| 31 | -} | ||
| 32 | -__aicore__ inline void SetAtomicnone() | ||
| 33 | -{ | ||
| 34 | - AscendC::SetAtomicNone(); | ||
| 35 | -} | ||
| 36 | -__aicore__ inline void SetMasknorm() | ||
| 37 | -{ | ||
| 38 | - | ||
| 39 | - return; | ||
| 40 | - | ||
| 41 | - AscendC::SetMaskNorm(); | ||
| 42 | -} | ||
| 43 | -__aicore__ inline void SetNdpara(uint16_t ndNum, uint16_t srcNdStride, uint16_t dstNdStride) | ||
| 44 | -{ | ||
| 45 | - AscendC::SetFixpipeNz2ndFlag(ndNum, srcNdStride, dstNdStride); | ||
| 46 | -} | ||
| 47 | -template <typename IN_DTYPE> | ||
| 48 | -__aicore__ inline void SetVectorMask(const uint64_t maskHigh, const uint64_t maskLow) | ||
| 49 | -{ | ||
| 50 | - AscendC::SetVectorMask<IN_DTYPE>(maskHigh, maskLow); | ||
| 51 | -} | ||
| 52 | -__aicore__ inline int64_t GetSubBlockidx() | ||
| 53 | -{ | ||
| 54 | - return AscendC::GetSubBlockIdx(); | ||
| 55 | -} | ||
| 56 | -__aicore__ inline void WaitFlagDev(uint16_t flagId) | ||
| 57 | -{ | ||
| 58 | - AscendC::WaitEvent(flagId); | ||
| 59 | -} | ||
| 60 | -template <pipe_t pipe, uint8_t mode> | ||
| 61 | -__aicore__ inline void FftsCrossCoreSync(uint16_t flagId) | ||
| 62 | -{ | ||
| 63 | - AscendC::CrossCoreSetFlag<mode, pipe>(flagId); | ||
| 64 | -} | ||
| 65 | -template <typename IN_DTYPE, bool setRelu = false> | ||
| 66 | -__aicore__ inline void SetFpc(const AscendC::LocalTensor<IN_DTYPE> &preTensor, bool isUnitFlag = false) | ||
| 67 | -{ | ||
| 68 | - AscendC::SetFixPipeConfig<IN_DTYPE, setRelu>(preTensor, isUnitFlag); | ||
| 69 | -} | ||
| 70 | - | ||
| @@ -1,76 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | -/*! | ||
| 12 | - * \file matmul_v2_compress_dequant.cpp | ||
| 13 | - * \brief matmul_v2_compress_dequant kernel entry, ported from pp_matmul_i8_nz_compress.cce | ||
| 14 | - */ | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | -extern "C" __global__ __aicore__ void mat_mul_v2_compress_dequant( | ||
| 21 | - GM_ADDR x1, | ||
| 22 | - GM_ADDR x2, | ||
| 23 | - GM_ADDR compress_index, | ||
| 24 | - GM_ADDR deq_scale, | ||
| 25 | - GM_ADDR bias, | ||
| 26 | - GM_ADDR offset_w, | ||
| 27 | - GM_ADDR out, | ||
| 28 | - GM_ADDR workspace, | ||
| 29 | - GM_ADDR tiling) | ||
| 30 | -{ | ||
| 31 | - SetPadding<uint64_t>((uint64_t)0x0); | ||
| 32 | - SetVectorMask<int8_t>((uint64_t)-1, (uint64_t)-1); | ||
| 33 | - SetAtomicnone(); | ||
| 34 | - | ||
| 35 | - PpMatmulI8NzCompress<0, false, true, false, int8_t, uint64_t, int32_t> kernel; | ||
| 36 | - | ||
| 37 | - AscendC::GlobalTensor<int32_t> gm_tiling; | ||
| 38 | - gm_tiling.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t *>(tiling)); | ||
| 39 | - AscendC::LocalTensor<int32_t> ub_tiling = | ||
| 40 | - kernel.buf.GetBuffer<BufferType::ASCEND_UB, int32_t>(0); | ||
| 41 | - gm_to_ub<ArchType::ASCEND_V200, int32_t>(ub_tiling, gm_tiling, | ||
| 42 | - 0, | ||
| 43 | - 1, | ||
| 44 | - 512 / 32, | ||
| 45 | - 0, | ||
| 46 | - 0); | ||
| 47 | - SET_FLAG(MTE2, S, EVENT_ID0); | ||
| 48 | - WAIT_FLAG(MTE2, S, EVENT_ID0); | ||
| 49 | - | ||
| 50 | - kernel.Init( | ||
| 51 | - reinterpret_cast<__gm__ uint8_t *>(x1), | ||
| 52 | - reinterpret_cast<__gm__ uint8_t *>(x2), | ||
| 53 | - reinterpret_cast<__gm__ uint8_t *>(bias), | ||
| 54 | - reinterpret_cast<__gm__ uint8_t *>(deq_scale), | ||
| 55 | - reinterpret_cast<__gm__ uint8_t *>(compress_index), | ||
| 56 | - reinterpret_cast<__gm__ uint8_t *>(out), | ||
| 57 | - ub_tiling.GetValue(0), // batchSize | ||
| 58 | - ub_tiling.GetValue(1), // m | ||
| 59 | - ub_tiling.GetValue(2), // k | ||
| 60 | - ub_tiling.GetValue(3), // n | ||
| 61 | - ub_tiling.GetValue(4), // m0 | ||
| 62 | - ub_tiling.GetValue(5), // k0 | ||
| 63 | - ub_tiling.GetValue(6), // n0 | ||
| 64 | - ub_tiling.GetValue(7), // mLoop | ||
| 65 | - ub_tiling.GetValue(8), // kLoop | ||
| 66 | - ub_tiling.GetValue(9), // nLoop | ||
| 67 | - ub_tiling.GetValue(10), // coreLoop | ||
| 68 | - ub_tiling.GetValue(11), // swizzlCount | ||
| 69 | - ub_tiling.GetValue(12), // tilingK | ||
| 70 | - ub_tiling.GetValue(13), // tilingN | ||
| 71 | - ub_tiling.GetValue(14)); // compressOverlapN | ||
| 72 | - | ||
| 73 | - kernel.Process(); | ||
| 74 | -} | ||
| 75 | - | ||
| 76 | - | ||
| @@ -1,412 +0,0 @@ | |||
| 1 | -/* | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | - | ||
| 15 | - | ||
| 16 | - | ||
| 17 | - | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | -constexpr uint32_t BLOCK_SIZE_32 = 32; | ||
| 24 | -constexpr uint32_t CUBE_BLOCK_SIZE_INT8 = 512; | ||
| 25 | -constexpr uint32_t BLOCK_SIZE_16 = 16; | ||
| 26 | -constexpr uint32_t BLOCK_SIZE_K = 32; | ||
| 27 | -constexpr uint32_t M_TILE_LIMIT = 64; | ||
| 28 | -constexpr uint32_t K_TILE_LIMIT = 256; | ||
| 29 | -constexpr uint32_t L0AB_PINGPONG_BUFFER_LEN = 32 * 1024; | ||
| 30 | - | ||
| 31 | -__aicore__ __force_inline__ uint32_t CeilDivKernel(const uint32_t dividend, const uint32_t divisor) | ||
| 32 | -{ | ||
| 33 | - if (divisor == 0) { | ||
| 34 | - return UINT32_MAX; | ||
| 35 | - } | ||
| 36 | - return (dividend + divisor - 1) / divisor; | ||
| 37 | -} | ||
| 38 | - | ||
| 39 | -__aicore__ __force_inline__ uint32_t RoundUpKernel(const uint32_t val, const uint32_t align) | ||
| 40 | -{ | ||
| 41 | - if (align == 0) { | ||
| 42 | - return 0; | ||
| 43 | - } | ||
| 44 | - return (val + align - 1) / align * align; | ||
| 45 | -} | ||
| 46 | - | ||
| 47 | -__aicore__ __force_inline__ uint32_t MinKernel(const uint32_t a, const uint32_t b) | ||
| 48 | -{ | ||
| 49 | - return a < b ? a : b; | ||
| 50 | -} | ||
| 51 | - | ||
| 52 | - | ||
| 53 | - | ||
| 54 | -template <uint32_t SWIZZL_DIR, bool TRANSPOSE_A, bool TRANSPOSE_B, bool SPLIT_K = false, | ||
| 55 | - typename IN_DTYPE = int8_t, typename DESCALE_TYPE = uint64_t, typename BIAS_TYPE = int32_t> | ||
| 56 | -class PpMatmulI8NzCompress { | ||
| 57 | -public: | ||
| 58 | - __aicore__ explicit PpMatmulI8NzCompress() | ||
| 59 | - { | ||
| 60 | - SetPadding<uint64_t>((uint64_t)0x0); | ||
| 61 | - SetAtomicnone(); | ||
| 62 | - set_ctrl(sbitset1(get_ctrl(), 62)); | ||
| 63 | - SetMasknorm(); | ||
| 64 | - } | ||
| 65 | - | ||
| 66 | - __aicore__ __force_inline__ void Init( | ||
| 67 | - __gm__ uint8_t *__restrict__ A, __gm__ uint8_t *__restrict__ B, | ||
| 68 | - __gm__ uint8_t *__restrict__ bias, __gm__ uint8_t *__restrict__ scale, | ||
| 69 | - __gm__ uint8_t *__restrict__ compress_index, __gm__ uint8_t *__restrict__ C, | ||
| 70 | - int32_t b, int32_t m, int32_t k, int32_t n, | ||
| 71 | - int32_t m0, int32_t k0, int32_t n0, | ||
| 72 | - int32_t m_loop, int32_t k_loop, int32_t n_loop, | ||
| 73 | - int32_t core_loop, int32_t swizzl_cnt, | ||
| 74 | - int32_t copress_tiling_k, int32_t copress_tiling_n, int32_t compress_overlap_n) | ||
| 75 | - { | ||
| 76 | - gm_a.SetGlobalBuffer(reinterpret_cast<__gm__ IN_DTYPE *>(A)); | ||
| 77 | - gm_b.SetGlobalBuffer(reinterpret_cast<__gm__ IN_DTYPE *>(B)); | ||
| 78 | - gm_c.SetGlobalBuffer(reinterpret_cast<__gm__ half *>(C)); | ||
| 79 | - gm_bias.SetGlobalBuffer(reinterpret_cast<__gm__ BIAS_TYPE *>(bias)); | ||
| 80 | - gm_scale.SetGlobalBuffer(reinterpret_cast<__gm__ DESCALE_TYPE *>(scale)); | ||
| 81 | - gm_compress_index.SetGlobalBuffer(reinterpret_cast<__gm__ IN_DTYPE *>(compress_index)); | ||
| 82 | - b_ = b; | ||
| 83 | - m_ = m; | ||
| 84 | - k_ = k; | ||
| 85 | - n_ = n; | ||
| 86 | - m0_ = m0; | ||
| 87 | - k0_ = m0 > M_TILE_LIMIT ? MinKernel(k0, K_TILE_LIMIT) : k0; | ||
| 88 | - n0_ = n0; | ||
| 89 | - m_org_ = m; | ||
| 90 | - n_org_ = n; | ||
| 91 | - m_loop_ = m_loop; | ||
| 92 | - k_loop_ = CeilDivKernel(k, k0_); | ||
| 93 | - n_loop_ = n_loop; | ||
| 94 | - core_loop_ = core_loop; | ||
| 95 | - swizzl_cnt_ = swizzl_cnt; | ||
| 96 | - copress_tiling_k_ = copress_tiling_k; | ||
| 97 | - copress_tiling_n_ = copress_tiling_n; | ||
| 98 | - compress_overlap_n_ = compress_overlap_n; | ||
| 99 | - n_compress_num = CeilDivKernel(n_, copress_tiling_n_ * 16); | ||
| 100 | - k_compress_num = CeilDivKernel(k_, copress_tiling_k_ * 32); | ||
| 101 | - compress_size = copress_tiling_k_ * copress_tiling_n_ * CUBE_BLOCK_SIZE_INT8; | ||
| 102 | - | ||
| 103 | - l1_a_ping = buf.template GetBuffer<BufferType::ASCEND_CB, IN_DTYPE>(0); | ||
| 104 | - l1_a_pong = buf.template GetBuffer<BufferType::ASCEND_CB, IN_DTYPE>(256 * 1024); | ||
| 105 | - l1_b_ping = buf.template GetBuffer<BufferType::ASCEND_CB, IN_DTYPE>(512 * 1024); | ||
| 106 | - l1_b_pong = buf.template GetBuffer<BufferType::ASCEND_CB, IN_DTYPE>(768 * 1024); | ||
| 107 | - l0_a_ping = buf.template GetBuffer<BufferType::ASCEND_L0A, IN_DTYPE>(0); | ||
| 108 | - l0_a_pong = buf.template GetBuffer<BufferType::ASCEND_L0A, IN_DTYPE>(32768); | ||
| 109 | - l0_b_ping = buf.template GetBuffer<BufferType::ASCEND_L0B, IN_DTYPE>(0); | ||
| 110 | - l0_b_pong = buf.template GetBuffer<BufferType::ASCEND_L0B, IN_DTYPE>(32768); | ||
| 111 | - ub_c = buf.template GetBuffer<BufferType::ASCEND_UB, half>(0 + 8 * 1024); | ||
| 112 | - ub_bias = buf.template GetBuffer<BufferType::ASCEND_UB, BIAS_TYPE>(131072 + 8 * 1024); | ||
| 113 | - ub_scale = buf.template GetBuffer<BufferType::ASCEND_UB, DESCALE_TYPE>(0); | ||
| 114 | - } | ||
| 115 | - | ||
| 116 | - __aicore__ __force_inline__ void GetIdx(uint32_t loop_idx, uint32_t &m_idx, uint32_t &n_idx) | ||
| 117 | - { | ||
| 118 | - uint32_t in_batch_idx = loop_idx % (m_loop_ * n_loop_); | ||
| 119 | - if constexpr (SWIZZL_DIR == 0) { | ||
| 120 | - int32_t tile_block_loop = (m_loop_ + swizzl_cnt_ - 1) / swizzl_cnt_; | ||
| 121 | - int32_t tile_block_idx = in_batch_idx / (swizzl_cnt_ * n_loop_); | ||
| 122 | - int32_t in_tile_block_idx = in_batch_idx % (swizzl_cnt_ * n_loop_); | ||
| 123 | - | ||
| 124 | - int32_t n_row = swizzl_cnt_; | ||
| 125 | - if (tile_block_idx == tile_block_loop - 1) { | ||
| 126 | - n_row = m_loop_ - swizzl_cnt_ * tile_block_idx; | ||
| 127 | - } | ||
| 128 | - m_idx = tile_block_idx * swizzl_cnt_ + in_tile_block_idx % n_row; | ||
| 129 | - n_idx = in_tile_block_idx / n_row; | ||
| 130 | - if (tile_block_idx % 2 != 0) { | ||
| 131 | - n_idx = n_loop_ - n_idx - 1; | ||
| 132 | - } | ||
| 133 | - } else if constexpr (SWIZZL_DIR == 1) { | ||
| 134 | - int32_t tile_block_loop = (n_loop_ + swizzl_cnt_ - 1) / swizzl_cnt_; | ||
| 135 | - int32_t tile_block_idx = in_batch_idx / (swizzl_cnt_ * m_loop_); | ||
| 136 | - int32_t in_tile_block_idx = in_batch_idx % (swizzl_cnt_ * m_loop_); | ||
| 137 | - | ||
| 138 | - int32_t n_col = swizzl_cnt_; | ||
| 139 | - if (tile_block_idx == tile_block_loop - 1) { | ||
| 140 | - n_col = n_loop_ - swizzl_cnt_ * tile_block_idx; | ||
| 141 | - } | ||
| 142 | - m_idx = in_tile_block_idx / n_col; | ||
| 143 | - n_idx = tile_block_idx * swizzl_cnt_ + in_tile_block_idx % n_col; | ||
| 144 | - if (tile_block_idx % 2 != 0) { | ||
| 145 | - m_idx = m_loop_ - m_idx - 1; | ||
| 146 | - } | ||
| 147 | - } | ||
| 148 | - } | ||
| 149 | - | ||
| 150 | - __aicore__ __force_inline__ void Process() | ||
| 151 | - { | ||
| 152 | - SET_FLAG(MTE1, MTE2, EVENT_ID0); | ||
| 153 | - SET_FLAG(MTE1, MTE2, EVENT_ID1); | ||
| 154 | - SET_FLAG(MTE1, MTE2, EVENT_ID2); | ||
| 155 | - SET_FLAG(MTE1, MTE2, EVENT_ID3); | ||
| 156 | - SET_FLAG(M, MTE1, EVENT_ID0); | ||
| 157 | - SET_FLAG(M, MTE1, EVENT_ID1); | ||
| 158 | - SET_FLAG(V, M, EVENT_ID0); | ||
| 159 | - SET_FLAG(MTE3, V, EVENT_ID0); | ||
| 160 | - uint32_t curr_core_loop = 0; | ||
| 161 | - uint32_t l1_ping_pong = 1; | ||
| 162 | - for (uint32_t loop_idx = 0; loop_idx < core_loop_; ++loop_idx) { | ||
| 163 | - if (loop_idx % block_num != block_idx) { | ||
| 164 | - continue; | ||
| 165 | - } | ||
| 166 | - uint64_t b_idx = loop_idx / (m_loop_ * n_loop_); | ||
| 167 | - uint32_t m_idx = 0, n_idx = 0; | ||
| 168 | - GetIdx(loop_idx, m_idx, n_idx); | ||
| 169 | - uint32_t m_actual = (m_idx == (m_loop_ - 1)) ? (m_ - m_idx * m0_) : m0_; | ||
| 170 | - uint32_t n_actual = (n_idx == (n_loop_ - 1)) ? (n_ - n_idx * n0_) : n0_; | ||
| 171 | - uint32_t m_round = RoundUpKernel(m_actual, 16); | ||
| 172 | - uint32_t n_round = RoundUpKernel(n_actual, 16); | ||
| 173 | - uint32_t k_actual = (k_loop_ == 1) ? k_ : k0_; | ||
| 174 | - uint32_t k_round = RoundUpKernel(k_actual, BLOCK_SIZE_32); | ||
| 175 | - uint64_t src_offset = 0, dst_offset = 0; | ||
| 176 | - uint64_t src_offset_index = 0, src_offset_compress_L1_size = 0; | ||
| 177 | - uint32_t mn_max = m_round > n_round ? m_round : n_round; | ||
| 178 | - uint32_t k_part_len = L0AB_PINGPONG_BUFFER_LEN / mn_max / 32 * 32; | ||
| 179 | - uint32_t repeat_n = copress_tiling_n_; | ||
| 180 | - uint32_t m_org_up = RoundUpKernel(m_, 16); | ||
| 181 | - uint32_t n_org_up = RoundUpKernel(n_, 16); | ||
| 182 | - uint32_t k_org_up = RoundUpKernel(k_, 32); | ||
| 183 | - | ||
| 184 | - uint32_t index_k_all = CeilDivKernel(k_, copress_tiling_k_ * BLOCK_SIZE_32); | ||
| 185 | - src_offset_index = (b_idx * n_compress_num * k_compress_num + | ||
| 186 | - n_idx * CeilDivKernel(n0_, copress_tiling_n_ * BLOCK_SIZE_16) * k_compress_num) * | ||
| 187 | - 8; | ||
| 188 | - load_unzip_index_from_gm(((__gm__ IN_DTYPE *)(gm_compress_index.GetPhyAddr() + src_offset_index)), | ||
| 189 | - (uint64_t)index_k_all); | ||
| 190 | - for (uint32_t k_idx = 0; k_idx < k_loop_; ++k_idx) { | ||
| 191 | - uint32_t k_actual = k_idx == k_loop_ - 1 ? k_ - k_idx * k0_ : k0_; | ||
| 192 | - uint32_t k_round = RoundUpKernel(k_actual, 32); | ||
| 193 | - | ||
| 194 | - AscendC::LocalTensor<int8_t> l1_a = l1_ping_pong ? l1_a_ping : l1_a_pong; | ||
| 195 | - AscendC::LocalTensor<int8_t> l1_b = l1_ping_pong ? l1_b_ping : l1_b_pong; | ||
| 196 | - | ||
| 197 | - event_t l1_a_event = l1_ping_pong ? EVENT_ID0 : EVENT_ID1; | ||
| 198 | - event_t l1_b_event = l1_ping_pong ? EVENT_ID2 : EVENT_ID3; | ||
| 199 | - uint32_t index_num_n = CeilDivKernel(n_round, copress_tiling_n_ * BLOCK_SIZE_16); | ||
| 200 | - uint32_t index_num_k = CeilDivKernel(k_round, copress_tiling_k_ * BLOCK_SIZE_32); | ||
| 201 | - WAIT_FLAG(MTE1, MTE2, l1_b_event); | ||
| 202 | - for (uint32_t _n_idx = 0; _n_idx < index_num_n; _n_idx++) { | ||
| 203 | - if (k_idx == k_loop_ - 1) { | ||
| 204 | - uint32_t copress_k_tile = | ||
| 205 | - (k_round - (index_num_k - 1) * (copress_tiling_k_ * BLOCK_SIZE_32)) / BLOCK_SIZE_32; | ||
| 206 | - uint32_t compress_tile_size = copress_k_tile * copress_tiling_n_ * CUBE_BLOCK_SIZE_INT8; | ||
| 207 | - src_offset_compress_L1_size = _n_idx * ((index_num_k - 1) * compress_size + compress_tile_size); | ||
| 208 | - } else { | ||
| 209 | - src_offset_compress_L1_size = _n_idx * index_num_k * compress_size; | ||
| 210 | - } | ||
| 211 | - | ||
| 212 | - for (uint32_t _k_idx = 0; _k_idx < index_num_k; _k_idx++) { | ||
| 213 | - load_gm_to_cbuf_unzip((__cbuf__ half *)(((__cbuf__ IN_DTYPE *)l1_b.GetPhyAddr() + | ||
| 214 | - src_offset_compress_L1_size + _k_idx * compress_size)), | ||
| 215 | - (__gm__ half *)(((__gm__ IN_DTYPE *)gm_b.GetPhyAddr()))); | ||
| 216 | - } | ||
| 217 | - } | ||
| 218 | - SET_FLAG(MTE2, MTE1, l1_b_event); | ||
| 219 | - WAIT_FLAG(MTE1, MTE2, l1_a_event); | ||
| 220 | - | ||
| 221 | - src_offset = b_idx * m_org_up * k_org_up + k_idx * k0_ * m_org_up + m_idx * m0_ * BLOCK_SIZE_32; | ||
| 222 | - gm_to_l1<ArchType::ASCEND_V200, IN_DTYPE, DataFormat::NZ, DataFormat::NZ>( | ||
| 223 | - l1_a, gm_a[src_offset], m_actual, m_round, m_org_up, k_actual, k_round, k_org_up); | ||
| 224 | - | ||
| 225 | - SET_FLAG(MTE2, MTE1, l1_a_event); | ||
| 226 | - AscendC::PipeBarrier<PIPE_MTE2>(); | ||
| 227 | - | ||
| 228 | - uint32_t k_part_loop = CeilDivKernel(k_actual, k_part_len); | ||
| 229 | - for (uint32_t k_part_idx = 0; k_part_idx < k_part_loop; ++k_part_idx) { | ||
| 230 | - uint32_t k0_round = k_part_idx < k_part_loop - 1 | ||
| 231 | - ? k_part_len | ||
| 232 | - : k_round - k_part_idx * k_part_len; | ||
| 233 | - uint32_t k0_actual = k_part_idx < k_part_loop - 1 | ||
| 234 | - ? k_part_len | ||
| 235 | - : k_actual - k_part_idx * k_part_len; | ||
| 236 | - uint32_t l0_ping_pong = 1 - k_part_idx % 2; | ||
| 237 | - event_t l0_event = l0_ping_pong ? EVENT_ID0 : EVENT_ID1; | ||
| 238 | - AscendC::LocalTensor<IN_DTYPE> l0_a = l0_ping_pong ? l0_a_ping : l0_a_pong; | ||
| 239 | - AscendC::LocalTensor<IN_DTYPE> l0_b = l0_ping_pong ? l0_b_ping : l0_b_pong; | ||
| 240 | - if (k_part_idx == 0) { | ||
| 241 | - WAIT_FLAG(MTE2, MTE1, l1_a_event); | ||
| 242 | - } | ||
| 243 | - WAIT_FLAG(M, MTE1, l0_event); | ||
| 244 | - l1_to_l0_a<ArchType::ASCEND_V200, IN_DTYPE, false, DataFormat::ZN, DataFormat::ZZ>( | ||
| 245 | - l0_a, l1_a[k_part_idx * k_part_len * m_round], m_round, k0_round, 1, m_round / 16, | ||
| 246 | - k0_round / 32, 1); | ||
| 247 | - if (k_part_idx == k_part_loop - 1) { | ||
| 248 | - SET_FLAG(MTE1, MTE2, l1_a_event); | ||
| 249 | - } | ||
| 250 | - if (k_part_idx == 0) { | ||
| 251 | - WAIT_FLAG(MTE2, MTE1, l1_b_event); | ||
| 252 | - } | ||
| 253 | - | ||
| 254 | - if (n_idx == (n_loop_ - 1) && (compress_overlap_n_ > 0)) { | ||
| 255 | - for (uint32_t i = 0; i < k0_round / 32; i++) { | ||
| 256 | - dst_offset = n_round * i * 32; | ||
| 257 | - src_offset = k_part_idx * k_part_len * n0_ + i * n0_ * 32 + compress_overlap_n_ * 512; | ||
| 258 | - l1_to_l0_b<ArchType::ASCEND_V200, IN_DTYPE, true, DataFormat::ZN, DataFormat::NZ>( | ||
| 259 | - l0_b[dst_offset], | ||
| 260 | - l1_b[src_offset], | ||
| 261 | - n_round, BLOCK_SIZE_32, 1, n0_ / 16, 1, n_round / 16); | ||
| 262 | - } | ||
| 263 | - } else { | ||
| 264 | - src_offset = k_part_idx * k_part_len * n_round; | ||
| 265 | - l1_to_l0_b<ArchType::ASCEND_V200, IN_DTYPE, false, DataFormat::VECTOR, DataFormat::VECTOR>( | ||
| 266 | - l0_b, | ||
| 267 | - l1_b[src_offset], | ||
| 268 | - 0, | ||
| 269 | - k0_round * n_round / CUBE_BLOCK_SIZE_INT8, | ||
| 270 | - 0, | ||
| 271 | - 1, | ||
| 272 | - 0, | ||
| 273 | - 0); | ||
| 274 | - } | ||
| 275 | - if (k_part_idx == k_part_loop - 1) { | ||
| 276 | - SET_FLAG(MTE1, MTE2, l1_b_event); | ||
| 277 | - } | ||
| 278 | - SET_FLAG(MTE1, M, l0_event); | ||
| 279 | - WAIT_FLAG(MTE1, M, l0_event); | ||
| 280 | - if (k_idx == 0 && k_part_idx == 0) { | ||
| 281 | - AscendC::PipeBarrier<PIPE_MTE2>(); | ||
| 282 | - gm_to_ub<ArchType::ASCEND_V200, BIAS_TYPE>(ub_bias, gm_bias[n_idx * n0_], | ||
| 283 | - 0, | ||
| 284 | - 1, | ||
| 285 | - n_round / 8, | ||
| 286 | - 0, | ||
| 287 | - 0); | ||
| 288 | - SET_FLAG(MTE2, V, EVENT_ID0); | ||
| 289 | - WAIT_FLAG(MTE2, V, EVENT_ID0); | ||
| 290 | - for (uint32_t i = 0; i < n_round / BLOCK_SIZE_16; i++) { | ||
| 291 | - for (uint32_t j = 0; j < m_round / BLOCK_SIZE_16; j++) { | ||
| 292 | - AscendC::BroadCastVecToMM( | ||
| 293 | - l0c[i * m_round * BLOCK_SIZE_16 + j * BLOCK_SIZE_16 * BLOCK_SIZE_16], | ||
| 294 | - ub_bias[i * BLOCK_SIZE_16], 1, 1, 0, 0); | ||
| 295 | - } | ||
| 296 | - } | ||
| 297 | - SET_FLAG(V, M, EVENT_ID1); | ||
| 298 | - WAIT_FLAG(V, M, EVENT_ID1); | ||
| 299 | - WAIT_FLAG(V, M, EVENT_ID0); | ||
| 300 | - } | ||
| 301 | - uint32_t m_mad_actual = (m_actual == 1) ? 2 : m_actual; | ||
| 302 | - AscendC::PipeBarrier<PIPE_M>(); | ||
| 303 | - mmad<ArchType::ASCEND_V200, IN_DTYPE, IN_DTYPE, int32_t, false>( | ||
| 304 | - l0c, l0_a, l0_b, | ||
| 305 | - m_mad_actual, | ||
| 306 | - n_actual, | ||
| 307 | - k0_actual, | ||
| 308 | - 0); | ||
| 309 | - SET_FLAG(M, MTE1, l0_event); | ||
| 310 | - } | ||
| 311 | - l1_ping_pong = 1 - l1_ping_pong; | ||
| 312 | - } | ||
| 313 | - curr_core_loop = curr_core_loop + 1; | ||
| 314 | - AscendC::PipeBarrier<PIPE_MTE2>(); | ||
| 315 | - gm_to_ub<ArchType::ASCEND_V200, DESCALE_TYPE>(ub_scale, gm_scale[n_idx * n0_], | ||
| 316 | - 0, | ||
| 317 | - 1, | ||
| 318 | - n_round / 4, | ||
| 319 | - 0, | ||
| 320 | - 0); | ||
| 321 | - SET_FLAG(M, V, EVENT_ID0); | ||
| 322 | - WAIT_FLAG(M, V, EVENT_ID0); | ||
| 323 | - WAIT_FLAG(MTE3, V, EVENT_ID0); | ||
| 324 | - SET_FLAG(MTE2, V, EVENT_ID1); | ||
| 325 | - WAIT_FLAG(MTE2, V, EVENT_ID1); | ||
| 326 | - l0c_to_ub<ArchType::ASCEND_V200, int32_t, half>(ub_c, l0c, | ||
| 327 | - (uint16_t)(n_round / BLOCK_SIZE_16), | ||
| 328 | - (uint16_t)(m_round / BLOCK_SIZE_16), | ||
| 329 | - (uint16_t)0, | ||
| 330 | - (uint16_t)0); | ||
| 331 | - AscendC::PipeBarrier<PIPE_V>(); | ||
| 332 | - if (m_actual == 1) { | ||
| 333 | - SetVectorMask<int8_t>((uint64_t)0x0, (uint64_t)0xffff); | ||
| 334 | - half zero = 0; | ||
| 335 | - for (uint32_t i = 0; i < n_round / 16; i++) { | ||
| 336 | - uint64_t curr_offset_c = i * m_round * BLOCK_SIZE_16 + m_actual * BLOCK_SIZE_16; | ||
| 337 | - muls_v<ArchType::ASCEND_V200, half>(ub_c[curr_offset_c], ub_c[curr_offset_c], | ||
| 338 | - zero, | ||
| 339 | - 1, | ||
| 340 | - 1, | ||
| 341 | - 1, | ||
| 342 | - 2, | ||
| 343 | - 2); | ||
| 344 | - } | ||
| 345 | - } | ||
| 346 | - SET_FLAG(V, M, EVENT_ID0); | ||
| 347 | - SET_FLAG(V, MTE3, EVENT_ID0); | ||
| 348 | - WAIT_FLAG(V, MTE3, EVENT_ID0); | ||
| 349 | - dst_offset = b_idx * m_org_up * n_org_up + n_idx * n0_ * m_org_up + m_idx * m0_ * 16; | ||
| 350 | - ub_to_gm<ArchType::ASCEND_V200, half, DataFormat::NZ, DataFormat::NZ>( | ||
| 351 | - gm_c[dst_offset], ub_c, m_round, m_round, m_org_up, n_round, n_round, n_org_up); | ||
| 352 | - SET_FLAG(MTE3, V, EVENT_ID0); | ||
| 353 | - } | ||
| 354 | - WAIT_FLAG(MTE1, MTE2, EVENT_ID0); | ||
| 355 | - WAIT_FLAG(MTE1, MTE2, EVENT_ID1); | ||
| 356 | - WAIT_FLAG(MTE1, MTE2, EVENT_ID2); | ||
| 357 | - WAIT_FLAG(MTE1, MTE2, EVENT_ID3); | ||
| 358 | - WAIT_FLAG(M, MTE1, EVENT_ID0); | ||
| 359 | - WAIT_FLAG(M, MTE1, EVENT_ID1); | ||
| 360 | - WAIT_FLAG(V, M, EVENT_ID0); | ||
| 361 | - WAIT_FLAG(MTE3, V, EVENT_ID0); | ||
| 362 | - AscendC::PipeBarrier<PIPE_ALL>(); | ||
| 363 | - } | ||
| 364 | - | ||
| 365 | -public: | ||
| 366 | - AsdopsBuffer<ArchType::ASCEND_V200> buf; | ||
| 367 | - | ||
| 368 | -private: | ||
| 369 | - AscendC::GlobalTensor<IN_DTYPE> gm_a; | ||
| 370 | - AscendC::GlobalTensor<IN_DTYPE> gm_b; | ||
| 371 | - AscendC::GlobalTensor<half> gm_c; | ||
| 372 | - AscendC::GlobalTensor<BIAS_TYPE> gm_bias; | ||
| 373 | - AscendC::GlobalTensor<DESCALE_TYPE> gm_scale; | ||
| 374 | - AscendC::GlobalTensor<IN_DTYPE> gm_compress_index; | ||
| 375 | - AscendC::LocalTensor<IN_DTYPE> l1_a_ping; | ||
| 376 | - AscendC::LocalTensor<IN_DTYPE> l1_a_pong; | ||
| 377 | - AscendC::LocalTensor<IN_DTYPE> l1_b_ping; | ||
| 378 | - AscendC::LocalTensor<IN_DTYPE> l1_b_pong; | ||
| 379 | - AscendC::LocalTensor<IN_DTYPE> l0_a_ping; | ||
| 380 | - AscendC::LocalTensor<IN_DTYPE> l0_a_pong; | ||
| 381 | - AscendC::LocalTensor<IN_DTYPE> l0_b_ping; | ||
| 382 | - AscendC::LocalTensor<IN_DTYPE> l0_b_pong; | ||
| 383 | - AscendC::LocalTensor<int32_t> l0c = buf.template GetBuffer<BufferType::ASCEND_L0C, int32_t>(0); | ||
| 384 | - AscendC::LocalTensor<half> ub_c; | ||
| 385 | - AscendC::LocalTensor<DESCALE_TYPE> ub_scale; | ||
| 386 | - AscendC::LocalTensor<BIAS_TYPE> ub_bias; | ||
| 387 | - | ||
| 388 | - uint32_t b_{0}; | ||
| 389 | - uint32_t m_{0}; | ||
| 390 | - uint32_t k_{0}; | ||
| 391 | - uint32_t n_{0}; | ||
| 392 | - uint32_t m0_{0}; | ||
| 393 | - uint32_t k0_{0}; | ||
| 394 | - uint32_t n0_{0}; | ||
| 395 | - uint32_t m_loop_{0}; | ||
| 396 | - uint32_t k_loop_{0}; | ||
| 397 | - uint32_t n_loop_{0}; | ||
| 398 | - uint32_t m_org_{0}; | ||
| 399 | - uint32_t n_org_{0}; | ||
| 400 | - uint32_t core_loop_{0}; | ||
| 401 | - uint32_t swizzl_cnt_{0}; | ||
| 402 | - uint32_t copress_tiling_k_{0}; | ||
| 403 | - uint32_t copress_tiling_n_{0}; | ||
| 404 | - uint32_t compress_size{0}; | ||
| 405 | - uint32_t compress_overlap_n_{0}; | ||
| 406 | - uint32_t n_compress_num{0}; | ||
| 407 | - uint32_t k_compress_num{0}; | ||
| 408 | -}; | ||
| 409 | - | ||
| 410 | - | ||
| 411 | - | ||
| 412 | - | ||
| @@ -1,17 +0,0 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | ||
| 2 | -# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | -# This file is a part of the CANN Open Software. | ||
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | -# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | -# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | -# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | -# ---------------------------------------------------------------------------- | ||
| 11 | - | ||
| 12 | -file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 13 | -foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 14 | - if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 15 | - add_subdirectory(${SUB_DIR}) | ||
| 16 | - endif() | ||
| 17 | -endforeach() | ||
| @@ -1,17 +0,0 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | ||
| 2 | -# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | -# This file is a part of the CANN Open Software. | ||
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | -# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | -# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | -# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | -# ---------------------------------------------------------------------------- | ||
| 11 | - | ||
| 12 | -file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 13 | -foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 14 | - if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 15 | - add_subdirectory(${SUB_DIR}) | ||
| 16 | - endif() | ||
| 17 | -endforeach() | ||
| @@ -1,22 +0,0 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | ||
| 2 | -# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | -# This file is a part of the CANN Open Software. | ||
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | -# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | -# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | -# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | -# ---------------------------------------------------------------------------- | ||
| 11 | - | ||
| 12 | -# if(UT_TEST_ALL OR OP_HOST_UT) | ||
| 13 | -# add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}) | ||
| 14 | -# add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}) | ||
| 15 | -# endif() | ||
| 16 | - | ||
| 17 | -file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 18 | -foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 19 | - if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 20 | - add_subdirectory(${SUB_DIR}) | ||
| 21 | - endif() | ||
| 22 | -endforeach() | ||
| @@ -1,14 +0,0 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | ||
| 2 | -# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | -# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | -# This file is a part of the CANN Open Software. | ||
| 5 | -# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | -# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | -# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | -# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | -# ---------------------------------------------------------------------------- | ||
| 11 | - | ||
| 12 | -if(UT_TEST_ALL OR OP_API_UT) | ||
| 13 | - # add_modules_ut_sources(UT_NAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}) | ||
| 14 | -endif() | ||
Dconversion/matmul_v2_compress_dequant/tests/ut/op_host/op_api/test_aclnn_matmul_compress_dequant.cpp+0-76
| @@ -1,76 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2025 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 | -using namespace std; | ||
| 21 | - | ||
| 22 | -class l2_matmul_compress_dequant_test : public testing::Test { | ||
| 23 | -protected: | ||
| 24 | - static void SetUpTestCase() | ||
| 25 | - { | ||
| 26 | - cout << "l2_matmul_compress_dequant_test SetUp" << endl; | ||
| 27 | - } | ||
| 28 | - static void TearDownTestCase() | ||
| 29 | - { | ||
| 30 | - cout << "l2_matmul_compress_dequant_test TearDown" << endl; | ||
| 31 | - } | ||
| 32 | - static void MatMulCompressDequantCommonTest( | ||
| 33 | - TensorDesc a_desc, TensorDesc b_desc, TensorDesc index_desc, TensorDesc bias_desc, TensorDesc deqScale_desc, | ||
| 34 | - IntArrayDesc compress_info_desc, TensorDesc out_desc, aclnnStatus expect_status) | ||
| 35 | - { | ||
| 36 | - auto ut = OP_API_UT( | ||
| 37 | - aclnnMatmulCompressDequant, | ||
| 38 | - INPUT(a_desc, b_desc, index_desc, bias_desc, deqScale_desc, nullptr, 1, compress_info_desc), | ||
| 39 | - OUTPUT(out_desc)); | ||
| 40 | - | ||
| 41 | - // SAMPLE: only test GetWorkspaceSize | ||
| 42 | - uint64_t workspace_size = 0; | ||
| 43 | - aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size); | ||
| 44 | - EXPECT_EQ(aclRet, expect_status); | ||
| 45 | - } | ||
| 46 | -}; | ||
| 47 | - | ||
| 48 | -TEST_F(l2_matmul_compress_dequant_test, ascend310P_test_normal_input) | ||
| 49 | -{ | ||
| 50 | - // 使用**Desc描述host api输入输出 | ||
| 51 | - TensorDesc a_desc = TensorDesc({16, 32}, ACL_INT8, ACL_FORMAT_ND); | ||
| 52 | - TensorDesc b_desc = TensorDesc({64}, ACL_INT8, ACL_FORMAT_ND); | ||
| 53 | - TensorDesc index_desc = TensorDesc({8}, ACL_INT8, ACL_FORMAT_ND); | ||
| 54 | - TensorDesc bias_desc = TensorDesc({16, 16}, ACL_INT32, ACL_FORMAT_ND); | ||
| 55 | - TensorDesc deqScale_desc = TensorDesc({16, 32}, ACL_UINT64, ACL_FORMAT_NC1HWC0); | ||
| 56 | - vector<int64_t> compress_info = {8, 8, 32, 16, 1}; | ||
| 57 | - IntArrayDesc compress_info_desc = IntArrayDesc(compress_info); | ||
| 58 | - TensorDesc out_desc = TensorDesc({16, 16}, ACL_FLOAT16, ACL_FORMAT_ND); | ||
| 59 | - MatMulCompressDequantCommonTest( | ||
| 60 | - a_desc, b_desc, index_desc, bias_desc, deqScale_desc, compress_info_desc, out_desc, ACLNN_SUCCESS); | ||
| 61 | -} | ||
| 62 | - | ||
| 63 | -TEST_F(l2_matmul_compress_dequant_test, ascend310P_test_empty) | ||
| 64 | -{ | ||
| 65 | - // 使用**Desc描述host api输入输出 | ||
| 66 | - TensorDesc a_desc = TensorDesc({16, 0}, ACL_INT8, ACL_FORMAT_ND); | ||
| 67 | - TensorDesc b_desc = TensorDesc({0}, ACL_INT8, ACL_FORMAT_ND); | ||
| 68 | - TensorDesc index_desc = TensorDesc({8}, ACL_INT8, ACL_FORMAT_ND); | ||
| 69 | - TensorDesc bias_desc = TensorDesc({16, 16}, ACL_INT32, ACL_FORMAT_ND); | ||
| 70 | - TensorDesc deqScale_desc = TensorDesc({16, 32}, ACL_UINT64, ACL_FORMAT_NC1HWC0); | ||
| 71 | - vector<int64_t> compress_info = {8, 8, 0, 16, 1}; | ||
| 72 | - IntArrayDesc compress_info_desc = IntArrayDesc(compress_info); | ||
| 73 | - TensorDesc out_desc = TensorDesc({16, 16}, ACL_FLOAT16, ACL_FORMAT_ND); | ||
| 74 | - MatMulCompressDequantCommonTest( | ||
| 75 | - a_desc, b_desc, index_desc, bias_desc, deqScale_desc, compress_info_desc, out_desc, ACLNN_SUCCESS); | ||
| 76 | -} | ||
| @@ -198,7 +198,6 @@ ${INSTALL_DIR}表示CANN安装后文件路径；\$\{ops\_project\}表示算子 | |||
| 198 | | [aclnnLtTensor&aclnnInplaceLtTensor](../../math/less/docs/aclnnLtTensor&aclnnInplaceLtTensor.md) | 判断输入self中的每个元素是否小于输入other中的元素，返回一个Bool类型的Tensor。 | 默认确定性实现| 默认确定性实现 | | 198 | | [aclnnLtTensor&aclnnInplaceLtTensor](../../math/less/docs/aclnnLtTensor&aclnnInplaceLtTensor.md) | 判断输入self中的每个元素是否小于输入other中的元素，返回一个Bool类型的Tensor。 | 默认确定性实现| 默认确定性实现 | |
| 199 | | [aclnnMaskedScale](../../math/masked_scale/docs/aclnnMaskedScale.md) | 完成elementwise计算。 | 默认确定性实现| 默认确定性实现| | 199 | | [aclnnMaskedScale](../../math/masked_scale/docs/aclnnMaskedScale.md) | 完成elementwise计算。 | 默认确定性实现| 默认确定性实现| |
| 200 | | [aclnnMaskedSelect](../../conversion/masked_select_v3/docs/aclnnMaskedSelect.md) | 根据一个布尔掩码张量（mask）中的值选择输入张量（self）中的元素作为输出，形成一个新的一维张量。 | 默认确定性实现| 默认确定性实现| | 200 | | [aclnnMaskedSelect](../../conversion/masked_select_v3/docs/aclnnMaskedSelect.md) | 根据一个布尔掩码张量（mask）中的值选择输入张量（self）中的元素作为输出，形成一个新的一维张量。 | 默认确定性实现| 默认确定性实现| |
| 201 | -| [aclnnMatmulCompressDequant](../../conversion/matmul_v2_compress_dequant/docs/aclnnMatmulCompressDequant.md) | 进行l@r矩阵乘计算时，可先通过msModelSlim工具对r矩阵进行无损压缩，减少r矩阵的内存占用大小，然后通过本接口完成无损解压缩、矩阵乘、反量化计算。 | - | - | | ||
| 202 | | [aclnnMaxDim](../../math/arg_max_with_value/docs/aclnnMaxDim.md) | 返回Tensor指定维度的最大值及其索引位置。 | 默认确定性实现| 默认确定性实现| | 201 | | [aclnnMaxDim](../../math/arg_max_with_value/docs/aclnnMaxDim.md) | 返回Tensor指定维度的最大值及其索引位置。 | 默认确定性实现| 默认确定性实现| |
| 203 | | [aclnnMax](../../math/reduce_max/docs/aclnnMax.md) | 返回Tensor所有元素中的最大值。 | 默认确定性实现| 默认确定性实现| | 202 | | [aclnnMax](../../math/reduce_max/docs/aclnnMax.md) | 返回Tensor所有元素中的最大值。 | 默认确定性实现| 默认确定性实现| |
| 204 | | [aclnnMaximum](../../math/maximum/docs/aclnnMaximum.md) | 计算两个张量中每个元素的最大值，并返回一个新的张量。 | 默认确定性实现| 默认确定性实现 | | 203 | | [aclnnMaximum](../../math/maximum/docs/aclnnMaximum.md) | 计算两个张量中每个元素的最大值，并返回一个新的张量。 | 默认确定性实现| 默认确定性实现 | |