Star686
996
代码
代码
Issues102
Pull Requests187
Actions
项目讨论
Wiki
分析
Star686
996
已合并
新增 SIMT 协作组 grid_group #4795
louyuxuan创建于 22 天前
新增 SIMT 协作组 grid_group #4795
已合并
louyuxuan创建于 22 天前
从louyuxuan/asc-devkitmaster合入到cann/asc-devkitmaster
共 24 个文件变更+1113-3
Mdocs/zh/api/README.md+15-0
@@ -2729,6 +2729,20 @@
2729 2729 
2730 - [协作组](SIMT-API/cooperative_groups/cooperative_groups.md)2730 - [协作组](SIMT-API/cooperative_groups/cooperative_groups.md)
2731 - [协作组简介](SIMT-API/cooperative_groups/cooperative_groups_intro.md)2731 - [协作组简介](SIMT-API/cooperative_groups/cooperative_groups_intro.md)
2732+ - [grid_group](SIMT-API/cooperative_groups/grid_group/grid_group.md)
2733+ - [grid_group简介](SIMT-API/cooperative_groups/grid_group/grid_group_intro.md)
2734+ - [grid_group_constructor](SIMT-API/cooperative_groups/grid_group/grid_group_constructor.md)
2735+ - [is_valid](SIMT-API/cooperative_groups/grid_group/is_valid.md)
2736+ - [sync](SIMT-API/cooperative_groups/grid_group/sync.md)
2737+ - [thread_rank](SIMT-API/cooperative_groups/grid_group/thread_rank.md)
2738+ - [block_rank](SIMT-API/cooperative_groups/grid_group/block_rank.md)
2739+ - [num_threads](SIMT-API/cooperative_groups/grid_group/num_threads.md)
2740+ - [num_blocks](SIMT-API/cooperative_groups/grid_group/num_blocks.md)
2741+ - [dim_blocks](SIMT-API/cooperative_groups/grid_group/dim_blocks.md)
2742+ - [block_index](SIMT-API/cooperative_groups/grid_group/block_index.md)
2743+ - [size](SIMT-API/cooperative_groups/grid_group/size.md)
2744+ - [group_dim](SIMT-API/cooperative_groups/grid_group/group_dim.md)
2745+ 
2732 - [thread_block](SIMT-API/cooperative_groups/thread_block/thread_block.md)2746 - [thread_block](SIMT-API/cooperative_groups/thread_block/thread_block.md)
2733 - [thread_block简介](SIMT-API/cooperative_groups/thread_block/thread_block_intro.md)2747 - [thread_block简介](SIMT-API/cooperative_groups/thread_block/thread_block_intro.md)
2734 - [thread_block构造函数](SIMT-API/cooperative_groups/thread_block/thread_block_constructor.md)2748 - [thread_block构造函数](SIMT-API/cooperative_groups/thread_block/thread_block_constructor.md)
@@ -2774,6 +2788,7 @@
2774 - [any](SIMT-API/cooperative_groups/thread_block_tile/any.md)2788 - [any](SIMT-API/cooperative_groups/thread_block_tile/any.md)
2775 - [all](SIMT-API/cooperative_groups/thread_block_tile/all.md)2789 - [all](SIMT-API/cooperative_groups/thread_block_tile/all.md)
2776 2790 
2791+ - [sync](SIMT-API/cooperative_groups/sync.md)
2777 - [tiled_partition](SIMT-API/cooperative_groups/tiled_partition.md)2792 - [tiled_partition](SIMT-API/cooperative_groups/tiled_partition.md)
2778 - [binary_partition](SIMT-API/cooperative_groups/binary_partition.md)2793 - [binary_partition](SIMT-API/cooperative_groups/binary_partition.md)
2779 2794 
Mdocs/zh/api/SIMT-API/SIMD_SIMT_hybrid_programming_intro/api_list.md+2-1
@@ -818,9 +818,10 @@
818| [any](../cooperative_groups/thread_block_tile/any.md) | 判断是否有组内线程的输入不为0。 |818| [any](../cooperative_groups/thread_block_tile/any.md) | 判断是否有组内线程的输入不为0。 |
819| [all](../cooperative_groups/thread_block_tile/all.md) | 判断是否所有组内线程的输入均不为0。 |819| [all](../cooperative_groups/thread_block_tile/all.md) | 判断是否所有组内线程的输入均不为0。 |
820 820 
821-**表30** 协作组划分接口821+**表30** 协作组通用与划分接口
822 822 
823| 接口名 | 功能描述 |823| 接口名 | 功能描述 |
824| --- | --- |824| --- | --- |
825+| [sync](../cooperative_groups/sync.md) | 同步指定协作组内的所有线程。 |
825| [tiled_partition](../cooperative_groups/tiled_partition.md) | 将一个线程组按指定大小划分为多个子组。 |826| [tiled_partition](../cooperative_groups/tiled_partition.md) | 将一个线程组按指定大小划分为多个子组。 |
826| [binary_partition](../cooperative_groups/binary_partition.md) | 根据一个标签（0或1）将父组划分为两个子组。 |827| [binary_partition](../cooperative_groups/binary_partition.md) | 根据一个标签（0或1）将父组划分为两个子组。 |
Mdocs/zh/api/SIMT-API/SIMT_programming_intro/api_list.md+18-1
@@ -823,9 +823,26 @@
823| [any](../cooperative_groups/thread_block_tile/any.md) | 判断是否有组内线程的输入不为0。 |823| [any](../cooperative_groups/thread_block_tile/any.md) | 判断是否有组内线程的输入不为0。 |
824| [all](../cooperative_groups/thread_block_tile/all.md) | 判断是否所有组内线程的输入均不为0。 |824| [all](../cooperative_groups/thread_block_tile/all.md) | 判断是否所有组内线程的输入均不为0。 |
825 825 
826-**表31** 协作组划分接口826+**表31** grid_group接口
827 827 
828| 接口名 | 功能描述 |828| 接口名 | 功能描述 |
829| --- | --- |829| --- | --- |
830+| [this_grid](../cooperative_groups/grid_group/grid_group_constructor.md) | 获取当前Grid协作组对象。 |
831+| [is_valid](../cooperative_groups/grid_group/is_valid.md) | 判断当前Grid协作组是否可使用`sync`接口进行跨线程块同步。 |
832+| [sync](../cooperative_groups/grid_group/sync.md) | 同步Grid内所有线程。 |
833+| [thread_rank](../cooperative_groups/grid_group/thread_rank.md) | 获取当前线程在Grid内的排名。 |
834+| [block_rank](../cooperative_groups/grid_group/block_rank.md) | 获取当前线程所属线程块在Grid内的排名。 |
835+| [num_threads](../cooperative_groups/grid_group/num_threads.md) | 获取Grid内线程总数。 |
836+| [num_blocks](../cooperative_groups/grid_group/num_blocks.md) | 获取Grid内线程块总数。 |
837+| [dim_blocks](../cooperative_groups/grid_group/dim_blocks.md) | 获取Grid的线程块维度配置。 |
838+| [block_index](../cooperative_groups/grid_group/block_index.md) | 获取当前线程所属线程块在Grid中的三维索引。 |
839+| [size](../cooperative_groups/grid_group/size.md) | 获取Grid内线程总数。 |
840+| [group_dim](../cooperative_groups/grid_group/group_dim.md) | 获取Grid的线程块维度配置。 |
841+ 
842+**表32** 协作组通用与划分接口
843+ 
844+| 接口名 | 功能描述 |
845+| --- | --- |
846+| [sync](../cooperative_groups/sync.md) | 同步指定协作组内的所有线程。 |
830| [tiled_partition](../cooperative_groups/tiled_partition.md) | 将一个线程组按指定大小划分为多个子组。 |847| [tiled_partition](../cooperative_groups/tiled_partition.md) | 将一个线程组按指定大小划分为多个子组。 |
831| [binary_partition](../cooperative_groups/binary_partition.md) | 根据一个标签（0或1）将父组划分为两个子组。 |848| [binary_partition](../cooperative_groups/binary_partition.md) | 根据一个标签（0或1）将父组划分为两个子组。 |
Mdocs/zh/api/SIMT-API/cooperative_groups/cooperative_groups.md+4-0
@@ -2,12 +2,16 @@
2 2 
3- **[协作组简介](cooperative_groups_intro.md)**3- **[协作组简介](cooperative_groups_intro.md)**
4 4 
5+- **[grid_group](grid_group/grid_group.md)**
A
Aai_xin15 天前

api/README.md未修改

likedislike
6+ 
5- **[thread_block](thread_block/thread_block.md)**7- **[thread_block](thread_block/thread_block.md)**
6 8 
7- **[coalesced_group](coalesced_group/coalesced_group.md)**9- **[coalesced_group](coalesced_group/coalesced_group.md)**
8 10 
9- **[thread_block_tile](thread_block_tile/thread_block_tile.md)**11- **[thread_block_tile](thread_block_tile/thread_block_tile.md)**
10 12 
13+- **[sync](sync.md)**
14+ 
11- **[tiled_partition](tiled_partition.md)**15- **[tiled_partition](tiled_partition.md)**
12 16 
13- **[binary_partition](binary_partition.md)**17- **[binary_partition](binary_partition.md)**
Mdocs/zh/api/SIMT-API/cooperative_groups/cooperative_groups_intro.md+7-1
@@ -7,12 +7,17 @@
7- [thread_block](thread_block/thread_block_intro.md)：对线程块的抽象，提供统一的线程管理接口。7- [thread_block](thread_block/thread_block_intro.md)：对线程块的抽象，提供统一的线程管理接口。
8- [coalesced_group](coalesced_group/coalesced_group_intro.md)：Warp中的活跃线程集合。8- [coalesced_group](coalesced_group/coalesced_group_intro.md)：Warp中的活跃线程集合。
9- [thread_block_tile](thread_block_tile/thread_block_tile_intro.md)：手动划分出的线程子组。9- [thread_block_tile](thread_block_tile/thread_block_tile_intro.md)：手动划分出的线程子组。
10+- [grid_group](grid_group/grid_group_intro.md)：对一个核函数启动的所有SIMT线程的抽象，提供跨线程块的查询和同步接口。
10 11 
11基于以上协作组类型，提供以下划分操作，支持将一个协作组划分为多个子组：12基于以上协作组类型，提供以下划分操作，支持将一个协作组划分为多个子组：
12 13 
13- [tiled_partition](tiled_partition.md)：将父组划分成指定大小的子组。14- [tiled_partition](tiled_partition.md)：将父组划分成指定大小的子组。
14- [binary_partition](binary_partition.md)：根据自定义的0、1标签将父组划分为两个子组。15- [binary_partition](binary_partition.md)：根据自定义的0、1标签将父组划分为两个子组。
15 16 
17+基于以上协作组类型，提供以下通用操作：
18+ 
19+- [sync](sync.md)：同步指定协作组内的所有线程，功能等价于调用该协作组对象的成员函数`sync()`。
20+ 
16## 需包含的头文件21## 需包含的头文件
17 22 
18```cpp23```cpp
@@ -20,4 +25,5 @@
20```25```
21 26 
22> [!NOTE]说明 27> [!NOTE]说明
23-> 所有协作组API均在`cooperative_groups`命名空间下，调用时需添加命名空间。28+> 所有协作组API均在`cooperative_groups`命名空间下，调用时需添加命名空间。
29+> `grid_group`仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/block_index.md+60-0
@@ -0,0 +1,60 @@
1+# block_index
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+获取当前线程所属线程块在Grid中的三维索引。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+static dim3 block_index()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前线程所属线程块在Grid中的三维索引。
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ dim3 block = g.block_index();
58+ ...
59+ }
60+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/block_rank.md+60-0
@@ -0,0 +1,60 @@
1+# block_rank
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+获取当前线程所属线程块在Grid内所有线程块中的排名，排名从0开始计数。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+static unsigned long long block_rank()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前线程所属线程块在Grid内的排名。
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ unsigned long long block = g.block_rank();
58+ ...
59+ }
60+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/dim_blocks.md+60-0
@@ -0,0 +1,60 @@
1+# dim_blocks
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+获取当前Grid的线程块维度配置。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+static dim3 dim_blocks()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前Grid的线程块维度配置。
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ dim3 grid_dim = g.dim_blocks();
58+ ...
59+ }
60+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/grid_group.md+25-0
@@ -0,0 +1,25 @@
1+# grid_group
2+ 
3+- **[grid_group简介](grid_group_intro.md)**
4+ 
5+- **[grid_group_constructor](grid_group_constructor.md)**
6+ 
7+- **[is_valid](is_valid.md)**
8+ 
9+- **[sync](sync.md)**
10+ 
11+- **[thread_rank](thread_rank.md)**
12+ 
13+- **[block_rank](block_rank.md)**
14+ 
15+- **[num_threads](num_threads.md)**
16+ 
17+- **[num_blocks](num_blocks.md)**
18+ 
19+- **[dim_blocks](dim_blocks.md)**
20+ 
21+- **[block_index](block_index.md)**
22+ 
23+- **[size](size.md)**
24+ 
25+- **[group_dim](group_dim.md)**
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/grid_group_constructor.md+49-0
@@ -0,0 +1,49 @@
1+# grid_group_constructor
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+`grid_group`不提供默认的构造函数，用户使用`this_grid`函数获取当前Grid对应的`grid_group`对象。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+grid_group this_grid()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前Grid对应的`grid_group`对象。
42+ 
43+## 约束说明
44+ 
45+- 仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+参考[调用示例](sync.md#调用示例)
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/grid_group_intro.md+28-0
@@ -0,0 +1,28 @@
1+# grid_group简介
2+ 
3+`grid_group`是对一个核函数启动的所有SIMT线程的抽象，线程范围覆盖当前Grid中的所有线程块及其线程。通过`grid_group`可以查询Grid维度、线程块排名、全局线程排名，并通过`sync()`在Grid范围内进行跨线程块同步。
4+ 
5+## Public成员函数
6+ 
7+```cpp
8+bool is_valid() const;
9+void sync() const;
10+static unsigned long long thread_rank();
11+static unsigned long long block_rank();
12+static unsigned long long num_threads();
13+static unsigned long long size();
14+static unsigned long long num_blocks();
15+static dim3 dim_blocks();
16+static dim3 group_dim();
17+static dim3 block_index();
18+```
19+ 
20+## 需包含的头文件
21+ 
22+```cpp
23+#include "simt_api/cooperative_groups.h"
24+```
25+ 
26+## 约束说明
27+ 
28+`grid_group`仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/group_dim.md+60-0
@@ -0,0 +1,60 @@
1+# group_dim
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+获取当前Grid的线程块维度配置，与接口[dim_blocks](dim_blocks.md)功能相同。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+static dim3 group_dim()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前Grid的线程块维度配置。
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ dim3 grid_dim = g.group_dim();
58+ ...
59+ }
60+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/is_valid.md+62-0
@@ -0,0 +1,62 @@
1+# is_valid
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+判断当前`grid_group`是否可使用`sync()`接口进行跨线程块同步。`grid_group`默认支持同步功能，因此接口固定返回`true`。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+bool is_valid() const
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+固定返回`true`
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ if (g.is_valid()) {
58+ g.sync();
59+ }
60+ ...
61+ }
62+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/num_blocks.md+60-0
@@ -0,0 +1,60 @@
1+# num_blocks
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+获取当前Grid内的线程块总数。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+static unsigned long long num_blocks()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前Grid内的线程块总数。
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ unsigned long long block_num = g.num_blocks();
58+ ...
59+ }
60+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/num_threads.md+60-0
@@ -0,0 +1,60 @@
1+# num_threads
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+获取当前Grid内的SIMT线程总数。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+static unsigned long long num_threads()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前Grid内的SIMT线程总数。
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ unsigned long long thread_num = g.num_threads();
58+ ...
59+ }
60+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/size.md+60-0
@@ -0,0 +1,60 @@
1+# size
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+获取当前Grid内的SIMT线程总数，与接口[num_threads](num_threads.md)功能相同。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+static unsigned long long size()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前Grid内的SIMT线程总数。
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ unsigned long long thread_num = g.size();
58+ ...
59+ }
60+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/sync.md+75-0
@@ -0,0 +1,75 @@
1+# sync
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+同步当前Grid内所有SIMT线程。所有线程块中的所有线程都执行到该同步点位置后，Grid内线程才能继续执行。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+void sync() const
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+无
42+ 
43+## 约束说明
44+ 
45+- 仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+- 必须保证Grid内所有线程都能执行到同一个`sync()`调用，否则是未定义行为。
47+- `sync()`接口仅保证线程指令执行的同步，不保证各线程块的Data Cache和Global Memory间的[缓存一致性](../../../../guide/编程指南/高级编程/内存模型/缓存一致性.md)。
48+ 
49+## 调用示例
50+ 
51+- SIMT编程场景：
52+ 
53+ ```cpp
54+ using namespace cooperative_groups;
55+ __global__ void simt_kernel(uint64_t* output)
56+ {
57+ grid_group g = this_grid(); // 获取grid_group对象
58+ unsigned int block = static_cast<unsigned int>(g.block_rank());
59+ 
60+ if (threadIdx.x == 0) {
61+ output[block] = block + 1;
62+ }
63+ 
64+ g.sync(); // 同步Grid内所有线程
65+ 
66+ if (threadIdx.x == 0) {
67+ uint64_t sum = 0;
68+ asc_dcci_entire(output); // 刷新Cache
69+ for (unsigned int i = 0; i < g.num_blocks(); ++i) {
70+ sum += output[i];
71+ }
72+ output[g.num_blocks() + block] = sum;
73+ }
74+ }
75+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/grid_group/thread_rank.md+60-0
@@ -0,0 +1,60 @@
1+# thread_rank
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+获取当前线程在Grid内所有线程中的排名，排名从0开始计数。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+static unsigned long long thread_rank()
33+```
34+ 
35+## 参数说明
36+ 
37+无
38+ 
39+## 返回值说明
40+ 
41+当前线程在Grid内所有SIMT线程中的排名。
42+ 
43+## 约束说明
44+ 
45+仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
46+ 
47+## 调用示例
48+ 
49+- SIMT编程场景：
50+ 
51+ ```cpp
52+ using namespace cooperative_groups;
53+ __global__ void simt_kernel(...)
54+ {
55+ ...
56+ grid_group g = this_grid();
57+ unsigned long long rank = g.thread_rank();
58+ ...
59+ }
60+ ```
Adocs/zh/api/SIMT-API/cooperative_groups/sync.md+88-0
@@ -0,0 +1,88 @@
1+# sync
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+- Atlas 推理系列产品Vector Core：不支持
20+<!-- end id5 -->
21+<!-- npu="910" id6 -->
22+- Atlas 训练系列产品：不支持
23+<!-- end id6 -->
24+ 
25+## 功能说明
26+ 
27+同步指定协作组内的所有线程，功能等价于调用该协作组对象的成员函数`sync()`。
28+ 
29+## 函数原型
30+ 
31+```cpp
32+template <typename GroupType>
33+void sync(const GroupType& g)
34+```
35+ 
36+## 参数说明
37+ 
38+**表1** 模板参数说明
39+ 
40+| 参数名 | 描述 |
41+| --- | --- |
42+| GroupType | 协作组对象的类型。 |
43+ 
44+**表2** 参数说明
45+ 
46+| 参数名 | 输入/输出 | 描述 |
47+| --- | --- | --- |
48+| g | 输入 | 待同步的协作组对象。 |
A
Aai_xin15 天前

模板参数 GroupType 未说明含义

likedislike
49+ 
50+## 返回值说明
51+ 
52+无
53+ 
54+## 约束说明
55+ 
56+- 对于`grid_group`，该接口仅支持SIMT编程场景，不支持SIMD与SIMT混合编程场景。
57+- 必须保证协作组内所有线程都能执行到同一个`sync(g)`调用，否则是未定义行为。
58+ 
59+## 调用示例
60+ 
61+- SIMT编程场景：
62+ 
63+ ```cpp
64+ using namespace cooperative_groups;
65+ __global__ void simt_kernel(...)
66+ {
67+ ...
68+ thread_block block = this_thread_block();
69+ sync(block); // 等价于block.sync()
70+ 
71+ grid_group grid = this_grid();
72+ sync(grid); // 等价于grid.sync()
73+ ...
74+ }
75+ ```
76+ 
77+- SIMD与SIMT混合编程场景：
78+ 
79+ ```cpp
80+ using namespace cooperative_groups;
81+ __simt_vf__ inline void simt_kernel(...)
82+ {
83+ ...
84+ thread_block block = this_thread_block();
85+ sync(block); // 等价于block.sync()
86+ ...
87+ }
88+ ```
Mimpl/simt_api/cooperative_groups_impl.h+98-0
@@ -67,6 +67,21 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int __fns_internal(unsigned int m
67}67}
68namespace details {68namespace details {
69 69 
70+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
71+static constexpr uint64_t grid_sync_counter_offset_bytes = 4U * 8U;
72+static constexpr uint64_t grid_sync_counter_phase_bit = 1ULL << 63;
73+ 
74+__SIMT_DEVICE_FUNCTIONS_DECL__ inline uint64_t* get_grid_sync_counter()
75+{
76+ return reinterpret_cast<uint64_t*>(__cce_simt_get_para_base_imp() - grid_sync_counter_offset_bytes);
77+}
78+ 
79+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bool has_grid_sync_phase_flipped(uint64_t old_counter, uint64_t current_counter)
80+{
81+ return ((old_counter ^ current_counter) & grid_sync_counter_phase_bit) != 0U;
82+}
83+#endif
84+ 
70__SIMT_DEVICE_FUNCTIONS_DECL__ inline void wait_warp_fully_active()85__SIMT_DEVICE_FUNCTIONS_DECL__ inline void wait_warp_fully_active()
71{86{
72 while (asc_activemask() != 0xFFFFFFFFU) {87 while (asc_activemask() != 0xFFFFFFFFU) {
@@ -190,6 +205,11 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long thread_group::size() co
190 case group_type::coalesced_group_type: {205 case group_type::coalesced_group_type: {
191 return static_cast<const coalesced_group*>(this)->size();206 return static_cast<const coalesced_group*>(this)->size();
192 }207 }
208+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
209+ case group_type::grid_group_type: {
210+ return static_cast<const grid_group*>(this)->size();
211+ }
212+#endif
193 case group_type::tiled_group_type: {213 case group_type::tiled_group_type: {
194 return static_cast<const tiled_group*>(this)->size();214 return static_cast<const tiled_group*>(this)->size();
195 }215 }
@@ -210,6 +230,11 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long thread_group::thread_ra
210 case group_type::coalesced_group_type: {230 case group_type::coalesced_group_type: {
211 return static_cast<const coalesced_group*>(this)->thread_rank();231 return static_cast<const coalesced_group*>(this)->thread_rank();
212 }232 }
233+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
234+ case group_type::grid_group_type: {
235+ return static_cast<const grid_group*>(this)->thread_rank();
236+ }
237+#endif
213 case group_type::tiled_group_type: {238 case group_type::tiled_group_type: {
214 return static_cast<const tiled_group*>(this)->thread_rank();239 return static_cast<const tiled_group*>(this)->thread_rank();
215 }240 }
@@ -230,6 +255,12 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline void thread_group::sync() const
230 static_cast<const coalesced_group*>(this)->sync();255 static_cast<const coalesced_group*>(this)->sync();
231 break;256 break;
232 }257 }
258+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
259+ case group_type::grid_group_type: {
260+ static_cast<const grid_group*>(this)->sync();
261+ break;
262+ }
263+#endif
233 case group_type::tiled_group_type: {264 case group_type::tiled_group_type: {
234 static_cast<const tiled_group*>(this)->sync();265 static_cast<const tiled_group*>(this)->sync();
235 break;266 break;
@@ -350,6 +381,67 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_group thread_block::create_tiled_gr
350 381 
351__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block this_thread_block() { return thread_block(); }382__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block this_thread_block() { return thread_block(); }
352 383 
384+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
385+__SIMT_DEVICE_FUNCTIONS_DECL__ inline grid_group::grid_group() : thread_group(group_type::grid_group_type) {}
386+ 
387+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bool grid_group::is_valid() const { return true; }
388+ 
389+__SIMT_DEVICE_FUNCTIONS_DECL__ inline void grid_group::sync() const
390+{
391+ asc_syncthreads();
392+ unsigned long long block_count = num_blocks();
393+ if (thread_block::thread_rank() == 0U) {
394+ auto sync_counter = details::get_grid_sync_counter();
395+ uint64_t add_value = 1ULL;
396+ if (block_rank() == 0U) {
397+ add_value = details::grid_sync_counter_phase_bit - (static_cast<uint64_t>(block_count) - 1ULL);
398+ }
399+ uint64_t old_counter = __asc_simt_vf::asc_atomic_add(sync_counter, add_value);
W
Wwulinyu18 天前

这里为什么要加__asc_simt_vf::

likedislike
L
louyuxuan
18 天前 评论：
400+ auto volatile_counter = reinterpret_cast<volatile uint64_t*>(sync_counter);
401+ while (!details::has_grid_sync_phase_flipped(old_counter, *volatile_counter)) {
402+#ifndef ASCENDC_CPU_DEBUG
403+ asc_nop();
zhanghao_0689
zhanghao_0689zhanghao_068914 天前

这个地方为啥要隔离一下，不如做一个cpu的装函数算了？

likedislike
L
louyuxuan
14 天前 评论：
404+#endif
405+ }
406+ }
407+ asc_syncthreads();
zhanghao_0689
zhanghao_0689zhanghao_068914 天前

这个syncthreads是否有必要？

likedislike
L
louyuxuan
14 天前 评论：
408+}
409+ 
410+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long grid_group::num_blocks()
411+{
412+ return static_cast<unsigned long long>(gridDim.x) * gridDim.y * gridDim.z;
413+}
414+ 
415+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long grid_group::num_threads()
416+{
417+ return num_blocks() * static_cast<unsigned long long>(thread_block::num_threads());
418+}
419+ 
420+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long grid_group::size() { return num_threads(); }
421+ 
422+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long grid_group::block_rank()
423+{
424+ return blockIdx.x + static_cast<unsigned long long>(blockIdx.y) * gridDim.x +
425+ static_cast<unsigned long long>(blockIdx.z) * gridDim.x * gridDim.y;
426+}
427+ 
428+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long grid_group::thread_rank()
429+{
430+ return block_rank() * static_cast<unsigned long long>(thread_block::num_threads()) + thread_block::thread_rank();
431+}
432+ 
433+__SIMT_DEVICE_FUNCTIONS_DECL__ inline dim3 grid_group::dim_blocks() { return dim3(gridDim.x, gridDim.y, gridDim.z); }
434+ 
435+__SIMT_DEVICE_FUNCTIONS_DECL__ inline dim3 grid_group::group_dim() { return dim_blocks(); }
436+ 
437+__SIMT_DEVICE_FUNCTIONS_DECL__ inline dim3 grid_group::block_index()
438+{
439+ return dim3(blockIdx.x, blockIdx.y, blockIdx.z);
440+}
441+ 
442+__SIMT_DEVICE_FUNCTIONS_DECL__ inline grid_group this_grid() { return grid_group(); }
443+#endif
444+ 
353#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__) || defined(ASCENDC_CPU_DEBUG)445#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__) || defined(ASCENDC_CPU_DEBUG)
354template <unsigned int MaxBlockSize>446template <unsigned int MaxBlockSize>
355__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block this_thread_block(block_tile_memory<MaxBlockSize>& scratch)447__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block this_thread_block(block_tile_memory<MaxBlockSize>& scratch)
@@ -1141,6 +1233,12 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline coalesced_group __binary_partition_interna
1141 return result;1233 return result;
1142}1234}
1143 1235 
1236+template <typename GroupType>
1237+__SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync(const GroupType& g)
1238+{
1239+ g.sync();
1240+}
1241+ 
1144__SIMT_DEVICE_FUNCTIONS_DECL__ inline coalesced_group binary_partition(const coalesced_group& g, bool pred)1242__SIMT_DEVICE_FUNCTIONS_DECL__ inline coalesced_group binary_partition(const coalesced_group& g, bool pred)
1145{1243{
1146 return __binary_partition_internal(g, pred);1244 return __binary_partition_internal(g, pred);
Minclude/simt_api/cooperative_groups.h+35-0
@@ -97,11 +97,17 @@ struct alignas(details::multi_warp_scratch::communication_size) block_tile_memor
97 97 
98class _coalesced_group_data_access;98class _coalesced_group_data_access;
99class thread_block;99class thread_block;
100+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
101+class grid_group;
102+#endif
100 103 
101enum class group_type : unsigned int {104enum class group_type : unsigned int {
102 thread_block_type,105 thread_block_type,
103 tiled_group_type,106 tiled_group_type,
104 coalesced_group_type,107 coalesced_group_type,
108+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
109+ grid_group_type,
110+#endif
105};111};
106 112 
107class thread_group {113class thread_group {
@@ -124,6 +130,9 @@ protected:
124 } _tiled_info;130 } _tiled_info;
125 131 
126 friend class thread_block;132 friend class thread_block;
133+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
134+ friend class grid_group;
135+#endif
127 136 
128 __SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_group(group_type type);137 __SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_group(group_type type);
129};138};
@@ -189,6 +198,29 @@ private:
189 198 
190__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block this_thread_block();199__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block this_thread_block();
191 200 
201+#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)
202+class grid_group : public thread_group {
203+ friend __SIMT_DEVICE_FUNCTIONS_DECL__ inline grid_group this_grid();
204+ 
205+public:
206+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline bool is_valid() const;
207+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;
208+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline static unsigned long long thread_rank();
209+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline static unsigned long long block_rank();
210+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline static unsigned long long num_threads();
211+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline static unsigned long long size();
212+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline static unsigned long long num_blocks();
213+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline static dim3 dim_blocks();
214+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline static dim3 group_dim();
215+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline static dim3 block_index();
216+ 
217+private:
218+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline grid_group();
219+};
220+ 
221+__SIMT_DEVICE_FUNCTIONS_DECL__ inline grid_group this_grid();
222+#endif
223+ 
192#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__) || defined(ASCENDC_CPU_DEBUG)224#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__) || defined(ASCENDC_CPU_DEBUG)
193template <unsigned int MaxBlockSize>225template <unsigned int MaxBlockSize>
194__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block this_thread_block(block_tile_memory<MaxBlockSize>& scratch);226__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block this_thread_block(block_tile_memory<MaxBlockSize>& scratch);
@@ -519,6 +551,9 @@ struct tiled_partition_impl<Size, thread_block_tile<ParentSize, GrandParent>>
519template <unsigned int Size, typename ParentT>551template <unsigned int Size, typename ParentT>
520__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block_tile<Size, ParentT> tiled_partition(const ParentT& g);552__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block_tile<Size, ParentT> tiled_partition(const ParentT& g);
521 553 
554+template <typename GroupType>
555+__SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync(const GroupType& g);
556+ 
522__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_group tiled_partition(const thread_group& parent, unsigned int tilesz);557__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_group tiled_partition(const thread_group& parent, unsigned int tilesz);
523__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_group tiled_partition(const thread_block& parent, unsigned int tilesz);558__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_group tiled_partition(const thread_block& parent, unsigned int tilesz);
524__SIMT_DEVICE_FUNCTIONS_DECL__ inline coalesced_group tiled_partition(559__SIMT_DEVICE_FUNCTIONS_DECL__ inline coalesced_group tiled_partition(
Mtests/api/simt_api/ascendc_case_ascend950pr_9599_simt/test_operator_simt_c_cooperative_groups.cpp+119-0
@@ -11,12 +11,24 @@
11#include <mockcpp/mockcpp.hpp>11#include <mockcpp/mockcpp.hpp>
12#include <type_traits>12#include <type_traits>
13#include "simt_compiler_stub.h"13#include "simt_compiler_stub.h"
14+#include "simt_api/device_functions.h"
15+#include "simt_api/asc_fp16.h"
16+#include "simt_api/asc_bf16.h"
17+#include "utils/debug/asc_assert.h"
18+#define __NPU_COMPILER_INTERNAL_PURE_SIMT__
14#include "simt_api/cooperative_groups.h"19#include "simt_api/cooperative_groups.h"
20+#undef __NPU_COMPILER_INTERNAL_PURE_SIMT__
15 21 
16using namespace std;22using namespace std;
17using namespace AscendC;23using namespace AscendC;
18using namespace cooperative_groups;24using namespace cooperative_groups;
19 25 
26+namespace {
27+uint64_t* GetGridSyncCounter() { return reinterpret_cast<uint64_t*>(__cce_simt_get_para_base_imp() - 4U * 8U); }
28+ 
29+constexpr uint64_t GridSyncPhaseBit = 1ULL << 63;
30+} // namespace
31+ 
20class CooperativeGroupsTestsuite : public testing::Test {32class CooperativeGroupsTestsuite : public testing::Test {
21protected:33protected:
22 void SetUp() {}34 void SetUp() {}
@@ -56,6 +68,7 @@ TEST_F(CooperativeGroupsTestsuite, GroupTypeTest)
56 EXPECT_EQ(static_cast<unsigned int>(group_type::thread_block_type), 0u);68 EXPECT_EQ(static_cast<unsigned int>(group_type::thread_block_type), 0u);
57 EXPECT_EQ(static_cast<unsigned int>(group_type::tiled_group_type), 1u);69 EXPECT_EQ(static_cast<unsigned int>(group_type::tiled_group_type), 1u);
58 EXPECT_EQ(static_cast<unsigned int>(group_type::coalesced_group_type), 2u);70 EXPECT_EQ(static_cast<unsigned int>(group_type::coalesced_group_type), 2u);
71+ EXPECT_EQ(static_cast<unsigned int>(group_type::grid_group_type), 3u);
59}72}
60 73 
61TEST_F(CooperativeGroupsTestsuite, ThreadGroupGetTypeTest)74TEST_F(CooperativeGroupsTestsuite, ThreadGroupGetTypeTest)
@@ -70,6 +83,106 @@ TEST_F(CooperativeGroupsTestsuite, ThreadGroupGetTypeTest)
70 EXPECT_EQ(tg3.get_type(), group_type::tiled_group_type);83 EXPECT_EQ(tg3.get_type(), group_type::tiled_group_type);
71}84}
72 85 
86+TEST_F(CooperativeGroupsTestsuite, GridGroupThisGridTest)
87+{
88+ grid_group gg = this_grid();
89+ EXPECT_EQ(gg.get_type(), group_type::grid_group_type);
90+ EXPECT_TRUE(gg.is_valid());
91+}
92+ 
93+TEST_F(CooperativeGroupsTestsuite, ThreadGroupBaseClassGridGroupDispatchTest)
94+{
95+ SimtDimGuard guard(cce::dim3(8u, 4u, 2u), cce::dim3(7u, 3u, 1u), cce::dim3(3u, 2u, 1u), cce::dim3(4u, 3u, 2u));
96+ grid_group gg = this_grid();
97+ thread_group& tg = gg;
98+ 
99+ EXPECT_EQ(tg.get_type(), group_type::grid_group_type);
100+ EXPECT_EQ(tg.size(), grid_group::size());
101+ EXPECT_EQ(tg.num_threads(), grid_group::num_threads());
102+ EXPECT_EQ(tg.thread_rank(), grid_group::thread_rank());
103+}
104+ 
105+TEST_F(CooperativeGroupsTestsuite, GridGroupNumBlocksThreadsAndRank3DTest)
106+{
107+ SimtDimGuard guard(cce::dim3(8u, 4u, 2u), cce::dim3(7u, 3u, 1u), cce::dim3(3u, 2u, 1u), cce::dim3(4u, 3u, 2u));
108+ 
109+ EXPECT_EQ(grid_group::num_blocks(), 24ull);
110+ EXPECT_EQ(grid_group::num_threads(), 1536ull);
111+ EXPECT_EQ(grid_group::size(), 1536ull);
112+ EXPECT_EQ(grid_group::block_rank(), 23ull);
113+ EXPECT_EQ(grid_group::thread_rank(), 1535ull);
114+}
115+ 
116+TEST_F(CooperativeGroupsTestsuite, GridGroupDimBlocksAndBlockIndex3DTest)
117+{
118+ SimtDimGuard guard(cce::dim3(8u, 4u, 2u), cce::dim3(0u, 0u, 0u), cce::dim3(3u, 2u, 1u), cce::dim3(4u, 3u, 2u));
119+ 
120+ dim3 dimBlocks = grid_group::dim_blocks();
121+ EXPECT_EQ(dimBlocks.x, 4u);
122+ EXPECT_EQ(dimBlocks.y, 3u);
123+ EXPECT_EQ(dimBlocks.z, 2u);
124+ 
125+ dim3 groupDim = grid_group::group_dim();
126+ EXPECT_EQ(groupDim.x, 4u);
127+ EXPECT_EQ(groupDim.y, 3u);
128+ EXPECT_EQ(groupDim.z, 2u);
129+ 
130+ dim3 blockIndex = grid_group::block_index();
131+ EXPECT_EQ(blockIndex.x, 3u);
132+ EXPECT_EQ(blockIndex.y, 2u);
133+ EXPECT_EQ(blockIndex.z, 1u);
134+}
135+ 
136+TEST_F(CooperativeGroupsTestsuite, GridGroupSyncSingleBlockFlipsCounterPhaseTest)
137+{
138+ SimtDimGuard guard(cce::dim3(8u, 1u, 1u), cce::dim3(0u, 0u, 0u), cce::dim3(0u, 0u, 0u), cce::dim3(1u, 1u, 1u));
139+ GetGridSyncCounter()[0] = 123u;
140+ 
141+ sync(this_grid());
142+ 
143+ EXPECT_EQ(GetGridSyncCounter()[0], GridSyncPhaseBit + 123u);
144+}
145+ 
146+TEST_F(CooperativeGroupsTestsuite, GridGroupSyncBlockZeroReleasesCounterTest)
147+{
148+ SimtDimGuard guard(cce::dim3(8u, 1u, 1u), cce::dim3(0u, 0u, 0u), cce::dim3(0u, 0u, 0u), cce::dim3(2u, 1u, 1u));
149+ GetGridSyncCounter()[0] = 1u;
150+ 
151+ this_grid().sync();
152+ 
153+ EXPECT_EQ(GetGridSyncCounter()[0], GridSyncPhaseBit);
154+}
155+ 
156+TEST_F(CooperativeGroupsTestsuite, GridGroupSyncNonZeroBlockReleasesCounterTest)
157+{
158+ SimtDimGuard guard(cce::dim3(8u, 1u, 1u), cce::dim3(0u, 0u, 0u), cce::dim3(1u, 0u, 0u), cce::dim3(2u, 1u, 1u));
159+ GetGridSyncCounter()[0] = GridSyncPhaseBit - 1ULL;
160+ 
161+ this_grid().sync();
162+ 
163+ EXPECT_EQ(GetGridSyncCounter()[0], GridSyncPhaseBit);
164+}
165+ 
166+TEST_F(CooperativeGroupsTestsuite, GridGroupSyncBlockZeroReleasesCounterForThreeBlocksTest)
167+{
168+ SimtDimGuard guard(cce::dim3(8u, 1u, 1u), cce::dim3(0u, 0u, 0u), cce::dim3(0u, 0u, 0u), cce::dim3(3u, 1u, 1u));
169+ GetGridSyncCounter()[0] = 2u;
170+ 
171+ this_grid().sync();
172+ 
173+ EXPECT_EQ(GetGridSyncCounter()[0], GridSyncPhaseBit);
174+}
175+ 
176+TEST_F(CooperativeGroupsTestsuite, GridGroupSyncBlockZeroReleasesNextPhaseTest)
177+{
178+ SimtDimGuard guard(cce::dim3(8u, 1u, 1u), cce::dim3(0u, 0u, 0u), cce::dim3(0u, 0u, 0u), cce::dim3(2u, 1u, 1u));
179+ GetGridSyncCounter()[0] = GridSyncPhaseBit + 1ULL;
180+ 
181+ this_grid().sync();
182+ 
183+ EXPECT_EQ(GetGridSyncCounter()[0], 0u);
184+}
185+ 
73TEST_F(CooperativeGroupsTestsuite, ThreadBlockThisThreadBlockTest)186TEST_F(CooperativeGroupsTestsuite, ThreadBlockThisThreadBlockTest)
74{187{
75 thread_block tb = this_thread_block();188 thread_block tb = this_thread_block();
@@ -78,6 +191,12 @@ TEST_F(CooperativeGroupsTestsuite, ThreadBlockThisThreadBlockTest)
78 191 
79TEST_F(CooperativeGroupsTestsuite, ThreadBlockSyncTest) { thread_block::sync(); }192TEST_F(CooperativeGroupsTestsuite, ThreadBlockSyncTest) { thread_block::sync(); }
80 193 
194+TEST_F(CooperativeGroupsTestsuite, OuterSyncThreadBlockTest)
195+{
196+ thread_block tb = this_thread_block();
197+ sync(tb);
198+}
199+ 
81TEST_F(CooperativeGroupsTestsuite, ThreadBlockThreadRankTest)200TEST_F(CooperativeGroupsTestsuite, ThreadBlockThreadRankTest)
82{201{
83 unsigned int rank = thread_block::thread_rank();202 unsigned int rank = thread_block::thread_rank();
Mtests/api/simt_api/common_simt/simt_compiler_stub.cpp+6-0
@@ -58,3 +58,9 @@ bool CheckFunVecBinaryScalarImpl(VecBinaryScalarApiParams& chkParams, const uint
58bool CheckFunVecBinaryScalarImpl(VecBinaryScalarApiParams& chkParams, const char* intriName) { return true; }58bool CheckFunVecBinaryScalarImpl(VecBinaryScalarApiParams& chkParams, const char* intriName) { return true; }
59} // namespace check59} // namespace check
60} // namespace AscendC60} // namespace AscendC
61+ 
62+namespace {
63+static uint64_t g_gridSyncArgs[5] = {0U, 0U, 0U, 0U, 0U};
64+} // namespace
65+ 
66+uint64_t __cce_simt_get_para_base_imp() { return reinterpret_cast<uint64_t>(&g_gridSyncArgs[4]); }
Mtests/api/simt_api/common_simt/simt_compiler_stub.h+2-0
@@ -27,4 +27,6 @@ static int32_t laneid() { return 0; }
27 27 
28static int32_t lanemask_lt() { return 0xFFFFFFFF; }28static int32_t lanemask_lt() { return 0xFFFFFFFF; }
29 29 
30+uint64_t __cce_simt_get_para_base_imp();
31+ 
30#endif32#endif