已合并
capi文档更新 #4983
chenmyk创建于 19 天前
capi文档更新 #4983
已合并
共 6 个文件变更+423-132
| @@ -1,3 +1,3 @@ | |||
| 1 | version https://git-lfs.github.com/spec/v1 | 1 | version https://git-lfs.github.com/spec/v1 |
| 2 | -oid sha256:f6606e78ec7b0848205e93a01661918ed76fc9ed5836f26a49f8a532f6ff37cd | 2 | +oid sha256:8645fbc05a7ea8caae867d739acacf8fec2e33997c0a66bf3531c2dd2167934b |
| 3 | -size 11764 | 3 | +size 11074 |
| @@ -1,3 +1,3 @@ | |||
| 1 | version https://git-lfs.github.com/spec/v1 | 1 | version https://git-lfs.github.com/spec/v1 |
| 2 | -oid sha256:bbb001232beaf6775fa35af5f696fb2a10f24fc7dce1204b552099982918dc95 | 2 | +oid sha256:a04ef54b14ecac1d2c025c544ee9af5d45298185bec30a2c07276b51e63ce993 |
| 3 | -size 10140 | 3 | +size 9532 |
| @@ -34,16 +34,27 @@ $$ | |||
| 34 | dst_i = src0_i \times src1_i | 34 | dst_i = src0_i \times src1_i |
| 35 | $$ | 35 | $$ |
| 36 | 36 | ||
| 37 | +本接口仅在AIV上生效。 | ||
| 38 | + | ||
| 37 | ## 函数原型 | 39 | ## 函数原型 |
| 38 | 40 | ||
| 39 | -```cpp | 41 | +```c |
| 40 | -__simd_callee__ inline void asc_mul(vector_int16_t& dst, vector_int16_t src0, vector_int16_t src1, vector_bool mask) | 42 | +__simd_callee__ inline void asc_mul(vector_<dtype>& dst, |
| 41 | -__simd_callee__ inline void asc_mul(vector_uint16_t& dst, vector_uint16_t src0, vector_uint16_t src1, vector_bool mask) | 43 | + vector_<dtype> src0, |
| 42 | -__simd_callee__ inline void asc_mul(vector_half& dst, vector_half src0, vector_half src1, vector_bool mask) | 44 | + vector_<dtype> src1, |
| 43 | -__simd_callee__ inline void asc_mul(vector_bfloat16_t& dst, vector_bfloat16_t src0, vector_bfloat16_t src1, vector_bool mask) | 45 | + vector_bool mask) |
| 44 | -__simd_callee__ inline void asc_mul(vector_int32_t& dst, vector_int32_t src0, vector_int32_t src1, vector_bool mask) | 46 | +``` |
| 45 | -__simd_callee__ inline void asc_mul(vector_uint32_t& dst, vector_uint32_t src0, vector_uint32_t src1, vector_bool mask) | 47 | + |
| 46 | -__simd_callee__ inline void asc_mul(vector_float& dst, vector_float src0, vector_float src1, vector_bool mask) | 48 | +dtype可取的数据类型为`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。 |
| 49 | + | ||
| 50 | +### 典型示例 | ||
| 51 | + | ||
| 52 | +```c | ||
| 53 | +// 示例：对half矢量数据寄存器执行逐元素乘法 | ||
| 54 | +__simd_callee__ inline void asc_mul(vector_half& dst, | ||
| 55 | + vector_half src0, | ||
| 56 | + vector_half src1, | ||
| 57 | + vector_bool mask) | ||
| 47 | ``` | 58 | ``` |
| 48 | 59 | ||
| 49 | ## 参数说明 | 60 | ## 参数说明 |
| @@ -65,23 +76,132 @@ __simd_callee__ inline void asc_mul(vector_float& dst, vector_float src0, vector | |||
| 65 | 76 | ||
| 66 | ## 约束说明 | 77 | ## 约束说明 |
| 67 | 78 | ||
| 68 | -mask控制源操作数是否参与计算，源操作数不参与计算的元素在输出对应位置置零。 | 79 | +- 本接口在非AIV上调用直接返回。 |
| 80 | +- mask需通过掩码设置接口预先赋值后再传入，未赋值的掩码寄存器内容不确定，会导致有效元素位置错误。 | ||
| 69 | 81 | ||
| 70 | ## 调用示例 | 82 | ## 调用示例 |
| 71 | 83 | ||
| 72 | -```cpp | 84 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 73 | -__simd_vf__ inline void mul_vf(__ubuf__ half* dst_addr, __ubuf__ half* src0_addr, __ubuf__ half* src1_addr, uint32_t count, int32_t one_repeat_size, uint16_t repeat_time) | 85 | + |
| 86 | +<!-- npu="950" id8 --> | ||
| 87 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 88 | + | ||
| 89 | +```bash | ||
| 90 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 91 | +``` | ||
| 92 | +<!-- end id8 --> | ||
| 93 | + | ||
| 94 | +```c | ||
| 95 | +#include <algorithm> | ||
| 96 | +#include <cmath> | ||
| 97 | +#include <cstdint> | ||
| 98 | +#include <iostream> | ||
| 99 | +#include <vector> | ||
| 100 | + | ||
| 101 | +#include "c_api/asc_simd.h" | ||
| 102 | +#include "acl/acl.h" | ||
| 103 | + | ||
| 104 | +namespace { | ||
| 105 | +template <typename T> | ||
| 106 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 74 | { | 107 | { |
| 75 | - vector_half src0; | 108 | + std::cout << label << ":"; |
| 76 | - vector_half src1; | 109 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 77 | - vector_half dst; | 110 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 78 | - vector_bool mask; | 111 | + if (values.size() > count) std::cout << " ..."; |
| 79 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 112 | + std::cout << std::endl; |
| 80 | - mask = asc_update_mask_b16(count); | 113 | +} |
| 81 | - asc_loadalign_postupdate(src0, src0_addr, one_repeat_size); | 114 | + |
| 82 | - asc_loadalign_postupdate(src1, src1_addr, one_repeat_size); | 115 | +template <typename T> |
| 83 | - asc_mul(dst, src0, src1, mask); | 116 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) |
| 84 | - asc_storealign_postupdate(dst_addr, dst, one_repeat_size, mask); | 117 | +{ |
| 118 | + if (actual.size() != expected.size()) return false; | ||
| 119 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 120 | + if (actual[i] == expected[i]) continue; | ||
| 121 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 122 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 85 | } | 123 | } |
| 124 | + return true; | ||
| 125 | +} | ||
| 126 | + | ||
| 127 | +constexpr uint32_t ELEMENT_COUNT = 64; | ||
| 128 | + | ||
| 129 | +__simd_vf__ inline void mul_vf(__ubuf__ float* dst, __ubuf__ float* src0, __ubuf__ float* src1) | ||
Y | |||
| 130 | +{ | ||
| 131 | + vector_float dst_reg; | ||
| 132 | + vector_float src0_reg; | ||
| 133 | + vector_float src1_reg; | ||
| 134 | + uint32_t count = ELEMENT_COUNT; | ||
| 135 | + vector_bool mask = asc_update_mask_b32(count); | ||
| 136 | + asc_loadalign(src0_reg, src0); | ||
| 137 | + asc_loadalign(src1_reg, src1); | ||
| 138 | + asc_mul(dst_reg, src0_reg, src1_reg, mask); | ||
| 139 | + asc_storealign(dst, dst_reg, mask); | ||
| 140 | +} | ||
| 141 | + | ||
| 142 | +__global__ __vector__ void asc_mul_kernel(__gm__ float* dst, __gm__ float* src0, __gm__ float* src1) | ||
| 143 | +{ | ||
| 144 | + asc_init(); | ||
| 145 | + __ubuf__ float dst_local[ELEMENT_COUNT]; | ||
| 146 | + __ubuf__ float src0_local[ELEMENT_COUNT]; | ||
| 147 | + __ubuf__ float src1_local[ELEMENT_COUNT]; | ||
| 148 | + asc_copy_gm2ub_align(src0_local, src0, ELEMENT_COUNT * sizeof(float)); | ||
| 149 | + asc_copy_gm2ub_align(src1_local, src1, ELEMENT_COUNT * sizeof(float)); | ||
| 150 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 151 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 152 | + mul_vf(dst_local, src0_local, src1_local); | ||
| 153 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 154 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 155 | + asc_copy_ub2gm_align(dst, dst_local, ELEMENT_COUNT * sizeof(float)); | ||
| 156 | + asc_sync(); | ||
| 157 | +} | ||
| 158 | + | ||
| 159 | +} // namespace | ||
| 160 | + | ||
| 161 | +int main() | ||
| 162 | +{ | ||
| 163 | + std::vector<float> src0(ELEMENT_COUNT); | ||
| 164 | + std::vector<float> src1(ELEMENT_COUNT); | ||
| 165 | + std::vector<float> output(ELEMENT_COUNT, 0.0f); | ||
| 166 | + std::vector<float> golden(ELEMENT_COUNT); | ||
| 167 | + for (uint32_t i = 0; i < ELEMENT_COUNT; ++i) { | ||
| 168 | + src0[i] = static_cast<float>(i) * 0.25f; | ||
| 169 | + src1[i] = static_cast<float>(i % 8) * 0.5f; | ||
| 170 | + golden[i] = src0[i] * src1[i]; | ||
| 171 | + } | ||
| 172 | + | ||
| 173 | + aclInit(nullptr); | ||
| 174 | + aclrtSetDevice(0); | ||
| 175 | + float* src0_device = nullptr; | ||
| 176 | + aclrtMalloc(reinterpret_cast<void**>(&src0_device), (ELEMENT_COUNT) * sizeof(float), | ||
| 177 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 178 | + float* src1_device = nullptr; | ||
| 179 | + aclrtMalloc(reinterpret_cast<void**>(&src1_device), (ELEMENT_COUNT) * sizeof(float), | ||
| 180 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 181 | + float* dst_device = nullptr; | ||
| 182 | + aclrtMalloc(reinterpret_cast<void**>(&dst_device), (ELEMENT_COUNT) * sizeof(float), | ||
| 183 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 184 | + aclrtMemcpy(src0_device, src0.size() * sizeof(float), src0.data(), src0.size() * sizeof(float), | ||
| 185 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 186 | + aclrtMemcpy(src1_device, src1.size() * sizeof(float), src1.data(), src1.size() * sizeof(float), | ||
| 187 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 188 | + | ||
| 189 | + asc_mul_kernel<<<1, 0>>>(dst_device, src0_device, src1_device); | ||
| 190 | + aclrtSynchronizeDevice(); | ||
| 191 | + aclrtMemcpy(output.data(), output.size() * sizeof(float), dst_device, output.size() * sizeof(float), | ||
| 192 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 193 | + | ||
| 194 | + print_data("Input src0", src0); | ||
| 195 | + print_data("Input src1", src1); | ||
| 196 | + print_data("Output", output); | ||
| 197 | + print_data("Golden", golden); | ||
| 198 | + const bool passed = compare_data(output, golden, 1e-6); | ||
| 199 | + std::cout << (passed ? "[Success] asc_mul passed." : "[Failed] asc_mul failed.") << std::endl; | ||
| 200 | + aclrtFree(dst_device); | ||
| 201 | + aclrtFree(src0_device); | ||
| 202 | + aclrtFree(src1_device); | ||
| 203 | + aclrtResetDevice(0); | ||
| 204 | + aclFinalize(); | ||
| 205 | + return passed ? 0 : 1; | ||
| 86 | } | 206 | } |
| 87 | ``` | 207 | ``` |
| @@ -86,7 +86,7 @@ __simd_vf__ inline void half2hif8_vf(__ubuf__ hifloat8_t* dst_addr, __ubuf__ hal | |||
| 86 | for (uint16_t i = 0; i < repeat_time; ++i) { | 86 | for (uint16_t i = 0; i < repeat_time; ++i) { |
| 87 | asc_loadalign_postupdate(src, src_addr, src_repeat_size); | 87 | asc_loadalign_postupdate(src, src_addr, src_repeat_size); |
| 88 | asc_half2hif8_rna(dst, src, mask); | 88 | asc_half2hif8_rna(dst, src, mask); |
| 89 | - asc_storealign_postupdate(dst_addr, dst, dst_repeat_size, mask); | 89 | + asc_storealign_postupdate(reinterpret_cast<__ubuf__ uint8_t*&>(dst_addr), reinterpret_cast<vector_uint8_t&>(dst), dst_repeat_size, mask); |
| 90 | } | 90 | } |
| 91 | } | 91 | } |
| 92 | ``` | 92 | ``` |
| @@ -26,7 +26,8 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -以传入的value为起始值，生成递增/递减的索引，并将生成的索引保存在dst中。算法逻辑表示如下： | 29 | +以传入的`value`为起始值，生成递增/递减的索引，并将生成的索引保存在`dst`中，[Vector Length (VL)](../reg_data_types/data_type_definition.md)表示矢量数据寄存器的位宽，`VL_T`表示该寄存器可存储的元素数量。算法逻辑表示如下： |
| 30 | + | ||
| 30 | ```cpp | 31 | ```cpp |
| 31 | // 递增 | 32 | // 递增 |
| 32 | {value, value + 1, value + 2, ... value + VL_T - 2, value + VL_T - 1} | 33 | {value, value + 1, value + 2, ... value + VL_T - 2, value + VL_T - 1} |
| @@ -34,37 +35,66 @@ | |||
| 34 | {value + VL_T - 1, value + VL_T - 2, value + VL_T - 3, ... value + 1, value} | 35 | {value + VL_T - 1, value + VL_T - 2, value + VL_T - 3, ... value + 1, value} |
| 35 | ``` | 36 | ``` |
| 36 | 37 | ||
| 37 | -以int16_t数据类型，起始值value=10为例： | 38 | +以int16_t数据类型，起始值`value=10`为例： |
| 38 | 递增索引为{10, 11, 12, 13, ... 135, 136, 137}，递减索引为{137, 136, 135, 134, ... 12, 11, 10}。 | 39 | 递增索引为{10, 11, 12, 13, ... 135, 136, 137}，递减索引为{137, 136, 135, 134, ... 12, 11, 10}。 |
| 39 | 40 | ||
| 41 | +本接口仅在AIV上生效。 | ||
| 42 | + | ||
| 40 | ## 函数原型 | 43 | ## 函数原型 |
| 41 | 44 | ||
| 42 | -- 递增模式 | 45 | +### 递增模式 |
| 43 | - ```cpp | ||
| 44 | - __simd_callee__ inline void asc_arange(vector_int8_t& dst, int8_t value) | ||
| 45 | - __simd_callee__ inline void asc_arange(vector_int16_t& dst, int16_t value) | ||
| 46 | - __simd_callee__ inline void asc_arange(vector_half& dst, half value) | ||
| 47 | - __simd_callee__ inline void asc_arange(vector_int32_t& dst, int32_t value) | ||
| 48 | - __simd_callee__ inline void asc_arange(vector_float& dst, float value) | ||
| 49 | - ``` | ||
| 50 | 46 | ||
| 51 | -- 递减模式 | 47 | +```c |
| 52 | - ```cpp | 48 | +__simd_callee__ inline void asc_arange(vector_<dtype>& dst, |
| 53 | - __simd_callee__ inline void asc_arange_descend(vector_int8_t& dst, int8_t value) | 49 | + <dtype> value) |
| 54 | - __simd_callee__ inline void asc_arange_descend(vector_int16_t& dst, int16_t value) | 50 | +``` |
| 55 | - __simd_callee__ inline void asc_arange_descend(vector_half& dst, half value) | 51 | + |
| 56 | - __simd_callee__ inline void asc_arange_descend(vector_int32_t& dst, int32_t value) | 52 | +dtype可取的数据类型为`int8_t`、`int16_t`、`half`、`int32_t`、`float`。 |
| 57 | - __simd_callee__ inline void asc_arange_descend(vector_float& dst, float value) | 53 | + |
| 58 | - ``` | 54 | +#### 典型示例 |
| 55 | + | ||
| 56 | +```c | ||
| 57 | +// 示例：以half标量为基值生成递增序列 | ||
| 58 | +__simd_callee__ inline void asc_arange(vector_half& dst, | ||
| 59 | + half value) | ||
| 60 | +``` | ||
| 61 | + | ||
| 62 | +### 递减模式 | ||
| 63 | + | ||
| 64 | +```c | ||
| 65 | +__simd_callee__ inline void asc_arange_descend(vector_<dtype>& dst, | ||
| 66 | + <dtype> value) | ||
| 67 | +``` | ||
| 68 | + | ||
| 69 | +dtype可取的数据类型为`int8_t`、`int16_t`、`half`、`int32_t`、`float`。 | ||
| 70 | + | ||
| 71 | +#### 典型示例 | ||
| 72 | + | ||
| 73 | +```c | ||
| 74 | +// 示例：以half标量为基值生成递减序列 | ||
| 75 | +__simd_callee__ inline void asc_arange_descend(vector_half& dst, | ||
| 76 | + half value) | ||
| 77 | +``` | ||
| 59 | 78 | ||
| 60 | ## 参数说明 | 79 | ## 参数说明 |
| 61 | 80 | ||
| 81 | +### 递增模式 | ||
| 82 | + | ||
| 62 | **表1** 参数说明 | 83 | **表1** 参数说明 |
| 63 | 84 | ||
| 64 | | 参数名 | 输入/输出 | 描述 | | 85 | | 参数名 | 输入/输出 | 描述 | |
| 65 | | --------- | ----- | ----------------- | | 86 | | --------- | ----- | ----------------- | |
| 66 | | dst | 输出 | 目的操作数（矢量数据寄存器）。 | | 87 | | dst | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 67 | -| value | 输入 | 源操作数（标量）。 | | 88 | +| value | 输入 | 源操作数（标量），dtype须与dst一致。作为递增序列的起点，序列第0个元素等于value，后续元素按1递增。取值范围为该dtype的可表示范围。 | |
| 89 | + | ||
| 90 | +### 递减模式 | ||
| 91 | + | ||
| 92 | +**表2** 参数说明 | ||
| 93 | + | ||
| 94 | +| 参数名 | 输入/输出 | 描述 | | ||
| 95 | +| --------- | ----- | ----------------- | | ||
| 96 | +| dst | 输出 | 目的操作数（矢量数据寄存器）。 | | ||
| 97 | +| value | 输入 | 源操作数（标量），dtype须与dst一致。作为递减序列的起点，序列第0个元素等于`value + VL_T - 1`，后续元素按1递减。取值范围为该dtype的可表示范围。 | | ||
| 68 | 98 | ||
| 69 | 矢量数据寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 99 | 矢量数据寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 70 | 100 | ||
| @@ -74,19 +104,94 @@ | |||
| 74 | 104 | ||
| 75 | ## 约束说明 | 105 | ## 约束说明 |
| 76 | 106 | ||
| 77 | -对于整型数据类型，如果生成的索引发生溢出，结果将环绕。 | 107 | +- 本接口在非AIV上调用直接返回。 |
| 108 | +- 整型dtype（int8_t、int16_t、int32_t）结果在超出该dtype可表示范围时回绕（wrap-around），不触发异常。例如int8_t取value=127时，序列前128个元素依次为127、−128、−127、…、-2；value=−128时，序列前128个元素依次为−128、−127、…、−1。 | ||
| 78 | 109 | ||
| 79 | ## 调用示例 | 110 | ## 调用示例 |
| 80 | 111 | ||
| 81 | -```cpp | 112 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 82 | -__simd_vf__ inline void arange_vf(__ubuf__ int8_t* dst_addr, int8_t value, uint32_t count, int32_t one_repeat_size, uint16_t repeat_time) | 113 | + |
| 114 | +<!-- npu="950" id8 --> | ||
| 115 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 116 | + | ||
| 117 | +```bash | ||
| 118 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 119 | +``` | ||
| 120 | +<!-- end id8 --> | ||
| 121 | + | ||
| 122 | +```c | ||
| 123 | +#include <cstdint> | ||
| 124 | +#include <iostream> | ||
| 125 | +#include <vector> | ||
| 126 | + | ||
| 127 | +#include "c_api/asc_simd.h" | ||
| 128 | +#include "acl/acl.h" | ||
| 129 | + | ||
| 130 | +namespace { | ||
| 131 | +template <typename T> | ||
| 132 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 83 | { | 133 | { |
| 84 | - vector_int8_t dst; | 134 | + std::cout << label << ":"; |
| 85 | - vector_bool mask; | 135 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 86 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 136 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 87 | - mask = asc_update_mask_b8(count); | 137 | + if (values.size() > count) std::cout << " ..."; |
| 88 | - asc_arange(dst, value); | 138 | + std::cout << std::endl; |
| 89 | - asc_storealign_postupdate(dst_addr, dst, one_repeat_size, mask); | 139 | +} |
| 90 | - } | 140 | + |
| 141 | +constexpr uint32_t ELEMENT_COUNT = 64; | ||
| 142 | +constexpr int32_t START_VALUE = 10; | ||
| 143 | + | ||
| 144 | +__simd_vf__ inline void arange_vf(__ubuf__ int32_t* ascending, __ubuf__ int32_t* descending) | ||
| 145 | +{ | ||
| 146 | + vector_int32_t ascending_reg; | ||
| 147 | + vector_int32_t descending_reg; | ||
| 148 | + uint32_t count = ELEMENT_COUNT; | ||
| 149 | + vector_bool mask = asc_update_mask_b32(count); | ||
| 150 | + asc_arange(ascending_reg, START_VALUE); | ||
Y 需要给出注释体现输入与输出。   | |||
| 151 | + asc_arange_descend(descending_reg, START_VALUE); | ||
| 152 | + asc_storealign(ascending, ascending_reg, mask); | ||
| 153 | + asc_storealign(descending, descending_reg, mask); | ||
| 154 | +} | ||
| 155 | + | ||
| 156 | +__global__ __vector__ void asc_arange_kernel(__gm__ int32_t* ascending, __gm__ int32_t* descending) | ||
| 157 | +{ | ||
| 158 | + asc_init(); | ||
| 159 | + __ubuf__ int32_t ascending_local[ELEMENT_COUNT]; | ||
| 160 | + __ubuf__ int32_t descending_local[ELEMENT_COUNT]; | ||
| 161 | + arange_vf(ascending_local, descending_local); | ||
| 162 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 163 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 164 | + asc_copy_ub2gm_align(ascending, ascending_local, ELEMENT_COUNT * sizeof(int32_t)); | ||
| 165 | + asc_copy_ub2gm_align(descending, descending_local, ELEMENT_COUNT * sizeof(int32_t)); | ||
| 166 | + asc_sync(); | ||
| 167 | +} | ||
| 168 | +} // namespace | ||
| 169 | + | ||
| 170 | +int main() | ||
| 171 | +{ | ||
| 172 | + std::vector<int32_t> input = {START_VALUE}; | ||
| 173 | + std::vector<int32_t> ascending(ELEMENT_COUNT, 0), descending(ELEMENT_COUNT, 0); | ||
| 174 | + aclInit(nullptr); | ||
| 175 | + aclrtSetDevice(0); | ||
| 176 | + int32_t* ascending_device = nullptr; | ||
| 177 | + aclrtMalloc(reinterpret_cast<void**>(&ascending_device), (ELEMENT_COUNT) * sizeof(int32_t), | ||
| 178 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 179 | + int32_t* descending_device = nullptr; | ||
| 180 | + aclrtMalloc(reinterpret_cast<void**>(&descending_device), (ELEMENT_COUNT) * sizeof(int32_t), | ||
| 181 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 182 | + asc_arange_kernel<<<1, 0>>>(ascending_device, descending_device); | ||
| 183 | + aclrtSynchronizeDevice(); | ||
| 184 | + aclrtMemcpy(ascending.data(), ascending.size() * sizeof(int32_t), ascending_device, ascending.size() * sizeof(int32_t), | ||
| 185 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 186 | + aclrtMemcpy(descending.data(), descending.size() * sizeof(int32_t), descending_device, descending.size() * sizeof(int32_t), | ||
| 187 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 188 | + print_data("Input start", input); | ||
| 189 | + print_data("Ascending output", ascending); | ||
| 190 | + print_data("Descending output", descending); | ||
| 191 | + aclrtFree(ascending_device); | ||
| 192 | + aclrtFree(descending_device); | ||
| 193 | + aclrtResetDevice(0); | ||
| 194 | + aclFinalize(); | ||
| 195 | + return 0; | ||
| 91 | } | 196 | } |
| 92 | ``` | 197 | ``` |
Mdocs/zh/api/SIMD-API/c_api/vector_data_move/asc_copy_gm2ub_align/asc_copy_gm2ub_align_arch_3510.md+143-77
| @@ -30,14 +30,14 @@ | |||
| 30 | 30 | ||
| 31 | 本接口支持以下两种数据搬运方式： | 31 | 本接口支持以下两种数据搬运方式： |
| 32 | 32 | ||
| 33 | -- 前n个数据搬运 | 33 | +- 连续数据搬运 |
| 34 | 34 | ||
| 35 | 若搬运数据长度非32字节对齐，搬运数据会补齐至32字节对齐，支持以下两种填充方式： | 35 | 若搬运数据长度非32字节对齐，搬运数据会补齐至32字节对齐，支持以下两种填充方式： |
| 36 | 36 | ||
| 37 | - 手动填充：搬运前调用[asc_set_copy_pad_val](../asc_set_copy_pad_val.md)配置填充值。 | 37 | - 手动填充：搬运前调用[asc_set_copy_pad_val](../asc_set_copy_pad_val.md)配置填充值。 |
| 38 | - 自动填充：由硬件自动填充dummy假数据，dummy假数据的值为数据块的第一个元素的值。 | 38 | - 自动填充：由硬件自动填充dummy假数据，dummy假数据的值为数据块的第一个元素的值。 |
| 39 | 39 | ||
| 40 | -- 高维切分搬运 | 40 | +- 高维切分数据搬运 |
| 41 | 41 | ||
| 42 | 若搬运数据长度非32字节对齐，会将搬运数据补齐至32字节对齐。可通过配置参数`dst_stride`选择Normal模式或Compact模式。非32字节对齐场景支持以下两种填充方式： | 42 | 若搬运数据长度非32字节对齐，会将搬运数据补齐至32字节对齐。可通过配置参数`dst_stride`选择Normal模式或Compact模式。非32字节对齐场景支持以下两种填充方式： |
| 43 | 43 | ||
| @@ -57,94 +57,90 @@ | |||
| 57 | 57 | ||
| 58 | 当只搬运1个数据块，或`len_burst`已经32字节对齐且无左右Padding时，两种模式的搬运结果相同。 | 58 | 当只搬运1个数据块，或`len_burst`已经32字节对齐且无左右Padding时，两种模式的搬运结果相同。 |
| 59 | 59 | ||
| 60 | +本接口仅在AIV上生效。 | ||
| 61 | + | ||
| 60 | ## 函数原型 | 62 | ## 函数原型 |
| 61 | 63 | ||
| 62 | -- 前n个数据搬运 | 64 | +### 连续数据搬运 |
| 63 | 65 | ||
| 64 | - ```cpp | 66 | +```c |
| 65 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int8_t* dst, __gm__ int8_t* src, uint32_t size) | 67 | +__aicore__ inline void asc_copy_gm2ub_align(__ubuf__ <dtype>* dst, |
| 66 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint8_t* dst, __gm__ uint8_t* src, uint32_t size) | 68 | + __gm__ <dtype>* src, |
| 67 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ hifloat8_t* dst, __gm__ hifloat8_t* src, uint32_t size) | 69 | + uint32_t size) |
| 68 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ fp8_e5m2_t* dst, __gm__ fp8_e5m2_t* src, uint32_t size) | 70 | +``` |
| 69 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ fp8_e4m3fn_t* dst, __gm__ fp8_e4m3fn_t* src, uint32_t size) | ||
| 70 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int16_t* dst, __gm__ int16_t* src, uint32_t size) | ||
| 71 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint16_t* dst, __gm__ uint16_t* src, uint32_t size) | ||
| 72 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ half* dst, __gm__ half* src, uint32_t size) | ||
| 73 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ bfloat16_t* dst, __gm__ bfloat16_t* src, uint32_t size) | ||
| 74 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int32_t* dst, __gm__ int32_t* src, uint32_t size) | ||
| 75 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint32_t* dst, __gm__ uint32_t* src, uint32_t size) | ||
| 76 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ float* dst, __gm__ float* src, uint32_t size) | ||
| 77 | - ``` | ||
| 78 | 71 | ||
| 79 | -- 同步搬运 | 72 | +dtype可取的数据类型为`int8_t`、`uint8_t`、`hifloat8_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。 |
| 80 | 73 | ||
| 81 | - ```cpp | 74 | +#### 典型示例 |
| 82 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ int8_t* dst, __gm__ int8_t* src, uint32_t size) | ||
| 83 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ uint8_t* dst, __gm__ uint8_t* src, uint32_t size) | ||
| 84 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ hifloat8_t* dst, __gm__ hifloat8_t* src, uint32_t size) | ||
| 85 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ fp8_e5m2_t* dst, __gm__ fp8_e5m2_t* src, uint32_t size) | ||
| 86 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ fp8_e4m3fn_t* dst, __gm__ fp8_e4m3fn_t* src, uint32_t size) | ||
| 87 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ int16_t* dst, __gm__ int16_t* src, uint32_t size) | ||
| 88 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ uint16_t* dst, __gm__ uint16_t* src, uint32_t size) | ||
| 89 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ half* dst, __gm__ half* src, uint32_t size) | ||
| 90 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ bfloat16_t* dst, __gm__ bfloat16_t* src, uint32_t size) | ||
| 91 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ int32_t* dst, __gm__ int32_t* src, uint32_t size) | ||
| 92 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ uint32_t* dst, __gm__ uint32_t* src, uint32_t size) | ||
| 93 | - __aicore__ inline void asc_copy_gm2ub_align_sync(__ubuf__ float* dst, __gm__ float* src, uint32_t size) | ||
| 94 | - ``` | ||
| 95 | 75 | ||
| 96 | -- 高维切分搬运 | 76 | +```c |
| 77 | +// 示例：源与目的数据类型为bfloat16_t | ||
| 78 | +__aicore__ inline void asc_copy_gm2ub_align(__ubuf__ bfloat16_t* dst, | ||
| 79 | + __gm__ bfloat16_t* src, | ||
| 80 | + uint32_t size) | ||
| 81 | +``` | ||
| 97 | 82 | ||
| 98 | - ```cpp | 83 | +### 高维切分数据搬运 |
| 99 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int8_t* dst, __gm__ int8_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 100 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint8_t* dst, __gm__ uint8_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 101 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ hifloat8_t* dst, __gm__ hifloat8_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 102 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ fp8_e5m2_t* dst, __gm__ fp8_e5m2_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 103 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ fp8_e4m3fn_t* dst, __gm__ fp8_e4m3fn_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 104 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int16_t* dst, __gm__ int16_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 105 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint16_t* dst, __gm__ uint16_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 106 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ half* dst, __gm__ half* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 107 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ bfloat16_t* dst, __gm__ bfloat16_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 108 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int32_t* dst, __gm__ int32_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 109 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint32_t* dst, __gm__ uint32_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 110 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ float* dst, __gm__ float* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, asc_load_l2_cache_mode l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 111 | - ``` | ||
| 112 | 84 | ||
| 113 | -- **以下函数原型已废弃，请使用`asc_load_l2_cache_mode`类型枚举值进行L2 Cache管理策略配置。** | 85 | +```c |
| 86 | +__aicore__ inline void asc_copy_gm2ub_align(__ubuf__ <dtype>* dst, | ||
| 87 | + __gm__ <dtype>* src, | ||
| 88 | + uint16_t n_burst, | ||
| 89 | + uint32_t len_burst, | ||
| 90 | + uint8_t left_padding_num, | ||
| 91 | + uint8_t right_padding_num, | ||
| 92 | + bool enable_constant_pad, | ||
| 93 | + asc_load_l2_cache_mode l2_cache_mode, | ||
| 94 | + uint64_t src_stride, | ||
| 95 | + uint32_t dst_stride) | ||
| 96 | +``` | ||
| 114 | 97 | ||
| 115 | - ```cpp | 98 | +dtype可取的数据类型为`int8_t`、`uint8_t`、`hifloat8_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。 |
| 116 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int8_t* dst, __gm__ int8_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 117 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint8_t* dst, __gm__ uint8_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 118 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ hifloat8_t* dst, __gm__ hifloat8_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 119 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ fp8_e5m2_t* dst, __gm__ fp8_e5m2_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 120 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ fp8_e4m3fn_t* dst, __gm__ fp8_e4m3fn_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 121 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int16_t* dst, __gm__ int16_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 122 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint16_t* dst, __gm__ uint16_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 123 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ half* dst, __gm__ half* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 124 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ bfloat16_t* dst, __gm__ bfloat16_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 125 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ int32_t* dst, __gm__ int32_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 126 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ uint32_t* dst, __gm__ uint32_t* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 127 | - __aicore__ inline void asc_copy_gm2ub_align(__ubuf__ float* dst, __gm__ float* src, uint16_t n_burst, uint32_t len_burst, uint8_t left_padding_num, uint8_t right_padding_num, bool enable_constant_pad, uint8_t l2_cache_mode, uint64_t src_stride, uint32_t dst_stride) | ||
| 128 | - ``` | ||
| 129 | 99 | ||
| 100 | +#### 典型示例 | ||
| 101 | + | ||
| 102 | +```c | ||
| 103 | +// 示例：源与目的数据类型为bfloat16_t | ||
| 104 | +__aicore__ inline void asc_copy_gm2ub_align(__ubuf__ bfloat16_t* dst, | ||
| 105 | + __gm__ bfloat16_t* src, | ||
| 106 | + uint16_t n_burst, | ||
| 107 | + uint32_t len_burst, | ||
| 108 | + uint8_t left_padding_num, | ||
| 109 | + uint8_t right_padding_num, | ||
| 110 | + bool enable_constant_pad, | ||
| 111 | + asc_load_l2_cache_mode l2_cache_mode, | ||
| 112 | + uint64_t src_stride, | ||
| 113 | + uint32_t dst_stride) | ||
| 114 | +``` | ||
| 130 | 115 | ||
| 131 | ## 参数说明 | 116 | ## 参数说明 |
| 132 | 117 | ||
| 118 | +### 连续数据搬运 | ||
| 119 | + | ||
| 133 | **表1** 参数说明 | 120 | **表1** 参数说明 |
| 134 | 121 | ||
| 135 | | 参数名 | 输入/输出 | 描述 | | 122 | | 参数名 | 输入/输出 | 描述 | |
| 136 | | :--- | :--- | :--- | | 123 | | :--- | :--- | :--- | |
| 137 | | dst | 输出 | 目的UB的起始地址。需要32字节对齐。 | | 124 | | dst | 输出 | 目的UB的起始地址。需要32字节对齐。 | |
| 138 | | src | 输入 | 源GM的起始地址。需要1字节对齐。 | | 125 | | src | 输入 | 源GM的起始地址。需要1字节对齐。 | |
| 139 | -| size | 输入 | 搬运数据大小，单位为字节。取值范围：[0, 2097151]。 | | 126 | +| size | 输入 | 搬运数据大小，单位为字节。取值范围：[1, $2^{21}−1$]。 | |
| 140 | -| n_burst | 输入 | 待搬运的连续传输数据块个数。取值范围：[0, 4095]。 | | 127 | + |
| 141 | -| len_burst | 输入 | 待搬运的每个连续传输数据块的长度，单位为字节。取值范围：[0, 2097151]。 | | 128 | +### 高维切分数据搬运 |
| 129 | + | ||
| 130 | +**表2** 参数说明 | ||
| 131 | + | ||
| 132 | +| 参数名 | 输入/输出 | 描述 | | ||
| 133 | +| :--- | :--- | :--- | | ||
| 134 | +| dst | 输出 | 目的UB的起始地址。需要32字节对齐。 | | ||
| 135 | +| src | 输入 | 源GM的起始地址。需要1字节对齐。 | | ||
| 136 | +| n_burst | 输入 | 待搬运的连续传输数据块个数。取值范围：[1, 4095]。 | | ||
| 137 | +| len_burst | 输入 | 待搬运的每个连续传输数据块的长度，单位为字节。取值范围：[1, $2^{21}−1$]。 | | ||
| 142 | | left_padding_num | 输入 | 连续搬运数据块左侧需要补充的元素个数。该参数对应的填充数据大小不能超过32字节。Compact模式下需要设置为0。 | | 138 | | left_padding_num | 输入 | 连续搬运数据块左侧需要补充的元素个数。该参数对应的填充数据大小不能超过32字节。Compact模式下需要设置为0。 | |
| 143 | | right_padding_num | 输入 | 连续搬运数据块右侧需要补充的元素个数。该参数对应的填充数据大小不能超过32字节。Compact模式下需要设置为0。 | | 139 | | right_padding_num | 输入 | 连续搬运数据块右侧需要补充的元素个数。该参数对应的填充数据大小不能超过32字节。Compact模式下需要设置为0。 | |
| 144 | | enable_constant_pad | 输入 | 当`left_padding_num`和`right_padding_num`均为0时，配置非对齐场景的填充方式。取值说明如下： <br>• `true`：手动填充，填充值为接口`asc_set_copy_pad_val`设置的值。 <br>• `false`：自动填充，由硬件填充dummy假数据，dummy假数据的值为数据块的第一个元素的值。<br>当`left_padding_num`或`right_padding_num`非0时，该参数不生效。 | | 140 | | enable_constant_pad | 输入 | 当`left_padding_num`和`right_padding_num`均为0时，配置非对齐场景的填充方式。取值说明如下： <br>• `true`：手动填充，填充值为接口`asc_set_copy_pad_val`设置的值。 <br>• `false`：自动填充，由硬件填充dummy假数据，dummy假数据的值为数据块的第一个元素的值。<br>当`left_padding_num`或`right_padding_num`非0时，该参数不生效。 | |
| 145 | | l2_cache_mode | 输入 | [asc_load_l2_cache_mode](../../enum/asc_load_l2_cache_mode.md)类型的枚举值，配置数据在L2 Cache中的管理策略。 | | 141 | | l2_cache_mode | 输入 | [asc_load_l2_cache_mode](../../enum/asc_load_l2_cache_mode.md)类型的枚举值，配置数据在L2 Cache中的管理策略。 | |
| 146 | -| src_stride | 输入 | 源操作数相邻连续数据块的距离（前面一个数据块的头与后面一个数据块的头的间隔），单位为字节。<br>只搬运1个数据块，即`n_burst`设置为1时，可以将此参数设置为0。 | | 142 | +| src_stride | 输入 | 源操作数相邻连续数据块的距离（前面一个数据块的头与后面一个数据块的头的间隔），单位为字节。取值范围：[0, $2^{40}−1$]。<br>只搬运1个数据块，即`n_burst`设置为1时，可以将此参数设置为0。 | |
| 147 | -| dst_stride | 输入 | 目的操作数相邻连续数据块的距离（前面一个数据块的头与后面一个数据块的头的间隔），单位为字节，用于选择数据搬运模式。<br>• 等于`len_burst`：Compact模式，目的数据块在UB中紧密排列，`dst_stride`支持字节对齐。<br>• 不等于`len_burst`：Normal模式，`dst_stride`需要满足32字节对齐要求。<br>只搬运1个数据块，即`n_burst`设置为1时，可以将此参数设置为0。 | | 143 | +| dst_stride | 输入 | 目的操作数相邻连续数据块的距离（前面一个数据块的头与后面一个数据块的头的间隔），单位为字节，用于选择数据搬运模式。取值范围：[0, $2^{21}−1$]。<br>• 等于`len_burst`：Compact模式，目的数据块在UB中紧密排列，`dst_stride`支持字节对齐。<br>• 不等于`len_burst`：Normal模式，`dst_stride`需要满足32字节对齐要求。<br>只搬运1个数据块，即`n_burst`设置为1时，可以将此参数设置为0。 | |
| 148 | 144 | ||
| 149 | ## 返回值说明 | 145 | ## 返回值说明 |
| 150 | 146 | ||
| @@ -156,20 +152,90 @@ PIPE_MTE2 | |||
| 156 | 152 | ||
| 157 | ## 约束说明 | 153 | ## 约束说明 |
| 158 | 154 | ||
| 155 | +### 通用约束 | ||
| 156 | + | ||
| 157 | +- 本接口在非AIV上调用直接返回。 | ||
| 159 | - 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。 | 158 | - 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。 |
| 160 | -- 当`n_burst`、`len_burst`中任意一个值为0时，该接口被视为NOP（空操作）。 | 159 | +- 如果本指令与其他指令的目的地址存在重叠，需要插入同步指令（[asc_sync_notify](../../sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/asc_sync_wait.md)），保证多个指令的串行化，防止出现异常数据。 |
| 161 | -- 当`size`值为0时，该接口被视为NOP（空操作）。 | 160 | + |
| 162 | -- 如果需要执行多条`asc_copy_gm2ub_align`指令，且`asc_copy_gm2ub_align`指令的目的地址存在重叠，需要插入同步指令（[asc_sync_notify](../../sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/asc_sync_wait.md)），保证多个`asc_copy_gm2ub_align`指令的串行化，防止出现异常数据。 | 161 | +### 连续数据搬运约束 |
| 162 | + | ||
| 163 | +- 若`size`非32字节对齐，搬运数据会补齐至32字节对齐，目的UB需要预留补齐后的空间。手动填充时，调用`asc_set_copy_pad_val`配置填充值；自动填充时，由硬件填充dummy假数据，dummy假数据的值为数据块的第一个元素的值。 | ||
| 164 | + | ||
| 165 | +### 高维切分数据搬运约束 | ||
| 166 | + | ||
| 163 | - 当`left_padding_num`或`right_padding_num`非0时，`enable_constant_pad`不生效，必须在搬运前调用`asc_set_copy_pad_val`配置填充值。`left_padding_num`、`right_padding_num`对应的填充数据大小均不能超过32字节。 | 167 | - 当`left_padding_num`或`right_padding_num`非0时，`enable_constant_pad`不生效，必须在搬运前调用`asc_set_copy_pad_val`配置填充值。`left_padding_num`、`right_padding_num`对应的填充数据大小均不能超过32字节。 |
| 164 | -- 前n个数据搬运接口：若`size`非32字节对齐，搬运数据会补齐至32字节对齐，目的UB需要预留补齐后的空间。手动填充时，调用`asc_set_copy_pad_val`配置填充值；自动填充时，由硬件填充dummy假数据，dummy假数据的值为数据块的第一个元素的值。 | ||
| 165 | - 当`dst_stride`不等于`len_burst`时，`dst_stride`要求32字节对齐。 | 168 | - 当`dst_stride`不等于`len_burst`时，`dst_stride`要求32字节对齐。 |
| 166 | 169 | ||
| 167 | ## 调用示例 | 170 | ## 调用示例 |
| 168 | 171 | ||
| 169 | -```cpp | 172 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 170 | -asc_set_gm2ub_loop_size(2, 2); | 173 | + |
| 171 | -asc_set_gm2ub_loop1_stride(96, 128); | 174 | +<!-- npu="950" id8 --> |
| 172 | -asc_set_gm2ub_loop2_stride(192, 288); | 175 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： |
| 173 | -asc_copy_gm2ub_align(dst, src, 2, 48 * sizeof(int8_t), 0, 0, false, asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, 48 * sizeof(int8_t), 48 * sizeof(int8_t)); | 176 | + |
| 174 | -asc_set_gm2ub_loop_size(1, 1); | 177 | +```bash |
| 178 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 179 | +``` | ||
| 180 | +<!-- end id8 --> | ||
| 181 | + | ||
| 182 | +```c | ||
| 183 | +#include <cstdint> | ||
| 184 | +#include <iostream> | ||
| 185 | +#include <vector> | ||
| 186 | +#include "c_api/asc_simd.h" | ||
| 187 | +#include "acl/acl.h" | ||
| 188 | + | ||
| 189 | +namespace { | ||
| 190 | + | ||
| 191 | +constexpr uint32_t INPUT_BYTES = 256; | ||
| 192 | +constexpr uint32_t OUTPUT_BYTES = 256; | ||
| 193 | + | ||
| 194 | +__global__ __vector__ void asc_copy_gm2ub_align_arch3510_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input) | ||
| 195 | +{ | ||
| 196 | + asc_init(); | ||
| 197 | + __ubuf__ uint8_t local[INPUT_BYTES]; | ||
| 198 | + // Copy INPUT_BYTES from GM to UB, then wait only for PIPE_MTE2. | ||
| 199 | + asc_copy_gm2ub_align(local, input, INPUT_BYTES); | ||
| 200 | + asc_sync_mte2(0); | ||
| 201 | + asc_copy_ub2gm_align(output, local, INPUT_BYTES); | ||
| 202 | + asc_sync_mte3(0); | ||
| 203 | +} | ||
| 204 | + | ||
| 205 | +void print_data(const char* name, const std::vector<uint8_t>& data) | ||
| 206 | +{ | ||
| 207 | + std::cout << name << ":"; | ||
| 208 | + const uint32_t count = data.size() < 32 ? data.size() : 32; | ||
| 209 | + for (uint32_t i = 0; i < count; ++i) std::cout << ' ' << +data[i]; | ||
| 210 | + if (data.size() > count) std::cout << " ..."; | ||
| 211 | + std::cout << std::endl; | ||
| 212 | +} | ||
| 213 | +} // namespace | ||
| 214 | + | ||
| 215 | +int main() | ||
| 216 | +{ | ||
| 217 | + std::vector<uint8_t> input(INPUT_BYTES), output(OUTPUT_BYTES, 0), golden(OUTPUT_BYTES, 0); | ||
| 218 | + for (uint32_t i = 0; i < INPUT_BYTES; ++i) input[i] = static_cast<uint8_t>(i + 1); | ||
| 219 | + for (uint32_t i = 0; i < 256; ++i) golden[i] = input[i]; | ||
| 220 | + aclInit(nullptr); | ||
| 221 | + aclrtSetDevice(0); | ||
| 222 | + uint8_t *input_device = nullptr, *output_device = nullptr; | ||
| 223 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), INPUT_BYTES, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 224 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), OUTPUT_BYTES, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 225 | + aclrtMemcpy(input_device, INPUT_BYTES, input.data(), INPUT_BYTES, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 226 | + aclrtMemcpy(output_device, OUTPUT_BYTES, output.data(), OUTPUT_BYTES, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 227 | + asc_copy_gm2ub_align_arch3510_kernel<<<1, 0>>>(output_device, input_device); | ||
| 228 | + aclrtSynchronizeDevice(); | ||
| 229 | + aclrtMemcpy(output.data(), OUTPUT_BYTES, output_device, OUTPUT_BYTES, ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 230 | + print_data("Input", input); | ||
| 231 | + print_data("Output", output); | ||
| 232 | + print_data("Golden", golden); | ||
| 233 | + const bool passed = output == golden; | ||
| 234 | + std::cout << (passed ? "[Success] asc_copy_gm2ub_align passed." : "[Failed] asc_copy_gm2ub_align failed.") << std::endl; | ||
| 235 | + aclrtFree(input_device); | ||
| 236 | + aclrtFree(output_device); | ||
| 237 | + aclrtResetDevice(0); | ||
| 238 | + aclFinalize(); | ||
| 239 | + return passed ? 0 : 1; | ||
| 240 | +} | ||
| 175 | ``` | 241 | ``` |
需要结合上下文添加注释说明API功能