已合并
完善初始化、比较选择、类型转换、系统变量和规约类capi资料 #5042
Cruiter创建于 27 天前
完善初始化、比较选择、类型转换、系统变量和规约类capi资料 #5042
已合并
共 12 个文件变更+1670-315
| @@ -26,7 +26,45 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -初始化NPU状态。 | 29 | +本接口将当前AI Core上可能由前序算子遗留的全局状态恢复为接口约定的初始状态，避免这些状态影响当前算子的计算结果。本接口仅初始化控制状态，不会清零Global Memory、Local Memory或用户数据。 |
| 30 | + | ||
| 31 | +不同产品的初始化范围和初始化后的状态如下。 | ||
| 32 | + | ||
| 33 | +<!-- npu="950" id8 --> | ||
| 34 | +**表 1** Ascend 950PR/Ascend 950DT初始化结果 | ||
| 35 | + | ||
| 36 | +| 状态类别 | 初始化项 | 初始化后的状态 | | ||
| 37 | +| -------- | -------------------------------------------------- | ------------------------------------------------------------------------------------------------------------------------------ | | ||
| 38 | +| 公共状态 | 原子累加状态 | 清空原子累加状态，后续数据搬运不执行原子累加。 | | ||
| 39 | +| 公共状态 | `Mask`工作模式 | 设置为`Normal`模式。 | | ||
| 40 | +| 公共状态 | `CTRL`寄存器 | 将`CTRL`寄存器恢复为接口约定的初始状态，其中保持位的值保持调用本接口前的值。 | | ||
| 41 | +| 公共状态 | 原子操作开启位与原子操作类型 | 寄存器配置值设置为`0b00100100`，其中数据类型字段为`half`，操作类型字段为求和。原子累加状态已清空，因此初始化后不执行原子累加。 | | ||
| 42 | +| AIC | `asc_copy_l12l0a`进行3D格式搬运时的`Padding`填充值 | 设置为0。 | | ||
| 43 | +| AIV | `Mask` | 高位`Mask`和低位`Mask`均设置为`0xffffffffffffffff`。在`Normal`模式下，单次迭代内的所有元素均参与计算。 | | ||
| 44 | +| AIV | `GM->UB`与`UB->GM`搬运的循环次数 | 内层循环次数和外层循环次数均设置为1，恢复为普通搬运模式。 | | ||
| 45 | +<!-- end id8 --> | ||
| 46 | + | ||
| 47 | +<!-- npu="A3" id9 --> | ||
| 48 | +**表 2** Atlas A3 训练系列产品/Atlas A3 推理系列产品初始化结果 | ||
| 49 | + | ||
| 50 | +| 状态类别 | 初始化项 | 初始化后的状态 | | ||
| 51 | +| -------- | -------------- | ------------------------------------------------------------------------------------------------------ | | ||
| 52 | +| AIC | 全局状态 | 不修改Cube Core的状态。 | | ||
| 53 | +| AIV | 原子累加状态 | 清空原子累加状态，后续数据搬运不执行原子累加。 | | ||
| 54 | +| AIV | `Mask`工作模式 | 设置为`Normal`模式。 | | ||
| 55 | +| AIV | `Mask` | 高位`Mask`和低位`Mask`均设置为`0xffffffffffffffff`。在`Normal`模式下，单次迭代内的所有元素均参与计算。 | | ||
| 56 | +<!-- end id9 --> | ||
| 57 | + | ||
| 58 | +<!-- npu="910b" id10 --> | ||
| 59 | +**表 3** Atlas A2 训练系列产品/Atlas A2 推理系列产品初始化结果 | ||
| 60 | + | ||
| 61 | +| 状态类别 | 初始化项 | 初始化后的状态 | | ||
| 62 | +| -------- | -------------- | ------------------------------------------------------------------------------------------------------ | | ||
| 63 | +| AIC | 全局状态 | 不修改Cube Core的状态。 | | ||
| 64 | +| AIV | 原子累加状态 | 清空原子累加状态，后续数据搬运不执行原子累加。 | | ||
| 65 | +| AIV | `Mask`工作模式 | 设置为`Normal`模式。 | | ||
| 66 | +| AIV | `Mask` | 高位`Mask`和低位`Mask`均设置为`0xffffffffffffffff`。在`Normal`模式下，单次迭代内的所有元素均参与计算。 | | ||
| 67 | +<!-- end id10 --> | ||
| 30 | 68 | ||
| 31 | ## 函数原型 | 69 | ## 函数原型 |
| 32 | 70 | ||
| @@ -44,12 +82,85 @@ __aicore__ inline void asc_init() | |||
| 44 | 82 | ||
| 45 | ## 流水类型 | 83 | ## 流水类型 |
| 46 | 84 | ||
| 47 | -PIPE_S | 85 | +`PIPE_S` |
| 48 | 86 | ||
| 49 | ## 约束说明 | 87 | ## 约束说明 |
| 50 | 88 | ||
| 51 | -无 | 89 | +在实际运行中，这些值可能被前序执行的算子修改，若不调用该接口进行初始化，非预期的值可能导致计算结果出现精度错误。 |
| 90 | + | ||
| 91 | +例如前序算子使用`Counter`模式但未重置为`Normal`模式，当前算子以默认的`Normal`模式设置`Mask`时，会导致`Mask`设置不符合预期，进而引发精度错误。 | ||
| 52 | 92 | ||
| 53 | ## 调用示例 | 93 | ## 调用示例 |
| 54 | 94 | ||
| 55 | -无 | 95 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 96 | + | ||
| 97 | +<!-- npu="950" id11 --> | ||
| 98 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 99 | + | ||
| 100 | +```bash | ||
| 101 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 102 | +``` | ||
| 103 | +<!-- end id11 --> | ||
| 104 | + | ||
| 105 | +```c | ||
| 106 | +#include <cstdint> | ||
| 107 | +#include <iostream> | ||
| 108 | +#include <vector> | ||
| 109 | +#include "c_api/asc_simd.h" | ||
| 110 | +#include "acl/acl.h" | ||
| 111 | + | ||
| 112 | +namespace { | ||
| 113 | +constexpr uint32_t ELEMENTS = 16; | ||
| 114 | +constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t); | ||
| 115 | + | ||
| 116 | +void print_data(const char* label, const std::vector<uint64_t>& data) | ||
| 117 | +{ | ||
| 118 | + std::cout << label << ":"; | ||
| 119 | + const size_t count = data.size() < 8 ? data.size() : 8; | ||
| 120 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << data[i]; | ||
| 121 | + if (data.size() > count) std::cout << " ..."; | ||
| 122 | + std::cout << std::endl; | ||
| 123 | +} | ||
| 124 | + | ||
| 125 | +__global__ __vector__ void asc_init_kernel(__gm__ uint64_t* output, __gm__ uint64_t* input) | ||
| 126 | +{ | ||
| 127 | + asc_init(); | ||
| 128 | + __ubuf__ uint32_t local[BYTES / sizeof(uint32_t)]; | ||
| 129 | + asc_copy_gm2ub_align(local, reinterpret_cast<__gm__ uint32_t*>(input), BYTES); | ||
| 130 | + asc_sync_mte2(0); | ||
| 131 | + asc_set_atomic_add_int(); | ||
| 132 | + // Reinitialize the core and clear persistent state such as the atomic mode. | ||
| 133 | + asc_init(); | ||
| 134 | + // asc_init resets the atomic mode, so this is an ordinary UB-to-GM copy. | ||
| 135 | + asc_copy_ub2gm(reinterpret_cast<__gm__ uint32_t*>(output), local, BYTES); | ||
| 136 | + asc_sync_mte3(0); | ||
| 137 | +} | ||
| 138 | + | ||
| 139 | +} // namespace | ||
| 140 | + | ||
| 141 | +int main() | ||
| 142 | +{ | ||
| 143 | + std::vector<uint64_t> input(ELEMENTS), output(ELEMENTS, 0); | ||
| 144 | + for (uint32_t i = 0; i < ELEMENTS; ++i) { input[i] = i + 1; output[i] = 10; } | ||
| 145 | + aclInit(nullptr); | ||
| 146 | + aclrtSetDevice(0); | ||
| 147 | + uint64_t *input_device = nullptr, *output_device = nullptr; | ||
| 148 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 149 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 150 | + aclrtMemcpy(input_device, BYTES, input.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 151 | + aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 152 | + asc_init_kernel<<<1, 0>>>(output_device, input_device); | ||
| 153 | + aclrtSynchronizeDevice(); | ||
| 154 | + aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 155 | + print_data("Input", input); | ||
| 156 | + print_data("Output after asc_init reset", output); | ||
| 157 | + print_data("Golden ordinary copy", input); | ||
| 158 | + const bool passed = output == input; | ||
| 159 | + std::cout << (passed ? "[Success] asc_init passed." : "[Failed] asc_init failed.") << std::endl; | ||
| 160 | + aclrtFree(input_device); | ||
| 161 | + aclrtFree(output_device); | ||
| 162 | + aclrtResetDevice(0); | ||
| 163 | + aclFinalize(); | ||
| 164 | + return passed ? 0 : 1; | ||
| 165 | +} | ||
| 166 | +``` | ||
| @@ -25,41 +25,51 @@ | |||
| 25 | <!-- end id7 --> | 25 | <!-- end id7 --> |
| 26 | ## 功能说明 | 26 | ## 功能说明 |
| 27 | 27 | ||
| 28 | -根据mask的比特位值，从源操作数src0，src1中选择元素，得到目的操作数dst。选择的规则为：当mask的比特位为1时，从src0中选取对应位置的数；当mask的比特位为0时，从src1中选取对应位置的数。 | 28 | +根据`mask`从源操作数`src0`、`src1`中选择元素，得到目的操作数`dst`。选择的规则为：当`mask`的比特位为1时，从`src0`中选取对应位置的数；当`mask`的比特位为0时，从`src1`中选取对应位置的数。计算公式如下： |
| 29 | - | ||
| 30 | -计算公式如下： | ||
| 31 | 29 | ||
| 32 | $$ | 30 | $$ |
| 33 | -dst_i = (mask_i == 1)\ ?\ src0_i : src1_i | 31 | +dst_i = |
| 32 | +\begin{cases} | ||
| 33 | + src0_i, & mask_i = 1 \\ | ||
| 34 | + src1_i, & mask_i = 0 \\ | ||
| 35 | +\end{cases} | ||
| 34 | $$ | 36 | $$ |
| 35 | 37 | ||
| 36 | ## 函数原型 | 38 | ## 函数原型 |
| 37 | 39 | ||
| 38 | -```cpp | 40 | +### 模板原型（占位符形式） |
| 39 | -__simd_callee__ inline void asc_select(vector_bool& dst, vector_bool src0, vector_bool src1, vector_bool mask) | 41 | + |
| 40 | -__simd_callee__ inline void asc_select(vector_uint8_t& dst, vector_uint8_t src0, vector_uint8_t src1, vector_bool mask) | 42 | +```c |
| 41 | -__simd_callee__ inline void asc_select(vector_int8_t& dst, vector_int8_t src0, vector_int8_t src1, vector_bool mask) | 43 | +__simd_callee__ inline void asc_select(vector_<dtype>& dst, |
| 42 | -__simd_callee__ inline void asc_select(vector_fp8_e4m3fn_t& dst, vector_fp8_e4m3fn_t src0, vector_fp8_e4m3fn_t src1, vector_bool mask) | 44 | + vector_<dtype> src0, |
| 43 | -__simd_callee__ inline void asc_select(vector_hifloat8_t& dst, vector_hifloat8_t src0, vector_hifloat8_t src1, vector_bool mask) | 45 | + vector_<dtype> src1, |
| 44 | -__simd_callee__ inline void asc_select(vector_fp8_e5m2_t& dst, vector_fp8_e5m2_t src0, vector_fp8_e5m2_t src1, vector_bool mask) | 46 | + vector_bool mask) |
| 45 | -__simd_callee__ inline void asc_select(vector_uint16_t& dst, vector_uint16_t src0, vector_uint16_t src1, vector_bool mask) | 47 | +``` |
| 46 | -__simd_callee__ inline void asc_select(vector_int16_t& dst, vector_int16_t src0, vector_int16_t src1, vector_bool mask) | 48 | + |
| 47 | -__simd_callee__ inline void asc_select(vector_half& dst, vector_half src0, vector_half src1, vector_bool mask) | 49 | +### dtype支持数据类型 |
| 48 | -__simd_callee__ inline void asc_select(vector_bfloat16_t& dst, vector_bfloat16_t src0, vector_bfloat16_t src1, vector_bool mask) | 50 | + |
| 49 | -__simd_callee__ inline void asc_select(vector_uint32_t& dst, vector_uint32_t src0, vector_uint32_t src1, vector_bool mask) | 51 | +`dtype`支持的数据类型为`bool`、`int8_t`、`uint8_t`、`hifloat8_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。 |
| 50 | -__simd_callee__ inline void asc_select(vector_int32_t& dst, vector_int32_t src0, vector_int32_t src1, vector_bool mask) | 52 | + |
| 51 | -__simd_callee__ inline void asc_select(vector_float& dst, vector_float src0, vector_float src1, vector_bool mask) | 53 | +### 典型示例 |
| 54 | + | ||
| 55 | +```c | ||
| 56 | +// 示例：对half矢量数据寄存器执行按掩码选择 | ||
| 57 | +__simd_callee__ inline void asc_select(vector_half& dst, | ||
| 58 | + vector_half src0, | ||
| 59 | + vector_half src1, | ||
| 60 | + vector_bool mask) | ||
| 52 | ``` | 61 | ``` |
| 53 | 62 | ||
| 54 | ## 参数说明 | 63 | ## 参数说明 |
| 55 | 64 | ||
| 56 | **表1** 参数说明 | 65 | **表1** 参数说明 |
| 57 | -| 参数名 | 输入/输出 | 描述 | | 66 | + |
| 58 | -| :----- | :------- | :------- | | 67 | +| 参数名 | 输入/输出 | 描述 | |
| 59 | -| dst | 输出 | 目的操作数（掩码寄存器或矢量数据寄存器）。 | | 68 | +| :----- | :-------- | :------------------------------------------------------------------------------------------------------------------------------- | |
| 60 | -| src0 | 输入 | 源操作数0（掩码寄存器或矢量数据寄存器）。 | | 69 | +| `dst` | 输出 | 目的操作数（掩码寄存器或矢量数据寄存器）。 | |
| 61 | -| src1 | 输入 | 源操作数1（掩码寄存器或矢量数据寄存器）。 | | 70 | +| `src0` | 输入 | 源操作数0（掩码寄存器或矢量数据寄存器）。 | |
| 62 | -| mask | 输入 | 源操作数掩码（掩码寄存器）。指定选择src0或src1为有效数据。mask的比特位为1时，选取src0；mask的比特位为0时，选取src1。 | | 71 | +| `src1` | 输入 | 源操作数1（掩码寄存器或矢量数据寄存器）。 | |
| 72 | +| `mask` | 输入 | 源操作数掩码（掩码寄存器）。指定选择`src0`或`src1`为有效数据。`mask`的比特位为1时，选取`src0`；`mask`的比特位为0时，选取`src1`。 | | ||
| 63 | 73 | ||
| 64 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 74 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 65 | 75 | ||
| @@ -69,21 +79,128 @@ __simd_callee__ inline void asc_select(vector_float& dst, vector_float src0, vec | |||
| 69 | 79 | ||
| 70 | ## 约束说明 | 80 | ## 约束说明 |
| 71 | 81 | ||
| 72 | -通过mask参数控制的未选中的元素保持源操作数的值而非清零。 | 82 | +- 通过`mask`参数控制的未选中的元素保持源操作数的值而非清零。 |
| 83 | +- `src0`、`src1`和`dst`的数据类型需要保持一致。 | ||
| 73 | 84 | ||
| 74 | ## 调用示例 | 85 | ## 调用示例 |
| 75 | 86 | ||
| 76 | -```cpp | 87 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 77 | -__simd_vf__ inline void select_vf(__ubuf__ half* src0_addr, __ubuf__ half* src1_addr, __ubuf__ half* dst_addr, uint32_t count, uint16_t one_repeat_size, uint16_t repeat_time) | 88 | + |
| 89 | +<!-- npu="950" id8 --> | ||
| 90 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 91 | + | ||
| 92 | +```bash | ||
| 93 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 94 | +``` | ||
| 95 | +<!-- end id8 --> | ||
| 96 | + | ||
| 97 | +```c | ||
| 98 | +#include <cstdint> | ||
| 99 | +#include <iostream> | ||
| 100 | +#include <vector> | ||
| 101 | +#include "c_api/asc_simd.h" | ||
| 102 | +#include "acl/acl.h" | ||
| 103 | +namespace { | ||
| 104 | +template <typename T> | ||
| 105 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 78 | { | 106 | { |
| 79 | - vector_half src0, src1, dst; | 107 | + std::cout << label << ":"; |
| 80 | - vector_bool mask; | 108 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 81 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 109 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 82 | - mask = asc_update_mask_b16(count); | 110 | + if (values.size() > count) std::cout << " ..."; |
| 83 | - asc_loadalign_postupdate(src0, src0_addr, one_repeat_size); | 111 | + std::cout << std::endl; |
| 84 | - asc_loadalign_postupdate(src1, src1_addr, one_repeat_size); | ||
| 85 | - asc_select(dst, src0, src1, mask); | ||
| 86 | - asc_storealign_postupdate(dst_addr, dst, one_repeat_size, mask); | ||
| 87 | - } | ||
| 88 | } | 112 | } |
| 89 | -``` | 113 | + |
| 114 | +template <typename T> | ||
| 115 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 116 | +{ | ||
| 117 | + if (actual.size() != expected.size()) return false; | ||
| 118 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 119 | + if (actual[i] == expected[i]) continue; | ||
| 120 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 121 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 122 | + } | ||
| 123 | + return true; | ||
| 124 | +} | ||
| 125 | + | ||
| 126 | +constexpr uint32_t ELEMENT_COUNT = 64; | ||
| 127 | + | ||
| 128 | +__simd_vf__ inline void select(__ubuf__ float* dst, __ubuf__ float* src0, __ubuf__ float* src1) | ||
| 129 | +{ | ||
| 130 | + vector_float dst_reg; | ||
| 131 | + vector_float src0_reg; | ||
| 132 | + vector_float src1_reg; | ||
| 133 | + vector_bool predicate; | ||
| 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_gt(predicate, src0_reg, src1_reg, mask); | ||
| 139 | + asc_select(dst_reg, src0_reg, src1_reg, predicate); | ||
| 140 | + asc_storealign(dst, dst_reg, mask); | ||
| 141 | +} | ||
| 142 | + | ||
| 143 | +__global__ __vector__ void asc_select_kernel(__gm__ float* dst, __gm__ float* src0, __gm__ float* src1) | ||
| 144 | +{ | ||
| 145 | + asc_init(); | ||
| 146 | + __ubuf__ float dst_local[ELEMENT_COUNT]; | ||
| 147 | + __ubuf__ float src0_local[ELEMENT_COUNT]; | ||
| 148 | + __ubuf__ float src1_local[ELEMENT_COUNT]; | ||
| 149 | + asc_copy_gm2ub_align(src0_local, src0, ELEMENT_COUNT * sizeof(float)); | ||
| 150 | + asc_copy_gm2ub_align(src1_local, src1, ELEMENT_COUNT * sizeof(float)); | ||
| 151 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 152 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 153 | + select(dst_local, src0_local, src1_local); | ||
| 154 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 155 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 156 | + asc_copy_ub2gm_align(dst, dst_local, ELEMENT_COUNT * sizeof(float)); | ||
| 157 | + asc_sync(); | ||
| 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); | ||
| 169 | + src1[i] = static_cast<float>(ELEMENT_COUNT - i); | ||
| 170 | + golden[i] = src0[i] > src1[i] ? src0[i] : src1[i]; | ||
| 171 | + } | ||
| 172 | + aclInit(nullptr); | ||
| 173 | + aclrtSetDevice(0); | ||
| 174 | + float* src0_device = nullptr; | ||
| 175 | + aclrtMalloc(reinterpret_cast<void**>(&src0_device), (ELEMENT_COUNT) * sizeof(float), | ||
| 176 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 177 | + float* src1_device = nullptr; | ||
| 178 | + aclrtMalloc(reinterpret_cast<void**>(&src1_device), (ELEMENT_COUNT) * sizeof(float), | ||
| 179 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 180 | + float* dst_device = nullptr; | ||
| 181 | + aclrtMalloc(reinterpret_cast<void**>(&dst_device), (ELEMENT_COUNT) * sizeof(float), | ||
| 182 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 183 | + aclrtMemcpy(src0_device, src0.size() * sizeof(float), src0.data(), src0.size() * sizeof(float), | ||
| 184 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 185 | + aclrtMemcpy(src1_device, src1.size() * sizeof(float), src1.data(), src1.size() * sizeof(float), | ||
| 186 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 187 | + asc_select_kernel<<<1, 0>>>(dst_device, src0_device, src1_device); | ||
| 188 | + aclrtSynchronizeDevice(); | ||
| 189 | + aclrtMemcpy(output.data(), output.size() * sizeof(float), dst_device, output.size() * sizeof(float), | ||
| 190 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 191 | + print_data("Input src0", src0); | ||
| 192 | + print_data("Input src1", src1); | ||
| 193 | + print_data("Output", output); | ||
| 194 | + print_data("Golden", golden); | ||
| 195 | + const bool passed = compare_data(output, golden); | ||
| 196 | + std::cout << (passed ? "[Success] asc_select passed." | ||
| 197 | + : "[Failed] asc_select failed.") | ||
| 198 | + << std::endl; | ||
| 199 | + aclrtFree(dst_device); | ||
| 200 | + aclrtFree(src0_device); | ||
| 201 | + aclrtFree(src1_device); | ||
| 202 | + aclrtResetDevice(0); | ||
| 203 | + aclFinalize(); | ||
| 204 | + return passed ? 0 : 1; | ||
| 205 | +} | ||
| 206 | +``` | ||
| @@ -26,66 +26,64 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将float类型数据转换成bfloat16_t类型，并支持多种舍入模式与饱和模式。 | 29 | +根据`mask`将`src`中的每个`float`类型数据转换成`bfloat16_t`类型，并支持多种舍入模式与饱和模式，结果写入`dst`。 |
| 30 | +由于源操作数与目的操作数类型位宽比为2:1，写入数据时需要将一个`VL`大小的数据分为两部分，根据不同接口选择数据写入索引为奇数的位置或偶数的位置。伪代码如下： | ||
| 30 | 31 | ||
| 31 | -关于舍入模式和饱和/非饱和模式的详细说明，请参见[舍入模式与饱和模式](rounding_mode.md)。 | 32 | +```python |
| 33 | +def asc_float2bfloat16_rn(dst, src, mask): | ||
| 34 | + for i in range(64): | ||
| 35 | + if mask[i]: | ||
| 36 | + dst[2 * i] = bfloat16(src[i]) # float -> bfloat16, RINT舍入, 写入偶数索引 | ||
| 37 | + else: | ||
| 38 | + dst[2 * i] = 0 | ||
| 39 | + for i in range(64): | ||
| 40 | + dst[2 * i + 1] = 0 # 奇数索引置零 | ||
| 32 | 41 | ||
| 33 | -由于源操作数与目的操作数类型位宽比为2:1，写入数据时需要将一个VL大小的数据分为两部分，根据不同接口选择数据写入索引为奇数的位置或偶数的位置。 | 42 | +def asc_float2bfloat16_rn_v2(dst, src, mask): |
| 43 | + for i in range(64): | ||
| 44 | + dst[2 * i] = 0 # 偶数索引置零 | ||
| 45 | + if mask[i]: | ||
| 46 | + dst[2 * i + 1] = bfloat16(src[i]) # float -> bfloat16, RINT舍入, 写入奇数索引 | ||
| 47 | + else: | ||
| 48 | + dst[2 * i + 1] = 0 | ||
| 49 | +``` | ||
| 50 | + | ||
| 51 | +关于舍入模式的详细说明，请参见[舍入模式](../data_type_convert/rounding_mode.md)。 | ||
| 34 | 52 | ||
| 35 | ## 函数原型 | 53 | ## 函数原型 |
| 36 | 54 | ||
| 37 | -```cpp | 55 | +### 模板原型（占位符形式） |
| 38 | -// FLOOR舍入模式，非饱和模式，数据写入索引为偶数的位置 | 56 | + |
| 39 | -__simd_callee__ inline void asc_float2bfloat16_rd(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | 57 | +```c |
| 40 | -// FLOOR舍入模式，饱和模式，数据写入索引为偶数的位置 | 58 | +__simd_callee__ inline void asc_float2bfloat16_<round_mode><sat_mode><position_mode>(vector_bfloat16_t& dst, |
| 41 | -__simd_callee__ inline void asc_float2bfloat16_rd_sat(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | 59 | + vector_float src, |
| 42 | -// FLOOR舍入模式，非饱和模式，数据写入索引为奇数的位置 | 60 | + vector_bool mask) |
| 43 | -__simd_callee__ inline void asc_float2bfloat16_rd_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | 61 | +``` |
| 44 | -// FLOOR舍入模式，饱和模式，数据写入索引为奇数的位置 | 62 | + |
| 45 | -__simd_callee__ inline void asc_float2bfloat16_rd_sat_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | 63 | +### 占位符说明 |
| 64 | + | ||
| 65 | +- `<round_mode>`表示支持的舍入模式，支持`rd`（`FLOOR`）、`rn`（`RINT`）、`rna`（`ROUND`）、`ru`（`CEIL`）和`rz`（`TRUNC`）。 | ||
| 66 | +- `<sat_mode>`表示饱和非饱和模式，为空时表示非饱和模式，取值为`_sat`时表示饱和模式。 | ||
| 67 | +- `<position_mode>`表示`dst`写入的位置，为空时，结果写入`dst`的偶数索引位置；取值为`_v2`时，结果写入`dst`的奇数索引位置。 | ||
| 68 | + | ||
| 69 | +### 典型示例 | ||
| 70 | + | ||
| 71 | +```c | ||
| 46 | // RINT舍入模式，非饱和模式，数据写入索引为偶数的位置 | 72 | // RINT舍入模式，非饱和模式，数据写入索引为偶数的位置 |
| 47 | -__simd_callee__ inline void asc_float2bfloat16_rn(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | 73 | +__simd_callee__ inline void asc_float2bfloat16_rn(vector_bfloat16_t& dst, |
| 48 | -// RINT舍入模式，饱和模式，数据写入索引为偶数的位置 | 74 | + vector_float src, |
| 49 | -__simd_callee__ inline void asc_float2bfloat16_rn_sat(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | 75 | + vector_bool mask) |
| 50 | -// RINT舍入模式，非饱和模式，数据写入索引为奇数的位置 | ||
| 51 | -__simd_callee__ inline void asc_float2bfloat16_rn_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 52 | -// RINT舍入模式，饱和模式，数据写入索引为奇数的位置 | ||
| 53 | -__simd_callee__ inline void asc_float2bfloat16_rn_sat_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 54 | -// ROUND舍入模式，非饱和模式，数据写入索引为偶数的位置 | ||
| 55 | -__simd_callee__ inline void asc_float2bfloat16_rna(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 56 | -// ROUND舍入模式，饱和模式，数据写入索引为偶数的位置 | ||
| 57 | -__simd_callee__ inline void asc_float2bfloat16_rna_sat(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 58 | -// ROUND舍入模式，非饱和模式，数据写入索引为奇数的位置 | ||
| 59 | -__simd_callee__ inline void asc_float2bfloat16_rna_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 60 | -// ROUND舍入模式，饱和模式，数据写入索引为奇数的位置 | ||
| 61 | -__simd_callee__ inline void asc_float2bfloat16_rna_sat_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 62 | -// CEIL舍入模式，非饱和模式，数据写入索引为偶数的位置 | ||
| 63 | -__simd_callee__ inline void asc_float2bfloat16_ru(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 64 | -// CEIL舍入模式，饱和模式，数据写入索引为偶数的位置 | ||
| 65 | -__simd_callee__ inline void asc_float2bfloat16_ru_sat(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 66 | -// CEIL舍入模式，非饱和模式，数据写入索引为奇数的位置 | ||
| 67 | -__simd_callee__ inline void asc_float2bfloat16_ru_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 68 | -// CEIL舍入模式，饱和模式，数据写入索引为奇数的位置 | ||
| 69 | -__simd_callee__ inline void asc_float2bfloat16_ru_sat_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 70 | -// TRUNC舍入模式，非饱和模式，数据写入索引为偶数的位置 | ||
| 71 | -__simd_callee__ inline void asc_float2bfloat16_rz(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 72 | -// TRUNC舍入模式，饱和模式，数据写入索引为偶数的位置 | ||
| 73 | -__simd_callee__ inline void asc_float2bfloat16_rz_sat(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 74 | -// TRUNC舍入模式，非饱和模式，数据写入索引为奇数的位置 | ||
| 75 | -__simd_callee__ inline void asc_float2bfloat16_rz_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 76 | -// TRUNC舍入模式，饱和模式，数据写入索引为奇数的位置 | ||
| 77 | -__simd_callee__ inline void asc_float2bfloat16_rz_sat_v2(vector_bfloat16_t& dst, vector_float src, vector_bool mask) | ||
| 78 | ``` | 76 | ``` |
| 79 | 77 | ||
| 80 | ## 参数说明 | 78 | ## 参数说明 |
| 81 | 79 | ||
| 82 | **表1** 参数说明 | 80 | **表1** 参数说明 |
| 83 | 81 | ||
| 84 | -| 参数名 | 输入/输出 | 描述 | | 82 | +| 参数名 | 输入/输出 | 描述 | |
| 85 | -| --- | --- | --- | | 83 | +| ------ | --------- | ------------------------------------------------------------------------------------------------------------------ | |
| 86 | -| dst | 输出 | 目的操作数（矢量数据寄存器）。| | 84 | +| `dst` | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 87 | -| src | 输入 | 源操作数（矢量数据寄存器）。| | 85 | +| `src` | 输入 | 源操作数（矢量数据寄存器）。 | |
| 88 | -| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在输出中置零。 | | 86 | +| `mask` | 输入 | 掩码寄存器，用于控制各元素是否参与计算。`mask`中与元素对应的比特位为1时，该元素参与计算；为0时，该元素不参与计算。 | |
| 89 | 87 | ||
| 90 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 88 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 91 | 89 | ||
| @@ -95,19 +93,249 @@ __simd_callee__ inline void asc_float2bfloat16_rz_sat_v2(vector_bfloat16_t& dst, | |||
| 95 | 93 | ||
| 96 | ## 约束说明 | 94 | ## 约束说明 |
| 97 | 95 | ||
| 98 | -- mask控制源操作数是否参与计算，源操作数不参与计算的元素在输出对应位置置零。 | 96 | +- `src`与`dst`的数据类型需要与函数原型匹配。 |
| 99 | -- 当数据写入目的操作数索引为奇数的位置，索引为偶数位置的数值置零。反之亦然。 | 97 | +- `mask`掩码位为0时，`dst`对应元素置0。 |
| 100 | -- 开启饱和模式和非饱和模式时，需配置ctrl寄存器，ctrl寄存器的详细说明请参见[asc_set_ctrl.md](../../sys_var/asc_set_ctrl.md)。 | 98 | +- 结果写入`dst`的奇数索引位置时，偶数索引位置置0；结果写入偶数索引位置时，奇数索引位置置0。 |
| 99 | +- 使用饱和或非饱和模式时，需要配置`ctrl`寄存器，详细说明请参见[asc_set_ctrl](../../sys_var/asc_set_ctrl.md)。 | ||
| 101 | 100 | ||
| 102 | ## 调用示例 | 101 | ## 调用示例 |
| 103 | 102 | ||
| 104 | -```cpp | 103 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 105 | -vector_bfloat16_t dst; | 104 | + |
| 106 | -vector_float src; | 105 | +<!-- npu="950" id8 --> |
| 107 | -vector_bool mask; | 106 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： |
| 108 | -mask = asc_create_mask_b32(PAT_ALL); | 107 | + |
| 109 | -asc_float2bfloat16_rd(dst, src, mask); // FLOOR舍入模式，不饱和模式，将src转换成vector_bfloat16_t类型并写入dst偶数索引位置 | 108 | +```bash |
| 110 | -asc_float2bfloat16_rd_sat(dst, src, mask); // FLOOR舍入模式，饱和模式，将src转换成vector_bfloat16_t类型并写入dst偶数索引位置 | 109 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main |
| 111 | -asc_float2bfloat16_rd_v2(dst, src, mask); // FLOOR舍入模式，不饱和模式，将src转换成vector_bfloat16_t类型并写入dst奇数索引位置 | 110 | +``` |
| 112 | -asc_float2bfloat16_rd_sat_v2(dst, src, mask); // FLOOR舍入模式，饱和模式，将src转换成vector_bfloat16_t类型并写入dst奇数索引位置 | 111 | + |
| 113 | -``` | 112 | +<!-- end id8 --> |
| 113 | + | ||
| 114 | +```c | ||
| 115 | +#include <cstdint> | ||
| 116 | +#include <iostream> | ||
| 117 | +#include <vector> | ||
| 118 | +#include "c_api/asc_simd.h" | ||
| 119 | +#include "acl/acl.h" | ||
| 120 | +namespace { | ||
| 121 | +template <typename T> | ||
| 122 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 123 | +{ | ||
| 124 | + std::cout << label << ":"; | ||
| 125 | + const size_t count = values.size() < 8 ? values.size() : 8; | ||
| 126 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; | ||
| 127 | + if (values.size() > count) std::cout << " ..."; | ||
| 128 | + std::cout << std::endl; | ||
| 129 | +} | ||
| 130 | + | ||
| 131 | +template <typename T> | ||
| 132 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 133 | +{ | ||
| 134 | + if (actual.size() != expected.size()) return false; | ||
| 135 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 136 | + if (actual[i] == expected[i]) continue; | ||
| 137 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 138 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 139 | + } | ||
| 140 | + return true; | ||
| 141 | +} | ||
| 142 | + | ||
| 143 | +constexpr uint32_t BUFFER_BYTES = 256; | ||
| 144 | +__simd_vf__ inline void convert(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input) | ||
| 145 | +{ | ||
| 146 | + vector_bool mask = asc_create_mask_b8(PAT_ALL); | ||
| 147 | + vector_bfloat16_t dst_0; | ||
| 148 | + vector_float src_0; | ||
| 149 | + asc_loadalign(dst_0, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 150 | + asc_loadalign(src_0, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 151 | + asc_float2bfloat16_rd_sat(dst_0, src_0, mask); | ||
| 152 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_0, mask); | ||
| 153 | + vector_bfloat16_t dst_1; | ||
| 154 | + vector_float src_1; | ||
| 155 | + asc_loadalign(dst_1, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 156 | + asc_loadalign(src_1, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 157 | + asc_float2bfloat16_rd_v2(dst_1, src_1, mask); | ||
| 158 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_1, mask); | ||
| 159 | + vector_bfloat16_t dst_2; | ||
| 160 | + vector_float src_2; | ||
| 161 | + asc_loadalign(dst_2, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 162 | + asc_loadalign(src_2, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 163 | + asc_float2bfloat16_rd_sat_v2(dst_2, src_2, mask); | ||
| 164 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_2, mask); | ||
| 165 | + vector_bfloat16_t dst_3; | ||
| 166 | + vector_float src_3; | ||
| 167 | + asc_loadalign(dst_3, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 168 | + asc_loadalign(src_3, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 169 | + asc_float2bfloat16_rn(dst_3, src_3, mask); | ||
| 170 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_3, mask); | ||
| 171 | + vector_bfloat16_t dst_4; | ||
| 172 | + vector_float src_4; | ||
| 173 | + asc_loadalign(dst_4, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 174 | + asc_loadalign(src_4, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 175 | + asc_float2bfloat16_rn_sat(dst_4, src_4, mask); | ||
| 176 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_4, mask); | ||
| 177 | + vector_bfloat16_t dst_5; | ||
| 178 | + vector_float src_5; | ||
| 179 | + asc_loadalign(dst_5, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 180 | + asc_loadalign(src_5, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 181 | + asc_float2bfloat16_rn_v2(dst_5, src_5, mask); | ||
| 182 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_5, mask); | ||
| 183 | + vector_bfloat16_t dst_6; | ||
| 184 | + vector_float src_6; | ||
| 185 | + asc_loadalign(dst_6, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 186 | + asc_loadalign(src_6, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 187 | + asc_float2bfloat16_rn_sat_v2(dst_6, src_6, mask); | ||
| 188 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_6, mask); | ||
| 189 | + vector_bfloat16_t dst_7; | ||
| 190 | + vector_float src_7; | ||
| 191 | + asc_loadalign(dst_7, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 192 | + asc_loadalign(src_7, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 193 | + asc_float2bfloat16_rna(dst_7, src_7, mask); | ||
| 194 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_7, mask); | ||
| 195 | + vector_bfloat16_t dst_8; | ||
| 196 | + vector_float src_8; | ||
| 197 | + asc_loadalign(dst_8, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 198 | + asc_loadalign(src_8, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 199 | + asc_float2bfloat16_rna_sat(dst_8, src_8, mask); | ||
| 200 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_8, mask); | ||
| 201 | + vector_bfloat16_t dst_9; | ||
| 202 | + vector_float src_9; | ||
| 203 | + asc_loadalign(dst_9, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 204 | + asc_loadalign(src_9, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 205 | + asc_float2bfloat16_rna_v2(dst_9, src_9, mask); | ||
| 206 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_9, mask); | ||
| 207 | + vector_bfloat16_t dst_10; | ||
| 208 | + vector_float src_10; | ||
| 209 | + asc_loadalign(dst_10, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 210 | + asc_loadalign(src_10, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 211 | + asc_float2bfloat16_rna_sat_v2(dst_10, src_10, mask); | ||
| 212 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_10, mask); | ||
| 213 | + vector_bfloat16_t dst_11; | ||
| 214 | + vector_float src_11; | ||
| 215 | + asc_loadalign(dst_11, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 216 | + asc_loadalign(src_11, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 217 | + asc_float2bfloat16_ru(dst_11, src_11, mask); | ||
| 218 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_11, mask); | ||
| 219 | + vector_bfloat16_t dst_12; | ||
| 220 | + vector_float src_12; | ||
| 221 | + asc_loadalign(dst_12, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 222 | + asc_loadalign(src_12, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 223 | + asc_float2bfloat16_ru_sat(dst_12, src_12, mask); | ||
| 224 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_12, mask); | ||
| 225 | + vector_bfloat16_t dst_13; | ||
| 226 | + vector_float src_13; | ||
| 227 | + asc_loadalign(dst_13, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 228 | + asc_loadalign(src_13, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 229 | + asc_float2bfloat16_ru_v2(dst_13, src_13, mask); | ||
| 230 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_13, mask); | ||
| 231 | + vector_bfloat16_t dst_14; | ||
| 232 | + vector_float src_14; | ||
| 233 | + asc_loadalign(dst_14, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 234 | + asc_loadalign(src_14, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 235 | + asc_float2bfloat16_ru_sat_v2(dst_14, src_14, mask); | ||
| 236 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_14, mask); | ||
| 237 | + vector_bfloat16_t dst_15; | ||
| 238 | + vector_float src_15; | ||
| 239 | + asc_loadalign(dst_15, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 240 | + asc_loadalign(src_15, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 241 | + asc_float2bfloat16_rz(dst_15, src_15, mask); | ||
| 242 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_15, mask); | ||
| 243 | + vector_bfloat16_t dst_16; | ||
| 244 | + vector_float src_16; | ||
| 245 | + asc_loadalign(dst_16, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 246 | + asc_loadalign(src_16, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 247 | + asc_float2bfloat16_rz_sat(dst_16, src_16, mask); | ||
| 248 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_16, mask); | ||
| 249 | + vector_bfloat16_t dst_17; | ||
| 250 | + vector_float src_17; | ||
| 251 | + asc_loadalign(dst_17, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 252 | + asc_loadalign(src_17, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 253 | + asc_float2bfloat16_rz_v2(dst_17, src_17, mask); | ||
| 254 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_17, mask); | ||
| 255 | + vector_bfloat16_t dst_18; | ||
| 256 | + vector_float src_18; | ||
| 257 | + asc_loadalign(dst_18, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 258 | + asc_loadalign(src_18, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 259 | + asc_float2bfloat16_rz_sat_v2(dst_18, src_18, mask); | ||
| 260 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_18, mask); | ||
| 261 | + vector_bfloat16_t dst_19; | ||
| 262 | + vector_float src_19; | ||
| 263 | + asc_loadalign(dst_19, reinterpret_cast<__ubuf__ bfloat16_t*>(output)); | ||
| 264 | + asc_loadalign(src_19, reinterpret_cast<__ubuf__ float*>(input)); | ||
| 265 | + asc_float2bfloat16_rd(dst_19, src_19, mask); | ||
| 266 | + asc_storealign(reinterpret_cast<__ubuf__ bfloat16_t*>(output), dst_19, mask); | ||
| 267 | +} | ||
| 268 | +__global__ __vector__ void asc_float2bfloat16_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input) | ||
| 269 | +{ | ||
| 270 | + asc_init(); | ||
| 271 | + __ubuf__ uint8_t output_local[BUFFER_BYTES]; | ||
| 272 | + __ubuf__ uint8_t input_local[BUFFER_BYTES]; | ||
| 273 | + asc_copy_gm2ub_align(input_local, input, BUFFER_BYTES); | ||
| 274 | + asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES); | ||
| 275 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 276 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 277 | + convert(output_local, input_local); | ||
| 278 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 279 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 280 | + asc_copy_ub2gm_align(output, output_local, BUFFER_BYTES); | ||
| 281 | + asc_sync(); | ||
| 282 | +} | ||
| 283 | +} // namespace | ||
| 284 | +int main() | ||
| 285 | +{ | ||
| 286 | + constexpr uint8_t input_pattern[] = {0x00, 0x00, 0x80, 0x3f}; | ||
| 287 | + constexpr uint8_t golden_pattern[] = {0x80, 0x3f, 0x00, 0x00}; | ||
| 288 | + std::vector<uint8_t> input(BUFFER_BYTES); | ||
| 289 | + std::vector<uint8_t> output(BUFFER_BYTES, 0xff); | ||
| 290 | + std::vector<uint8_t> golden(BUFFER_BYTES); | ||
| 291 | + for (size_t i = 0; i < BUFFER_BYTES; ++i) { | ||
| 292 | + input[i] = input_pattern[i % sizeof(input_pattern)]; | ||
| 293 | + golden[i] = golden_pattern[i % sizeof(golden_pattern)]; | ||
| 294 | + } | ||
| 295 | + aclInit(nullptr); | ||
| 296 | + aclrtSetDevice(0); | ||
| 297 | + uint8_t* input_device = nullptr; | ||
| 298 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 299 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 300 | + uint8_t* output_device = nullptr; | ||
| 301 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 302 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 303 | + aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t), | ||
| 304 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 305 | + asc_float2bfloat16_kernel<<<1, 0>>>(output_device, input_device); | ||
| 306 | + aclrtSynchronizeDevice(); | ||
| 307 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t), | ||
| 308 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 309 | + std::cout << "Conversion: vector_float -> vector_bfloat16_t" << std::endl; | ||
| 310 | + print_data("Input 1.0 (raw bytes)", input); | ||
| 311 | + print_data("Output (raw bytes)", output); | ||
| 312 | + print_data("Golden 1.0 (raw bytes)", golden); | ||
| 313 | + const bool passed = compare_data(output, golden); | ||
| 314 | + std::cout << (passed ? "[Success] asc_float2bfloat16_rd/asc_float2bfloat16_rd_sat/" | ||
| 315 | + "asc_float2bfloat16_rd_v2/asc_float2bfloat16_rd_sat_v2/" | ||
| 316 | + "asc_float2bfloat16_rn/asc_float2bfloat16_rn_sat/" | ||
| 317 | + "asc_float2bfloat16_rn_v2/asc_float2bfloat16_rn_sat_v2/" | ||
| 318 | + "asc_float2bfloat16_rna/asc_float2bfloat16_rna_sat/" | ||
| 319 | + "asc_float2bfloat16_rna_v2/asc_float2bfloat16_rna_sat_v2/" | ||
| 320 | + "asc_float2bfloat16_ru/asc_float2bfloat16_ru_sat/" | ||
| 321 | + "asc_float2bfloat16_ru_v2/asc_float2bfloat16_ru_sat_v2/" | ||
| 322 | + "asc_float2bfloat16_rz/asc_float2bfloat16_rz_sat/" | ||
| 323 | + "asc_float2bfloat16_rz_v2/asc_float2bfloat16_rz_sat_v2 passed." | ||
| 324 | + : "[Failed] asc_float2bfloat16_rd/asc_float2bfloat16_rd_sat/" | ||
| 325 | + "asc_float2bfloat16_rd_v2/asc_float2bfloat16_rd_sat_v2/" | ||
| 326 | + "asc_float2bfloat16_rn/asc_float2bfloat16_rn_sat/" | ||
| 327 | + "asc_float2bfloat16_rn_v2/asc_float2bfloat16_rn_sat_v2/" | ||
| 328 | + "asc_float2bfloat16_rna/asc_float2bfloat16_rna_sat/" | ||
| 329 | + "asc_float2bfloat16_rna_v2/asc_float2bfloat16_rna_sat_v2/" | ||
| 330 | + "asc_float2bfloat16_ru/asc_float2bfloat16_ru_sat/" | ||
| 331 | + "asc_float2bfloat16_ru_v2/asc_float2bfloat16_ru_sat_v2/" | ||
| 332 | + "asc_float2bfloat16_rz/asc_float2bfloat16_rz_sat/" | ||
| 333 | + "asc_float2bfloat16_rz_v2/asc_float2bfloat16_rz_sat_v2 failed.") | ||
| 334 | + << std::endl; | ||
| 335 | + aclrtFree(input_device); | ||
| 336 | + aclrtFree(output_device); | ||
| 337 | + aclrtResetDevice(0); | ||
| 338 | + aclFinalize(); | ||
| 339 | + return passed ? 0 : 1; | ||
| 340 | +} | ||
| 341 | +``` | ||
| @@ -26,36 +26,51 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将int32_t类型转化为float类型，并支持多种舍入模式。 | 29 | +根据`mask`将`src`中的每个`int32_t`类型数据转换成`float`类型，并支持多种舍入模式，结果写入`dst`。参考伪代码如下： |
| 30 | -$$ | 30 | + |
| 31 | -dst_i = (float)src_i | 31 | +```python |
| 32 | -$$ | 32 | +def asc_int322float_rn(dst, src, mask): |
| 33 | + for i in range(64): | ||
| 34 | + if mask[i]: | ||
| 35 | + dst[i] = float(src[i]) # int32 -> float, RINT舍入 | ||
| 36 | + else: | ||
| 37 | + dst[i] = 0.0 # mask未选中, 置零 | ||
| 38 | +``` | ||
| 33 | 39 | ||
| 34 | 关于舍入模式的详细说明，请参见[舍入模式](../data_type_convert/rounding_mode.md)。 | 40 | 关于舍入模式的详细说明，请参见[舍入模式](../data_type_convert/rounding_mode.md)。 |
| 35 | 41 | ||
| 36 | ## 函数原型 | 42 | ## 函数原型 |
| 37 | 43 | ||
| 38 | -```cpp | 44 | +### 模板原型（占位符形式） |
| 45 | + | ||
| 46 | +```c | ||
| 47 | +__simd_callee__ inline void asc_int322float_<round_mode>(vector_float& dst, | ||
| 48 | + vector_int32_t src, | ||
| 49 | + vector_bool mask) | ||
| 50 | +``` | ||
| 51 | + | ||
| 52 | +### 占位符说明 | ||
| 53 | + | ||
| 54 | +- `<round_mode>`表示支持的舍入模式，支持`rd`（`FLOOR`）、`rn`（`RINT`）、`rna`（`ROUND`）、`ru`（`CEIL`）和`rz`（`TRUNC`）。 | ||
| 55 | + | ||
| 56 | +### 典型示例 | ||
| 57 | + | ||
| 58 | +```c | ||
| 39 | // RINT舍入模式 | 59 | // RINT舍入模式 |
| 40 | -__simd_callee__ inline void asc_int322float_rn(vector_float& dst, vector_int32_t src, vector_bool mask) | 60 | +__simd_callee__ inline void asc_int322float_rn(vector_float& dst, |
| 41 | -// ROUND舍入模式 | 61 | + vector_int32_t src, |
| 42 | -__simd_callee__ inline void asc_int322float_rna(vector_float& dst, vector_int32_t src, vector_bool mask) | 62 | + vector_bool mask) |
| 43 | -// FLOOR舍入模式 | ||
| 44 | -__simd_callee__ inline void asc_int322float_rd(vector_float& dst, vector_int32_t src, vector_bool mask) | ||
| 45 | -// CEIL舍入模式 | ||
| 46 | -__simd_callee__ inline void asc_int322float_ru(vector_float& dst, vector_int32_t src, vector_bool mask) | ||
| 47 | -// TRUNC舍入模式 | ||
| 48 | -__simd_callee__ inline void asc_int322float_rz(vector_float& dst, vector_int32_t src, vector_bool mask) | ||
| 49 | ``` | 63 | ``` |
| 50 | 64 | ||
| 51 | ## 参数说明 | 65 | ## 参数说明 |
| 52 | 66 | ||
| 53 | **表1** 参数说明 | 67 | **表1** 参数说明 |
| 54 | -| 参数名 | 输入/输出 | 描述 | | 68 | + |
| 55 | -| :------ | :----- | :----------------- | | 69 | +| 参数名 | 输入/输出 | 描述 | |
| 56 | -| dst | 输出 | 目的操作数（矢量数据寄存器）。 | | 70 | +| ------ | --------- | ------------------------------------------------------------------------------------------------------------------ | |
| 57 | -| src | 输入 | 源操作数（矢量数据寄存器）。 | | 71 | +| `dst` | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 58 | -| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在输出中置零。 | | 72 | +| `src` | 输入 | 源操作数（矢量数据寄存器）。 | |
| 73 | +| `mask` | 输入 | 掩码寄存器，用于控制各元素是否参与计算。`mask`中与元素对应的比特位为1时，该元素参与计算；为0时，该元素不参与计算。 | | ||
| 59 | 74 | ||
| 60 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 75 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 61 | 76 | ||
| @@ -65,22 +80,141 @@ __simd_callee__ inline void asc_int322float_rz(vector_float& dst, vector_int32_t | |||
| 65 | 80 | ||
| 66 | ## 约束说明 | 81 | ## 约束说明 |
| 67 | 82 | ||
| 68 | -- 当int32_t的值超出float可精确表示的整数范围（绝对值大于2^24）时，按函数后缀指定的舍入模式进行舍入。 | 83 | +- `src`与`dst`的数据类型需要与函数原型匹配。 |
| 69 | -- mask未筛选的元素在输出中置零。 | 84 | +- 掩码位为0时，`dst`对应元素置0。 |
| 85 | +- 当`int32_t`的绝对值大于$2^{24}$时，按函数后缀指定的舍入模式进行舍入。 | ||
| 70 | 86 | ||
| 71 | ## 调用示例 | 87 | ## 调用示例 |
| 72 | 88 | ||
| 73 | -```cpp | 89 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 74 | -__simd_vf__ inline void int322float_vf(__ubuf__ int32_t* src_addr, __ubuf__ float* dst_addr, uint32_t count, uint16_t one_repeat_size, uint16_t repeat_time) | 90 | + |
| 91 | +<!-- npu="950" id8 --> | ||
| 92 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 93 | + | ||
| 94 | +```bash | ||
| 95 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 96 | +``` | ||
| 97 | +<!-- end id8 --> | ||
| 98 | + | ||
| 99 | +```c | ||
| 100 | +#include <cstdint> | ||
| 101 | +#include <iostream> | ||
| 102 | +#include <vector> | ||
| 103 | +#include "c_api/asc_simd.h" | ||
| 104 | +#include "acl/acl.h" | ||
| 105 | +namespace { | ||
| 106 | +template <typename T> | ||
| 107 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 75 | { | 108 | { |
| 76 | - vector_int32_t src; | 109 | + std::cout << label << ":"; |
| 77 | - vector_float dst; | 110 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 78 | - vector_bool mask; | 111 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 79 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 112 | + if (values.size() > count) std::cout << " ..."; |
| 80 | - mask = asc_update_mask_b32(count); | 113 | + std::cout << std::endl; |
| 81 | - asc_loadalign_postupdate(src, src_addr, one_repeat_size); | 114 | +} |
| 82 | - asc_int322float_rn(dst, src, mask); | 115 | + |
| 83 | - asc_storealign_postupdate(dst_addr, dst, one_repeat_size, mask); | 116 | +template <typename T> |
| 117 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 118 | +{ | ||
| 119 | + if (actual.size() != expected.size()) return false; | ||
| 120 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 121 | + if (actual[i] == expected[i]) continue; | ||
| 122 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 123 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 84 | } | 124 | } |
| 125 | + return true; | ||
| 126 | +} | ||
| 127 | + | ||
| 128 | +constexpr uint32_t BUFFER_BYTES = 256; | ||
| 129 | +__simd_vf__ inline void convert(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input) | ||
| 130 | +{ | ||
| 131 | + vector_bool mask = asc_create_mask_b8(PAT_ALL); | ||
| 132 | + vector_float dst_0; | ||
| 133 | + vector_int32_t src_0; | ||
| 134 | + asc_loadalign(dst_0, reinterpret_cast<__ubuf__ float*>(output)); | ||
| 135 | + asc_loadalign(src_0, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 136 | + asc_int322float_rna(dst_0, src_0, mask); | ||
| 137 | + asc_storealign(reinterpret_cast<__ubuf__ float*>(output), dst_0, mask); | ||
| 138 | + vector_float dst_1; | ||
| 139 | + vector_int32_t src_1; | ||
| 140 | + asc_loadalign(dst_1, reinterpret_cast<__ubuf__ float*>(output)); | ||
| 141 | + asc_loadalign(src_1, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 142 | + asc_int322float_rd(dst_1, src_1, mask); | ||
| 143 | + asc_storealign(reinterpret_cast<__ubuf__ float*>(output), dst_1, mask); | ||
| 144 | + vector_float dst_2; | ||
| 145 | + vector_int32_t src_2; | ||
| 146 | + asc_loadalign(dst_2, reinterpret_cast<__ubuf__ float*>(output)); | ||
| 147 | + asc_loadalign(src_2, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 148 | + asc_int322float_ru(dst_2, src_2, mask); | ||
| 149 | + asc_storealign(reinterpret_cast<__ubuf__ float*>(output), dst_2, mask); | ||
| 150 | + vector_float dst_3; | ||
| 151 | + vector_int32_t src_3; | ||
| 152 | + asc_loadalign(dst_3, reinterpret_cast<__ubuf__ float*>(output)); | ||
| 153 | + asc_loadalign(src_3, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 154 | + asc_int322float_rz(dst_3, src_3, mask); | ||
| 155 | + asc_storealign(reinterpret_cast<__ubuf__ float*>(output), dst_3, mask); | ||
| 156 | + vector_float dst_4; | ||
| 157 | + vector_int32_t src_4; | ||
| 158 | + asc_loadalign(dst_4, reinterpret_cast<__ubuf__ float*>(output)); | ||
| 159 | + asc_loadalign(src_4, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 160 | + asc_int322float_rn(dst_4, src_4, mask); | ||
| 161 | + asc_storealign(reinterpret_cast<__ubuf__ float*>(output), dst_4, mask); | ||
| 162 | +} | ||
| 163 | +__global__ __vector__ void asc_int322float_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input) | ||
| 164 | +{ | ||
| 165 | + asc_init(); | ||
| 166 | + __ubuf__ uint8_t output_local[BUFFER_BYTES]; | ||
| 167 | + __ubuf__ uint8_t input_local[BUFFER_BYTES]; | ||
| 168 | + asc_copy_gm2ub_align(input_local, input, BUFFER_BYTES); | ||
| 169 | + asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES); | ||
| 170 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 171 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 172 | + convert(output_local, input_local); | ||
| 173 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 174 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 175 | + asc_copy_ub2gm_align(output, output_local, BUFFER_BYTES); | ||
| 176 | + asc_sync(); | ||
| 177 | +} | ||
| 178 | +} // namespace | ||
| 179 | +int main() | ||
| 180 | +{ | ||
| 181 | + constexpr uint8_t input_pattern[] = {0x01, 0x00, 0x00, 0x00}; | ||
| 182 | + constexpr uint8_t golden_pattern[] = {0x00, 0x00, 0x80, 0x3f}; | ||
| 183 | + std::vector<uint8_t> input(BUFFER_BYTES); | ||
| 184 | + std::vector<uint8_t> output(BUFFER_BYTES, 0xff); | ||
| 185 | + std::vector<uint8_t> golden(BUFFER_BYTES); | ||
| 186 | + for (size_t i = 0; i < BUFFER_BYTES; ++i) { | ||
| 187 | + input[i] = input_pattern[i % sizeof(input_pattern)]; | ||
| 188 | + golden[i] = golden_pattern[i % sizeof(golden_pattern)]; | ||
| 189 | + } | ||
| 190 | + aclInit(nullptr); | ||
| 191 | + aclrtSetDevice(0); | ||
| 192 | + uint8_t* input_device = nullptr; | ||
| 193 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 194 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 195 | + uint8_t* output_device = nullptr; | ||
| 196 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 197 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 198 | + aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t), | ||
| 199 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 200 | + asc_int322float_kernel<<<1, 0>>>(output_device, input_device); | ||
| 201 | + aclrtSynchronizeDevice(); | ||
| 202 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t), | ||
| 203 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 204 | + std::cout << "Conversion: vector_int32_t -> vector_float" << std::endl; | ||
| 205 | + print_data("Input 1.0 (raw bytes)", input); | ||
| 206 | + print_data("Output (raw bytes)", output); | ||
| 207 | + print_data("Golden 1.0 (raw bytes)", golden); | ||
| 208 | + const bool passed = compare_data(output, golden); | ||
| 209 | + std::cout << (passed ? "[Success] asc_int322float_rn/asc_int322float_rna/" | ||
| 210 | + "asc_int322float_rd/asc_int322float_ru/asc_int322float_rz passed." | ||
| 211 | + : "[Failed] asc_int322float_rn/asc_int322float_rna/" | ||
| 212 | + "asc_int322float_rd/asc_int322float_ru/asc_int322float_rz failed.") | ||
| 213 | + << std::endl; | ||
| 214 | + aclrtFree(input_device); | ||
| 215 | + aclrtFree(output_device); | ||
| 216 | + aclrtResetDevice(0); | ||
| 217 | + aclFinalize(); | ||
| 218 | + return passed ? 0 : 1; | ||
| 85 | } | 219 | } |
| 86 | ``` | 220 | ``` |
| @@ -26,41 +26,62 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将int32_t类型转化为int16_t类型，支持饱和/非饱和模式。 | 29 | +根据`mask`将`src`中的每个`int32_t`类型元素转化为`int16_t`类型，支持饱和/非饱和模式，结果写入`dst`。由于源操作数与目的操作数类型位宽比为2:1，写入数据时需要将一个`VL`大小的数据分为两部分，根据不同接口选择数据写入索引为奇数的位置或偶数的位置。参考伪代码如下： |
| 30 | + | ||
| 31 | +```python | ||
| 32 | +def asc_int322int16(dst, src, mask): | ||
| 33 | + for i in range(64): | ||
| 34 | + if mask[i]: | ||
| 35 | + dst[2 * i] = int16(src[i]) # int32 -> int16, 写入偶数索引 | ||
| 36 | + else: | ||
| 37 | + dst[2 * i] = 0 | ||
| 38 | + for i in range(64): | ||
| 39 | + dst[2 * i + 1] = 0 # 奇数索引置零 | ||
| 40 | + | ||
| 41 | +def asc_int322int16_v2(dst, src, mask): | ||
| 42 | + for i in range(64): | ||
| 43 | + dst[2 * i] = 0 # 偶数索引置零 | ||
| 44 | + if mask[i]: | ||
| 45 | + dst[2 * i + 1] = int16(src[i]) # int32 -> int16, 写入奇数索引 | ||
| 46 | + else: | ||
| 47 | + dst[2 * i + 1] = 0 | ||
| 48 | +``` | ||
| 30 | 49 | ||
| 31 | 关于舍入模式和饱和/非饱和模式的详细说明，请参见[舍入模式](../data_type_convert/rounding_mode.md)。 | 50 | 关于舍入模式和饱和/非饱和模式的详细说明，请参见[舍入模式](../data_type_convert/rounding_mode.md)。 |
| 32 | 51 | ||
| 33 | -由于源操作数与目的操作数类型位宽比为2:1，写入数据时需要将一个VL大小的数据分为两部分，根据不同接口选择数据写入索引为奇数的位置或偶数的位置。 | ||
| 34 | - | ||
| 35 | -- asc_int322int16：非饱和模式，数据写入索引为偶数的位置。 | ||
| 36 | - | ||
| 37 | -- asc_int322int16_sat：饱和模式，数据写入索引为偶数的位置。 | ||
| 38 | - | ||
| 39 | -- asc_int322int16_v2：非饱和模式，数据写入索引为奇数的位置。 | ||
| 40 | - | ||
| 41 | -- asc_int322int16_sat_v2：饱和模式，数据写入索引为奇数的位置。 | ||
| 42 | - | ||
| 43 | ## 函数原型 | 52 | ## 函数原型 |
| 44 | 53 | ||
| 45 | -```cpp | 54 | +### 模板原型（占位符形式） |
| 55 | + | ||
| 56 | +```c | ||
| 57 | +__simd_callee__ inline void asc_int322int16<sat_mode><position_mode>(vector_int16_t& dst, | ||
| 58 | + vector_int32_t src, | ||
| 59 | + vector_bool mask) | ||
| 60 | +``` | ||
| 61 | + | ||
| 62 | +### 占位符说明 | ||
| 63 | + | ||
| 64 | +- `<sat_mode>`表示饱和非饱和模式，为空时表示非饱和模式，取值为`_sat`时表示饱和模式。 | ||
| 65 | +- `<position_mode>`表示`dst`写入的位置，为空时，结果写入`dst`的偶数索引位置；取值为`_v2`时，结果写入`dst`的奇数索引位置。 | ||
| 66 | + | ||
| 67 | +### 典型示例 | ||
| 68 | + | ||
| 69 | +```c | ||
| 46 | // 非饱和模式，数据写入索引为偶数的位置 | 70 | // 非饱和模式，数据写入索引为偶数的位置 |
| 47 | -__simd_callee__ inline void asc_int322int16(vector_int16_t& dst, vector_int32_t src, vector_bool mask) | 71 | +__simd_callee__ inline void asc_int322int16(vector_int16_t& dst, |
| 48 | -// 饱和模式，数据写入索引为偶数的位置 | 72 | + vector_int32_t src, |
| 49 | -__simd_callee__ inline void asc_int322int16_sat(vector_int16_t& dst, vector_int32_t src, vector_bool mask) | 73 | + vector_bool mask) |
| 50 | -// 非饱和模式，数据写入索引为奇数的位置 | ||
| 51 | -__simd_callee__ inline void asc_int322int16_v2(vector_int16_t& dst, vector_int32_t src, vector_bool mask) | ||
| 52 | -// 饱和模式，数据写入索引为奇数的位置 | ||
| 53 | -__simd_callee__ inline void asc_int322int16_sat_v2(vector_int16_t& dst, vector_int32_t src, vector_bool mask) | ||
| 54 | ``` | 74 | ``` |
| 55 | 75 | ||
| 56 | ## 参数说明 | 76 | ## 参数说明 |
| 57 | 77 | ||
| 58 | **表1** 参数说明 | 78 | **表1** 参数说明 |
| 59 | -| 参数名 | 输入/输出 | 描述 | | 79 | + |
| 60 | -| :------ | :----- | :----------------- | | 80 | +| 参数名 | 输入/输出 | 描述 | |
| 61 | -| dst | 输出 | 目的操作数（矢量数据寄存器）。 | | 81 | +| ------ | --------- | ------------------------------------------------------------------------------------------------------------------ | |
| 62 | -| src | 输入 | 源操作数（矢量数据寄存器）。 | | 82 | +| `dst` | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 63 | -| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在输出中置零。 | | 83 | +| `src` | 输入 | 源操作数（矢量数据寄存器）。 | |
| 84 | +| `mask` | 输入 | 掩码寄存器，用于控制各元素是否参与计算。`mask`中与元素对应的比特位为1时，该元素参与计算；为0时，该元素不参与计算。 | | ||
| 64 | 85 | ||
| 65 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 86 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 66 | 87 | ||
| @@ -70,21 +91,134 @@ __simd_callee__ inline void asc_int322int16_sat_v2(vector_int16_t& dst, vector_i | |||
| 70 | 91 | ||
| 71 | ## 约束说明 | 92 | ## 约束说明 |
| 72 | 93 | ||
| 73 | -mask未筛选的元素在输出中置零。 | 94 | +- `src`与`dst`的数据类型需要与函数原型匹配。 |
| 95 | +- `mask`掩码位为0时，`dst`对应元素置0。 | ||
| 74 | 96 | ||
| 75 | ## 调用示例 | 97 | ## 调用示例 |
| 76 | 98 | ||
| 77 | -```cpp | 99 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 78 | -__simd_vf__ inline void int322int16_vf(__ubuf__ int32_t* src_addr, __ubuf__ int16_t* dst_addr, uint32_t count, uint16_t one_repeat_size, uint16_t one_block_size, uint16_t repeat_time) | 100 | + |
| 101 | +<!-- npu="950" id8 --> | ||
| 102 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 103 | + | ||
| 104 | +```bash | ||
| 105 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 106 | +``` | ||
| 107 | +<!-- end id8 --> | ||
| 108 | + | ||
| 109 | +```c | ||
| 110 | +#include <cstdint> | ||
| 111 | +#include <iostream> | ||
| 112 | +#include <vector> | ||
| 113 | +#include "c_api/asc_simd.h" | ||
| 114 | +#include "acl/acl.h" | ||
| 115 | +namespace { | ||
| 116 | +template <typename T> | ||
| 117 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 79 | { | 118 | { |
| 80 | - vector_int32_t src; | 119 | + std::cout << label << ":"; |
| 81 | - vector_int16_t dst; | 120 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 82 | - vector_bool mask; | 121 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 83 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 122 | + if (values.size() > count) std::cout << " ..."; |
| 84 | - mask = asc_update_mask_b32(count); | 123 | + std::cout << std::endl; |
| 85 | - asc_loadalign_postupdate(src, src_addr, one_repeat_size); | 124 | +} |
| 86 | - asc_int322int16(dst, src, mask); | 125 | + |
| 87 | - asc_storealign_pack_postupdate(dst_addr, dst, one_block_size, mask); | 126 | +template <typename T> |
| 127 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 128 | +{ | ||
| 129 | + if (actual.size() != expected.size()) return false; | ||
| 130 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 131 | + if (actual[i] == expected[i]) continue; | ||
| 132 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 133 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 88 | } | 134 | } |
| 135 | + return true; | ||
| 136 | +} | ||
| 137 | + | ||
| 138 | +constexpr uint32_t BUFFER_BYTES = 256; | ||
| 139 | +__simd_vf__ inline void convert(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input) | ||
| 140 | +{ | ||
| 141 | + vector_bool mask = asc_create_mask_b8(PAT_ALL); | ||
| 142 | + vector_int16_t dst_0; | ||
| 143 | + vector_int32_t src_0; | ||
| 144 | + asc_loadalign(dst_0, reinterpret_cast<__ubuf__ int16_t*>(output)); | ||
| 145 | + asc_loadalign(src_0, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 146 | + asc_int322int16_sat(dst_0, src_0, mask); | ||
| 147 | + asc_storealign(reinterpret_cast<__ubuf__ int16_t*>(output), dst_0, mask); | ||
| 148 | + vector_int16_t dst_1; | ||
| 149 | + vector_int32_t src_1; | ||
| 150 | + asc_loadalign(dst_1, reinterpret_cast<__ubuf__ int16_t*>(output)); | ||
| 151 | + asc_loadalign(src_1, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 152 | + asc_int322int16_v2(dst_1, src_1, mask); | ||
| 153 | + asc_storealign(reinterpret_cast<__ubuf__ int16_t*>(output), dst_1, mask); | ||
| 154 | + vector_int16_t dst_2; | ||
| 155 | + vector_int32_t src_2; | ||
| 156 | + asc_loadalign(dst_2, reinterpret_cast<__ubuf__ int16_t*>(output)); | ||
| 157 | + asc_loadalign(src_2, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 158 | + asc_int322int16_sat_v2(dst_2, src_2, mask); | ||
| 159 | + asc_storealign(reinterpret_cast<__ubuf__ int16_t*>(output), dst_2, mask); | ||
| 160 | + vector_int16_t dst_3; | ||
| 161 | + vector_int32_t src_3; | ||
| 162 | + asc_loadalign(dst_3, reinterpret_cast<__ubuf__ int16_t*>(output)); | ||
| 163 | + asc_loadalign(src_3, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 164 | + asc_int322int16(dst_3, src_3, mask); | ||
| 165 | + asc_storealign(reinterpret_cast<__ubuf__ int16_t*>(output), dst_3, mask); | ||
| 166 | +} | ||
| 167 | +__global__ __vector__ void asc_int322int16_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input) | ||
| 168 | +{ | ||
| 169 | + asc_init(); | ||
| 170 | + __ubuf__ uint8_t output_local[BUFFER_BYTES]; | ||
| 171 | + __ubuf__ uint8_t input_local[BUFFER_BYTES]; | ||
| 172 | + asc_copy_gm2ub_align(input_local, input, BUFFER_BYTES); | ||
| 173 | + asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES); | ||
| 174 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 175 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 176 | + convert(output_local, input_local); | ||
| 177 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 178 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 179 | + asc_copy_ub2gm_align(output, output_local, BUFFER_BYTES); | ||
| 180 | + asc_sync(); | ||
| 181 | +} | ||
| 182 | +} // namespace | ||
| 183 | +int main() | ||
| 184 | +{ | ||
| 185 | + constexpr uint8_t input_pattern[] = {0x01, 0x00, 0x00, 0x00}; | ||
| 186 | + constexpr uint8_t golden_pattern[] = {0x01, 0x00, 0x00, 0x00}; | ||
| 187 | + std::vector<uint8_t> input(BUFFER_BYTES); | ||
| 188 | + std::vector<uint8_t> output(BUFFER_BYTES, 0xff); | ||
| 189 | + std::vector<uint8_t> golden(BUFFER_BYTES); | ||
| 190 | + for (size_t i = 0; i < BUFFER_BYTES; ++i) { | ||
| 191 | + input[i] = input_pattern[i % sizeof(input_pattern)]; | ||
| 192 | + golden[i] = golden_pattern[i % sizeof(golden_pattern)]; | ||
| 193 | + } | ||
| 194 | + aclInit(nullptr); | ||
| 195 | + aclrtSetDevice(0); | ||
| 196 | + uint8_t* input_device = nullptr; | ||
| 197 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 198 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 199 | + uint8_t* output_device = nullptr; | ||
| 200 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 201 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 202 | + aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t), | ||
| 203 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 204 | + asc_int322int16_kernel<<<1, 0>>>(output_device, input_device); | ||
| 205 | + aclrtSynchronizeDevice(); | ||
| 206 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t), | ||
| 207 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 208 | + std::cout << "Conversion: vector_int32_t -> vector_int16_t" << std::endl; | ||
| 209 | + print_data("Input 1.0 (raw bytes)", input); | ||
| 210 | + print_data("Output (raw bytes)", output); | ||
| 211 | + print_data("Golden 1.0 (raw bytes)", golden); | ||
| 212 | + const bool passed = compare_data(output, golden); | ||
| 213 | + std::cout << (passed ? "[Success] asc_int322int16/asc_int322int16_sat/asc_int322int16_v2/" | ||
| 214 | + "asc_int322int16_sat_v2 passed." | ||
| 215 | + : "[Failed] asc_int322int16/asc_int322int16_sat/asc_int322int16_v2/" | ||
| 216 | + "asc_int322int16_sat_v2 failed.") | ||
| 217 | + << std::endl; | ||
| 218 | + aclrtFree(input_device); | ||
| 219 | + aclrtFree(output_device); | ||
| 220 | + aclrtResetDevice(0); | ||
| 221 | + aclFinalize(); | ||
| 222 | + return passed ? 0 : 1; | ||
| 89 | } | 223 | } |
| 90 | ``` | 224 | ``` |
| @@ -26,36 +26,57 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将vector_int32_t类型的源操作数转换为vector_int64_t类型，并写入目的操作数。 | 29 | +根据`mask`将`src`中的每个`int32_t`类型元素转换为`int64_t`类型，结果写入`dst`。由于源操作数与目的操作数类型位宽比为1:2，读取数据时需要将一个`VL`大小的数据分为两部分，根据不同接口选择输入数据索引为奇数的位置或偶数的位置。伪代码如下： |
| 30 | 30 | ||
| 31 | -由于源操作数与目的操作数类型位宽比为1:2，读取数据时需要将一个VL大小的数据分为两部分，根据不同接口选择输入数据索引为奇数的位置或偶数的位置。 | 31 | +```python |
| 32 | +def asc_int322int64(dst, src, mask): | ||
| 33 | + for i in range(32): | ||
| 34 | + if mask[2 * i]: # 偶数索引位置参与计算 | ||
| 35 | + dst[i] = int64(src[2 * i]) # int32 -> int64 扩展转换 | ||
| 36 | + else: | ||
| 37 | + dst[i] = 0 # mask未选中, 置零 | ||
| 32 | 38 | ||
| 33 | -- asc_int322int64：取src的偶数索引元素（索引0, 2, 4, ...）进行转换。 | 39 | +def asc_int322int64_v2(dst, src, mask): |
| 34 | - | 40 | + for i in range(32): |
| 35 | -- asc_int322int64_v2：取src的奇数索引元素（索引1, 3, 5, ...）进行转换。 | 41 | + if mask[2 * i + 1]: # 奇数索引位置参与计算 |
| 36 | - | 42 | + dst[i] = int64(src[2 * i + 1]) |
| 37 | - | 43 | + else: |
| 38 | -$$ | 44 | + dst[i] = 0 |
| 39 | -dst_i = (int64\_t)src_i | 45 | +``` |
| 40 | -$$ | ||
| 41 | 46 | ||
| 42 | ## 函数原型 | 47 | ## 函数原型 |
| 43 | 48 | ||
| 44 | -```cpp | 49 | +### 模板原型（占位符形式） |
| 45 | -// 取src的偶数索引元素（索引0, 2, 4, ...）进行转换 | 50 | + |
| 46 | -__simd_callee__ inline void asc_int322int64(vector_int64_t& dst, vector_int32_t src, vector_bool mask) | 51 | +```c |
| 47 | -// 取src的奇数索引元素（索引1, 3, 5, ...）进行转换 | 52 | +__simd_callee__ inline void asc_int322int64<position_mode>(vector_int64_t& dst, |
| 48 | -__simd_callee__ inline void asc_int322int64_v2(vector_int64_t& dst, vector_int32_t src, vector_bool mask) | 53 | + vector_int32_t src, |
| 54 | + vector_bool mask) | ||
| 55 | +``` | ||
| 56 | + | ||
| 57 | +### 占位符说明 | ||
| 58 | + | ||
| 59 | +- `<position_mode>`表示`src`读取的位置，为空时，读取`src`的偶数索引位置元素；取值为`_v2`时，读取`src`的奇数索引位置元素。 | ||
| 60 | + | ||
| 61 | + | ||
| 62 | +### 典型示例 | ||
| 63 | + | ||
| 64 | +```c | ||
| 65 | +// 数据写入索引为偶数的位置 | ||
| 66 | +__simd_callee__ inline void asc_int322int64(vector_int64_t& dst, | ||
| 67 | + vector_int32_t src, | ||
| 68 | + vector_bool mask) | ||
| 49 | ``` | 69 | ``` |
| 50 | 70 | ||
| 51 | ## 参数说明 | 71 | ## 参数说明 |
| 52 | 72 | ||
| 53 | **表1** 参数说明 | 73 | **表1** 参数说明 |
| 54 | -| 参数名 | 输入/输出 | 描述 | | 74 | + |
| 55 | -| --------- | ----- | ---------------- | | 75 | +| 参数名 | 输入/输出 | 描述 | |
| 56 | -| dst | 输出 | 目的操作数（矢量数据寄存器）。 | | 76 | +| ------ | --------- | ------------------------------------------------------------------------------------------------------------------ | |
| 57 | -| src | 输入 | 源操作数（矢量数据寄存器）。 | | 77 | +| `dst` | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 58 | -| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在输出中置零。 | | 78 | +| `src` | 输入 | 源操作数（矢量数据寄存器）。 | |
| 79 | +| `mask` | 输入 | 掩码寄存器，用于控制各元素是否参与计算。`mask`中与元素对应的比特位为1时，该元素参与计算；为0时，该元素不参与计算。 | | ||
| 59 | 80 | ||
| 60 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 81 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 61 | 82 | ||
| @@ -65,21 +86,121 @@ __simd_callee__ inline void asc_int322int64_v2(vector_int64_t& dst, vector_int32 | |||
| 65 | 86 | ||
| 66 | ## 约束说明 | 87 | ## 约束说明 |
| 67 | 88 | ||
| 68 | -mask未筛选的元素在输出中置零。 | 89 | +- `src`与`dst`的数据类型需要与函数原型匹配。 |
| 90 | +- `mask`掩码位为0时，`dst`对应元素置0。 | ||
| 69 | 91 | ||
| 70 | ## 调用示例 | 92 | ## 调用示例 |
| 71 | 93 | ||
| 72 | -```cpp | 94 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 73 | -__simd_vf__ inline void int322int64_vf(__ubuf__ int32_t* src_addr, __ubuf__ int64_t* dst_addr, uint32_t count, uint16_t one_repeat_size, uint16_t one_block_size, uint16_t repeat_time) | 95 | + |
| 96 | +<!-- npu="950" id8 --> | ||
| 97 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 98 | + | ||
| 99 | +```bash | ||
| 100 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 101 | +``` | ||
| 102 | + | ||
| 103 | +<!-- end id8 --> | ||
| 104 | + | ||
| 105 | +```c | ||
| 106 | +#include <cstdint> | ||
| 107 | +#include <iostream> | ||
| 108 | +#include <vector> | ||
| 109 | +#include "c_api/asc_simd.h" | ||
| 110 | +#include "acl/acl.h" | ||
| 111 | +namespace { | ||
| 112 | +template <typename T> | ||
| 113 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 74 | { | 114 | { |
| 75 | - vector_int32_t src; | 115 | + std::cout << label << ":"; |
| 76 | - vector_int64_t dst; | 116 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 77 | - vector_bool mask; | 117 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 78 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 118 | + if (values.size() > count) std::cout << " ..."; |
| 79 | - mask = asc_update_mask_b32(count); | 119 | + std::cout << std::endl; |
| 80 | - asc_loadalign_postupdate(src, src_addr, one_repeat_size); | ||
| 81 | - asc_int322int64(dst, src, mask); | ||
| 82 | - asc_storealign_pack_postupdate(dst_addr, dst, one_block_size, mask); | ||
| 83 | - } | ||
| 84 | } | 120 | } |
| 85 | -``` | 121 | + |
| 122 | +template <typename T> | ||
| 123 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 124 | +{ | ||
| 125 | + if (actual.size() != expected.size()) return false; | ||
| 126 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 127 | + if (actual[i] == expected[i]) continue; | ||
| 128 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 129 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 130 | + } | ||
| 131 | + return true; | ||
| 132 | +} | ||
| 133 | + | ||
| 134 | +constexpr uint32_t BUFFER_BYTES = 256; | ||
| 135 | +__simd_vf__ inline void convert(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input) | ||
| 136 | +{ | ||
| 137 | + vector_bool mask = asc_create_mask_b8(PAT_ALL); | ||
| 138 | + vector_int64_t dst_0; | ||
| 139 | + vector_int32_t src_0; | ||
| 140 | + asc_loadalign(dst_0, reinterpret_cast<__ubuf__ int64_t*>(output)); | ||
| 141 | + asc_loadalign(src_0, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 142 | + asc_int322int64_v2(dst_0, src_0, mask); | ||
| 143 | + asc_store(reinterpret_cast<__ubuf__ int64_t*>(output), dst_0); | ||
| 144 | + vector_int64_t dst_1; | ||
| 145 | + vector_int32_t src_1; | ||
| 146 | + asc_loadalign(dst_1, reinterpret_cast<__ubuf__ int64_t*>(output)); | ||
| 147 | + asc_loadalign(src_1, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 148 | + asc_int322int64(dst_1, src_1, mask); | ||
| 149 | + asc_store(reinterpret_cast<__ubuf__ int64_t*>(output), dst_1); | ||
| 150 | +} | ||
| 151 | +__global__ __vector__ void asc_int322int64_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input) | ||
| 152 | +{ | ||
| 153 | + asc_init(); | ||
| 154 | + __ubuf__ uint8_t output_local[BUFFER_BYTES]; | ||
| 155 | + __ubuf__ uint8_t input_local[BUFFER_BYTES]; | ||
| 156 | + asc_copy_gm2ub_align(input_local, input, BUFFER_BYTES); | ||
| 157 | + asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES); | ||
| 158 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 159 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 160 | + convert(output_local, input_local); | ||
| 161 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 162 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 163 | + asc_copy_ub2gm_align(output, output_local, BUFFER_BYTES); | ||
| 164 | + asc_sync(); | ||
| 165 | +} | ||
| 166 | +} // namespace | ||
| 167 | +int main() | ||
| 168 | +{ | ||
| 169 | + constexpr uint8_t input_pattern[] = {0x01, 0x00, 0x00, 0x00}; | ||
| 170 | + constexpr uint8_t golden_pattern[] = {0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00}; | ||
| 171 | + std::vector<uint8_t> input(BUFFER_BYTES); | ||
| 172 | + std::vector<uint8_t> output(BUFFER_BYTES, 0xff); | ||
| 173 | + std::vector<uint8_t> golden(BUFFER_BYTES); | ||
| 174 | + for (size_t i = 0; i < BUFFER_BYTES; ++i) { | ||
| 175 | + input[i] = input_pattern[i % sizeof(input_pattern)]; | ||
| 176 | + golden[i] = golden_pattern[i % sizeof(golden_pattern)]; | ||
| 177 | + } | ||
| 178 | + aclInit(nullptr); | ||
| 179 | + aclrtSetDevice(0); | ||
| 180 | + uint8_t* input_device = nullptr; | ||
| 181 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 182 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 183 | + uint8_t* output_device = nullptr; | ||
| 184 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 185 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 186 | + aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t), | ||
| 187 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 188 | + asc_int322int64_kernel<<<1, 0>>>(output_device, input_device); | ||
| 189 | + aclrtSynchronizeDevice(); | ||
| 190 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t), | ||
| 191 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 192 | + std::cout << "Conversion: vector_int32_t -> vector_int64_t" << std::endl; | ||
| 193 | + print_data("Input 1.0 (raw bytes)", input); | ||
| 194 | + print_data("Output (raw bytes)", output); | ||
| 195 | + print_data("Golden 1.0 (raw bytes)", golden); | ||
| 196 | + const bool passed = compare_data(output, golden); | ||
| 197 | + std::cout << (passed ? "[Success] asc_int322int64/asc_int322int64_v2 passed." | ||
| 198 | + : "[Failed] asc_int322int64/asc_int322int64_v2 failed.") | ||
| 199 | + << std::endl; | ||
| 200 | + aclrtFree(input_device); | ||
| 201 | + aclrtFree(output_device); | ||
| 202 | + aclrtResetDevice(0); | ||
| 203 | + aclFinalize(); | ||
| 204 | + return passed ? 0 : 1; | ||
| 205 | +} | ||
| 206 | +``` | ||
| @@ -26,41 +26,62 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将vector_int32_t类型转换成vector_uint16_t类型，并支持不同的饱和模式。 | 29 | +根据`mask`将`src`中的每个`int32_t`类型元素转换为`uint16_t`类型，并支持不同的饱和模式，结果写入`dst`。由于源操作数与目的操作数类型位宽比为2:1，写入数据时需要将一个`VL`大小的数据分为两部分，根据不同接口选择数据写入索引为奇数的位置或偶数的位置。参考伪代码： |
| 30 | + | ||
| 31 | +```python | ||
| 32 | +def asc_int322uint16(dst, src, mask): | ||
| 33 | + for i in range(64): | ||
| 34 | + if mask[i]: | ||
| 35 | + dst[2 * i] = uint16(src[i]) # int32 -> uint16, 写入偶数索引 | ||
| 36 | + else: | ||
| 37 | + dst[2 * i] = 0 | ||
| 38 | + for i in range(64): | ||
| 39 | + dst[2 * i + 1] = 0 # 奇数索引置零 | ||
| 40 | + | ||
| 41 | +def asc_int322uint16_v2(dst, src, mask): | ||
| 42 | + for i in range(64): | ||
| 43 | + dst[2 * i] = 0 # 偶数索引置零 | ||
| 44 | + if mask[i]: | ||
| 45 | + dst[2 * i + 1] = uint16(src[i]) # int32 -> uint16, 写入奇数索引 | ||
| 46 | + else: | ||
| 47 | + dst[2 * i + 1] = 0 | ||
| 48 | +``` | ||
| 30 | 49 | ||
| 31 | 关于舍入模式和饱和/非饱和模式的详细说明，请参见[舍入模式](../data_type_convert/rounding_mode.md)。 | 50 | 关于舍入模式和饱和/非饱和模式的详细说明，请参见[舍入模式](../data_type_convert/rounding_mode.md)。 |
| 32 | 51 | ||
| 33 | -由于源操作数与目的操作数类型位宽比为2:1，写入数据时需要将一个VL大小的数据分为两部分，根据不同接口选择数据写入索引为奇数的位置或偶数的位置。 | ||
| 34 | - | ||
| 35 | -- asc_int322uint16：不饱和模式，数据写入索引为偶数的位置。 | ||
| 36 | - | ||
| 37 | -- asc_int322uint16_sat：饱和模式，数据写入索引为偶数的位置。 | ||
| 38 | - | ||
| 39 | -- asc_int322uint16_v2：不饱和模式，数据写入索引为奇数的位置。 | ||
| 40 | - | ||
| 41 | -- asc_int322uint16_sat_v2：饱和模式，数据写入索引为奇数的位置。 | ||
| 42 | - | ||
| 43 | ## 函数原型 | 52 | ## 函数原型 |
| 44 | 53 | ||
| 45 | -```cpp | 54 | +### 模板原型（占位符形式） |
| 55 | + | ||
| 56 | +```c | ||
| 57 | +__simd_callee__ inline void asc_int322uint16<sat_mode><position_mode>(vector_uint16_t& dst, | ||
| 58 | + vector_int32_t src, | ||
| 59 | + vector_bool mask) | ||
| 60 | +``` | ||
| 61 | + | ||
| 62 | +### 占位符说明 | ||
| 63 | + | ||
| 64 | +- `<sat_mode>`表示饱和非饱和模式，为空时表示非饱和模式，取值为`_sat`时表示饱和模式。 | ||
| 65 | +- `<position_mode>`表示`dst`写入的位置，为空时，结果写入`dst`的偶数索引位置；取值为`_v2`时，结果写入`dst`的奇数索引位置。 | ||
| 66 | + | ||
| 67 | +### 典型示例 | ||
| 68 | + | ||
| 69 | +```c | ||
| 46 | // 非饱和模式，数据写入索引为偶数的位置 | 70 | // 非饱和模式，数据写入索引为偶数的位置 |
| 47 | -__simd_callee__ inline void asc_int322uint16(vector_uint16_t& dst, vector_int32_t src, vector_bool mask) | 71 | +__simd_callee__ inline void asc_int322uint16(vector_uint16_t& dst, |
| 48 | -// 饱和模式，数据写入索引为偶数的位置 | 72 | + vector_int32_t src, |
| 49 | -__simd_callee__ inline void asc_int322uint16_sat(vector_uint16_t& dst, vector_int32_t src, vector_bool mask) | 73 | + vector_bool mask) |
| 50 | -// 非饱和模式，数据写入索引为奇数的位置 | ||
| 51 | -__simd_callee__ inline void asc_int322uint16_v2(vector_uint16_t& dst, vector_int32_t src, vector_bool mask) | ||
| 52 | -// 饱和模式，数据写入索引为奇数的位置 | ||
| 53 | -__simd_callee__ inline void asc_int322uint16_sat_v2(vector_uint16_t& dst, vector_int32_t src, vector_bool mask) | ||
| 54 | ``` | 74 | ``` |
| 55 | 75 | ||
| 56 | ## 参数说明 | 76 | ## 参数说明 |
| 57 | 77 | ||
| 58 | **表1** 参数说明 | 78 | **表1** 参数说明 |
| 59 | -| 参数名 | 输入/输出 | 描述 | | 79 | + |
| 60 | -| --------- | ----- | ---------------- | | 80 | +| 参数名 | 输入/输出 | 描述 | |
| 61 | -| dst | 输出 | 目的操作数（矢量数据寄存器）。 | | 81 | +| ------ | --------- | ------------------------------------------------------------------------------------------------------------------ | |
| 62 | -| src | 输入 | 源操作数（矢量数据寄存器）。 | | 82 | +| `dst` | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 63 | -| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在输出中置零。 | | 83 | +| `src` | 输入 | 源操作数（矢量数据寄存器）。 | |
| 84 | +| `mask` | 输入 | 掩码寄存器，用于控制各元素是否参与计算。`mask`中与元素对应的比特位为1时，该元素参与计算；为0时，该元素不参与计算。 | | ||
| 64 | 85 | ||
| 65 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 86 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 66 | 87 | ||
| @@ -70,21 +91,134 @@ __simd_callee__ inline void asc_int322uint16_sat_v2(vector_uint16_t& dst, vector | |||
| 70 | 91 | ||
| 71 | ## 约束说明 | 92 | ## 约束说明 |
| 72 | 93 | ||
| 73 | -mask未筛选的元素在输出中置零。 | 94 | +- `src`与`dst`的数据类型需要与函数原型匹配。 |
| 95 | +- `mask`掩码位为0时，`dst`对应元素置0。 | ||
| 74 | 96 | ||
| 75 | ## 调用示例 | 97 | ## 调用示例 |
| 76 | 98 | ||
| 77 | -```cpp | 99 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 78 | -__simd_vf__ inline void int322uint16_vf(__ubuf__ int32_t* src_addr, __ubuf__ uint16_t* dst_addr, uint32_t count, uint16_t one_repeat_size, uint16_t one_block_size, uint16_t repeat_time) | 100 | + |
| 101 | +<!-- npu="950" id8 --> | ||
| 102 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 103 | + | ||
| 104 | +```bash | ||
| 105 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 106 | +``` | ||
| 107 | +<!-- end id8 --> | ||
| 108 | + | ||
| 109 | +```c | ||
| 110 | +#include <cstdint> | ||
| 111 | +#include <iostream> | ||
| 112 | +#include <vector> | ||
| 113 | +#include "c_api/asc_simd.h" | ||
| 114 | +#include "acl/acl.h" | ||
| 115 | +namespace { | ||
| 116 | +template <typename T> | ||
| 117 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 79 | { | 118 | { |
| 80 | - vector_int32_t src; | 119 | + std::cout << label << ":"; |
| 81 | - vector_uint16_t dst; | 120 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 82 | - vector_bool mask; | 121 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 83 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 122 | + if (values.size() > count) std::cout << " ..."; |
| 84 | - mask = asc_update_mask_b32(count); | 123 | + std::cout << std::endl; |
| 85 | - asc_loadalign_postupdate(src, src_addr, one_repeat_size); | ||
| 86 | - asc_int322uint16(dst, src, mask); | ||
| 87 | - asc_storealign_pack_postupdate(dst_addr, dst, one_block_size, mask); | ||
| 88 | - } | ||
| 89 | } | 124 | } |
| 90 | -``` | 125 | + |
| 126 | +template <typename T> | ||
| 127 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 128 | +{ | ||
| 129 | + if (actual.size() != expected.size()) return false; | ||
| 130 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 131 | + if (actual[i] == expected[i]) continue; | ||
| 132 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 133 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 134 | + } | ||
| 135 | + return true; | ||
| 136 | +} | ||
| 137 | + | ||
| 138 | +constexpr uint32_t BUFFER_BYTES = 256; | ||
| 139 | +__simd_vf__ inline void convert(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input) | ||
| 140 | +{ | ||
| 141 | + vector_bool mask = asc_create_mask_b8(PAT_ALL); | ||
| 142 | + vector_uint16_t dst_0; | ||
| 143 | + vector_int32_t src_0; | ||
| 144 | + asc_loadalign(dst_0, reinterpret_cast<__ubuf__ uint16_t*>(output)); | ||
| 145 | + asc_loadalign(src_0, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 146 | + asc_int322uint16_sat(dst_0, src_0, mask); | ||
| 147 | + asc_storealign(reinterpret_cast<__ubuf__ uint16_t*>(output), dst_0, mask); | ||
| 148 | + vector_uint16_t dst_1; | ||
| 149 | + vector_int32_t src_1; | ||
| 150 | + asc_loadalign(dst_1, reinterpret_cast<__ubuf__ uint16_t*>(output)); | ||
| 151 | + asc_loadalign(src_1, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 152 | + asc_int322uint16_v2(dst_1, src_1, mask); | ||
| 153 | + asc_storealign(reinterpret_cast<__ubuf__ uint16_t*>(output), dst_1, mask); | ||
| 154 | + vector_uint16_t dst_2; | ||
| 155 | + vector_int32_t src_2; | ||
| 156 | + asc_loadalign(dst_2, reinterpret_cast<__ubuf__ uint16_t*>(output)); | ||
| 157 | + asc_loadalign(src_2, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 158 | + asc_int322uint16_sat_v2(dst_2, src_2, mask); | ||
| 159 | + asc_storealign(reinterpret_cast<__ubuf__ uint16_t*>(output), dst_2, mask); | ||
| 160 | + vector_uint16_t dst_3; | ||
| 161 | + vector_int32_t src_3; | ||
| 162 | + asc_loadalign(dst_3, reinterpret_cast<__ubuf__ uint16_t*>(output)); | ||
| 163 | + asc_loadalign(src_3, reinterpret_cast<__ubuf__ int32_t*>(input)); | ||
| 164 | + asc_int322uint16(dst_3, src_3, mask); | ||
| 165 | + asc_storealign(reinterpret_cast<__ubuf__ uint16_t*>(output), dst_3, mask); | ||
| 166 | +} | ||
| 167 | +__global__ __vector__ void asc_int322uint16_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input) | ||
| 168 | +{ | ||
| 169 | + asc_init(); | ||
| 170 | + __ubuf__ uint8_t output_local[BUFFER_BYTES]; | ||
| 171 | + __ubuf__ uint8_t input_local[BUFFER_BYTES]; | ||
| 172 | + asc_copy_gm2ub_align(input_local, input, BUFFER_BYTES); | ||
| 173 | + asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES); | ||
| 174 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 175 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 176 | + convert(output_local, input_local); | ||
| 177 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 178 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 179 | + asc_copy_ub2gm_align(output, output_local, BUFFER_BYTES); | ||
| 180 | + asc_sync(); | ||
| 181 | +} | ||
| 182 | +} // namespace | ||
| 183 | +int main() | ||
| 184 | +{ | ||
| 185 | + constexpr uint8_t input_pattern[] = {0x01, 0x00, 0x00, 0x00}; | ||
| 186 | + constexpr uint8_t golden_pattern[] = {0x01, 0x00, 0x00, 0x00}; | ||
| 187 | + std::vector<uint8_t> input(BUFFER_BYTES); | ||
| 188 | + std::vector<uint8_t> output(BUFFER_BYTES, 0xff); | ||
| 189 | + std::vector<uint8_t> golden(BUFFER_BYTES); | ||
| 190 | + for (size_t i = 0; i < BUFFER_BYTES; ++i) { | ||
| 191 | + input[i] = input_pattern[i % sizeof(input_pattern)]; | ||
| 192 | + golden[i] = golden_pattern[i % sizeof(golden_pattern)]; | ||
| 193 | + } | ||
| 194 | + aclInit(nullptr); | ||
| 195 | + aclrtSetDevice(0); | ||
| 196 | + uint8_t* input_device = nullptr; | ||
| 197 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 198 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 199 | + uint8_t* output_device = nullptr; | ||
| 200 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 201 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 202 | + aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t), | ||
| 203 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 204 | + asc_int322uint16_kernel<<<1, 0>>>(output_device, input_device); | ||
| 205 | + aclrtSynchronizeDevice(); | ||
| 206 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t), | ||
| 207 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 208 | + std::cout << "Conversion: vector_int32_t -> vector_uint16_t" << std::endl; | ||
| 209 | + print_data("Input 1.0 (raw bytes)", input); | ||
| 210 | + print_data("Output (raw bytes)", output); | ||
| 211 | + print_data("Golden 1.0 (raw bytes)", golden); | ||
| 212 | + const bool passed = compare_data(output, golden); | ||
| 213 | + std::cout << (passed ? "[Success] asc_int322uint16/asc_int322uint16_sat/" | ||
| 214 | + "asc_int322uint16_v2/asc_int322uint16_sat_v2 passed." | ||
| 215 | + : "[Failed] asc_int322uint16/asc_int322uint16_sat/" | ||
| 216 | + "asc_int322uint16_v2/asc_int322uint16_sat_v2 failed.") | ||
| 217 | + << std::endl; | ||
| 218 | + aclrtFree(input_device); | ||
| 219 | + aclrtFree(output_device); | ||
| 220 | + aclrtResetDevice(0); | ||
| 221 | + aclFinalize(); | ||
| 222 | + return passed ? 0 : 1; | ||
| 223 | +} | ||
| 224 | +``` | ||
| @@ -26,28 +26,58 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将uint8_t类型数据转为uint16_t类型。 | 29 | +根据`mask`将源操作数`src`中的每个`uint8_t`类型元素转为`uint16_t`类型，结果写入`dst`。由于源操作数与目的操作数类型位宽比为1:2，读取数据时需要将一个`VL`大小的数据分为两部分，根据不同接口选择输入数据索引为奇数的位置或偶数的位置。参考伪代码如下： |
| 30 | 30 | ||
| 31 | -由于源操作数与目的操作数类型位宽比为1:2，读取数据时需要将一个VL大小的数据分为两部分，根据不同接口选择输入数据索引为奇数的位置或偶数的位置。 | 31 | +```python |
| 32 | +def asc_uint82uint16(dst, src, mask): | ||
| 33 | + for i in range(128): | ||
| 34 | + if mask[2 * i]: # 偶数索引位置参与计算 | ||
| 35 | + dst[i] = uint16(src[2 * i]) # uint8 -> uint16 扩展转换 | ||
| 36 | + else: | ||
| 37 | + dst[i] = 0 # mask未选中, 置零 | ||
| 38 | + | ||
| 39 | +def asc_uint82uint16_v2(dst, src, mask): | ||
| 40 | + for i in range(128): | ||
| 41 | + if mask[2 * i + 1]: # 奇数索引位置参与计算 | ||
| 42 | + dst[i] = uint16(src[2 * i + 1]) | ||
| 43 | + else: | ||
| 44 | + dst[i] = 0 | ||
| 45 | +``` | ||
| 32 | 46 | ||
| 33 | ## 函数原型 | 47 | ## 函数原型 |
| 34 | 48 | ||
| 35 | -```cpp | 49 | +### 模板原型（占位符形式） |
| 50 | + | ||
| 51 | +本接口包含2种分块位置，模板原型如下： | ||
| 52 | + | ||
| 53 | +```c | ||
| 54 | +__simd_callee__ inline void asc_uint82uint16<position_mode>(vector_uint16_t& dst, | ||
| 55 | + vector_uint8_t src, | ||
| 56 | + vector_bool mask) | ||
| 57 | +``` | ||
| 58 | + | ||
| 59 | +### 占位符说明 | ||
| 60 | + | ||
| 61 | +- `<position_mode>`表示`src`读取的位置，为空时，读取`src`的偶数索引位置；取值为`_v2`时，读取`src`的奇数索引位置。 | ||
| 62 | + | ||
| 63 | +### 典型示例 | ||
| 64 | + | ||
| 65 | +```c | ||
| 36 | // 数据读取索引为偶数的位置 | 66 | // 数据读取索引为偶数的位置 |
| 37 | -__simd_callee__ inline void asc_uint82uint16(vector_uint16_t& dst, vector_uint8_t src, vector_bool mask) | 67 | +__simd_callee__ inline void asc_uint82uint16(vector_uint16_t& dst, |
| 38 | -// 数据读取索引为奇数的位置 | 68 | + vector_uint8_t src, |
| 39 | -__simd_callee__ inline void asc_uint82uint16_v2(vector_uint16_t& dst, vector_uint8_t src, vector_bool mask) | 69 | + vector_bool mask) |
| 40 | ``` | 70 | ``` |
| 41 | 71 | ||
| 42 | ## 参数说明 | 72 | ## 参数说明 |
| 43 | 73 | ||
| 44 | **表1** 参数说明 | 74 | **表1** 参数说明 |
| 45 | 75 | ||
| 46 | -| 参数名 | 输入/输出 | 描述 | | 76 | +| 参数名 | 输入/输出 | 描述 | |
| 47 | -| --- | --- | --- | | 77 | +| ------ | --------- | ------------------------------------------------------------------------------------------------------------------ | |
| 48 | -| dst | 输出 | 目的操作数（矢量数据寄存器）。| | 78 | +| `dst` | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 49 | -| src | 输入 | 源操作数（矢量数据寄存器）。| | 79 | +| `src` | 输入 | 源操作数（矢量数据寄存器）。 | |
| 50 | -| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在输出中置零。 | | 80 | +| `mask` | 输入 | 掩码寄存器，用于控制各元素是否参与计算。`mask`中与元素对应的比特位为1时，该元素参与计算；为0时，该元素不参与计算。 | |
| 51 | 81 | ||
| 52 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 82 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 53 | 83 | ||
| @@ -57,15 +87,121 @@ __simd_callee__ inline void asc_uint82uint16_v2(vector_uint16_t& dst, vector_uin | |||
| 57 | 87 | ||
| 58 | ## 约束说明 | 88 | ## 约束说明 |
| 59 | 89 | ||
| 60 | -- mask控制源操作数是否参与计算，源操作数不参与计算的元素在输出对应位置置零。 | 90 | +- `src`与`dst`的数据类型需要与函数原型匹配。 |
| 91 | +- `mask`掩码位为0时，`dst`对应元素置0。 | ||
| 61 | 92 | ||
| 62 | ## 调用示例 | 93 | ## 调用示例 |
| 63 | 94 | ||
| 64 | -```cpp | 95 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 65 | -vector_uint16_t dst; | 96 | + |
| 66 | -vector_uint8_t src; | 97 | +<!-- npu="950" id8 --> |
| 67 | -vector_bool mask; | 98 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： |
| 68 | -mask = asc_create_mask_b8(PAT_ALL); | 99 | + |
| 69 | -asc_uint82uint16(dst, src, mask); // 取src的偶数索引元素进行转换并写入dst | 100 | +```bash |
| 70 | -asc_uint82uint16_v2(dst, src, mask); // 取src的奇数索引元素进行转换并写入dst | 101 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main |
| 71 | -``` | 102 | +``` |
| 103 | + | ||
| 104 | +<!-- end id8 --> | ||
| 105 | + | ||
| 106 | +```c | ||
| 107 | +#include <cstdint> | ||
| 108 | +#include <iostream> | ||
| 109 | +#include <vector> | ||
| 110 | +#include "c_api/asc_simd.h" | ||
| 111 | +#include "acl/acl.h" | ||
| 112 | +namespace { | ||
| 113 | +template <typename T> | ||
| 114 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 115 | +{ | ||
| 116 | + std::cout << label << ":"; | ||
| 117 | + const size_t count = values.size() < 8 ? values.size() : 8; | ||
| 118 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; | ||
| 119 | + if (values.size() > count) std::cout << " ..."; | ||
| 120 | + std::cout << std::endl; | ||
| 121 | +} | ||
| 122 | + | ||
| 123 | +template <typename T> | ||
| 124 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 125 | +{ | ||
| 126 | + if (actual.size() != expected.size()) return false; | ||
| 127 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 128 | + if (actual[i] == expected[i]) continue; | ||
| 129 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 130 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 131 | + } | ||
| 132 | + return true; | ||
| 133 | +} | ||
| 134 | + | ||
| 135 | +constexpr uint32_t BUFFER_BYTES = 256; | ||
| 136 | +__simd_vf__ inline void convert(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input) | ||
| 137 | +{ | ||
| 138 | + vector_bool mask = asc_create_mask_b8(PAT_ALL); | ||
| 139 | + vector_uint16_t dst_0; | ||
| 140 | + vector_uint8_t src_0; | ||
| 141 | + asc_loadalign(dst_0, reinterpret_cast<__ubuf__ uint16_t*>(output)); | ||
| 142 | + asc_loadalign(src_0, reinterpret_cast<__ubuf__ uint8_t*>(input)); | ||
| 143 | + asc_uint82uint16_v2(dst_0, src_0, mask); | ||
| 144 | + asc_storealign(reinterpret_cast<__ubuf__ uint16_t*>(output), dst_0, mask); | ||
| 145 | + vector_uint16_t dst_1; | ||
| 146 | + vector_uint8_t src_1; | ||
| 147 | + asc_loadalign(dst_1, reinterpret_cast<__ubuf__ uint16_t*>(output)); | ||
| 148 | + asc_loadalign(src_1, reinterpret_cast<__ubuf__ uint8_t*>(input)); | ||
| 149 | + asc_uint82uint16(dst_1, src_1, mask); | ||
| 150 | + asc_storealign(reinterpret_cast<__ubuf__ uint16_t*>(output), dst_1, mask); | ||
| 151 | +} | ||
| 152 | +__global__ __vector__ void asc_uint82uint16_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input) | ||
| 153 | +{ | ||
| 154 | + asc_init(); | ||
| 155 | + __ubuf__ uint8_t output_local[BUFFER_BYTES]; | ||
| 156 | + __ubuf__ uint8_t input_local[BUFFER_BYTES]; | ||
| 157 | + asc_copy_gm2ub_align(input_local, input, BUFFER_BYTES); | ||
| 158 | + asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES); | ||
| 159 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 160 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 161 | + convert(output_local, input_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(output, output_local, BUFFER_BYTES); | ||
| 165 | + asc_sync(); | ||
| 166 | +} | ||
| 167 | +} // namespace | ||
| 168 | +int main() | ||
| 169 | +{ | ||
| 170 | + constexpr uint8_t input_pattern[] = {0x01}; | ||
| 171 | + constexpr uint8_t golden_pattern[] = {0x01, 0x00}; | ||
| 172 | + std::vector<uint8_t> input(BUFFER_BYTES); | ||
| 173 | + std::vector<uint8_t> output(BUFFER_BYTES, 0xff); | ||
| 174 | + std::vector<uint8_t> golden(BUFFER_BYTES); | ||
| 175 | + for (size_t i = 0; i < BUFFER_BYTES; ++i) { | ||
| 176 | + input[i] = input_pattern[i % sizeof(input_pattern)]; | ||
| 177 | + golden[i] = golden_pattern[i % sizeof(golden_pattern)]; | ||
| 178 | + } | ||
| 179 | + aclInit(nullptr); | ||
| 180 | + aclrtSetDevice(0); | ||
| 181 | + uint8_t* input_device = nullptr; | ||
| 182 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 183 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 184 | + uint8_t* output_device = nullptr; | ||
| 185 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 186 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 187 | + aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t), | ||
| 188 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 189 | + asc_uint82uint16_kernel<<<1, 0>>>(output_device, input_device); | ||
| 190 | + aclrtSynchronizeDevice(); | ||
| 191 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t), | ||
| 192 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 193 | + std::cout << "Conversion: vector_uint8_t -> vector_uint16_t" << std::endl; | ||
| 194 | + print_data("Input 1.0 (raw bytes)", input); | ||
| 195 | + print_data("Output (raw bytes)", output); | ||
| 196 | + print_data("Golden 1.0 (raw bytes)", golden); | ||
| 197 | + const bool passed = compare_data(output, golden); | ||
| 198 | + std::cout << (passed ? "[Success] asc_uint82uint16/asc_uint82uint16_v2 passed." | ||
| 199 | + : "[Failed] asc_uint82uint16/asc_uint82uint16_v2 failed.") | ||
| 200 | + << std::endl; | ||
| 201 | + aclrtFree(input_device); | ||
| 202 | + aclrtFree(output_device); | ||
| 203 | + aclrtResetDevice(0); | ||
| 204 | + aclFinalize(); | ||
| 205 | + return passed ? 0 : 1; | ||
| 206 | +} | ||
| 207 | +``` | ||
| @@ -26,27 +26,58 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -根据mask对源操作数src进行归约求和操作，用于将src中的所有参与计算的元素求和，得到的结果保存在dst中的0位置。 | 29 | +根据`mask`对源操作数`src`进行归约求和操作，用于将`src`中的所有参与计算的元素求和，得到的结果保存在`dst`中的0位置。参考伪代码： |
| 30 | + | ||
| 31 | +```python | ||
| 32 | +def asc_reduce_sum(dst, src, mask): | ||
| 33 | + values = [src[i] if mask[i] else 0 for i in range(len(src))] | ||
| 34 | + while len(values) > 1: | ||
| 35 | + values = [values[i] + values[i + 1] | ||
| 36 | + for i in range(0, len(values), 2)] | ||
| 37 | + dst[0] = values[0] # 求和值 | ||
| 38 | + for i in range(1, len(dst)): | ||
| 39 | + dst[i] = 0 # 其余位置置0 | ||
| 40 | +``` | ||
| 30 | 41 | ||
| 31 | ## 函数原型 | 42 | ## 函数原型 |
| 32 | 43 | ||
| 33 | -```cpp | 44 | +### 模板原型（占位符形式） |
| 34 | -__simd_callee__ inline void asc_reduce_sum(vector_int32_t& dst, vector_int16_t src, vector_bool mask) | 45 | + |
| 35 | -__simd_callee__ inline void asc_reduce_sum(vector_int32_t& dst, vector_int32_t src, vector_bool mask) | 46 | +```c |
| 36 | -__simd_callee__ inline void asc_reduce_sum(vector_uint32_t& dst, vector_uint16_t src, vector_bool mask) | 47 | +__simd_callee__ inline void asc_reduce_sum(vector_<dst_dtype>& dst, |
| 37 | -__simd_callee__ inline void asc_reduce_sum(vector_uint32_t& dst, vector_uint32_t src, vector_bool mask) | 48 | + vector_<src_dtype> src, |
| 38 | -__simd_callee__ inline void asc_reduce_sum(vector_half& dst, vector_half src, vector_bool mask) | 49 | + vector_bool mask) |
| 39 | -__simd_callee__ inline void asc_reduce_sum(vector_float& dst, vector_float src, vector_bool mask) | 50 | +``` |
| 51 | + | ||
| 52 | +### src_dtype与dst_dtype支持组合 | ||
| 53 | + | ||
| 54 | +| `src_dtype` | `dst_dtype` | | ||
| 55 | +| ----------- | ----------- | | ||
| 56 | +| `int16_t` | `int32_t` | | ||
| 57 | +| `uint16_t` | `uint32_t` | | ||
| 58 | +| `half` | `half` | | ||
| 59 | +| `int32_t` | `int32_t` | | ||
| 60 | +| `uint32_t` | `uint32_t` | | ||
| 61 | +| `float` | `float` | | ||
| 62 | + | ||
| 63 | +### 典型示例 | ||
| 64 | + | ||
| 65 | +```c | ||
| 66 | +// 示例：float类型归约求和 | ||
| 67 | +__simd_callee__ inline void asc_reduce_sum(vector_float& dst, | ||
| 68 | + vector_float src, | ||
| 69 | + vector_bool mask) | ||
| 40 | ``` | 70 | ``` |
| 41 | 71 | ||
| 42 | ## 参数说明 | 72 | ## 参数说明 |
| 43 | 73 | ||
| 44 | **表1** 参数说明 | 74 | **表1** 参数说明 |
| 45 | -| 参数名 | 输入/输出 | 描述 | | 75 | + |
| 46 | -| :----- | :------- | :------- | | 76 | +| 参数名 | 输入/输出 | 描述 | |
| 47 | -| dst | 输出 | 目的操作数（矢量数据寄存器）。 | | 77 | +| ------ | --------- | ---------------------------------------- | |
| 48 | -| src | 输入 | 源操作数（矢量数据寄存器）。 | | 78 | +| `dst` | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 49 | -| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。 | | 79 | +| `src` | 输入 | 源操作数（矢量数据寄存器）。 | |
| 80 | +| `mask` | 输入 | 掩码寄存器，用于控制各元素是否参与归约。 | | ||
| 50 | 81 | ||
| 51 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 82 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 52 | 83 | ||
| @@ -57,36 +88,120 @@ __simd_callee__ inline void asc_reduce_sum(vector_float& dst, vector_float src, | |||
| 57 | ## 约束说明 | 88 | ## 约束说明 |
| 58 | 89 | ||
| 59 | - 指令内累加顺序采用二叉树累加方式，结果连续写入到目的操作数，目的操作数中的其它元素置0。 | 90 | - 指令内累加顺序采用二叉树累加方式，结果连续写入到目的操作数，目的操作数中的其它元素置0。 |
| 60 | -- 当所有元素均不参与计算（mask全为0）时，将0写入目的操作数对应位置（特别的，对于浮点数为+0）。 | 91 | +- 当所有元素均不参与计算（`mask`全为0）时，将0写入目的操作数对应位置（特别的，对于浮点数为+0）。 |
| 61 | -- 对于输入为uint16_t/int16_t类型的情况，会提升精度到uint32_t/int32_t进行计算。 | 92 | +- 对于输入为`uint16_t`/`int16_t`类型的情况，会提升精度到`uint32_t`/`int32_t`进行计算。 |
| 62 | 93 | ||
| 63 | ## 关键特性 | 94 | ## 关键特性 |
| 64 | 95 | ||
| 65 | **asc_reduce_sum累加顺序**： | 96 | **asc_reduce_sum累加顺序**： |
| 66 | 97 | ||
| 67 | -以二叉树累加的方式计算源操作数src内有效元素的数据总和。 | 98 | +以二叉树累加的方式计算源操作数`src`内有效元素的数据总和。 |
| 68 | 99 | ||
| 69 | -以half类型的数据求和为例，在src内有128个数，通过二叉树的方式，两两相加，计算过程如下图所示： | 100 | +以`half`类型的数据求和为例，在`src`内有128个数，通过二叉树的方式，两两相加，计算过程如下图所示： |
| 70 | 101 | ||
| 71 | 1. data1和data2相加得到data01，data3和data4相加得到data02，……，data125和data126相加得到data63，data127和data128相加得到data64； | 102 | 1. data1和data2相加得到data01，data3和data4相加得到data02，……，data125和data126相加得到data63，data127和data128相加得到data64； |
| 72 | 2. data01和data02相加得到data001，data03和data04相加得到data002，……，data63和data64相加得到data032； | 103 | 2. data01和data02相加得到data001，data03和data04相加得到data002，……，data63和data64相加得到data032； |
| 73 | -3. 以此类推，得到目的操作数为1个half类型的数据sum。 | 104 | +3. 以此类推，得到目的操作数为1个`half`类型的数据sum。 |
| 74 | 105 | ||
| 75 | **图1** reg_reduce_index | 106 | **图1** reg_reduce_index |
| 107 | + | ||
| 76 |  | 108 |  |
| 77 | 109 | ||
| 78 | ## 调用示例 | 110 | ## 调用示例 |
| 79 | 111 | ||
| 80 | -```cpp | 112 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 81 | -__simd_vf__ inline void reduce_sum_vf(__ubuf__ half* src_addr, __ubuf__ half* dst_addr, uint32_t count, uint16_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 | +#include "c_api/asc_simd.h" | ||
| 127 | +#include "acl/acl.h" | ||
| 128 | +namespace { | ||
| 129 | +template <typename T> | ||
| 130 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 82 | { | 131 | { |
| 83 | - vector_half src, dst; | 132 | + std::cout << label << ":"; |
| 84 | - vector_bool mask; | 133 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 85 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 134 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 86 | - mask = asc_update_mask_b16(count); | 135 | + if (values.size() > count) std::cout << " ..."; |
| 87 | - asc_loadalign_postupdate(src, src_addr, one_repeat_size); | 136 | + std::cout << std::endl; |
| 88 | - asc_reduce_sum(dst, src, mask); | 137 | +} |
| 89 | - asc_storealign_postupdate(dst_addr, dst, one_repeat_size, mask); | 138 | + |
| 139 | +template <typename T> | ||
| 140 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 141 | +{ | ||
| 142 | + if (actual.size() != expected.size()) return false; | ||
| 143 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 144 | + if (actual[i] == expected[i]) continue; | ||
| 145 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 146 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 90 | } | 147 | } |
| 148 | + return true; | ||
| 149 | +} | ||
| 150 | + | ||
| 151 | +constexpr uint32_t ELEMENT_COUNT = 64; | ||
| 152 | +__simd_vf__ inline void reduce(__ubuf__ uint32_t* dst, __ubuf__ uint32_t* src) | ||
| 153 | +{ | ||
| 154 | + vector_uint32_t dst_reg, src_reg; | ||
| 155 | + uint32_t count = ELEMENT_COUNT; | ||
| 156 | + vector_bool mask = asc_update_mask_b32(count); | ||
| 157 | + asc_loadalign(src_reg, src); | ||
| 158 | + asc_reduce_sum(dst_reg, src_reg, mask); | ||
| 159 | + asc_storealign(dst, dst_reg, mask); | ||
| 160 | +} | ||
| 161 | +__global__ __vector__ void asc_reduce_sum_kernel(__gm__ uint32_t* dst, __gm__ uint32_t* src) | ||
| 162 | +{ | ||
| 163 | + asc_init(); | ||
| 164 | + __ubuf__ uint32_t dst_local[ELEMENT_COUNT], src_local[ELEMENT_COUNT]; | ||
| 165 | + asc_copy_gm2ub_align(src_local, src, ELEMENT_COUNT * sizeof(uint32_t)); | ||
| 166 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 167 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 168 | + reduce(dst_local, src_local); | ||
| 169 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 170 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 171 | + asc_copy_ub2gm_align(dst, dst_local, ELEMENT_COUNT * sizeof(uint32_t)); | ||
| 172 | + asc_sync(); | ||
| 173 | +} | ||
| 174 | +} // namespace | ||
| 175 | +int main() | ||
| 176 | +{ | ||
| 177 | + std::vector<uint32_t> src(ELEMENT_COUNT), output(ELEMENT_COUNT, 0), golden(ELEMENT_COUNT, 0); | ||
| 178 | + for (uint32_t i = 0; i < ELEMENT_COUNT; ++i) src[i] = i + 1; | ||
| 179 | + golden[0] = ELEMENT_COUNT * (ELEMENT_COUNT + 1) / 2; | ||
| 180 | + aclInit(nullptr); | ||
| 181 | + aclrtSetDevice(0); | ||
| 182 | + uint32_t* src_device = nullptr; | ||
| 183 | + aclrtMalloc(reinterpret_cast<void**>(&src_device), (ELEMENT_COUNT) * sizeof(uint32_t), | ||
| 184 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 185 | + uint32_t* dst_device = nullptr; | ||
| 186 | + aclrtMalloc(reinterpret_cast<void**>(&dst_device), (ELEMENT_COUNT) * sizeof(uint32_t), | ||
| 187 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 188 | + aclrtMemcpy(src_device, src.size() * sizeof(uint32_t), src.data(), src.size() * sizeof(uint32_t), | ||
| 189 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 190 | + asc_reduce_sum_kernel<<<1, 0>>>(dst_device, src_device); | ||
| 191 | + aclrtSynchronizeDevice(); | ||
| 192 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint32_t), dst_device, output.size() * sizeof(uint32_t), | ||
| 193 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 194 | + print_data("Input", src); | ||
| 195 | + print_data("Output", output); | ||
| 196 | + print_data("Golden", golden); | ||
| 197 | + const bool passed = compare_data(output, golden); | ||
| 198 | + std::cout << (passed ? "[Success] asc_reduce_sum passed." | ||
| 199 | + : "[Failed] asc_reduce_sum failed.") | ||
| 200 | + << std::endl; | ||
| 201 | + aclrtFree(dst_device); | ||
| 202 | + aclrtFree(src_device); | ||
| 203 | + aclrtResetDevice(0); | ||
| 204 | + aclFinalize(); | ||
| 205 | + return passed ? 0 : 1; | ||
| 91 | } | 206 | } |
| 92 | ``` | 207 | ``` |
| @@ -26,27 +26,64 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将每个DataBlock（32B）中参与计算的元素相加，并将最终的计算结果依次保存在dst的最低位。 | 29 | +根据`mask`将每个`DataBlock`（32B）中参与计算的元素相加，并将最终的计算结果依次保存在`dst`的最低位。参考伪代码： |
| 30 | + | ||
| 31 | +```python | ||
| 32 | +def asc_reduce_sum_datablock(dst, src, mask): | ||
| 33 | + num_blocks = 8 # 128元素 / 16元素每块 = 8块 | ||
| 34 | + elems_per_block = 16 # 32B / 2B(half) = 16 | ||
| 35 | + for b in range(num_blocks): | ||
| 36 | + values = [] | ||
| 37 | + for j in range(elems_per_block): | ||
| 38 | + i = b * elems_per_block + j | ||
| 39 | + values.append(src[i] if mask[i] else 0) | ||
| 40 | + while len(values) > 1: | ||
| 41 | + values = [values[j] + values[j + 1] | ||
| 42 | + for j in range(0, len(values), 2)] | ||
| 43 | + dst[b] = values[0] # 每块求和值连续写入dst低位 | ||
| 44 | + for i in range(num_blocks, 128): | ||
| 45 | + dst[i] = 0 # 其余位置置0 | ||
| 46 | +``` | ||
| 30 | 47 | ||
| 31 | ## 函数原型 | 48 | ## 函数原型 |
| 32 | 49 | ||
| 33 | -```cpp | 50 | +### 模板原型（占位符形式） |
| 34 | -__simd_callee__ inline void asc_reduce_sum_datablock(vector_half& dst, vector_half src, vector_bool mask) | 51 | + |
| 35 | -__simd_callee__ inline void asc_reduce_sum_datablock(vector_float& dst, vector_float src, vector_bool mask) | 52 | +```c |
| 36 | -__simd_callee__ inline void asc_reduce_sum_datablock(vector_uint32_t& dst, vector_uint16_t src, vector_bool mask) | 53 | +__simd_callee__ inline void asc_reduce_sum_datablock(vector_<dst_dtype>& dst, |
| 37 | -__simd_callee__ inline void asc_reduce_sum_datablock(vector_int32_t& dst, vector_int16_t src, vector_bool mask) | 54 | + vector_<src_dtype> src, |
| 38 | -__simd_callee__ inline void asc_reduce_sum_datablock(vector_uint32_t& dst, vector_uint32_t src, vector_bool mask) | 55 | + vector_bool mask) |
| 39 | -__simd_callee__ inline void asc_reduce_sum_datablock(vector_int32_t& dst, vector_int32_t src, vector_bool mask) | 56 | +``` |
| 57 | + | ||
| 58 | +### src_dtype与dst_dtype数据类型支持情况 | ||
| 59 | + | ||
| 60 | +| `src_dtype` | `dst_dtype` | | ||
| 61 | +| ----------- | ----------- | | ||
| 62 | +| `int16_t` | `int32_t` | | ||
| 63 | +| `uint16_t` | `uint32_t` | | ||
| 64 | +| `half` | `half` | | ||
| 65 | +| `int32_t` | `int32_t` | | ||
| 66 | +| `uint32_t` | `uint32_t` | | ||
| 67 | +| `float` | `float` | | ||
| 68 | + | ||
| 69 | +### 典型示例 | ||
| 70 | + | ||
| 71 | +```c | ||
| 72 | +// 示例：half类型树形归约求和 | ||
| 73 | +__simd_callee__ inline void asc_reduce_sum_datablock(vector_half& dst, | ||
| 74 | + vector_half src, | ||
| 75 | + vector_bool mask) | ||
| 40 | ``` | 76 | ``` |
| 41 | 77 | ||
| 42 | ## 参数说明 | 78 | ## 参数说明 |
| 43 | 79 | ||
| 44 | **表1** 参数说明 | 80 | **表1** 参数说明 |
| 45 | -| 参数名 | 输入/输出 | 描述 | | 81 | + |
| 46 | -| :----- | :------- | :------- | | 82 | +| 参数名 | 输入/输出 | 描述 | |
| 47 | -| dst | 输出 | 目的操作数（矢量数据寄存器）。 | | 83 | +| ------ | --------- | ---------------------------------------- | |
| 48 | -| src | 输入 | 源操作数（矢量数据寄存器）。 | | 84 | +| `dst` | 输出 | 目的操作数（矢量数据寄存器）。 | |
| 49 | -| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。 | | 85 | +| `src` | 输入 | 源操作数（矢量数据寄存器）。 | |
| 86 | +| `mask` | 输入 | 掩码寄存器，用于控制各元素是否参与归约。 | | ||
| 50 | 87 | ||
| 51 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 | 88 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。 |
| 52 | 89 | ||
| @@ -56,37 +93,121 @@ __simd_callee__ inline void asc_reduce_sum_datablock(vector_int32_t& dst, vector | |||
| 56 | 93 | ||
| 57 | ## 约束说明 | 94 | ## 约束说明 |
| 58 | 95 | ||
| 59 | -- 指令内累加顺序采用二叉树累加方式，在每个DataBlock（32B）内两两相加逐层归约求和，结果连续写入到目的操作数，目的操作数中的其它元素置0。 | 96 | +- 指令内累加顺序采用二叉树累加方式，在每个`DataBlock`（32B）内两两相加逐层归约求和，结果连续写入到目的操作数，目的操作数中的其它元素置0。 |
| 60 | -- 当DataBlock中的元素均不参与计算（mask全为0）时，将0写入dst对应位置（对于浮点数则为+0）。 | 97 | +- 当`DataBlock`中的元素均不参与计算（`mask`全为0）时，将0写入`dst`对应位置（对于浮点数则为+0）。 |
| 61 | -- 对于输入为uint16_t/int16_t类型的情况，会提升精度到uint32_t/int32_t进行计算。 | 98 | +- 对于输入为`uint16_t`/`int16_t`类型的情况，会提升精度到`uint32_t`/`int32_t`进行计算。 |
| 62 | 99 | ||
| 63 | ## 关键特性 | 100 | ## 关键特性 |
| 64 | 101 | ||
| 65 | **asc_reduce_sum_datablock累加顺序**： | 102 | **asc_reduce_sum_datablock累加顺序**： |
| 66 | 103 | ||
| 67 | -以二叉树累加的方式计算每个DataBlock内的和。 | 104 | +以二叉树累加的方式计算每个`DataBlock`内的和。 |
| 68 | 105 | ||
| 69 | -以half类型的数据求和为例，在每个DataBlock内有16个数，通过二叉树的方式，两两相加，计算过程如下图所示： | 106 | +以`half`类型的数据求和为例，在每个`DataBlock`内有16个数，通过二叉树的方式，两两相加，计算过程如下图所示： |
| 70 | 107 | ||
| 71 | 1. data1和data2相加得到data01，data3和data4相加得到data02，……，data13和data14相加得到data07，data15和data16相加得到data08； | 108 | 1. data1和data2相加得到data01，data3和data4相加得到data02，……，data13和data14相加得到data07，data15和data16相加得到data08； |
| 72 | 2. data01和data02相加得到data001，data03和data04相加得到data002，……，data07和data08相加得到data004； | 109 | 2. data01和data02相加得到data001，data03和data04相加得到data002，……，data07和data08相加得到data004； |
| 73 | -3. 以此类推，得到目的操作数为1个half类型的数据sum。 | 110 | +3. 以此类推，得到目的操作数为1个`half`类型的数据sum。 |
| 74 | 111 | ||
| 75 | **图1** asc_reduce_sum_datablock示意图 | 112 | **图1** asc_reduce_sum_datablock示意图 |
| 113 | + | ||
| 76 |  | 114 |  |
| 77 | 115 | ||
| 78 | ## 调用示例 | 116 | ## 调用示例 |
| 79 | 117 | ||
| 80 | -```cpp | 118 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。 |
| 81 | -__simd_vf__ inline void reduce_sum_datablock_vf(__ubuf__ half* src_addr, __ubuf__ half* dst_addr, uint32_t count, uint16_t one_repeat_size, uint16_t repeat_time) | 119 | + |
| 120 | +<!-- npu="950" id8 --> | ||
| 121 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 122 | + | ||
| 123 | +```bash | ||
| 124 | +bisheng example.asc -o main --npu-arch=dav-3510; ./main | ||
| 125 | +``` | ||
| 126 | +<!-- end id8 --> | ||
| 127 | + | ||
| 128 | +```c | ||
| 129 | +#include <cstdint> | ||
| 130 | +#include <iostream> | ||
| 131 | +#include <vector> | ||
| 132 | +#include "c_api/asc_simd.h" | ||
| 133 | +#include "acl/acl.h" | ||
| 134 | +namespace { | ||
| 135 | +template <typename T> | ||
| 136 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 82 | { | 137 | { |
| 83 | - vector_half src, dst; | 138 | + std::cout << label << ":"; |
| 84 | - vector_bool mask; | 139 | + const size_t count = values.size() < 8 ? values.size() : 8; |
| 85 | - for (uint16_t i = 0; i < repeat_time; ++i) { | 140 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; |
| 86 | - mask = asc_update_mask_b16(count); | 141 | + if (values.size() > count) std::cout << " ..."; |
| 87 | - asc_loadalign_postupdate(src, src_addr, one_repeat_size); | 142 | + std::cout << std::endl; |
| 88 | - asc_reduce_sum_datablock(dst, src, mask); | 143 | +} |
| 89 | - asc_storealign_postupdate(dst_addr, dst, one_repeat_size, mask); | 144 | + |
| 145 | +template <typename T> | ||
| 146 | +bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0) | ||
| 147 | +{ | ||
| 148 | + if (actual.size() != expected.size()) return false; | ||
| 149 | + for (size_t i = 0; i < actual.size(); ++i) { | ||
| 150 | + if (actual[i] == expected[i]) continue; | ||
| 151 | + const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]); | ||
| 152 | + if (diff > tolerance || diff < -tolerance) return false; | ||
| 90 | } | 153 | } |
| 154 | + return true; | ||
| 155 | +} | ||
| 156 | + | ||
| 157 | +constexpr uint32_t ELEMENT_COUNT = 64; | ||
| 158 | +__simd_vf__ inline void reduce(__ubuf__ uint32_t* dst, __ubuf__ uint32_t* src) | ||
| 159 | +{ | ||
| 160 | + vector_uint32_t dst_reg, src_reg; | ||
| 161 | + uint32_t count = ELEMENT_COUNT; | ||
| 162 | + vector_bool mask = asc_update_mask_b32(count); | ||
| 163 | + asc_loadalign(src_reg, src); | ||
| 164 | + asc_reduce_sum_datablock(dst_reg, src_reg, mask); | ||
| 165 | + asc_storealign(dst, dst_reg, mask); | ||
| 166 | +} | ||
| 167 | +__global__ __vector__ void asc_reduce_sum_datablock_kernel(__gm__ uint32_t* dst, __gm__ uint32_t* src) | ||
| 168 | +{ | ||
| 169 | + asc_init(); | ||
| 170 | + __ubuf__ uint32_t dst_local[ELEMENT_COUNT], src_local[ELEMENT_COUNT]; | ||
| 171 | + asc_copy_gm2ub_align(src_local, src, ELEMENT_COUNT * sizeof(uint32_t)); | ||
| 172 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 173 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 174 | + reduce(dst_local, src_local); | ||
| 175 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 176 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 177 | + asc_copy_ub2gm_align(dst, dst_local, ELEMENT_COUNT * sizeof(uint32_t)); | ||
| 178 | + asc_sync(); | ||
| 179 | +} | ||
| 180 | +} // namespace | ||
| 181 | +int main() | ||
| 182 | +{ | ||
| 183 | + std::vector<uint32_t> src(ELEMENT_COUNT), output(ELEMENT_COUNT, 0), golden(ELEMENT_COUNT, 0); | ||
| 184 | + for (uint32_t i = 0; i < ELEMENT_COUNT; ++i) src[i] = i + 1; | ||
| 185 | + for (uint32_t block = 0; block < 8; ++block) for (uint32_t i = 0; i < 8; ++i) golden[block] += src[block * 8 + i]; | ||
| 186 | + aclInit(nullptr); | ||
| 187 | + aclrtSetDevice(0); | ||
| 188 | + uint32_t* src_device = nullptr; | ||
| 189 | + aclrtMalloc(reinterpret_cast<void**>(&src_device), (ELEMENT_COUNT) * sizeof(uint32_t), | ||
| 190 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 191 | + uint32_t* dst_device = nullptr; | ||
| 192 | + aclrtMalloc(reinterpret_cast<void**>(&dst_device), (ELEMENT_COUNT) * sizeof(uint32_t), | ||
| 193 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 194 | + aclrtMemcpy(src_device, src.size() * sizeof(uint32_t), src.data(), src.size() * sizeof(uint32_t), | ||
| 195 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 196 | + asc_reduce_sum_datablock_kernel<<<1, 0>>>(dst_device, src_device); | ||
| 197 | + aclrtSynchronizeDevice(); | ||
| 198 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint32_t), dst_device, output.size() * sizeof(uint32_t), | ||
| 199 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 200 | + print_data("Input", src); | ||
| 201 | + print_data("Output", output); | ||
| 202 | + print_data("Golden", golden); | ||
| 203 | + const bool passed = compare_data(output, golden); | ||
| 204 | + std::cout << (passed ? "[Success] asc_reduce_sum_datablock passed." | ||
| 205 | + : "[Failed] asc_reduce_sum_datablock failed.") | ||
| 206 | + << std::endl; | ||
| 207 | + aclrtFree(dst_device); | ||
| 208 | + aclrtFree(src_device); | ||
| 209 | + aclrtResetDevice(0); | ||
| 210 | + aclFinalize(); | ||
| 211 | + return passed ? 0 : 1; | ||
| 91 | } | 212 | } |
| 92 | ``` | 213 | ``` |
| @@ -28,6 +28,8 @@ | |||
| 28 | 28 | ||
| 29 | 获取当前运行核的索引。 | 29 | 获取当前运行核的索引。 |
| 30 | 30 | ||
| 31 | +**此接口后续版本会废弃，请使用内置变量[block_idx](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#内置变量)。** | ||
| 32 | + | ||
| 31 | ## 函数原型 | 33 | ## 函数原型 |
| 32 | 34 | ||
| 33 | ```cpp | 35 | ```cpp |
| @@ -44,7 +46,7 @@ __aicore__ inline int64_t asc_get_block_idx() | |||
| 44 | 46 | ||
| 45 | ## 流水类型 | 47 | ## 流水类型 |
| 46 | 48 | ||
| 47 | -PIPE_S | 49 | +`PIPE_S` |
| 48 | 50 | ||
| 49 | ## 约束说明 | 51 | ## 约束说明 |
| 50 | 52 | ||
| @@ -28,6 +28,8 @@ | |||
| 28 | 28 | ||
| 29 | 获取当前任务配置的核数，用于代码内部的多核逻辑控制等。 | 29 | 获取当前任务配置的核数，用于代码内部的多核逻辑控制等。 |
| 30 | 30 | ||
| 31 | +**此接口后续版本会废弃，请使用内置变量[block_num](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#内置变量)。** | ||
| 32 | + | ||
| 31 | ## 函数原型 | 33 | ## 函数原型 |
| 32 | 34 | ||
| 33 | ```cpp | 35 | ```cpp |
| @@ -44,7 +46,7 @@ __aicore__ inline int64_t asc_get_block_num() | |||
| 44 | 46 | ||
| 45 | ## 流水类型 | 47 | ## 流水类型 |
| 46 | 48 | ||
| 47 | -PIPE_S | 49 | +`PIPE_S` |
| 48 | 50 | ||
| 49 | ## 约束说明 | 51 | ## 约束说明 |
| 50 | 52 | ||