refactor: use asc_uintdiv SIMT API in integer fast div sample-asc-devkit-AtomGit

已合并

refactor: use asc_uintdiv SIMT API in integer fast div sample #6277

zhangshunwei2创建于 22 天前从zhangshunwei2/asc-devkit： fastdiv_example_useapi 合入到 cann/asc-devkit： master

refactor: use asc_uintdiv SIMT API in integer fast div sample #6277

已合并

zhangshunwei2创建于 22 天前

从zhangshunwei2/asc-devkit fastdiv_example_useapi合入到cann/asc-devkit master

共 4 个文件变更+69-95

Mdocs/zh/guide/operator_practice/simd_simt_hybrid_optimization/instruction_optimization/fast_integer_division.md+33-19

	@@ -9,7 +9,7 @@
9		9
10	【优先级】低	10	【优先级】低
11		11
12	-【描述】整数除法指令在AI处理器上相比乘法和移位，执行开销显著更高。当除数固定时，可以使用基于乘法和移位的快速除法算法，在SIMD中根据除数预计算乘法魔数magic和移位量shift，在SIMT并行计算中使用乘法和移位操作，得到与普通整数除法一致的结果，从而降低每次除法的执行开销。	12	+【描述】整数除法指令在AI处理器上相比乘法和移位，执行开销显著更高。当除数固定时，可以使用基于乘法和移位的快速除法算法，在SIMD标量计算逻辑中调用[asc_get_uintdiv_magic_and_shift](../../../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md)接口根据除数预计算乘法魔数magic和移位量shift，在SIMT并行计算中调用[asc_uintdiv](../../../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md)接口使用乘法和移位操作，得到与普通整数除法一致的结果，从而降低每次除法的执行开销。
13		13
14	【样例介绍】以uint32类型整数除法算子为例，算子功能为将输入input的每个元素除以固定除数divisor，并将结果写入输出output。样例对比在SIMT中直接使用普通除法，以及针对固定除数预计算magic和shift后使用乘法和移位替代普通除法的性能差异。完整的算子实现代码请参考[SIMD与SIMT混合编程实现快速除法算子样例](../../../../../../examples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div)。	14	【样例介绍】以uint32类型整数除法算子为例，算子功能为将输入input的每个元素除以固定除数divisor，并将结果写入输出output。样例对比在SIMT中直接使用普通除法，以及针对固定除数预计算magic和shift后使用乘法和移位替代普通除法的性能差异。完整的算子实现代码请参考[SIMD与SIMT混合编程实现快速除法算子样例](../../../../../../examples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div)。
15		15
	@@ -54,7 +54,7 @@ __simt_vf__ __launch_bounds__(THREAD_LIMIT) inline void simt_normal_div(
54		54
55	【正例】	55	【正例】
56		56
57	-基于SIMD与SIMT混合编程的快速除法实现：对应样例中的场景1（SCENARIO\_NUM=1）。当除数固定时，可以在SIMD编程中预先计算所需的magic和shift，再在SIMT编程中使用乘法和移位组合运算的方式替代普通除法。	57	+基于SIMD与SIMT混合编程的快速除法实现：对应样例中的场景1（SCENARIO\_NUM=1）。当除数固定时，可以在SIMD标量计算逻辑中调用快除预计算接口[asc_get_uintdiv_magic_and_shift](../../../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md)计算所需的magic和shift，再在SIMT计算中调用快除接口[asc_uintdiv](../../../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md)以乘法和移位组合运算的方式替代普通除法。
58		58
59	对于uint32类型整数除法，可作如下变形，将除法计算转换为乘法和移位操作：	59	对于uint32类型整数除法，可作如下变形，将除法计算转换为乘法和移位操作：
60		60
	@@ -80,28 +80,43 @@ $$
80	\left\lfloor \left(\frac{n \times magic}{2^{32}} + n\right) >> shift \right\rfloor	80	\left\lfloor \left(\frac{n \times magic}{2^{32}} + n\right) >> shift \right\rfloor
81	$$	81	$$
82		82
83	-其中magic和shift即为快速除法所需的乘法魔数和移位数。magic和shift的计算只与固定除数有关，可以在SIMD中预计算，在SIMT计算过程中复用，相关代码如下。	83	+其中magic和shift即为快速除法所需的乘法魔数和移位数。magic和shift的计算只与固定除数有关，可以在SIMD标量计算逻辑中调用[asc_get_uintdiv_magic_and_shift](../../../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md)接口预计算，在SIMT计算过程中复用，相关代码如下。
84		84
85	```cpp	85	```cpp
86	-__aicore__ inline void calc_magic_shift(uint32_t& magic_out, uint32_t& shift_out, uint32_t divisor)	86	+__global__ __vector__ void integer_div_kernel(
		87	+ __gm__ uint32_t* input,
		88	+ __gm__ uint32_t* output,
		89	+ uint32_t divisor,
		90	+ uint32_t total_length)
87	{	91	{
88	- if (divisor == 0) {	92	+ // 在标量计算逻辑中，根据固定除数预计算magic和shift。
89	- magic_out = 0;	93	+ uint32_t magic = 0;
90	- shift_out = 0;	94	+ uint32_t shift = 0;
91	- return;	95	+ asc_get_uintdiv_magic_and_shift(&magic, &shift, divisor);
92	- }
93		96
94	- int64_t pos = BIT_64_LEN - count_leading_zero(divisor);	97	+ uint32_t block_offset = block_idx * THREAD_COUNT;
95	- int64_t cnt1 = get_bit_count1(divisor);
96	- uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;
97	- uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;
98		98
99	- magic_out = magic;	99	+ __ubuf__ uint32_t input_buf[THREAD_COUNT];
100	- shift_out = shift;	100	+ __ubuf__ uint32_t output_buf[THREAD_COUNT];
		101	+ uint32_t blk_length = THREAD_COUNT * sizeof(uint32_t);
		102	+ asc_copy_gm2ub_align(
		103	+ input_buf, input + block_offset, 1, blk_length, 0, 0, false, asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);
		104	+
		105	+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
		106	+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
		107	+
		108	+ // 将magic和shift传入SIMT VF，完成快速除法计算。
		109	+ asc_vf_call<simt_fast_div>(dim3(THREAD_COUNT), input_buf, output_buf, magic, shift, total_length);
		110	+
		111	+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
		112	+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
		113	+
		114	+ asc_copy_ub2gm_align(
		115	+ output + block_offset, output_buf, 1, blk_length, asc_store_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);
101	}	116	}
102	```	117	```
103		118
104	-在SIMT计算过程中，每个线程读取待处理的被除数value，调用__umulhi(value, magic)获取value与magic乘积的高32位结果q，再对value + q右移shift位，得到与value / divisor等价的结果。关键代码如下。	119	+在SIMT计算过程中，每个线程读取待处理的被除数value，调用[asc_uintdiv](../../../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md)接口完成乘法和移位组合运算，得到与value / divisor等价的结果。关键代码如下。
105		120
106	```cpp	121	```cpp
107	__simt_vf__ __launch_bounds__(THREAD_LIMIT) inline void simt_fast_div(	122	__simt_vf__ __launch_bounds__(THREAD_LIMIT) inline void simt_fast_div(
	@@ -119,8 +134,7 @@ __simt_vf__ __launch_bounds__(THREAD_LIMIT) inline void simt_fast_div(
119	}	134	}
120		135
121	uint32_t value = input[local_idx];	136	uint32_t value = input[local_idx];
122	- uint32_t q = __umulhi(value, magic);	137	+ uint32_t result = asc_uintdiv(value, magic, shift);
123	- uint32_t result = (value + q) >> shift;
124	output[local_idx] = result;	138	output[local_idx] = result;
125	}	139	}
126	```	140	```
	@@ -138,4 +152,4 @@ __simt_vf__ __launch_bounds__(THREAD_LIMIT) inline void simt_fast_div(
138		152
139	相比场景0，场景1使用乘法和移位替代普通整数除法，Task Duration从117.229μs降低至103.889μs，端到端耗时下降约11.4%。aiv\_vec\_time从0.497μs降低至0.270μs，下降约45.6%，说明将普通除法替换为乘法和移位后，计算指令耗时明显降低。aiv\_scalar\_time从0.128μs增加至0.142μs，主要原因是快速除法需要在Scalar计算单元中额外计算magic和shift；但Task Duration仍然下降，说明快速除法带来的计算收益可以覆盖Scalar计算开销。	153	相比场景0，场景1使用乘法和移位替代普通整数除法，Task Duration从117.229μs降低至103.889μs，端到端耗时下降约11.4%。aiv\_vec\_time从0.497μs降低至0.270μs，下降约45.6%，说明将普通除法替换为乘法和移位后，计算指令耗时明显降低。aiv\_scalar\_time从0.128μs增加至0.142μs，主要原因是快速除法需要在Scalar计算单元中额外计算magic和shift；但Task Duration仍然下降，说明快速除法带来的计算收益可以覆盖Scalar计算开销。
140		154
141	-【总结】对于除数固定的SIMT整数除法场景，可以在SIMD中预计算乘法魔数magic和移位量shift，并在SIMT计算过程中复用。通过将普通整数除法替换为乘法和移位操作，可以有效降低除法计算开销。	155	+【总结】对于除数固定的SIMT整数除法场景，可以在SIMD标量计算逻辑中调用[asc_get_uintdiv_magic_and_shift](../../../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md)接口预计算乘法魔数magic和移位量shift，并在SIMT计算过程中调用[asc_uintdiv](../../../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md)接口复用该结果完成快速除法。通过将普通整数除法替换为乘法和移位操作，可以有效降低除法计算开销。

Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div/README.md+15-21

	@@ -61,10 +61,8 @@ output[i] = input[i] / divisor
61	```cpp	61	```cpp
62	asc_copy_gm2ub_align(input_buf, input + block_offset, 1, blk_length, 0, 0, false, asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);	62	asc_copy_gm2ub_align(input_buf, input + block_offset, 1, blk_length, 0, 0, false, asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);
63		63
64	-if ASC_IS_AIV {	64	+asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
65	- asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);	65	+asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
66	- asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
67	-}
68		66
69	if constexpr (scenario == 0) {	67	if constexpr (scenario == 0) {
70	asc_vf_call<simt_normal_div>(dim3(THREAD_COUNT), input_buf, output_buf, divisor, total_length);	68	asc_vf_call<simt_normal_div>(dim3(THREAD_COUNT), input_buf, output_buf, divisor, total_length);
	@@ -72,10 +70,8 @@ if constexpr (scenario == 0) {
72	asc_vf_call<simt_fast_div>(dim3(THREAD_COUNT), input_buf, output_buf, magic, shift, total_length);	70	asc_vf_call<simt_fast_div>(dim3(THREAD_COUNT), input_buf, output_buf, magic, shift, total_length);
73	}	71	}
74		72
75	-if ASC_IS_AIV {	73	+asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
76	- asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);	74	+asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
77	- asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
78	-}
79		75
80	asc_copy_ub2gm_align(output + block_offset, output_buf, 1, blk_length, asc_store_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);	76	asc_copy_ub2gm_align(output + block_offset, output_buf, 1, blk_length, asc_store_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);
81	```	77	```
	@@ -132,8 +128,8 @@ uint32_t result = value / divisor;
132		128
133	核心实现：	129	核心实现：
134		130
135	-- 在SIMD侧根据固定除数计算快速除法所需的 `magic` 和 `shift`，计算结果可被SIMT线程复用。	131	+- 在SIMD标量计算逻辑中调用 [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) 接口，根据固定除数计算快速除法所需的 `magic` 和 `shift`，计算结果可被SIMT线程复用。
136	-- 每个SIMT线程从UB读取1个输入元素，线程内使用 `__umulhi` 和右移完成等价除法计算。	132	+- 每个SIMT线程从UB读取1个输入元素，线程内调用 [asc_uintdiv](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) 接口，使用乘法和移位完成等价除法计算。
137		133
138	#### 快速除法转换原理	134	#### 快速除法转换原理
139		135
	@@ -162,24 +158,22 @@ $$
162	$$	158	$$
163		159
164	其中 `magic` 和 `shift` 表示快速除法所需的乘法魔数和移位数。	160	其中 `magic` 和 `shift` 表示快速除法所需的乘法魔数和移位数。
165	-本样例中基于上述除法转换思路，在SIMD部分使用Scalar计算单元中的 `clz()`、`bcnt1()` 标量计算接口，获取快速除法所需的 `magic` 和 `shift`，供SIMT线程复用。	161	+本样例中基于上述除法转换思路，在SIMD标量计算逻辑中调用快速除法预计算接口 [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md)，获取快速除法所需的 `magic` 和 `shift`，供SIMT线程复用；SIMT线程内调用快速除法接口 [asc_uintdiv](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) 完成除法计算。
166		162
167	关键代码：	163	关键代码：
168		164
169	-`calc_magic_shift()` 函数中计算快速除法的 `magic` 和 `shift`：	165	+`integer_div_kernel()` 核函数中调用 [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) 接口，计算快速除法的 `magic` 和 `shift`：
170		166
171	```cpp	167	```cpp
172	-int64_t pos = BIT_64_LEN - count_leading_zero(divisor);	168	+uint32_t magic = 0;
173	-int64_t cnt1 = get_bit_count1(divisor);	169	+uint32_t shift = 0;
174	-uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;	170	+asc_get_uintdiv_magic_and_shift(&magic, &shift, divisor);
175	-uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;
176	```	171	```
177		172
178	-`simt_fast_div()` 函数中对当前线程处理的输入元素 `value` 执行快速除法：	173	+`simt_fast_div()` 函数中对当前线程处理的输入元素 `value` 调用 [asc_uintdiv](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) 接口执行快速除法：
179		174
180	```cpp	175	```cpp
181	-uint32_t q = __umulhi(value, magic);	176	+uint32_t result = asc_uintdiv(value, magic, shift);
182	-uint32_t result = (value + q) >> shift;
183	```	177	```
184		178
185	性能数据：	179	性能数据：
	@@ -212,8 +206,8 @@ uint32_t result = (value + q) >> shift;
212		206
213	\| 优化手段 \| 核心原理 \| 样例体现 \|	207	\| 优化手段 \| 核心原理 \| 样例体现 \|
214	\|:---\|:---\|:---\|	208	\|:---\|:---\|:---\|
215	-\| 固定除数预计算 \| 当除数在核函数执行期间固定时，可以提前计算 `magic` 和 `shift`，SIMT线程复用该结果完成快速除法。 \| Case 1在核函数中调用 `calc_magic_shift()`，SIMT线程内不再直接执行 `/`。 \|	209	+\| 固定除数预计算 \| 当除数在核函数执行期间固定时，可以在SIMD标量计算逻辑中调用 [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) 接口计算 `magic` 和 `shift`，SIMT线程复用该结果完成快速除法。 \| Case 1在核函数中调用 [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md)，SIMT线程内不再直接执行 `/`。 \|
216	-\| 乘法和移位替代普通除法 \| 普通整数除法指令开销较高，使用乘法和移位操作可降低逐元素计算开销。 \| Case 1相比Case 0，`Task Duration` 从117.229μs降低到103.889μs。 \|	210	+\| 乘法和移位替代普通除法 \| 普通整数除法指令开销较高，SIMT线程内调用 [asc_uintdiv](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) 接口，使用乘法和移位操作降低逐元素计算开销。 \| Case 1相比Case 0，`Task Duration` 从117.229μs降低到103.889μs。 \|
217		211
218	## 编译运行	212	## 编译运行
219		213

Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div/README_en.md+15-21

	@@ -61,10 +61,8 @@ The shared data transfer and SIMT invocation flow for both cases is as follows:
61	```cpp	61	```cpp
62	asc_copy_gm2ub_align(input_buf, input + block_offset, 1, blk_length, 0, 0, false, asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);	62	asc_copy_gm2ub_align(input_buf, input + block_offset, 1, blk_length, 0, 0, false, asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);
63		63
64	-if ASC_IS_AIV {	64	+asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
65	- asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);	65	+asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
66	- asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
67	-}
68		66
69	if constexpr (scenario == 0) {	67	if constexpr (scenario == 0) {
70	asc_vf_call<simt_normal_div>(dim3(THREAD_COUNT), input_buf, output_buf, divisor, total_length);	68	asc_vf_call<simt_normal_div>(dim3(THREAD_COUNT), input_buf, output_buf, divisor, total_length);
	@@ -72,10 +70,8 @@ if constexpr (scenario == 0) {
72	asc_vf_call<simt_fast_div>(dim3(THREAD_COUNT), input_buf, output_buf, magic, shift, total_length);	70	asc_vf_call<simt_fast_div>(dim3(THREAD_COUNT), input_buf, output_buf, magic, shift, total_length);
73	}	71	}
74		72
75	-if ASC_IS_AIV {	73	+asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
76	- asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);	74	+asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
77	- asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
78	-}
79		75
80	asc_copy_ub2gm_align(output + block_offset, output_buf, 1, blk_length, asc_store_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);	76	asc_copy_ub2gm_align(output + block_offset, output_buf, 1, blk_length, asc_store_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);
81	```	77	```
	@@ -132,8 +128,8 @@ uint32_t result = value / divisor;
132		128
133	Core Implementation:	129	Core Implementation:
134		130
135	-- On the SIMD side, compute the `magic` and `shift` required for fast division based on the fixed divisor. The computation results can be reused by SIMT threads.	131	+- In the SIMD scalar computation logic, call the [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) API to compute the `magic` and `shift` required for fast division based on the fixed divisor. The computation results can be reused by SIMT threads.
136	-- Each SIMT thread reads 1 input element from UB and uses `__umulhi` and right shift within the thread to complete the equivalent division computation.	132	+- Each SIMT thread reads 1 input element from UB and calls the [asc_uintdiv](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) API within the thread to complete the equivalent division computation using multiplication and shift.
137		133
138	#### Fast Division Conversion Principle	134	#### Fast Division Conversion Principle
139		135
	@@ -162,24 +158,22 @@ $$
162	$$	158	$$
163		159
164	Where `magic` and `shift` represent the multiplication magic number and shift bits required for fast division.	160	Where `magic` and `shift` represent the multiplication magic number and shift bits required for fast division.
165	-In this example, based on the above division conversion approach, the SIMD part uses the `clz()` and `bcnt1()` scalar computation interfaces in the Scalar computation unit to obtain the `magic` and `shift` required for fast division, for reuse by SIMT threads.	161	+In this example, based on the above division conversion approach, the SIMD scalar computation logic calls the fast division precomputation API [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) to obtain the `magic` and `shift` required for fast division, for reuse by SIMT threads; SIMT threads call the fast division API [asc_uintdiv](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) to complete the division computation.
166		162
167	Key Code:	163	Key Code:
168		164
169	-The `calc_magic_shift()` function computes `magic` and `shift` for fast division:	165	+The `integer_div_kernel()` kernel function calls the [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) API to compute `magic` and `shift` for fast division:
170		166
171	```cpp	167	```cpp
172	-int64_t pos = BIT_64_LEN - count_leading_zero(divisor);	168	+uint32_t magic = 0;
173	-int64_t cnt1 = get_bit_count1(divisor);	169	+uint32_t shift = 0;
174	-uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;	170	+asc_get_uintdiv_magic_and_shift(&magic, &shift, divisor);
175	-uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;
176	```	171	```
177		172
178	-The `simt_fast_div()` function performs fast division on the input element `value` processed by the current thread:	173	+The `simt_fast_div()` function calls the [asc_uintdiv](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) API to perform fast division on the input element `value` processed by the current thread:
179		174
180	```cpp	175	```cpp
181	-uint32_t q = __umulhi(value, magic);	176	+uint32_t result = asc_uintdiv(value, magic, shift);
182	-uint32_t result = (value + q) >> shift;
183	```	177	```
184		178
185	Performance Data:	179	Performance Data:
	@@ -212,8 +206,8 @@ uint32_t result = (value + q) >> shift;
212		206
213	\| Optimization Method \| Core Principle \| Example Demonstration \|	207	\| Optimization Method \| Core Principle \| Example Demonstration \|
214	\|:---\|:---\|:---\|	208	\|:---\|:---\|:---\|
215	-\| Fixed divisor precomputation \| When the divisor is fixed during kernel function execution, `magic` and `shift` can be computed in advance, and SIMT threads reuse these results for fast division. \| Case 1 calls `calc_magic_shift()` in the kernel function, and SIMT threads no longer directly execute `/`. \|	209	+\| Fixed divisor precomputation \| When the divisor is fixed during kernel function execution, the [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) API can be called in the SIMD scalar computation logic to compute `magic` and `shift`, and SIMT threads reuse these results for fast division. \| Case 1 calls [asc_get_uintdiv_magic_and_shift](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_get_uintdiv_magic_and_shift.md) in the kernel function, and SIMT threads no longer directly execute `/`. \|
216	-\| Multiplication and shift replacing standard division \| Standard integer division instructions have high overhead. Using multiplication and shift operations reduces per-element computation overhead. \| Case 1 reduces `Task Duration` from 117.229μs to 103.889μs compared to Case 0. \|	210	+\| Multiplication and shift replacing standard division \| Standard integer division instructions have high overhead. SIMT threads call the [asc_uintdiv](../../../../docs/zh/api/SIMT-API/math_functions/integer_math_functions/asc_uintdiv.md) API to use multiplication and shift operations to reduce per-element computation overhead. \| Case 1 reduces `Task Duration` from 117.229μs to 103.889μs compared to Case 0. \|
217		211
218	## Build and Run	212	## Build and Run
219		213

Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div/integer_div.asc+6-34

	@@ -38,34 +38,11 @@ namespace {
38	constexpr uint32_t THREAD_COUNT = 2048;	38	constexpr uint32_t THREAD_COUNT = 2048;
39	constexpr uint32_t THREAD_LIMIT = 2048;	39	constexpr uint32_t THREAD_LIMIT = 2048;
40	constexpr uint32_t MAX_UB_ELEMENTS = 32768;	40	constexpr uint32_t MAX_UB_ELEMENTS = 32768;
41	-constexpr uint32_t BIT_64_LEN = 64;
42	-constexpr uint8_t BIT_32_LEN = 32;
43		41
44	constexpr uint32_t DATA_SIZE = 8388608; // number of input elements	42	constexpr uint32_t DATA_SIZE = 8388608; // number of input elements
45	constexpr uint32_t DIVISOR = 7; // fixed divisor	43	constexpr uint32_t DIVISOR = 7; // fixed divisor
46	} // namespace	44	} // namespace
47		45
48	-__aicore__ inline int64_t count_leading_zero(uint64_t x) { return clz(x); }
49	-
50	-__aicore__ inline int64_t get_bit_count1(uint64_t x) { return bcnt1(x); }
51	-
52	-__aicore__ inline void calc_magic_shift(uint32_t& magic_out, uint32_t& shift_out, uint32_t divisor)
53	-{
54	- if (divisor == 0) {
55	- magic_out = 0;
56	- shift_out = 0;
57	- return;
58	- }
59	-
60	- int64_t pos = BIT_64_LEN - count_leading_zero(divisor);
61	- int64_t cnt1 = get_bit_count1(divisor);
62	- uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;
63	- uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;
64	-
65	- magic_out = magic;
66	- shift_out = shift;
67	-}
68	-
69	__simt_vf__ __launch_bounds__(THREAD_LIMIT) inline void simt_normal_div(	46	__simt_vf__ __launch_bounds__(THREAD_LIMIT) inline void simt_normal_div(
70	__ubuf__ uint32_t* input, __ubuf__ uint32_t* output, uint32_t divisor, uint32_t total_length)	47	__ubuf__ uint32_t* input, __ubuf__ uint32_t* output, uint32_t divisor, uint32_t total_length)
71	{	48	{
	@@ -92,8 +69,7 @@ __simt_vf__ __launch_bounds__(THREAD_LIMIT) inline void simt_fast_div(
92	}	69	}
93		70
94	uint32_t value = input[local_idx];	71	uint32_t value = input[local_idx];
95	- uint32_t q = __umulhi(value, magic);	72	+ uint32_t result = asc_uintdiv(value, magic, shift);
96	- uint32_t result = (value + q) >> shift;
97	output[local_idx] = result;	73	output[local_idx] = result;
98	}	74	}
99		75
	@@ -104,7 +80,7 @@ __global__ __vector__ void integer_div_kernel(
104	uint32_t magic = 0;	80	uint32_t magic = 0;
105	uint32_t shift = 0;	81	uint32_t shift = 0;
106	if constexpr (scenario == 1) {	82	if constexpr (scenario == 1) {
107	- calc_magic_shift(magic, shift, divisor);	83	+ asc_get_uintdiv_magic_and_shift(&magic, &shift, divisor);
108	}	84	}
109		85
110	uint32_t block_offset = block_idx * THREAD_COUNT;	86	uint32_t block_offset = block_idx * THREAD_COUNT;
	@@ -115,10 +91,8 @@ __global__ __vector__ void integer_div_kernel(
115	asc_copy_gm2ub_align(	91	asc_copy_gm2ub_align(
116	input_buf, input + block_offset, 1, blk_length, 0, 0, false, asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);	92	input_buf, input + block_offset, 1, blk_length, 0, 0, false, asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);
117		93
118	- if ASC_IS_AIV {	94	+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
119	- asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);	95	+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
120	- asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
121	- }
122		96
123	if constexpr (scenario == 0) {	97	if constexpr (scenario == 0) {
124	asc_vf_call<simt_normal_div>(dim3(THREAD_COUNT), input_buf, output_buf, divisor, total_length);	98	asc_vf_call<simt_normal_div>(dim3(THREAD_COUNT), input_buf, output_buf, divisor, total_length);
	@@ -126,10 +100,8 @@ __global__ __vector__ void integer_div_kernel(
126	asc_vf_call<simt_fast_div>(dim3(THREAD_COUNT), input_buf, output_buf, magic, shift, total_length);	100	asc_vf_call<simt_fast_div>(dim3(THREAD_COUNT), input_buf, output_buf, magic, shift, total_length);
127	}	101	}
128		102
129	- if ASC_IS_AIV {	103	+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
130	- asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);	104	+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
131	- asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
132	- }
133		105
134	asc_copy_ub2gm_align(	106	asc_copy_ub2gm_align(
135	output + block_offset, output_buf, 1, blk_length, asc_store_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);	107	output + block_offset, output_buf, 1, blk_length, asc_store_l2_cache_mode::NORMAL_FIRST_VICTIM, 0, 0);