Merged
带返回值接口资料修改 #5501
qwerwjyAdd on Aug 31
带返回值接口资料修改 #5501
Merged
Total: 8 files changed+415-22
| @@ -1077,6 +1077,8 @@ | |||
| 1077 | - [asc_scatter](SIMD-API/c_api/reg_compute/scatter/asc_scatter.md) | 1077 | - [asc_scatter](SIMD-API/c_api/reg_compute/scatter/asc_scatter.md) |
| 1078 | - [Reg间搬运](SIMD-API/c_api/reg_compute/reg_copy/reg_copy.md) | 1078 | - [Reg间搬运](SIMD-API/c_api/reg_compute/reg_copy/reg_copy.md) |
| 1079 | - [asc_copy](SIMD-API/c_api/reg_compute/reg_copy/asc_copy.md) | 1079 | - [asc_copy](SIMD-API/c_api/reg_compute/reg_copy/asc_copy.md) |
| 1080 | + - [asc_extract_mask](SIMD-API/c_api/reg_compute/reg_copy/asc_extract_mask.md) | ||
| 1081 | + - [asc_mask](SIMD-API/c_api/reg_compute/reg_copy/asc_mask.md) | ||
| 1080 | - [基础算术](SIMD-API/c_api/reg_compute/reg_arith/reg_arith.md) | 1082 | - [基础算术](SIMD-API/c_api/reg_compute/reg_arith/reg_arith.md) |
| 1081 | - [asc_abs](SIMD-API/c_api/reg_compute/reg_arith/asc_abs.md) | 1083 | - [asc_abs](SIMD-API/c_api/reg_compute/reg_arith/asc_abs.md) |
| 1082 | - [asc_add](SIMD-API/c_api/reg_compute/reg_arith/asc_add.md) | 1084 | - [asc_add](SIMD-API/c_api/reg_compute/reg_arith/asc_add.md) |
| @@ -394,6 +394,8 @@ C API文档按最细粒度公共头文件归类。除试验接口外，整体使 | |||
| 394 | |API名称|说明|最细粒度公共头文件| | 394 | |API名称|说明|最细粒度公共头文件| |
| 395 | |---|---|---| | 395 | |---|---|---| |
| 396 | |[asc_copy](reg_compute/reg_copy/asc_copy.md)|完成数据在寄存器内的搬运。|`c_api/reg_compute/reg_copy.h`| | 396 | |[asc_copy](reg_compute/reg_copy/asc_copy.md)|完成数据在寄存器内的搬运。|`c_api/reg_compute/reg_copy.h`| |
| 397 | +|[asc_extract_mask](reg_compute/reg_copy/asc_extract_mask.md)|从矢量数据寄存器中提取掩码寄存器，通过函数返回值返回结果。|`c_api/reg_compute/reg_copy.h`| | ||
| 398 | +|[asc_mask](reg_compute/reg_copy/asc_mask.md)|通过函数返回值返回矢量数据寄存器或掩码寄存器的掩码复制结果。|`c_api/reg_compute/reg_copy.h`| | ||
| 397 | 399 | ||
| 398 | ### 同步控制 | 400 | ### 同步控制 |
| 399 | 401 | ||
| @@ -0,0 +1,155 @@ | |||
| 1 | +# asc_extract_mask | ||
| 2 | + | ||
| 3 | +## 产品支持情况 | ||
| 4 | + | ||
| 5 | +<!-- npu="950" id1 --> | ||
| 6 | +- Ascend 950PR/Ascend 950DT：支持 | ||
| 7 | +<!-- end id1 --> | ||
| 8 | +<!-- npu="A3" id2 --> | ||
| 9 | +- Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持 | ||
| 10 | +<!-- end id2 --> | ||
| 11 | +<!-- npu="910b" id3 --> | ||
| 12 | +- Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持 | ||
| 13 | +<!-- end id3 --> | ||
| 14 | +<!-- npu="310b" id4 --> | ||
| 15 | +- Atlas 200I/500 A2 推理产品：不支持 | ||
| 16 | +<!-- end id4 --> | ||
| 17 | +<!-- npu="310p" id5 --> | ||
| 18 | +- Atlas 推理系列产品AI Core：不支持 | ||
| 19 | +<!-- end id5 --> | ||
| 20 | +<!-- npu="310p" id6 --> | ||
| 21 | +- Atlas 推理系列产品Vector Core：不支持 | ||
| 22 | +<!-- end id6 --> | ||
| 23 | +<!-- npu="910" id7 --> | ||
| 24 | +- Atlas 训练系列产品：不支持 | ||
| 25 | +<!-- end id7 --> | ||
| 26 | + | ||
| 27 | +## 功能说明 | ||
| 28 | + | ||
| 29 | +从矢量数据寄存器中提取掩码寄存器，并通过函数返回值返回结果。 | ||
| 30 | + | ||
| 31 | +- 数据类型为uint16_t时，矢量数据寄存器（256B）分成16个数据块（16B），用户通过参数可选择某一个数据块，数据块的每个bit 广播成2bit后变为32B，即输出的mask。 | ||
| 32 | +- 数据类型为uint32_t时，矢量数据寄存器（256B）分成32个数据块（8B），用户通过参数可选择某一个数据块，数据块的每个bit 广播成4bit后变为32B，即输出的mask。 | ||
| 33 | + | ||
| 34 | +具体搬运原理请参考[asc_copy关键特性说明](asc_copy.md#关键特性说明)。 | ||
| 35 | + | ||
| 36 | +本接口与[asc_copy](asc_copy.md)的矢量数据寄存器复制到掩码寄存器模式功能相同，区别在于本接口通过函数返回值返回结果。 | ||
| 37 | + | ||
| 38 | +本接口为Reg矢量计算接口，仅在AIV上生效。 | ||
| 39 | + | ||
| 40 | +## 函数原型 | ||
| 41 | + | ||
| 42 | +**part的值必须在编译期确定，可传入整数字面量或constexpr整数常量，不支持传入运行时变量。** | ||
| 43 | + | ||
| 44 | +```c | ||
| 45 | +__simd_callee__ inline vector_bool asc_extract_mask(vector_uint16_t src, | ||
| 46 | + int16_t part) | ||
| 47 | +__simd_callee__ inline vector_bool asc_extract_mask(vector_uint32_t src, | ||
| 48 | + int16_t part) | ||
| 49 | +``` | ||
| 50 | + | ||
| 51 | +## 参数说明 | ||
| 52 | + | ||
| 53 | +**表1** 参数说明 | ||
| 54 | + | ||
| 55 | +| 参数名 | 输入/输出 | 描述 | | ||
| 56 | +| --- | --- | --- | | ||
| 57 | +| src | 输入 | 源矢量数据寄存器。 | | ||
| 58 | +| part | 输入 | part决定了src中需要搬运的数据块，具体描述请参考[asc_copy关键特性说明](asc_copy.md#关键特性说明)。<br>• 当数据类型为uint16_t，part∈[0, 15]。<br>• 当数据类型为uint32_t，part∈[0, 31]。 | | ||
| 59 | + | ||
| 60 | +矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../../defs/type/data_type_definition.md)。 | ||
| 61 | + | ||
| 62 | +## 返回值说明 | ||
| 63 | + | ||
| 64 | +返回保存提取结果的掩码寄存器，类型为`vector_bool`。 | ||
| 65 | + | ||
| 66 | +## 约束说明 | ||
| 67 | + | ||
| 68 | +- 本接口仅在AIV上生效，非AIV调用直接返回。 | ||
| 69 | +- 本接口在Vector Function（`__simd_vf__`标记的函数）内调用，src为矢量数据寄存器。 | ||
| 70 | +- **part的值必须在编译期确定，可传入整数字面量或constexpr整数常量，不支持传入运行时变量。** | ||
| 71 | + - 当数据类型为uint16_t时，part∈[0, 15]。 | ||
| 72 | + - 当数据类型为uint32_t时，part∈[0, 31]。 | ||
| 73 | + | ||
| 74 | +## 调用示例 | ||
| 75 | + | ||
| 76 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。 | ||
| 77 | + | ||
| 78 | +<!-- npu="950" id8 --> | ||
| 79 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 80 | + | ||
| 81 | +```bash | ||
| 82 | +bisheng example.asc -o main --npu-arch=dav-3510&& ./main | ||
| 83 | +``` | ||
| 84 | +<!-- end id8 --> | ||
| 85 | + | ||
| 86 | +```cpp | ||
| 87 | +#include <cstdint> | ||
| 88 | +#include <iostream> | ||
| 89 | +#include <vector> | ||
| 90 | +#include "c_api/asc_simd.h" | ||
| 91 | +#include "acl/acl.h" | ||
| 92 | +namespace { | ||
| 93 | +template <typename T> | ||
| 94 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 95 | +{ | ||
| 96 | + std::cout << label << ":"; | ||
| 97 | + const size_t count = values.size() < 8 ? values.size() : 8; | ||
| 98 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; | ||
| 99 | + if (values.size() > count) std::cout << " ..."; | ||
| 100 | + std::cout << std::endl; | ||
| 101 | +} | ||
| 102 | + | ||
| 103 | +constexpr uint32_t ELEMENT_COUNT = 128; | ||
| 104 | +__simd_vf__ inline void extract_mask(__ubuf__ uint32_t* output, __ubuf__ uint16_t* input) | ||
| 105 | +{ | ||
| 106 | + vector_uint16_t src_reg; | ||
| 107 | + asc_loadalign(src_reg, input); | ||
| 108 | + vector_bool mask = asc_extract_mask(src_reg, 0); | ||
| 109 | + asc_storealign(output, mask); | ||
| 110 | +} | ||
| 111 | +__global__ __vector__ void asc_extract_mask_kernel(__gm__ uint32_t* output, __gm__ uint16_t* input) | ||
| 112 | +{ | ||
| 113 | + asc_init(); | ||
| 114 | + __ubuf__ uint32_t output_local[ELEMENT_COUNT]; | ||
| 115 | + __ubuf__ uint16_t input_local[ELEMENT_COUNT]; | ||
| 116 | + asc_copy_gm2ub_align(input_local, input, ELEMENT_COUNT * sizeof(uint16_t)); | ||
| 117 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 118 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 119 | + extract_mask(output_local, input_local); | ||
| 120 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 121 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 122 | + asc_copy_ub2gm_align(output, output_local, ELEMENT_COUNT * sizeof(uint32_t)); | ||
| 123 | + asc_sync(); | ||
| 124 | +} | ||
| 125 | +} // namespace | ||
| 126 | +int main() | ||
| 127 | +{ | ||
| 128 | + std::vector<uint16_t> input(ELEMENT_COUNT); | ||
| 129 | + std::vector<uint32_t> output(ELEMENT_COUNT, 0); | ||
| 130 | + for (uint32_t i = 0; i < ELEMENT_COUNT; ++i) input[i] = static_cast<uint16_t>(i); | ||
| 131 | + aclInit(nullptr); | ||
| 132 | + aclrtSetDevice(0); | ||
| 133 | + uint16_t* input_device = nullptr; | ||
| 134 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), (ELEMENT_COUNT) * sizeof(uint16_t), | ||
| 135 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 136 | + uint32_t* output_device = nullptr; | ||
| 137 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), (ELEMENT_COUNT) * sizeof(uint32_t), | ||
| 138 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 139 | + aclrtMemcpy(input_device, input.size() * sizeof(uint16_t), input.data(), input.size() * sizeof(uint16_t), | ||
| 140 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 141 | + asc_extract_mask_kernel<<<1, 0>>>(output_device, input_device); | ||
| 142 | + aclrtSynchronizeDevice(); | ||
| 143 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint32_t), output_device, output.size() * sizeof(uint32_t), | ||
| 144 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 145 | + print_data("Input", input); | ||
| 146 | + print_data("Output", output); | ||
| 147 | + const bool passed = true; | ||
| 148 | + std::cout << "[Success] asc_extract_mask completed." << std::endl; | ||
| 149 | + aclrtFree(input_device); | ||
| 150 | + aclrtFree(output_device); | ||
| 151 | + aclrtResetDevice(0); | ||
| 152 | + aclFinalize(); | ||
| 153 | + return passed ? 0 : 1; | ||
| 154 | +} | ||
| 155 | +``` | ||
| @@ -0,0 +1,179 @@ | |||
| 1 | +# asc_mask | ||
| 2 | + | ||
| 3 | +## 产品支持情况 | ||
| 4 | + | ||
| 5 | +<!-- npu="950" id1 --> | ||
| 6 | +- Ascend 950PR/Ascend 950DT：支持 | ||
| 7 | +<!-- end id1 --> | ||
| 8 | +<!-- npu="A3" id2 --> | ||
| 9 | +- Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持 | ||
| 10 | +<!-- end id2 --> | ||
| 11 | +<!-- npu="910b" id3 --> | ||
| 12 | +- Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持 | ||
| 13 | +<!-- end id3 --> | ||
| 14 | +<!-- npu="310b" id4 --> | ||
| 15 | +- Atlas 200I/500 A2 推理产品：不支持 | ||
| 16 | +<!-- end id4 --> | ||
| 17 | +<!-- npu="310p" id5 --> | ||
| 18 | +- Atlas 推理系列产品AI Core：不支持 | ||
| 19 | +<!-- end id5 --> | ||
| 20 | +<!-- npu="310p" id6 --> | ||
| 21 | +- Atlas 推理系列产品Vector Core：不支持 | ||
| 22 | +<!-- end id6 --> | ||
| 23 | +<!-- npu="910" id7 --> | ||
| 24 | +- Atlas 训练系列产品：不支持 | ||
| 25 | +<!-- end id7 --> | ||
| 26 | + | ||
| 27 | +## 功能说明 | ||
| 28 | + | ||
| 29 | +将src中被mask筛选的有效元素复制到返回值对应位置，未被mask筛选的位置置为0（ZEROING模式），并通过函数返回值返回结果。支持以下两种模式： | ||
| 30 | + | ||
| 31 | +- 矢量数据寄存器复制到矢量数据寄存器：将src中被mask筛选的有效元素复制到返回值对应位置，未被mask筛选的位置保持原值。 | ||
| 32 | +- 掩码寄存器复制到掩码寄存器：将src中被mask筛选的bit复制到返回值中，未被mask筛选的位置置为0。支持带mask和不带mask两种重载： | ||
| 33 | + - 带mask：将src中被mask筛选的bit复制到返回值。对于b8模式，mask的每个bit控制src的一个bit；对于b16模式，mask的每2bit为一组、仅LSB控制src的一个bit；对于b32模式，mask的每4bit为一组、仅LSB控制src的一个bit。 | ||
| 34 | + - 不带mask：将src的所有bit复制到返回值。 | ||
| 35 | + | ||
| 36 | +本接口与[asc_copy](asc_copy.md)的矢量数据寄存器复制到矢量数据寄存器模式及掩码寄存器复制到掩码寄存器模式功能相同，区别在于本接口通过函数返回值返回结果。 | ||
| 37 | + | ||
| 38 | +本接口为Reg矢量计算接口，仅在AIV上生效。 | ||
| 39 | + | ||
| 40 | +## 函数原型 | ||
| 41 | + | ||
| 42 | +### 矢量数据寄存器复制到矢量数据寄存器 | ||
| 43 | + | ||
| 44 | +```c | ||
| 45 | +__simd_callee__ inline vector_<dtype> asc_mask(vector_<dtype> src, | ||
| 46 | + vector_bool mask) | ||
| 47 | +``` | ||
| 48 | + | ||
| 49 | +#### dtype支持数据类型 | ||
| 50 | +dtype支持的数据类型：int8_t、uint8_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float。 | ||
| 51 | + | ||
| 52 | +#### 函数原型典型示例 | ||
| 53 | + | ||
| 54 | +```c | ||
| 55 | +// 示例：对half矢量数据寄存器执行掩码复制。 | ||
| 56 | +__simd_callee__ inline vector_half asc_mask(vector_half src, | ||
| 57 | + vector_bool mask) | ||
| 58 | +``` | ||
| 59 | + | ||
| 60 | +### 掩码寄存器复制到掩码寄存器 | ||
| 61 | + | ||
| 62 | +```c | ||
| 63 | +__simd_callee__ inline vector_bool asc_mask(vector_bool src, | ||
| 64 | + vector_bool mask) | ||
| 65 | +__simd_callee__ inline vector_bool asc_mask(vector_bool src) | ||
| 66 | +``` | ||
| 67 | + | ||
| 68 | +## 参数说明 | ||
| 69 | + | ||
| 70 | +**表1** 参数说明（矢量数据寄存器复制到矢量数据寄存器） | ||
| 71 | + | ||
| 72 | +| 参数名 | 输入/输出 | 描述 | | ||
| 73 | +| --- | --- | --- | | ||
| 74 | +| src | 输入 | 源矢量数据寄存器。 | | ||
| 75 | +| mask | 输入 | 源掩码寄存器，用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在返回值保持原值。 | | ||
| 76 | + | ||
| 77 | +**表2** 参数说明（掩码寄存器复制到掩码寄存器） | ||
| 78 | + | ||
| 79 | +| 参数名 | 输入/输出 | 描述 | | ||
| 80 | +| --- | --- | --- | | ||
| 81 | +| src | 输入 | 源掩码寄存器。 | | ||
| 82 | +| mask | 输入 | 源掩码寄存器，用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在返回值中置0。仅带mask的重载包含该参数。 | | ||
| 83 | + | ||
| 84 | +矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../../defs/type/data_type_definition.md)。 | ||
| 85 | + | ||
| 86 | +## 返回值说明 | ||
| 87 | + | ||
| 88 | +- 矢量数据寄存器复制到矢量数据寄存器：返回保存掩码复制结果的矢量数据寄存器，数据类型与`src`保持一致。 | ||
| 89 | +- 掩码寄存器复制到掩码寄存器：返回保存掩码复制结果的掩码寄存器，类型为`vector_bool`。 | ||
| 90 | + | ||
| 91 | +## 约束说明 | ||
| 92 | + | ||
| 93 | +- 本接口仅在AIV上生效，非AIV调用直接返回。 | ||
| 94 | +- 本接口在Vector Function（`__simd_vf__`标记的函数）内调用，src为矢量数据寄存器或掩码寄存器。 | ||
| 95 | + | ||
| 96 | +## 调用示例 | ||
| 97 | + | ||
| 98 | +将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。 | ||
| 99 | + | ||
| 100 | +<!-- npu="950" id8 --> | ||
| 101 | +以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下： | ||
| 102 | + | ||
| 103 | +```bash | ||
| 104 | +bisheng example.asc -o main --npu-arch=dav-3510&& ./main | ||
| 105 | +``` | ||
| 106 | +<!-- end id8 --> | ||
| 107 | + | ||
| 108 | +```cpp | ||
| 109 | +#include <cstdint> | ||
| 110 | +#include <iostream> | ||
| 111 | +#include <vector> | ||
| 112 | +#include "c_api/asc_simd.h" | ||
| 113 | +#include "acl/acl.h" | ||
| 114 | +namespace { | ||
| 115 | +template <typename T> | ||
| 116 | +void print_data(const char* label, const std::vector<T>& values) | ||
| 117 | +{ | ||
| 118 | + std::cout << label << ":"; | ||
| 119 | + const size_t count = values.size() < 8 ? values.size() : 8; | ||
| 120 | + for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i]; | ||
| 121 | + if (values.size() > count) std::cout << " ..."; | ||
| 122 | + std::cout << std::endl; | ||
| 123 | +} | ||
| 124 | + | ||
| 125 | +constexpr uint32_t BUFFER_BYTES = 256; | ||
| 126 | +__simd_vf__ inline void load(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input) | ||
| 127 | +{ | ||
| 128 | + vector_bool mask = asc_create_mask_b8(PAT_ALL); | ||
| 129 | + for (uint16_t repeat = 0; repeat < 1; ++repeat) { | ||
| 130 | + vector_int8_t arg_0_0; | ||
| 131 | + asc_loadalign(arg_0_0, reinterpret_cast<__ubuf__ int8_t*>(input)); | ||
| 132 | + vector_int8_t result = asc_mask(arg_0_0, mask); | ||
| 133 | + asc_storealign(reinterpret_cast<__ubuf__ int8_t*>(output), result, mask); | ||
| 134 | + } | ||
| 135 | +} | ||
| 136 | +__global__ __vector__ void asc_mask_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input) | ||
| 137 | +{ | ||
| 138 | + asc_init(); | ||
| 139 | + __ubuf__ uint8_t output_local[BUFFER_BYTES], input_local[BUFFER_BYTES]; | ||
| 140 | + asc_copy_gm2ub_align(input_local, input, BUFFER_BYTES); | ||
| 141 | + asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES); | ||
| 142 | + asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 143 | + asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); | ||
| 144 | + load(output_local, input_local); | ||
| 145 | + asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 146 | + asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); | ||
| 147 | + asc_copy_ub2gm_align(output, output_local, BUFFER_BYTES); | ||
| 148 | + asc_sync(); | ||
| 149 | +} | ||
| 150 | +} // namespace | ||
| 151 | +int main() | ||
| 152 | +{ | ||
| 153 | + std::vector<uint8_t> input(BUFFER_BYTES), output(BUFFER_BYTES, 0xff); | ||
| 154 | + for (uint32_t i = 0; i < BUFFER_BYTES; ++i) input[i] = static_cast<uint8_t>(i % 251); | ||
| 155 | + aclInit(nullptr); | ||
| 156 | + aclrtSetDevice(0); | ||
| 157 | + uint8_t* input_device = nullptr; | ||
| 158 | + aclrtMalloc(reinterpret_cast<void**>(&input_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 159 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 160 | + uint8_t* output_device = nullptr; | ||
| 161 | + aclrtMalloc(reinterpret_cast<void**>(&output_device), (BUFFER_BYTES) * sizeof(uint8_t), | ||
| 162 | + ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 163 | + aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t), | ||
| 164 | + ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 165 | + asc_mask_kernel<<<1, 0>>>(output_device, input_device); | ||
| 166 | + aclrtSynchronizeDevice(); | ||
| 167 | + aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t), | ||
| 168 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 169 | + print_data("Input bytes", input); | ||
| 170 | + print_data("Output bytes", output); | ||
| 171 | + const bool passed = true; | ||
| 172 | + std::cout << "[Success] asc_mask completed." << std::endl; | ||
| 173 | + aclrtFree(input_device); | ||
| 174 | + aclrtFree(output_device); | ||
| 175 | + aclrtResetDevice(0); | ||
| 176 | + aclFinalize(); | ||
| 177 | + return passed ? 0 : 1; | ||
| 178 | +} | ||
| 179 | +``` | ||
| @@ -1,3 +1,5 @@ | |||
| 1 | # Reg间搬运 | 1 | # Reg间搬运 |
| 2 | 2 | ||
| 3 | - **[asc_copy](asc_copy.md)** | 3 | - **[asc_copy](asc_copy.md)** |
| 4 | +- **[asc_extract_mask](asc_extract_mask.md)** | ||
| 5 | +- **[asc_mask](asc_mask.md)** | ||
| @@ -29,15 +29,17 @@ | |||
| 29 | asc_xor支持两种接口： | 29 | asc_xor支持两种接口： |
| 30 | 30 | ||
| 31 | - 对矢量数据寄存器操作： | 31 | - 对矢量数据寄存器操作： |
| 32 | - 根据`mask`对源矢量数据寄存器`src0`和`src1`进行按位异或（^）操作，将结果写入目的矢量数据寄存器`dst`。未被`mask`筛选的位置被置为0。 | 32 | + 根据`mask`对源矢量数据寄存器`src0`和`src1`进行按位异或（^）操作，将结果写入目的矢量数据寄存器`dst`或通过函数返回值返回。未被`mask`筛选的位置被置为0。 |
| 33 | - 对掩码寄存器操作： | 33 | - 对掩码寄存器操作： |
| 34 | - 根据`mask`对两个源掩码寄存器`src0`和`src1`进行按位异或（^）操作，将结果写入目的掩码数据寄存器`dst`。未被`mask`筛选的位置被置为0。 | 34 | + 根据`mask`对两个源掩码寄存器`src0`和`src1`进行按位异或（^）操作，将结果写入目的掩码数据寄存器`dst`或通过函数返回值返回。未被`mask`筛选的位置被置为0。 |
| 35 | 35 | ||
| 36 | 计算公式如下： | 36 | 计算公式如下： |
| 37 | 37 | ||
| 38 | -$$ | 38 | +```python |
| 39 | -dst_i = src0_i \oplus src1_i | 39 | +def asc_xor(dst, src0, src1, mask): |
| 40 | -$$ | 40 | + for i in range(len(src0)): |
| 41 | + dst[i] = src0[i] ^ src1[i] if mask[i] else 0 | ||
| 42 | +``` | ||
| 41 | 43 | ||
| 42 | 本接口为`Reg`矢量计算接口，仅在AIV上生效。 | 44 | 本接口为`Reg`矢量计算接口，仅在AIV上生效。 |
| 43 | 45 | ||
| @@ -46,6 +48,12 @@ $$ | |||
| 46 | ### 矢量数据寄存器按位异或 | 48 | ### 矢量数据寄存器按位异或 |
| 47 | 49 | ||
| 48 | ```c | 50 | ```c |
| 51 | +// 通过函数返回值返回结果。 | ||
| 52 | +__simd_callee__ inline vector_<dtype> asc_xor(vector_<dtype> src0, | ||
| 53 | + vector_<dtype> src1, | ||
| 54 | + vector_bool mask) | ||
| 55 | + | ||
| 56 | +// 通过引用参数输出结果。 | ||
| 49 | __simd_callee__ inline void asc_xor(vector_<dtype>& dst, | 57 | __simd_callee__ inline void asc_xor(vector_<dtype>& dst, |
| 50 | vector_<dtype> src0, | 58 | vector_<dtype> src0, |
| 51 | vector_<dtype> src1, | 59 | vector_<dtype> src1, |
| @@ -58,6 +66,11 @@ __simd_callee__ inline void asc_xor(vector_<dtype>& dst, | |||
| 58 | #### 函数原型典型示例 | 66 | #### 函数原型典型示例 |
| 59 | 67 | ||
| 60 | ```c | 68 | ```c |
| 69 | +// 示例：对int8_t矢量数据寄存器执行按位异或，通过函数返回值返回结果。 | ||
| 70 | +__simd_callee__ inline vector_int8_t asc_xor(vector_int8_t src0, | ||
| 71 | + vector_int8_t src1, | ||
| 72 | + vector_bool mask) | ||
| 73 | + | ||
| 61 | // 示例：对int8_t矢量数据寄存器执行按位异或。 | 74 | // 示例：对int8_t矢量数据寄存器执行按位异或。 |
| 62 | __simd_callee__ inline void asc_xor(vector_int8_t& dst, | 75 | __simd_callee__ inline void asc_xor(vector_int8_t& dst, |
| 63 | vector_int8_t src0, | 76 | vector_int8_t src0, |
| @@ -68,6 +81,12 @@ __simd_callee__ inline void asc_xor(vector_int8_t& dst, | |||
| 68 | ### 掩码寄存器按位异或 | 81 | ### 掩码寄存器按位异或 |
| 69 | 82 | ||
| 70 | ```c | 83 | ```c |
| 84 | +// 通过函数返回值返回结果。 | ||
| 85 | +__simd_callee__ inline vector_bool asc_xor(vector_bool src0, | ||
| 86 | + vector_bool src1, | ||
| 87 | + vector_bool mask) | ||
| 88 | + | ||
| 89 | +// 通过引用参数输出结果。 | ||
| 71 | __simd_callee__ inline void asc_xor(vector_bool& dst, | 90 | __simd_callee__ inline void asc_xor(vector_bool& dst, |
| 72 | vector_bool src0, | 91 | vector_bool src0, |
| 73 | vector_bool src1, | 92 | vector_bool src1, |
| @@ -80,7 +99,7 @@ __simd_callee__ inline void asc_xor(vector_bool& dst, | |||
| 80 | 99 | ||
| 81 | | 参数名 | 输入/输出 | 描述 | | 100 | | 参数名 | 输入/输出 | 描述 | |
| 82 | | --- | --- | --- | | 101 | | --- | --- | --- | |
| 83 | -| dst | 输出 | 目的矢量数据寄存器，写入按位异或结果。`dtype`须与`src0`、`src1`一致。 | | 102 | +| dst | 输出 | 目的矢量数据寄存器，仅无返回值类型接口包含该参数，写入按位异或结果。`dtype`须与`src0`、`src1`一致。 | |
| 84 | | src0 | 输入 | 源矢量数据寄存器，参与按位异或的操作数。`dtype`须与`dst`一致。 | | 103 | | src0 | 输入 | 源矢量数据寄存器，参与按位异或的操作数。`dtype`须与`dst`一致。 | |
| 85 | | src1 | 输入 | 源矢量数据寄存器，参与按位异或的操作数。`dtype`须与`dst`一致。 | | 104 | | src1 | 输入 | 源矢量数据寄存器，参与按位异或的操作数。`dtype`须与`dst`一致。 | |
| 86 | | mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。`mask`未筛选的元素在输出中置零。 | | 105 | | mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。`mask`未筛选的元素在输出中置零。 | |
| @@ -89,7 +108,7 @@ __simd_callee__ inline void asc_xor(vector_bool& dst, | |||
| 89 | 108 | ||
| 90 | | 参数名 | 输入/输出 | 描述 | | 109 | | 参数名 | 输入/输出 | 描述 | |
| 91 | | --- | --- | --- | | 110 | | --- | --- | --- | |
| 92 | -| dst | 输出 | 目的掩码寄存器，写入按位异或结果。 | | 111 | +| dst | 输出 | 目的掩码寄存器，仅无返回值类型接口包含该参数，写入按位异或结果。 | |
| 93 | | src0 | 输入 | 源掩码寄存器，参与按位异或的操作数。 | | 112 | | src0 | 输入 | 源掩码寄存器，参与按位异或的操作数。 | |
| 94 | | src1 | 输入 | 源掩码寄存器，参与按位异或的操作数。 | | 113 | | src1 | 输入 | 源掩码寄存器，参与按位异或的操作数。 | |
| 95 | | mask | 输入 | 源掩码寄存器，用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。`mask`未筛选的元素在输出中置零。 | | 114 | | mask | 输入 | 源掩码寄存器，用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。`mask`未筛选的元素在输出中置零。 | |
| @@ -98,7 +117,7 @@ __simd_callee__ inline void asc_xor(vector_bool& dst, | |||
| 98 | 117 | ||
| 99 | ## 返回值说明 | 118 | ## 返回值说明 |
| 100 | 119 | ||
| 101 | -无 | 120 | +对于返回值类型接口，返回保存按位异或结果的矢量数据寄存器或掩码寄存器，数据类型与`src0`保持一致。 |
| 102 | 121 | ||
| 103 | ## 约束说明 | 122 | ## 约束说明 |
| 104 | 123 | ||
| @@ -26,10 +26,10 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将src中低半部分或高半部分的元素以扩充位宽的方式写入dst，支持： | 29 | +将src中低半部分或高半部分的元素以扩充位宽的方式写入dst或通过函数返回值返回，支持： |
| 30 | 30 | ||
| 31 | -- asc_unpack_lower：取src的低半部分（LOWER），每个元素高位填0扩充位宽后写入dst。 | 31 | +- asc_unpack_lower：取src的低半部分（LOWER），每个元素高位填0扩充位宽后写入dst或通过函数返回值返回。 |
| 32 | -- asc_unpack_upper：取src的高半部分（HIGHER），每个元素高位填0扩充位宽后写入dst。 | 32 | +- asc_unpack_upper：取src的高半部分（HIGHER），每个元素高位填0扩充位宽后写入dst或通过函数返回值返回。 |
| 33 | 33 | ||
| 34 | 不同数据类型的扩充方式： | 34 | 不同数据类型的扩充方式： |
| 35 | 35 | ||
| @@ -56,10 +56,15 @@ | |||
| 56 | ### 矢量数据寄存器解压缩 | 56 | ### 矢量数据寄存器解压缩 |
| 57 | 57 | ||
| 58 | ```c | 58 | ```c |
| 59 | -// 解压缩低位 | 59 | +// 解压缩低位，通过函数返回值返回结果。 |
| 60 | +__simd_callee__ inline vector_<dtype_dst> asc_unpack_lower(vector_<dtype_src> src) | ||
| 61 | +// 解压缩高位，通过函数返回值返回结果。 | ||
| 62 | +__simd_callee__ inline vector_<dtype_dst> asc_unpack_upper(vector_<dtype_src> src) | ||
| 63 | + | ||
| 64 | +// 解压缩低位，通过引用参数输出结果。 | ||
| 60 | __simd_callee__ inline void asc_unpack_lower(vector_<dtype_dst>& dst, | 65 | __simd_callee__ inline void asc_unpack_lower(vector_<dtype_dst>& dst, |
| 61 | vector_<dtype_src> src) | 66 | vector_<dtype_src> src) |
| 62 | -// 解压缩高位 | 67 | +// 解压缩高位，通过引用参数输出结果。 |
| 63 | __simd_callee__ inline void asc_unpack_upper(vector_<dtype_dst>& dst, | 68 | __simd_callee__ inline void asc_unpack_upper(vector_<dtype_dst>& dst, |
| 64 | vector_<dtype_src> src) | 69 | vector_<dtype_src> src) |
| 65 | ``` | 70 | ``` |
| @@ -77,6 +82,9 @@ dtype_src与dtype_dst支持的数据类型对如下： | |||
| 77 | #### 函数原型典型示例 | 82 | #### 函数原型典型示例 |
| 78 | 83 | ||
| 79 | ```c | 84 | ```c |
| 85 | +// 示例：将uint16_t矢量数据寄存器解压缩为uint32_t，通过函数返回值返回结果。 | ||
| 86 | +__simd_callee__ inline vector_uint32_t asc_unpack_lower(vector_uint16_t src) | ||
| 87 | + | ||
| 80 | // 示例：将uint16_t矢量数据寄存器解压缩为uint32_t。 | 88 | // 示例：将uint16_t矢量数据寄存器解压缩为uint32_t。 |
| 81 | __simd_callee__ inline void asc_unpack_lower(vector_uint32_t& dst, | 89 | __simd_callee__ inline void asc_unpack_lower(vector_uint32_t& dst, |
| 82 | vector_uint16_t src) | 90 | vector_uint16_t src) |
| @@ -85,10 +93,15 @@ __simd_callee__ inline void asc_unpack_lower(vector_uint32_t& dst, | |||
| 85 | ### 掩码寄存器解压缩 | 93 | ### 掩码寄存器解压缩 |
| 86 | 94 | ||
| 87 | ```c | 95 | ```c |
| 88 | -// 解压缩低位 | 96 | +// 解压缩低位，通过函数返回值返回结果。 |
| 97 | +__simd_callee__ inline vector_bool asc_unpack_lower(vector_bool src) | ||
| 98 | +// 解压缩高位，通过函数返回值返回结果。 | ||
| 99 | +__simd_callee__ inline vector_bool asc_unpack_upper(vector_bool src) | ||
| 100 | + | ||
| 101 | +// 解压缩低位，通过引用参数输出结果。 | ||
| 89 | __simd_callee__ inline void asc_unpack_lower(vector_bool& dst, | 102 | __simd_callee__ inline void asc_unpack_lower(vector_bool& dst, |
| 90 | vector_bool src) | 103 | vector_bool src) |
| 91 | -// 解压缩高位 | 104 | +// 解压缩高位，通过引用参数输出结果。 |
| 92 | __simd_callee__ inline void asc_unpack_upper(vector_bool& dst, | 105 | __simd_callee__ inline void asc_unpack_upper(vector_bool& dst, |
| 93 | vector_bool src) | 106 | vector_bool src) |
| 94 | ``` | 107 | ``` |
| @@ -99,21 +112,21 @@ __simd_callee__ inline void asc_unpack_upper(vector_bool& dst, | |||
| 99 | 112 | ||
| 100 | | 参数名 | 输入/输出 | 描述 | | 113 | | 参数名 | 输入/输出 | 描述 | |
| 101 | | --- | --- | --- | | 114 | | --- | --- | --- | |
| 102 | -| dst | 输出 | 目的矢量数据寄存器，读取src低半段或高半段元素，扩展位宽后写入dst全部有效位置，dst全部有效位置均被改写。 | | 115 | +| dst | 输出 | 目的矢量数据寄存器，仅无返回值类型接口包含该参数，读取src低半段或高半段元素，扩展位宽后写入dst全部有效位置，dst全部有效位置均被改写。 | |
| 103 | | src | 输入 | 源矢量数据寄存器，取src低半段或高半段元素扩展位宽写入dst。 | | 116 | | src | 输入 | 源矢量数据寄存器，取src低半段或高半段元素扩展位宽写入dst。 | |
| 104 | 117 | ||
| 105 | **表2** 参数说明（操作数为掩码寄存器） | 118 | **表2** 参数说明（操作数为掩码寄存器） |
| 106 | 119 | ||
| 107 | | 参数名 | 输入/输出 | 描述 | | 120 | | 参数名 | 输入/输出 | 描述 | |
| 108 | | --- | --- | --- | | 121 | | --- | --- | --- | |
| 109 | -| dst | 输出 | 目的掩码寄存器，读取src低半段或高半段元素，扩展位宽后写入dst全部有效位置，dst全部有效位置均被改写。 | | 122 | +| dst | 输出 | 目的掩码寄存器，仅无返回值类型接口包含该参数，读取src低半段或高半段元素，扩展位宽后写入dst全部有效位置，dst全部有效位置均被改写。 | |
| 110 | | src | 输入 | 源掩码寄存器，取src低半段或高半段元素扩展位宽写入dst。 | | 123 | | src | 输入 | 源掩码寄存器，取src低半段或高半段元素扩展位宽写入dst。 | |
| 111 | 124 | ||
| 112 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../../defs/type/data_type_definition.md)。 | 125 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../../defs/type/data_type_definition.md)。 |
| 113 | 126 | ||
| 114 | ## 返回值说明 | 127 | ## 返回值说明 |
| 115 | 128 | ||
| 116 | -无 | 129 | +对于返回值类型接口，返回保存解压缩结果的矢量数据寄存器或掩码寄存器，数据类型与`dst`保持一致。 |
| 117 | 130 | ||
| 118 | ## 约束说明 | 131 | ## 约束说明 |
| 119 | 132 | ||
| @@ -26,7 +26,7 @@ | |||
| 26 | 26 | ||
| 27 | ## 功能说明 | 27 | ## 功能说明 |
| 28 | 28 | ||
| 29 | -将dst中数据根据mask进行解压缩。解压缩方式：dst中第0个元素置为0，dst中的第i个元素等于mask中从第0个到第(i-1)个元素中1的数量。mask最高位被忽略，不参与统计。 | 29 | +将dst中数据根据mask进行解压缩，或将解压缩结果通过函数返回值返回。解压缩方式：dst中第0个元素置为0，dst中的第i个元素等于mask中从第0个到第(i-1)个元素中1的数量。mask最高位被忽略，不参与统计。 |
| 30 | 30 | ||
| 31 | 具体算法如图1所示，dst的首位为0。对于后续元素，与dst[i-1]对应的有效mask位为1时，dst[i]的值为dst[i-1] + 1；对应的有效mask位为0时，dst[i]的值为dst[i-1]。 | 31 | 具体算法如图1所示，dst的首位为0。对于后续元素，与dst[i-1]对应的有效mask位为1时，dst[i]的值为dst[i-1] + 1；对应的有效mask位为0时，dst[i]的值为dst[i-1]。 |
| 32 | 32 | ||
| @@ -39,16 +39,37 @@ | |||
| 39 | ## 函数原型 | 39 | ## 函数原型 |
| 40 | 40 | ||
| 41 | ```c | 41 | ```c |
| 42 | +// 通过函数返回值返回结果。 | ||
| 43 | +__simd_callee__ inline vector_uint8_t asc_unsqueeze_u8(vector_bool mask) | ||
| 44 | +__simd_callee__ inline vector_int8_t asc_unsqueeze_s8(vector_bool mask) | ||
| 45 | +__simd_callee__ inline vector_uint16_t asc_unsqueeze_u16(vector_bool mask) | ||
| 46 | +__simd_callee__ inline vector_int16_t asc_unsqueeze_s16(vector_bool mask) | ||
| 47 | +__simd_callee__ inline vector_uint32_t asc_unsqueeze_u32(vector_bool mask) | ||
| 48 | +__simd_callee__ inline vector_int32_t asc_unsqueeze_s32(vector_bool mask) | ||
| 49 | + | ||
| 50 | +// 通过引用参数输出结果。 | ||
| 42 | __simd_callee__ inline void asc_unsqueeze(vector_<dtype>& dst, | 51 | __simd_callee__ inline void asc_unsqueeze(vector_<dtype>& dst, |
| 43 | vector_bool mask) | 52 | vector_bool mask) |
| 44 | ``` | 53 | ``` |
| 45 | 54 | ||
| 46 | #### dtype支持数据类型 | 55 | #### dtype支持数据类型 |
| 47 | -dtype支持的数据类型：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t。 | 56 | +dtype支持的数据类型：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t。返回值类型接口通过函数名后缀区分数据类型，对应关系如下： |
| 57 | + | ||
| 58 | +| 函数名后缀 | 数据类型 | | ||
| 59 | +|---|---| | ||
| 60 | +| u8 | uint8_t | | ||
| 61 | +| s8 | int8_t | | ||
| 62 | +| u16 | uint16_t | | ||
| 63 | +| s16 | int16_t | | ||
| 64 | +| u32 | uint32_t | | ||
| 65 | +| s32 | int32_t | | ||
| 48 | 66 | ||
| 49 | #### 函数原型典型示例 | 67 | #### 函数原型典型示例 |
| 50 | 68 | ||
| 51 | ```c | 69 | ```c |
| 70 | +// 示例：对int8_t矢量数据寄存器执行解压缩，通过函数返回值返回结果。 | ||
| 71 | +__simd_callee__ inline vector_int8_t asc_unsqueeze_s8(vector_bool mask) | ||
| 72 | + | ||
| 52 | // 示例：对int8_t矢量数据寄存器执行解压缩。 | 73 | // 示例：对int8_t矢量数据寄存器执行解压缩。 |
| 53 | __simd_callee__ inline void asc_unsqueeze(vector_int8_t& dst, | 74 | __simd_callee__ inline void asc_unsqueeze(vector_int8_t& dst, |
| 54 | vector_bool mask) | 75 | vector_bool mask) |
| @@ -60,14 +81,14 @@ __simd_callee__ inline void asc_unsqueeze(vector_int8_t& dst, | |||
| 60 | 81 | ||
| 61 | | 参数名 | 输入/输出 | 描述 | | 82 | | 参数名 | 输入/输出 | 描述 | |
| 62 | | --- | --- | --- | | 83 | | --- | --- | --- | |
| 63 | -| dst | 输入/输出 | 源和目的矢量数据寄存器，接口在掩码前缀计数过程中写入计数结果。 | | 84 | +| dst | 输入/输出 | 源和目的矢量数据寄存器，仅无返回值类型接口包含该参数，接口在掩码前缀计数过程中写入计数结果。 | |
| 64 | | mask | 输入 | 源掩码寄存器，用于提供dst解压缩信息。 | | 85 | | mask | 输入 | 源掩码寄存器，用于提供dst解压缩信息。 | |
| 65 | 86 | ||
| 66 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../../defs/type/data_type_definition.md)。 | 87 | 矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../../defs/type/data_type_definition.md)。 |
| 67 | 88 | ||
| 68 | ## 返回值说明 | 89 | ## 返回值说明 |
| 69 | 90 | ||
| 70 | -无 | 91 | +对于返回值类型接口，返回保存解压缩结果的矢量数据寄存器，数据类型由函数名后缀决定（参见函数原型中的对应关系表）。 |
| 71 | 92 | ||
| 72 | ## 约束说明 | 93 | ## 约束说明 |
| 73 | 94 | ||