Star776
1079
代码
代码
Issues72
Pull Requests213
Actions
项目讨论
Wiki
分析
Star776
1079
已合并
混合编程interger fast div样例去除c++风格接口 && SIMT编程样例支持版本修改 #4374
zhangshunwei2创建于 7月15日
混合编程interger fast div样例去除c++风格接口 && SIMT编程样例支持版本修改 #4374
已合并
zhangshunwei2创建于 7月15日
从zhangshunwei2/asc-devkitfastdiv_change合入到cann/asc-devkitmaster
共 34 个文件变更+40-42
Mdocs/zh/guide/算子实践参考/SIMD与SIMT混合算子性能优化/指令优化/整数除法的快速算法.md+2-2
@@ -90,8 +90,8 @@ __aicore__ inline void calc_magic_shift(uint32_t& magic_out, uint32_t& shift_out
90 return;90 return;
91 }91 }
92 92 
93- int64_t pos = BIT_64_LEN - CountLeadingZero(divisor);93+ int64_t pos = BIT_64_LEN - count_leading_zero(divisor);
94- int64_t cnt1 = GetBitCount1(divisor);94+ int64_t cnt1 = get_bit_count1(divisor);
95 uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;95 uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;
96 uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;96 uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;
97 97 
Mexamples/03_simt_api/00_introduction/00_quickstart/hello_world_simt/README.md+1-1
@@ -9,7 +9,7 @@
9- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
10 10 
11## 支持的CANN软件版本11## 支持的CANN软件版本
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## 目录结构介绍14## 目录结构介绍
15 15 
Mexamples/03_simt_api/00_introduction/00_quickstart/hello_world_simt/README_en.md+1-1
@@ -9,7 +9,7 @@ This example is a SIMT programming introductory example. It uses the `<<<>>>` ke
9- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
10 10 
11## Supported CANN Software Versions11## Supported CANN Software Versions
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## Directory Structure14## Directory Structure
15 15 
Mexamples/03_simt_api/01_utilities/00_printf/README.md+1-1
@@ -9,7 +9,7 @@
9- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
10 10 
11## 支持的CANN软件版本11## 支持的CANN软件版本
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## 目录结构介绍14## 目录结构介绍
15 15 
Mexamples/03_simt_api/01_utilities/00_printf/README_en.md+1-1
@@ -9,7 +9,7 @@ This example demonstrates how to use the `printf()` interface in SIMT programmin
9- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
10 10 
11## Supported CANN Software Version11## Supported CANN Software Version
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## Directory Structure14## Directory Structure
15 15 
Mexamples/03_simt_api/01_utilities/01_assert/README.md+1-1
@@ -9,7 +9,7 @@
9- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
10 10 
11## 支持的CANN软件版本11## 支持的CANN软件版本
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## 目录结构介绍14## 目录结构介绍
15 15 
Mexamples/03_simt_api/01_utilities/01_assert/README_en.md+1-1
@@ -9,7 +9,7 @@ This example demonstrates how to use the `assert()` interface in SIMT programmin
9- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
10 10 
11## Supported CANN Software Version11## Supported CANN Software Version
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## Directory Structure14## Directory Structure
15 15 
Mexamples/03_simt_api/01_utilities/02_clock/README.md+1-1
@@ -10,7 +10,7 @@
10 10 
11## 支持的CANN软件版本11## 支持的CANN软件版本
12 12 
13-- \> CANN 9.0.013+- \>= CANN 9.1.0
14 14 
15## 目录结构介绍15## 目录结构介绍
16 16 
Mexamples/03_simt_api/01_utilities/02_clock/README_en.md+1-1
@@ -10,7 +10,7 @@ This example demonstrates how to use the `clock()` interface in SIMT programming
10 10 
11## Supported CANN Software Version11## Supported CANN Software Version
12 12 
13-- \> CANN 9.0.013+- \>= CANN 9.1.0
14 14 
15## Directory Structure15## Directory Structure
16 16 
Mexamples/03_simt_api/02_features/00_framework/00_pytorch/torch_library_roll/README.md+1-1
@@ -9,7 +9,7 @@
9 9 
10## 支持的CANN软件版本10## 支持的CANN软件版本
11 11 
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## 目录结构介绍14## 目录结构介绍
15 15 
Mexamples/03_simt_api/02_features/00_framework/00_pytorch/torch_library_roll/README_en.md+1-1
@@ -9,7 +9,7 @@ This example uses the Roll operator to demonstrate how to register custom operat
9 9 
10## Supported CANN Software Version10## Supported CANN Software Version
11 11 
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## Directory Structure14## Directory Structure
15 15 
Mexamples/03_simt_api/02_features/01_api_features/00_memory_access/insert_hash_table/README.md+1-1
@@ -9,7 +9,7 @@
9- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
10 10 
11## 支持的CANN软件版本11## 支持的CANN软件版本
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## 目录结构介绍14## 目录结构介绍
15 15 
Mexamples/03_simt_api/02_features/01_api_features/00_memory_access/insert_hash_table/README_en.md+1-1
@@ -9,7 +9,7 @@ This example introduces the InsertHashTable operator, demonstrating a large-scal
9- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
10 10 
11## Supported CANN Software Version11## Supported CANN Software Version
12-- \> CANN 9.0.012+- \>= CANN 9.1.0
13 13 
14## Directory Structure14## Directory Structure
15 15 
Mexamples/03_simt_api/02_features/01_api_features/01_sync_instruction/memory_fence/README.md+1-1
@@ -11,7 +11,7 @@
11- Ascend 950PR/Ascend 950DT11- Ascend 950PR/Ascend 950DT
12 12 
13## 支持的CANN软件版本13## 支持的CANN软件版本
14-- \> CANN 9.0.014+- \>= CANN 9.1.0
15 15 
16## 目录结构介绍16## 目录结构介绍
17 17 
Mexamples/03_simt_api/02_features/01_api_features/01_sync_instruction/memory_fence/README_en.md+1-1
@@ -11,7 +11,7 @@ The 2 scenarios correspond to reduction sum for small shape input and large shap
11- Ascend 950PR/Ascend 950DT11- Ascend 950PR/Ascend 950DT
12 12 
13## Supported CANN Software Version13## Supported CANN Software Version
14-- \> CANN 9.0.014+- \>= CANN 9.1.0
15 15 
16## Directory Structure16## Directory Structure
17 17 
Mexamples/03_simt_api/02_features/01_api_features/02_atomic_operation/histogram/README.md+1-1
@@ -11,7 +11,7 @@
11 11 
12## 支持的CANN软件版本12## 支持的CANN软件版本
13 13 
14-- \> CANN 9.0.014+- \>= CANN 9.1.0
15 15 
16## 目录结构介绍16## 目录结构介绍
17 17 
Mexamples/03_simt_api/02_features/01_api_features/02_atomic_operation/histogram/README_en.md+1-1
@@ -11,7 +11,7 @@ This example demonstrates how to use the Ascend C SIMT [`asc_atomic_add()`](http
11 11 
12## Supported CANN Software Version12## Supported CANN Software Version
13 13 
14-- \> CANN 9.0.014+- \>= CANN 9.1.0
15 15 
16## Directory Structure16## Directory Structure
17 17 
Mexamples/03_simt_api/02_features/06_compile/00_basic_compile/README.md+1-1
@@ -17,7 +17,7 @@
17 17 
18## 支持的CANN软件版本18## 支持的CANN软件版本
19 19 
20-- \> CANN 9.0.020+- \>= CANN 9.1.0
21 21 
22## 目录结构介绍22## 目录结构介绍
23 23 
Mexamples/03_simt_api/02_features/06_compile/00_basic_compile/README_en.md+1-1
@@ -17,7 +17,7 @@ Figure 1: Code structure and call relationship
17 17 
18## Supported CANN Software Version18## Supported CANN Software Version
19 19 
20-- \> CANN 9.0.020+- \>= CANN 9.1.0
21 21 
22## Directory Structure22## Directory Structure
23 23 
Mexamples/03_simt_api/02_features/06_compile/01_separate_compile/README.md+1-1
@@ -17,7 +17,7 @@
17 17 
18## 支持的CANN软件版本18## 支持的CANN软件版本
19 19 
20-- \> CANN 9.0.020+- \>= CANN 9.1.0
21 21 
22## 目录结构介绍22## 目录结构介绍
23 23 
Mexamples/03_simt_api/02_features/06_compile/01_separate_compile/README_en.md+1-1
@@ -17,7 +17,7 @@ Figure 1: Code structure and call relationship
17 17 
18## Supported CANN Software Version18## Supported CANN Software Version
19 19 
20-- \> CANN 9.0.020+- \>= CANN 9.1.0
21 21 
22## Directory Structure22## Directory Structure
23 23 
Mexamples/03_simt_api/02_features/06_compile/02_dynamic_library_compile/README.md+1-1
@@ -17,7 +17,7 @@
17 17 
18## 支持的CANN软件版本18## 支持的CANN软件版本
19 19 
20-- \> CANN 9.0.020+- \>= CANN 9.1.0
21 21 
22## 目录结构介绍22## 目录结构介绍
23 23 
Mexamples/03_simt_api/02_features/06_compile/02_dynamic_library_compile/README_en.md+1-1
@@ -17,7 +17,7 @@ Figure 1: Code structure and call relationship
17 17 
18## Supported CANN Software Version18## Supported CANN Software Version
19 19 
20-- \> CANN 9.0.020+- \>= CANN 9.1.0
21 21 
22## Directory Structure22## Directory Structure
23 23 
Mexamples/03_simt_api/02_features/06_compile/03_static_library_compile/README.md+1-1
@@ -17,7 +17,7 @@
17 17 
18## 支持的CANN软件版本18## 支持的CANN软件版本
19 19 
20-- \> CANN 9.0.020+- \>= CANN 9.1.0
21 21 
22## 目录结构介绍22## 目录结构介绍
23 23 
Mexamples/03_simt_api/02_features/06_compile/03_static_library_compile/README_en.md+1-1
@@ -17,7 +17,7 @@ Figure 1: Code structure and call relationship
17 17 
18## Supported CANN Software Version18## Supported CANN Software Version
19 19 
20-- \> CANN 9.0.020+- \>= CANN 9.1.0
21 21 
22## Directory Structure22## Directory Structure
23 23 
Mexamples/03_simt_api/03_best_practices/00_memory_optimizations/aligned_types/README.md+1-1
@@ -10,7 +10,7 @@
10 10 
11## 支持的 CANN 软件版本11## 支持的 CANN 软件版本
12 12 
13- - \> CANN 9.0.013+ - \>= CANN 9.1.0
14 14 
15## 目录结构15## 目录结构
16 16 
Mexamples/03_simt_api/03_best_practices/00_memory_optimizations/aligned_types/README_en.md+1-1
@@ -10,7 +10,7 @@
10 10 
11## Supported CANN Software Version11## Supported CANN Software Version
12 12 
13- - \> CANN 9.0.013+ - \>= CANN 9.1.0
14 14 
15## Directory Structure15## Directory Structure
16 16 
Mexamples/03_simt_api/03_best_practices/00_memory_optimizations/cache_hint/README.md+1-1
@@ -10,7 +10,7 @@
10 10 
11## 支持的CANN软件版本11## 支持的CANN软件版本
12 12 
13-- \> CANN 9.0.013+- \>= CANN 9.1.0
14 14 
15## 目录结构介绍15## 目录结构介绍
16 16 
Mexamples/03_simt_api/03_best_practices/00_memory_optimizations/cache_hint/README_en.md+1-1
@@ -10,7 +10,7 @@ This example uses table-lookup sin computation to demonstrate data cache optimiz
10 10 
11## Supported CANN Software Version11## Supported CANN Software Version
12 12 
13-- \> CANN 9.0.013+- \>= CANN 9.1.0
14 14 
15## Directory Structure15## Directory Structure
16 16 
Mexamples/03_simt_api/03_best_practices/00_memory_optimizations/short_vector_add/README.md+1-1
@@ -16,7 +16,7 @@
16 16 
17## 支持的CANN软件版本17## 支持的CANN软件版本
18 18 
19-- \> CANN 9.0.019+- \>= CANN 9.1.0
20 20 
21## 目录结构介绍21## 目录结构介绍
22 22 
Mexamples/03_simt_api/03_best_practices/00_memory_optimizations/short_vector_add/README_en.md+1-1
@@ -15,7 +15,7 @@ This example uses vector addition to demonstrate performance tuning methods for
15 15 
16## Supported CANN Software Version16## Supported CANN Software Version
17 17 
18-- \> CANN 9.0.018+- \>= CANN 9.1.0
19 19 
20## Directory Structure20## Directory Structure
21 21 
Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div/README.md+2-2
@@ -168,8 +168,8 @@ $$
168`calc_magic_shift()` 函数中计算快速除法的 `magic` 和 `shift`：168`calc_magic_shift()` 函数中计算快速除法的 `magic` 和 `shift`：
169 169 
170```cpp170```cpp
171-int64_t pos = BIT_64_LEN - CountLeadingZero(divisor);171+int64_t pos = BIT_64_LEN - count_leading_zero(divisor);
172-int64_t cnt1 = GetBitCount1(divisor);172+int64_t cnt1 = get_bit_count1(divisor);
173uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;173uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;
174uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;174uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;
175```175```
Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div/README_en.md+2-2
@@ -168,8 +168,8 @@ In this example, based on the above division conversion approach, the SIMD part
168The `calc_magic_shift()` function computes `magic` and `shift` for fast division:168The `calc_magic_shift()` function computes `magic` and `shift` for fast division:
169 169 
170```cpp170```cpp
171-int64_t pos = BIT_64_LEN - CountLeadingZero(divisor);171+int64_t pos = BIT_64_LEN - count_leading_zero(divisor);
172-int64_t cnt1 = GetBitCount1(divisor);172+int64_t cnt1 = get_bit_count1(divisor);
173uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;173uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;
174uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;174uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;
175```175```
Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_integer_fast_div/integer_div.asc+4-6
@@ -23,7 +23,6 @@
23#include <random>23#include <random>
24#include <vector>24#include <vector>
25#include "acl/acl.h"25#include "acl/acl.h"
26-#include "kernel_operator.h"
27#include "simt_api/asc_simt.h"26#include "simt_api/asc_simt.h"
28#include "c_api/asc_simd.h"27#include "c_api/asc_simd.h"
29 28 
@@ -46,9 +45,9 @@ constexpr uint32_t DATA_SIZE = 8388608; // number of input elements
46constexpr uint32_t DIVISOR = 7; // fixed divisor45constexpr uint32_t DIVISOR = 7; // fixed divisor
47} // namespace46} // namespace
48 47 
49-__aicore__ inline int64_t CountLeadingZero(uint64_t x) { return clz(x); }48+__aicore__ inline int64_t count_leading_zero(uint64_t x) { return clz(x); }
50 49 
51-__aicore__ inline int64_t GetBitCount1(uint64_t x) { return bcnt1(x); }50+__aicore__ inline int64_t get_bit_count1(uint64_t x) { return bcnt1(x); }
52 51 
53__aicore__ inline void calc_magic_shift(uint32_t& magic_out, uint32_t& shift_out, uint32_t divisor)52__aicore__ inline void calc_magic_shift(uint32_t& magic_out, uint32_t& shift_out, uint32_t divisor)
54{53{
@@ -58,8 +57,8 @@ __aicore__ inline void calc_magic_shift(uint32_t& magic_out, uint32_t& shift_out
58 return;57 return;
59 }58 }
60 59 
61- int64_t pos = BIT_64_LEN - CountLeadingZero(divisor);60+ int64_t pos = BIT_64_LEN - count_leading_zero(divisor);
62- int64_t cnt1 = GetBitCount1(divisor);61+ int64_t cnt1 = get_bit_count1(divisor);
63 uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;62 uint32_t shift = (cnt1 == 1) ? (pos - 1) : pos;
64 uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;63 uint32_t magic = (1l << BIT_32_LEN) * ((1l << shift) - divisor) / divisor + 1;
65 64 
@@ -108,7 +107,6 @@ __global__ __vector__ void integer_div_kernel(
108 calc_magic_shift(magic, shift, divisor);107 calc_magic_shift(magic, shift, divisor);
109 }108 }
110 109 
111- uint32_t block_idx = AscendC::GetBlockIdx();
112 uint32_t block_offset = block_idx * THREAD_COUNT;110 uint32_t block_offset = block_idx * THREAD_COUNT;
atomgit-botL
atomgit-botatomgit-bot7月15日

🔴 Critical

该 diff 删除了 uint32_t block_idx = AscendC::GetBlockIdx();（C++ 风格接口），但紧接的第 111 行（新行号）仍然使用 block_idx 计算 block_offset：uint32_t block_offset = block_idx * THREAD_COUNT;。删除后 block_idx 在整个函数及文件中无任何声明或定义，会产生未定义标识符的编译错误。

该 PR 的目标是去掉 C++ 风格接口并替换为 C 风格接口。同类文件 matrix_transpose.asc 中已使用 get_block_idx()（底层内建函数），官方 C API 为 asc_get_block_idx()（定义于 include/c_api/sys_var/sys_var.h，返回 int64_t）。应补充 C 风格替代写法，而非仅删除该行。

建议：将删除的行替换为 C 风格 API 调用。可参考同仓库 matrix_transpose.asc 使用 get_block_idx() 或使用官方 C API asc_get_block_idx()。注意 asc_get_block_idx() 返回 int64_t，赋值给 uint32_t 时建议加上显式类型转换。

likedislike
不准确？
Z
zhangshunwei2
7月15日 评论：
Llouyuxuan7月15日

block_idx 定义删了，这里还在用，能编过吗？

likedislike
Z
zhangshunwei2
7月15日 评论：
113 111 
114 __ubuf__ uint32_t input_buf[THREAD_COUNT];112 __ubuf__ uint32_t input_buf[THREAD_COUNT];