Star744
1035
代码
代码
Issues57
Pull Requests194
Actions
项目讨论
Wiki
分析
Star744
1035
已合并
c-api minimal set #5021
DianaBlessYou创建于 28 天前
c-api minimal set #5021
已合并
DianaBlessYou创建于 28 天前
从lijunjie147/asc-devkit_9351perf合入到cann/asc-devkitmaster
共 4 个文件变更+609-107
Mdocs/zh/api/SIMD-API/c_api/reg/compare_and_select/asc_eq.md+161-27
@@ -26,7 +26,7 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-eq（equal to），该接口用于逐元素比较两个源操作数大小，将比较结果（$src0_i == src1_i$）写入目的操作数dst对应比特位，如果比较结果为真，则对应比特位为1，否则为0。29+eq表示相等（equal to），该接口用于逐元素比较两个源操作数大小，将比较结果（$src0_i == src1_i$）写入目的操作数dst对应比特位，如果比较结果为真，则对应比特位为1，否则为0。
30 30 
31计算公式如下：31计算公式如下：
32 32 
@@ -34,18 +34,27 @@ $$
34dst_i = (src0_i == src1_i)34dst_i = (src0_i == src1_i)
35$$35$$
36 36 
37+本接口仅在AIV上生效，非AIV调用直接返回。
M
Mmunanhw28 天前

和约束说明重复

likedislike
38+ 
37## 函数原型39## 函数原型
38 40 
39-```cpp41+```c
40-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_int8_t src0, vector_int8_t src1, vector_bool mask)42+__simd_callee__ inline void asc_eq(vector_bool& dst,
41-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_uint8_t src0, vector_uint8_t src1, vector_bool mask)43+ vector_<dtype> src0,
42-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_int16_t src0, vector_int16_t src1, vector_bool mask)44+ vector_<dtype> src1,
43-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_uint16_t src0, vector_uint16_t src1, vector_bool mask)45+ vector_bool mask)
44-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_half src0, vector_half src1, vector_bool mask)46+```
45-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_bfloat16_t src0, vector_bfloat16_t src1, vector_bool mask)47+ 
46-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_int32_t src0, vector_int32_t src1, vector_bool mask)48+dtype取值为`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。
47-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_uint32_t src0, vector_uint32_t src1, vector_bool mask)49+ 
48-__simd_callee__ inline void asc_eq(vector_bool& dst, vector_float src0, vector_float src1, vector_bool mask)50+### 典型示例
51+ 
52+```c
53+// 典型示例：half类型逐元素相等比较
54+__simd_callee__ inline void asc_eq(vector_bool& dst,
55+ vector_half src0,
56+ vector_half src1,
57+ vector_bool mask)
49```58```
50 59 
51## 参数说明60## 参数说明
@@ -57,7 +66,7 @@ __simd_callee__ inline void asc_eq(vector_bool& dst, vector_float src0, vector_f
57| dst | 输出 | 目的操作数（掩码寄存器）。 |66| dst | 输出 | 目的操作数（掩码寄存器）。 |
58| src0 | 输入 | 源操作数（矢量数据寄存器）。 |67| src0 | 输入 | 源操作数（矢量数据寄存器）。 |
59| src1 | 输入 | 源操作数（矢量数据寄存器）。 |68| src1 | 输入 | 源操作数（矢量数据寄存器）。 |
60-| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。对应位置为1时参与计算，为0时不参与计算。mask未筛选的元素在输出中置零。 |69+| mask | 输入 | 源操作数掩码（掩码寄存器），用于指示在计算过程中哪些元素参与计算。mask中与元素对应的比特位为1时，该元素参与计算；为0时，该元素不参与计算。 |
61 70 
62矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。71矢量数据寄存器和掩码寄存器的详细说明请参见[reg数据类型定义](../reg_data_types/data_type_definition.md)。
63 72 
@@ -67,26 +76,151 @@ __simd_callee__ inline void asc_eq(vector_bool& dst, vector_float src0, vector_f
67 76 
68## 约束说明77## 约束说明
69 78 
70-- 通过mask参数控制的未选中元素在目的操作数中被置零。79+- 本接口仅在AIV上生效，非AIV调用直接返回。
71-- 操作数重叠约束：src0和src1可以是同一个矢量数据寄存器。80+- 本接口在Vector Function（`__simd_vf__`标记的函数）内调用，dst为掩码寄存器，src0、src1为矢量数据寄存器。
72-- 浮点数据类型比较，涉及`nan`时比较结果为0。81+- mask为`vector_bool`类型掩码寄存器，需通过掩码设置接口预先赋值后再传入；未赋值的掩码寄存器内容不确定，会导致有效元素位置错误。
73-- 当源操作数为浮点数据类型时，+0与-0视为相等。82+- src0与src1的dtype需要保持一致。
83+- mask比特位为0时，dst对应比特位写0。
84+- 浮点比较时，+0与-0视为相等。
85+- 浮点比较输入含nan时，dst对应比特位写0。
74 86 
75## 调用示例87## 调用示例
76 88 
89+将代码保存为example.asc后，可通过bisheng命令编译运行，其中--npu-arch参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。
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+ 
77```cpp99```cpp
78-__simd_vf__ inline void eq_vf(__ubuf__ uint16_t* dst_addr, __ubuf__ half* src0_addr, __ubuf__ half* src1_addr, uint32_t count, uint16_t one_repeat_size, uint16_t one_block_size, uint16_t repeat_time)100+#include <cstdint>
L
LLycheeeee28 天前

增加bisheng编译命令。

likedislike
101+#include <iostream>
102+#include <vector>
103+ 
104+#include "c_api/asc_simd.h"
105+#include "acl/acl.h"
106+ 
107+namespace {
108+template <typename T>
109+void print_data(const char* label, const std::vector<T>& values)
79{110{
80- vector_half src0;111+ std::cout << label << ":";
81- vector_half src1;112+ const size_t count = values.size() < 8 ? values.size() : 8;
82- vector_bool dst;113+ for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i];
83- vector_bool mask;114+ if (values.size() > count) std::cout << " ...";
84- for (uint16_t i = 0; i < repeat_time; ++i) {115+ std::cout << std::endl;
85- mask = asc_update_mask_b16(count);116+}
86- asc_loadalign_postupdate(src0, src0_addr, one_repeat_size);117+ 
87- asc_loadalign_postupdate(src1, src1_addr, one_repeat_size);118+template <typename T>
88- asc_eq(dst, src0, src1, mask);119+bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0)
89- asc_storealign_postupdate(dst_addr, dst, one_block_size);120+{
121+ if (actual.size() != expected.size()) return false;
122+ for (size_t i = 0; i < actual.size(); ++i) {
123+ if (actual[i] == expected[i]) continue;
124+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
125+ if (diff > tolerance || diff < -tolerance) return false;
90 }126 }
127+ return true;
128+}
129+ 
130+template <typename T>
131+bool compare_range_data(const std::vector<T>& actual, const std::vector<T>& expected,
132+ size_t begin, size_t count, double tolerance = 0.0)
133+{
134+ if (begin + count > actual.size() || begin + count > expected.size()) return false;
135+ for (size_t i = begin; i < begin + count; ++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 ELEMENT_COUNT = 64;
144+constexpr float COMPARE_VALUE = 3.0f;
145+ 
146+__simd_vf__ inline void compare(__ubuf__ float* dst, __ubuf__ float* src0, __ubuf__ float* src1)
147+{
148+ vector_float dst_reg;
149+ vector_float ones;
150+ vector_float zeros;
151+ vector_float src0_reg;
152+ vector_float src1_reg;
153+ vector_bool predicate;
154+ uint32_t count = ELEMENT_COUNT;
155+ vector_bool mask = asc_update_mask_b32(count);
156+ asc_loadalign(src0_reg, src0);
157+ asc_loadalign(src1_reg, src1);
158+ asc_eq(predicate, src0_reg, src1_reg, mask);
159+ asc_duplicate_scalar(ones, 1.0f, mask);
160+ asc_duplicate_scalar(zeros, 0.0f, mask);
161+ asc_select(dst_reg, ones, zeros, predicate);
162+ asc_storealign(dst, dst_reg, mask);
163+}
164+ 
165+__global__ __vector__ void asc_eq_kernel(__gm__ float* dst, __gm__ float* src0, __gm__ float* src1)
166+{
167+ asc_init();
168+ __ubuf__ float dst_local[ELEMENT_COUNT];
169+ __ubuf__ float src0_local[ELEMENT_COUNT];
170+ __ubuf__ float src1_local[ELEMENT_COUNT];
171+ asc_copy_gm2ub_align(src0_local, src0, ELEMENT_COUNT * sizeof(float));
172+ asc_copy_gm2ub_align(src1_local, src1, ELEMENT_COUNT * sizeof(float));
173+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
174+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
175+ compare(dst_local, src0_local, src1_local);
176+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
177+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
178+ asc_copy_ub2gm_align(dst, dst_local, ELEMENT_COUNT * sizeof(float));
179+ asc_sync();
180+}
181+} // namespace
182+ 
183+int main()
184+{
185+ std::vector<float> src0(ELEMENT_COUNT);
186+ std::vector<float> src1(ELEMENT_COUNT);
187+ std::vector<float> output(ELEMENT_COUNT, 0.0f);
188+ std::vector<float> golden(ELEMENT_COUNT);
189+ for (uint32_t i = 0; i < ELEMENT_COUNT; ++i) {
190+ src0[i] = static_cast<float>(i % 8);
191+ src1[i] = static_cast<float>((i / 2) % 8);
192+ golden[i] = src0[i] == src1[i] ? 1.0f : 0.0f;
193+ }
194+ aclInit(nullptr);
195+ aclrtSetDevice(0);
196+ float* src0_device = nullptr;
197+ aclrtMalloc(reinterpret_cast<void**>(&src0_device), (ELEMENT_COUNT) * sizeof(float),
198+ ACL_MEM_MALLOC_HUGE_FIRST);
199+ float* src1_device = nullptr;
200+ aclrtMalloc(reinterpret_cast<void**>(&src1_device), (ELEMENT_COUNT) * sizeof(float),
201+ ACL_MEM_MALLOC_HUGE_FIRST);
202+ float* dst_device = nullptr;
203+ aclrtMalloc(reinterpret_cast<void**>(&dst_device), (ELEMENT_COUNT) * sizeof(float),
204+ ACL_MEM_MALLOC_HUGE_FIRST);
205+ aclrtMemcpy(src0_device, src0.size() * sizeof(float), src0.data(), src0.size() * sizeof(float),
206+ ACL_MEMCPY_HOST_TO_DEVICE);
207+ aclrtMemcpy(src1_device, src1.size() * sizeof(float), src1.data(), src1.size() * sizeof(float),
208+ ACL_MEMCPY_HOST_TO_DEVICE);
209+ asc_eq_kernel<<<1, 0>>>(dst_device, src0_device, src1_device);
210+ aclrtSynchronizeDevice();
211+ aclrtMemcpy(output.data(), output.size() * sizeof(float), dst_device, output.size() * sizeof(float),
212+ ACL_MEMCPY_DEVICE_TO_HOST);
213+ print_data("Input src0", src0);
214+ print_data("Input src1", src1);
215+ print_data("Output", output);
216+ print_data("Golden", golden);
217+ const bool passed = compare_data(output, golden);
218+ std::cout << (passed ? "[Success] asc_eq passed." : "[Failed] asc_eq failed.") << std::endl;
219+ aclrtFree(dst_device);
220+ aclrtFree(src0_device);
221+ aclrtFree(src1_device);
222+ aclrtResetDevice(0);
223+ aclFinalize();
224+ return passed ? 0 : 1;
91}225}
92```226```
Mdocs/zh/api/SIMD-API/c_api/reg/logic_compute/asc_and.md+169-29
@@ -26,11 +26,11 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-该接口可用于对矢量数据寄存器或掩码寄存器操作：29+根据mask对源操作数src0、src1执行按位与（&）操作，将结果写入目的操作数dst。
30 30 
31-- 对矢量数据寄存器操作：根据mask对源操作数src0、src1执行按位与（&）操作，将结果写入目的操作数dst。31+- 矢量数据寄存器按位与：对两个矢量数据寄存器执行按位与（&），结果写入目的矢量数据寄存器。
32 32 
33-- 对掩码寄存器操作：根据mask对源操作数src0、src1的有效bit执行按位与（&）操作，将结果写入目的操作数dst。33+- 掩码寄存器按位与：对两个掩码寄存器执行按位与（&），结果写入目的掩码寄存器。
34 34 
35计算公式如下：35计算公式如下：
36 36 
@@ -38,28 +38,45 @@ $$
38dst_i = src0_i \& src1_i38dst_i = src0_i \& src1_i
39$$39$$
40 40 
41+本接口仅在AIV上生效，非AIV调用直接返回。
42+ 
41## 函数原型43## 函数原型
42 44 
43-```cpp45+### 矢量数据寄存器按位与
44-__simd_callee__ inline void asc_and(vector_bool& dst, vector_bool src0, vector_bool src1, vector_bool mask)46+ 
45-__simd_callee__ inline void asc_and(vector_int8_t& dst, vector_int8_t src0, vector_int8_t src1, vector_bool mask)47+```c
46-__simd_callee__ inline void asc_and(vector_uint8_t& dst, vector_uint8_t src0, vector_uint8_t src1, vector_bool mask)48+__simd_callee__ inline void asc_and(vector_<dtype>& dst,
47-__simd_callee__ inline void asc_and(vector_fp8_e5m2_t& dst, vector_fp8_e5m2_t src0, vector_fp8_e5m2_t src1, vector_bool mask)49+ vector_<dtype> src0,
48-__simd_callee__ inline void asc_and(vector_fp8_e4m3fn_t& dst, vector_fp8_e4m3fn_t src0, vector_fp8_e4m3fn_t src1, vector_bool mask)50+ vector_<dtype> src1,
49-__simd_callee__ inline void asc_and(vector_int16_t& dst, vector_int16_t src0, vector_int16_t src1, vector_bool mask)51+ vector_bool mask)
M
Mmunanhw28 天前

多了一个空行 排查全文

likedislike
DianaBlessYou
DianaBlessYou
27 天前 评论：
50-__simd_callee__ inline void asc_and(vector_uint16_t& dst, vector_uint16_t src0, vector_uint16_t src1, vector_bool mask)52+```
51-__simd_callee__ inline void asc_and(vector_half& dst, vector_half src0, vector_half src1, vector_bool mask)53+ 
52-__simd_callee__ inline void asc_and(vector_bfloat16_t& dst, vector_bfloat16_t src0, vector_bfloat16_t src1, vector_bool mask)54+dtype取值为`int8_t`、`uint8_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。
53-__simd_callee__ inline void asc_and(vector_int32_t& dst, vector_int32_t src0, vector_int32_t src1, vector_bool mask)55+ 
54-__simd_callee__ inline void asc_and(vector_uint32_t& dst, vector_uint32_t src0, vector_uint32_t src1, vector_bool mask)56+#### 典型示例
55-__simd_callee__ inline void asc_and(vector_float& dst, vector_float src0, vector_float src1, vector_bool mask)57+ 
58+```c
59+// 示例：对half矢量数据寄存器执行按位与
60+__simd_callee__ inline void asc_and(vector_half& dst,
61+ vector_half src0,
62+ vector_half src1,
63+ vector_bool mask)
64+```
65+ 
66+### 掩码寄存器按位与
67+ 
68+```c
69+__simd_callee__ inline void asc_and(vector_bool& dst,
70+ vector_bool src0,
71+ vector_bool src1,
72+ vector_bool mask)
56```73```
57 74 
58## 参数说明75## 参数说明
59 76 
60**表1** 参数说明77**表1** 参数说明
61 78 
62-| 参数名 | 输入/输出 | 描述 |79+| 参数名 | 输入/输出 | 描述 |
63| :----- | :------- | :------- |80| :----- | :------- | :------- |
64| dst | 输出 | 目的操作数（矢量数据寄存器或掩码寄存器）。 |81| dst | 输出 | 目的操作数（矢量数据寄存器或掩码寄存器）。 |
65| src0 | 输入 | 源操作数（矢量数据寄存器或掩码寄存器）。 |82| src0 | 输入 | 源操作数（矢量数据寄存器或掩码寄存器）。 |
@@ -74,23 +91,146 @@ __simd_callee__ inline void asc_and(vector_float& dst, vector_float src0, vector
74 91 
75## 约束说明92## 约束说明
76 93 
77-无94+- 本接口仅在AIV上生效，非AIV调用直接返回。
95+- mask必须为`vector_bool`类型掩码寄存器，需通过掩码设置接口预先赋值后再传入；未赋值的掩码寄存器内容不确定，会导致有效元素位置错误。
96+- 参与计算的元素个数由矢量长度（VL）决定：
97+ - 矢量数据寄存器按位与中元素个数 = VL ÷ sizeof(dtype)；
98+ - 掩码寄存器按位与中比特个数 = VL。
99+- src0、src1、dst的dtype需要保持一致。
100+- mask比特位为0时，dst对应比特位写0。
78 101 
79## 调用示例102## 调用示例
80 103 
104+将代码保存为example.asc后，可通过bisheng命令编译运行，其中--npu-arch参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。
105+ 
106+<!-- npu="950" id8 -->
107+以Ascend 950PR/Ascend 950DT产品（对应NPU架构为dav-3510）为例，编译运行命令如下：
108+ 
109+```bash
110+bisheng example.asc -o main --npu-arch=dav-3510; ./main
111+```
112+<!-- end id8 -->
113+ 
81```cpp114```cpp
82-__simd_vf__ inline void and_vf(__ubuf__ half* dst_addr, __ubuf__ half* src0_addr, __ubuf__ half* src1_addr, uint32_t count, uint16_t one_repeat_size, uint16_t repeat_time)115+#include <cstdint>
116+#include <iostream>
117+#include <vector>
118+ 
119+#include "c_api/asc_simd.h"
120+#include "acl/acl.h"
121+ 
122+namespace {
123+template <typename T>
124+void print_data(const char* label, const std::vector<T>& values)
83{125{
84- vector_half dst;126+ std::cout << label << ":";
85- vector_half src0;127+ const size_t count = values.size() < 8 ? values.size() : 8;
86- vector_half src1;128+ for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i];
87- vector_bool mask;129+ if (values.size() > count) std::cout << " ...";
88- for (uint16_t i = 0; i < repeat_time; ++i) {130+ std::cout << std::endl;
89- mask = asc_update_mask_b16(count);131+}
90- asc_loadalign(src0, src0_addr + i * one_repeat_size);132+ 
91- asc_loadalign(src1, src1_addr + i * one_repeat_size);133+template <typename T>
92- asc_and(dst, src0, src1, mask);134+bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0)
93- asc_storealign(dst_addr + i * one_repeat_size, dst, mask);135+{
136+ if (actual.size() != expected.size()) return false;
137+ for (size_t i = 0; i < actual.size(); ++i) {
138+ if (actual[i] == expected[i]) continue;
139+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
140+ if (diff > tolerance || diff < -tolerance) return false;
94 }141 }
142+ return true;
143+}
144+ 
145+template <typename T>
146+bool compare_range_data(const std::vector<T>& actual, const std::vector<T>& expected,
147+ size_t begin, size_t count, double tolerance = 0.0)
148+{
149+ if (begin + count > actual.size() || begin + count > expected.size()) return false;
150+ for (size_t i = begin; i < begin + count; ++i) {
151+ if (actual[i] == expected[i]) continue;
152+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
153+ if (diff > tolerance || diff < -tolerance) return false;
154+ }
155+ return true;
156+}
157+ 
158+constexpr uint32_t ELEMENT_COUNT = 64;
159+ 
160+__simd_vf__ inline void compute(__ubuf__ int32_t* dst, __ubuf__ int32_t* src0, __ubuf__ int32_t* src1)
161+{
162+ vector_int32_t dst_reg;
163+ vector_int32_t src0_reg;
164+ vector_int32_t src1_reg;
165+ uint32_t count = ELEMENT_COUNT;
166+ vector_bool mask = asc_update_mask_b32(count);
167+ asc_loadalign(src0_reg, src0);
168+ asc_loadalign(src1_reg, src1);
169+ asc_and(dst_reg, src0_reg, src1_reg, mask);
170+ asc_storealign(dst, dst_reg, mask);
171+}
172+ 
173+__global__ __vector__ void asc_and_kernel(__gm__ int32_t* dst, __gm__ int32_t* src0, __gm__ int32_t* src1)
174+{
175+ asc_init();
176+ __ubuf__ int32_t dst_local[ELEMENT_COUNT];
177+ __ubuf__ int32_t src0_local[ELEMENT_COUNT];
178+ __ubuf__ int32_t src1_local[ELEMENT_COUNT];
179+ asc_copy_gm2ub_align(src0_local, src0, ELEMENT_COUNT * sizeof(int32_t));
180+ asc_copy_gm2ub_align(src1_local, src1, ELEMENT_COUNT * sizeof(int32_t));
181+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
182+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
183+ compute(dst_local, src0_local, src1_local);
184+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
185+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
186+ asc_copy_ub2gm_align(dst, dst_local, ELEMENT_COUNT * sizeof(int32_t));
187+ asc_sync();
188+}
189+} // namespace
190+ 
191+int main()
192+{
193+ std::vector<int32_t> src0(ELEMENT_COUNT);
194+ std::vector<int32_t> src1(ELEMENT_COUNT);
195+ std::vector<int32_t> output(ELEMENT_COUNT, 0);
196+ std::vector<int32_t> golden(ELEMENT_COUNT);
197+ for (uint32_t i = 0; i < ELEMENT_COUNT; ++i) {
198+ src0[i] = static_cast<int32_t>(0x101010U + i);
199+ src1[i] = static_cast<int32_t>(0x0f0f0000U + i);
200+ golden[i] = src0[i] & src1[i];
201+ }
202+ 
203+ aclInit(nullptr);
204+ aclrtSetDevice(0);
205+ int32_t* src0_device = nullptr;
206+ aclrtMalloc(reinterpret_cast<void**>(&src0_device), (ELEMENT_COUNT) * sizeof(int32_t),
207+ ACL_MEM_MALLOC_HUGE_FIRST);
208+ int32_t* src1_device = nullptr;
209+ aclrtMalloc(reinterpret_cast<void**>(&src1_device), (ELEMENT_COUNT) * sizeof(int32_t),
210+ ACL_MEM_MALLOC_HUGE_FIRST);
211+ int32_t* dst_device = nullptr;
212+ aclrtMalloc(reinterpret_cast<void**>(&dst_device), (ELEMENT_COUNT) * sizeof(int32_t),
213+ ACL_MEM_MALLOC_HUGE_FIRST);
214+ aclrtMemcpy(src0_device, src0.size() * sizeof(int32_t), src0.data(), src0.size() * sizeof(int32_t),
215+ ACL_MEMCPY_HOST_TO_DEVICE);
216+ aclrtMemcpy(src1_device, src1.size() * sizeof(int32_t), src1.data(), src1.size() * sizeof(int32_t),
217+ ACL_MEMCPY_HOST_TO_DEVICE);
218+ 
219+ asc_and_kernel<<<1, 0>>>(dst_device, src0_device, src1_device);
220+ aclrtSynchronizeDevice();
221+ aclrtMemcpy(output.data(), output.size() * sizeof(int32_t), dst_device, output.size() * sizeof(int32_t),
222+ ACL_MEMCPY_DEVICE_TO_HOST);
223+ print_data("Input src0", src0);
224+ print_data("Input src1", src1);
225+ print_data("Output", output);
226+ print_data("Golden", golden);
227+ const bool passed = compare_data(output, golden);
228+ std::cout << (passed ? "[Success] asc_and passed." : "[Failed] asc_and failed.") << std::endl;
229+ aclrtFree(dst_device);
230+ aclrtFree(src0_device);
231+ aclrtFree(src1_device);
232+ aclrtResetDevice(0);
233+ aclFinalize();
234+ return passed ? 0 : 1;
95}235}
96```236```
Mdocs/zh/api/SIMD-API/c_api/reg/logic_compute/asc_or.md+168-25
@@ -26,11 +26,11 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-该接口可用于对矢量数据寄存器或掩码寄存器操作：29+根据mask对源操作数src0、src1执行按位或（|）操作，将结果写入目的操作数dst。
30 30 
31-- 对矢量数据寄存器操作：根据mask对源操作数src0、src1执行按位或（|）操作，将结果写入目的操作数dst。31+- 矢量数据寄存器按位或：对两个矢量数据寄存器执行按位或（|），结果写入目的矢量数据寄存器。
32 32 
33-- 对掩码寄存器操作：根据mask对源操作数src0、src1的有效bit执行按位或（|）操作，将结果写入目的操作数dst。33+- 掩码寄存器按位或：对两个掩码寄存器执行按位或（|），结果写入目的掩码寄存器。
34 34 
35计算公式如下：35计算公式如下：
36 36 
@@ -38,18 +38,38 @@ $$
38dst_i = src0_i | src1_i38dst_i = src0_i | src1_i
39$$39$$
40 40 
41+本接口仅在AIV上生效，非AIV调用直接返回。
42+ 
41## 函数原型43## 函数原型
42 44 
43-```cpp45+### 矢量数据寄存器按位或
44-__simd_callee__ inline void asc_or(vector_bool& dst, vector_bool src0, vector_bool src1, vector_bool mask)46+ 
45-__simd_callee__ inline void asc_or(vector_int8_t& dst, vector_int8_t src0, vector_int8_t src1, vector_bool mask)47+```c
46-__simd_callee__ inline void asc_or(vector_uint8_t& dst, vector_uint8_t src0, vector_uint8_t src1, vector_bool mask)48+__simd_callee__ inline void asc_or(vector_<dtype>& dst,
47-__simd_callee__ inline void asc_or(vector_int16_t& dst, vector_int16_t src0, vector_int16_t src1, vector_bool mask)49+ vector_<dtype> src0,
48-__simd_callee__ inline void asc_or(vector_uint16_t& dst, vector_uint16_t src0, vector_uint16_t src1, vector_bool mask)50+ vector_<dtype> src1,
49-__simd_callee__ inline void asc_or(vector_half& dst, vector_half src0, vector_half src1, vector_bool mask)51+ vector_bool mask)
50-__simd_callee__ inline void asc_or(vector_int32_t& dst, vector_int32_t src0, vector_int32_t src1, vector_bool mask)52+```
51-__simd_callee__ inline void asc_or(vector_uint32_t& dst, vector_uint32_t src0, vector_uint32_t src1, vector_bool mask)53+ 
52-__simd_callee__ inline void asc_or(vector_float& dst, vector_float src0, vector_float src1, vector_bool mask)54+dtype取值为`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`half`、`int32_t`、`uint32_t`、`float`。
55+ 
56+#### 典型示例
57+ 
58+```c
59+// 示例：对half矢量数据寄存器执行按位或
60+__simd_callee__ inline void asc_or(vector_half& dst,
61+ vector_half src0,
62+ vector_half src1,
63+ vector_bool mask)
64+```
65+ 
66+### 掩码寄存器按位或
67+ 
68+```c
69+__simd_callee__ inline void asc_or(vector_bool& dst,
70+ vector_bool src0,
71+ vector_bool src1,
72+ vector_bool mask)
53```73```
54 74 
55## 参数说明75## 参数说明
@@ -71,23 +91,146 @@ __simd_callee__ inline void asc_or(vector_float& dst, vector_float src0, vector_
71 91 
72## 约束说明92## 约束说明
73 93 
74-无94+- 本接口仅在AIV上生效，非AIV调用直接返回。
95+- mask必须为`vector_bool`类型掩码寄存器，需通过掩码设置接口预先赋值后再传入；未赋值的掩码寄存器内容不确定，会导致有效元素位置错误。
96+- 参与计算的元素个数由矢量长度（VL）决定：
97+ - 矢量数据寄存器按位或中元素个数 = VL ÷ sizeof(dtype)；
98+ - 掩码寄存器按位或中比特个数 = VL。
99+- src0、src1、dst的dtype需要保持一致。
100+- mask比特位为0时，dst对应比特位写0。
75 101 
76## 调用示例102## 调用示例
77 103 
104+将代码保存为example.asc后，可通过bisheng命令编译运行，其中--npu-arch参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。
105+ 
106+<!-- npu="950" id8 -->
107+以Ascend 950PR/Ascend 950DT产品（对应NPU架构为dav-3510）为例，编译运行命令如下：
108+ 
109+```bash
110+bisheng example.asc -o main --npu-arch=dav-3510; ./main
111+```
112+<!-- end id8 -->
113+ 
78```cpp114```cpp
79-__simd_vf__ inline void or_vf(__ubuf__ half* dst_addr, __ubuf__ half* src0_addr, __ubuf__ half* src1_addr, uint32_t count, uint16_t one_repeat_size, uint16_t repeat_time)115+#include <cstdint>
116+#include <iostream>
117+#include <vector>
118+ 
119+#include "c_api/asc_simd.h"
120+#include "acl/acl.h"
121+ 
122+namespace {
123+template <typename T>
124+void print_data(const char* label, const std::vector<T>& values)
80{125{
81- vector_half dst;126+ std::cout << label << ":";
82- vector_half src0;127+ const size_t count = values.size() < 8 ? values.size() : 8;
83- vector_half src1;128+ for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i];
84- vector_bool mask;129+ if (values.size() > count) std::cout << " ...";
85- for (uint16_t i = 0; i < repeat_time; ++i) {130+ std::cout << std::endl;
86- mask = asc_update_mask_b16(count);131+}
87- asc_loadalign(src0, src0_addr + i * one_repeat_size);132+ 
88- asc_loadalign(src1, src1_addr + i * one_repeat_size);133+template <typename T>
89- asc_or(dst, src0, src1, mask);134+bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0)
90- asc_storealign(dst_addr + i * one_repeat_size, dst, mask);135+{
136+ if (actual.size() != expected.size()) return false;
137+ for (size_t i = 0; i < actual.size(); ++i) {
138+ if (actual[i] == expected[i]) continue;
139+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
140+ if (diff > tolerance || diff < -tolerance) return false;
91 }141 }
142+ return true;
143+}
144+ 
145+template <typename T>
146+bool compare_range_data(const std::vector<T>& actual, const std::vector<T>& expected,
147+ size_t begin, size_t count, double tolerance = 0.0)
148+{
149+ if (begin + count > actual.size() || begin + count > expected.size()) return false;
150+ for (size_t i = begin; i < begin + count; ++i) {
151+ if (actual[i] == expected[i]) continue;
152+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
153+ if (diff > tolerance || diff < -tolerance) return false;
154+ }
155+ return true;
156+}
157+ 
158+constexpr uint32_t ELEMENT_COUNT = 64;
159+ 
160+__simd_vf__ inline void compute(__ubuf__ int32_t* dst, __ubuf__ int32_t* src0, __ubuf__ int32_t* src1)
161+{
162+ vector_int32_t dst_reg;
163+ vector_int32_t src0_reg;
164+ vector_int32_t src1_reg;
165+ uint32_t count = ELEMENT_COUNT;
166+ vector_bool mask = asc_update_mask_b32(count);
167+ asc_loadalign(src0_reg, src0);
168+ asc_loadalign(src1_reg, src1);
169+ asc_or(dst_reg, src0_reg, src1_reg, mask);
170+ asc_storealign(dst, dst_reg, mask);
171+}
172+ 
173+__global__ __vector__ void asc_or_kernel(__gm__ int32_t* dst, __gm__ int32_t* src0, __gm__ int32_t* src1)
174+{
175+ asc_init();
176+ __ubuf__ int32_t dst_local[ELEMENT_COUNT];
177+ __ubuf__ int32_t src0_local[ELEMENT_COUNT];
178+ __ubuf__ int32_t src1_local[ELEMENT_COUNT];
179+ asc_copy_gm2ub_align(src0_local, src0, ELEMENT_COUNT * sizeof(int32_t));
180+ asc_copy_gm2ub_align(src1_local, src1, ELEMENT_COUNT * sizeof(int32_t));
181+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
182+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
183+ compute(dst_local, src0_local, src1_local);
184+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
185+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
186+ asc_copy_ub2gm_align(dst, dst_local, ELEMENT_COUNT * sizeof(int32_t));
187+ asc_sync();
188+}
189+} // namespace
190+ 
191+int main()
192+{
193+ std::vector<int32_t> src0(ELEMENT_COUNT);
194+ std::vector<int32_t> src1(ELEMENT_COUNT);
195+ std::vector<int32_t> output(ELEMENT_COUNT, 0);
196+ std::vector<int32_t> golden(ELEMENT_COUNT);
197+ for (uint32_t i = 0; i < ELEMENT_COUNT; ++i) {
198+ src0[i] = static_cast<int32_t>(0x101010U + i);
199+ src1[i] = static_cast<int32_t>(0x0f0f0000U + i);
200+ golden[i] = src0[i] | src1[i];
201+ }
202+ 
203+ aclInit(nullptr);
204+ aclrtSetDevice(0);
205+ int32_t* src0_device = nullptr;
206+ aclrtMalloc(reinterpret_cast<void**>(&src0_device), (ELEMENT_COUNT) * sizeof(int32_t),
207+ ACL_MEM_MALLOC_HUGE_FIRST);
208+ int32_t* src1_device = nullptr;
209+ aclrtMalloc(reinterpret_cast<void**>(&src1_device), (ELEMENT_COUNT) * sizeof(int32_t),
210+ ACL_MEM_MALLOC_HUGE_FIRST);
211+ int32_t* dst_device = nullptr;
212+ aclrtMalloc(reinterpret_cast<void**>(&dst_device), (ELEMENT_COUNT) * sizeof(int32_t),
213+ ACL_MEM_MALLOC_HUGE_FIRST);
214+ aclrtMemcpy(src0_device, src0.size() * sizeof(int32_t), src0.data(), src0.size() * sizeof(int32_t),
215+ ACL_MEMCPY_HOST_TO_DEVICE);
216+ aclrtMemcpy(src1_device, src1.size() * sizeof(int32_t), src1.data(), src1.size() * sizeof(int32_t),
217+ ACL_MEMCPY_HOST_TO_DEVICE);
218+ 
219+ asc_or_kernel<<<1, 0>>>(dst_device, src0_device, src1_device);
220+ aclrtSynchronizeDevice();
221+ aclrtMemcpy(output.data(), output.size() * sizeof(int32_t), dst_device, output.size() * sizeof(int32_t),
222+ ACL_MEMCPY_DEVICE_TO_HOST);
223+ print_data("Input src0", src0);
224+ print_data("Input src1", src1);
225+ print_data("Output", output);
226+ print_data("Golden", golden);
227+ const bool passed = compare_data(output, golden);
228+ std::cout << (passed ? "[Success] asc_or passed." : "[Failed] asc_or failed.") << std::endl;
229+ aclrtFree(dst_device);
230+ aclrtFree(src0_device);
231+ aclrtFree(src1_device);
232+ aclrtResetDevice(0);
233+ aclFinalize();
234+ return passed ? 0 : 1;
92}235}
93```236```
Mdocs/zh/api/SIMD-API/c_api/scalar_compute/asc_store_dev.md+111-26
@@ -31,49 +31,134 @@
31 31 
32## 函数原型32## 函数原型
33 33 
34-```cpp34+```c
35-__aicore__ inline void asc_store_dev(__gm__ int8_t* addr, int8_t value)35+__aicore__ inline void asc_store_dev(__gm__ <dtype>* addr,
36+ <dtype> value)
37+```
36 38 
37-__aicore__ inline void asc_store_dev(__gm__ uint8_t* addr, uint8_t value)39+dtype取值为`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`int32_t`、`uint32_t`、`int64_t`、`uint64_t`。
38 40 
39-__aicore__ inline void asc_store_dev(__gm__ int16_t* addr, int16_t value)41+### 典型示例
40 42 
41-__aicore__ inline void asc_store_dev(__gm__ uint16_t* addr, uint16_t value)43+```c
42- 44+// 示例：写入32bit无符号整数
43-__aicore__ inline void asc_store_dev(__gm__ int32_t* addr, int32_t value)45+__aicore__ inline void asc_store_dev(__gm__ uint32_t* addr,
44- 46+ uint32_t value)
45-__aicore__ inline void asc_store_dev(__gm__ uint32_t* addr, uint32_t value)
46- 
47-__aicore__ inline void asc_store_dev(__gm__ int64_t* addr, int64_t value)
48- 
49-__aicore__ inline void asc_store_dev(__gm__ uint64_t* addr, uint64_t value)
50```47```
51 48 
52## 参数说明49## 参数说明
53 50 
54**表1** 参数说明51**表1** 参数说明
55 52 
56-|参数名|输入/输出|描述|53+| 参数名 | 输入/输出 | 描述 |
57-|------------|------------|-----------|54+| ------------ | ------------ | ----------- |
58-| addr | 输出 | 目标GM地址。|55+| addr | 输出 | 目标GM地址。 |
59-| value | 输入 | 待写入目标的数据。|56+| value | 输入 | 待写入目标的数据。 |
60 57 
61## 返回值说明58## 返回值说明
62 59 
63无60无
64 61 
65-## 流水类型
66- 
67-PIPE_S
68- 
69## 约束说明62## 约束说明
70 63 
71-无64+- `addr`起始地址须按写入dtype字节数对齐。
65+- `addr`须落在GM可访问地址空间内。
66+- 本接口运行在标量流水上，与后续依赖该写入结果的指令之间存在标量数据依赖；如后续有读取同一GM地址的指令，须通过同步指令建立依赖顺序，标量流水本身的顺序执行不保证跨指令访存可见性。
67+- 本接口访问GM时绕过DCache，不维护缓存一致性。若其他核或其他通路通过缓存访问同一GM地址，调用方需使用[asc_dcci](../cache_ctrl/asc_dcci.md)清理或失效对应Cache Line，并使用[asc_sync_data_barrier](../sync/asc_sync_data_barrier.md)保证相关访存操作的执行顺序和数据可见性。详情可参考[Scalar原子操作与DCache一致性](../../../../guide/编程指南/高级编程/内存模型/缓存一致性.md#scalar原子操作与dcache一致性)。
72 68 
73## 调用示例69## 调用示例
74 70 
75-```cpp71+将代码保存为example.asc后，可通过bisheng命令编译运行，其中--npu-arch参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。
76-// addr是外部输入的GM地址，value是待写入GM内存的数据，类型为 int32_t72+ 
77-int32_t value = 2;73+<!-- npu="950" id8 -->
78-asc_store_dev(addr, value);74+以Ascend 950PR/Ascend 950DT产品（对应NPU架构为dav-3510）为例，编译运行命令如下：
75+ 
76+```bash
77+bisheng example.asc -o main --npu-arch=dav-3510; ./main
78+```
79+<!-- end id8 -->
80+ 
81+```cpp
82+#include <cstdint>
83+#include <iostream>
84+#include <vector>
85+ 
86+#include "c_api/asc_simd.h"
87+#include "acl/acl.h"
88+ 
89+namespace {
90+template <typename T>
91+void print_data(const char* label, const std::vector<T>& values)
92+{
93+ std::cout << label << ":";
94+ const size_t count = values.size() < 8 ? values.size() : 8;
95+ for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i];
96+ if (values.size() > count) std::cout << " ...";
97+ std::cout << std::endl;
98+}
99+ 
100+template <typename T>
101+bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0)
102+{
103+ if (actual.size() != expected.size()) return false;
104+ for (size_t i = 0; i < actual.size(); ++i) {
105+ if (actual[i] == expected[i]) continue;
106+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
107+ if (diff > tolerance || diff < -tolerance) return false;
108+ }
109+ return true;
110+}
111+ 
112+template <typename T>
113+bool compare_range_data(const std::vector<T>& actual, const std::vector<T>& expected,
114+ size_t begin, size_t count, double tolerance = 0.0)
115+{
116+ if (begin + count > actual.size() || begin + count > expected.size()) return false;
117+ for (size_t i = begin; i < begin + count; ++i) {
118+ if (actual[i] == expected[i]) continue;
119+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
120+ if (diff > tolerance || diff < -tolerance) return false;
121+ }
122+ return true;
123+}
124+ 
125+constexpr uint32_t ELEMENTS = 8;
126+ 
127+__global__ __vector__ void asc_store_dev_kernel(__gm__ int64_t* output)
128+{
129+ asc_init();
130+ asc_sync_data_barrier(mem_dsb_t::DSB_ALL);
131+ asc_store_dev(output, static_cast<int64_t>(42));
132+ asc_sync();
133+}
134+} // namespace
135+ 
136+int main()
137+{
138+ std::vector<int64_t> input = {0};
139+ std::vector<int64_t> golden = {42};
140+ input.resize(ELEMENTS, 0);
141+ golden.resize(ELEMENTS, 0);
142+ std::vector<int64_t> output(ELEMENTS, -1);
143+ aclInit(nullptr);
144+ aclrtSetDevice(0);
145+ int64_t* output_device = nullptr;
146+ aclrtMalloc(reinterpret_cast<void**>(&output_device), (ELEMENTS) * sizeof(int64_t),
147+ ACL_MEM_MALLOC_HUGE_FIRST);
148+ aclrtMemcpy(output_device, input.size() * sizeof(int64_t), input.data(), input.size() * sizeof(int64_t),
149+ ACL_MEMCPY_HOST_TO_DEVICE);
150+ asc_store_dev_kernel<<<1, 0>>>(output_device);
151+ aclrtSynchronizeDevice();
152+ aclrtMemcpy(output.data(), output.size() * sizeof(int64_t), output_device, output.size() * sizeof(int64_t),
153+ ACL_MEMCPY_DEVICE_TO_HOST);
154+ print_data("Input", input);
155+ print_data("Output", output);
156+ print_data("Golden", golden);
157+ const bool passed = compare_data(output, golden);
158+ std::cout << (passed ? "[Success] asc_store_dev passed." : "[Failed] asc_store_dev failed.") << std::endl;
159+ aclrtFree(output_device);
160+ aclrtResetDevice(0);
161+ aclFinalize();
162+ return passed ? 0 : 1;
163+}
79```164```