Star1166
1191
代码
代码
Issues147
Pull Requests386
Actions
项目讨论
Wiki
分析
Star1166
1191
已合并
修复SIMT协作组接口返回类型错误 #5329
louyuxuan创建于 8月25日
修复SIMT协作组接口返回类型错误 #5329
已合并
louyuxuan创建于 8月25日
从louyuxuan/asc-devkitdocs合入到cann/asc-devkitmaster
共 19 个文件变更+89-94
Mdocs/zh/api/SIMT-API/cooperative_groups/coalesced_group/all.md+2-2
@@ -60,7 +60,7 @@ int all(int predicate) const
60 ...60 ...
61 if (threadIdx.x % 2 == 0) {61 if (threadIdx.x % 2 == 0) {
62 coalesced_group active = coalesced_threads();62 coalesced_group active = coalesced_threads();
63- uint32_t result = active.all(1); // 返回163+ int result = active.all(1); // 返回1
W
Wwulinyu8月26日

修改涉及到API对外返回值的变动 需要和算子确认下，当前是否有使用

likedislike
L
louyuxuan
8月26日 评论：
64 }64 }
65 ...65 ...
66 }66 }
@@ -75,7 +75,7 @@ int all(int predicate) const
75 ...75 ...
76 if (threadIdx.x % 2 == 0) {76 if (threadIdx.x % 2 == 0) {
77 coalesced_group active = coalesced_threads();77 coalesced_group active = coalesced_threads();
78- uint32_t result = active.all(1); // 返回178+ int result = active.all(1); // 返回1
79 }79 }
80 ...80 ...
81 }81 }
Mdocs/zh/api/SIMT-API/cooperative_groups/coalesced_group/any.md+2-2
@@ -60,7 +60,7 @@ int any(int predicate) const
60 ...60 ...
61 if (threadIdx.x % 2 == 0) {61 if (threadIdx.x % 2 == 0) {
62 coalesced_group active = coalesced_threads();62 coalesced_group active = coalesced_threads();
63- uint32_t result = active.any(1); // 返回163+ int result = active.any(1); // 返回1
64 }64 }
65 ...65 ...
66 }66 }
@@ -75,7 +75,7 @@ int any(int predicate) const
75 ...75 ...
76 if (threadIdx.x % 2 == 0) {76 if (threadIdx.x % 2 == 0) {
77 coalesced_group active = coalesced_threads();77 coalesced_group active = coalesced_threads();
78- uint32_t result = active.any(1); // 返回178+ int result = active.any(1); // 返回1
79 }79 }
80 ...80 ...
81 }81 }
Mdocs/zh/api/SIMT-API/cooperative_groups/coalesced_group/coalesced_group_intro.md+5-5
@@ -9,10 +9,10 @@
9 9 
10```cpp10```cpp
11void sync() const;11void sync() const;
12-unsigned long long num_threads() const;12+unsigned int num_threads() const;
13-unsigned long long thread_rank() const;13+unsigned int thread_rank() const;
14-unsigned long long meta_group_size() const;14+unsigned int meta_group_size() const;
15-unsigned long long meta_group_rank() const;15+unsigned int meta_group_rank() const;
16template <typename T>16template <typename T>
17T shfl(T var, int src_rank) const;17T shfl(T var, int src_rank) const;
18template <typename T>18template <typename T>
@@ -22,5 +22,5 @@ T shfl_down(T var, unsigned int delta) const;
22int any(int predicate) const;22int any(int predicate) const;
23int all(int predicate) const;23int all(int predicate) const;
24unsigned int ballot(int predicate) const;24unsigned int ballot(int predicate) const;
25-unsigned long long size() const;25+unsigned int size() const;
26```26```
Mdocs/zh/api/SIMT-API/cooperative_groups/coalesced_group/meta_group_rank.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long meta_group_rank() const33+unsigned int meta_group_rank() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -58,7 +58,7 @@ unsigned long long meta_group_rank() const
58 ...58 ...
59 if (threadIdx.x % 2 == 0) {59 if (threadIdx.x % 2 == 0) {
60 coalesced_group active = coalesced_threads();60 coalesced_group active = coalesced_threads();
61- unsigned long long thread_num = active.meta_group_rank(); // 返回061+ unsigned int group_rank = active.meta_group_rank(); // 返回0
62 }62 }
63 ...63 ...
64 }64 }
@@ -73,7 +73,7 @@ unsigned long long meta_group_rank() const
73 ...73 ...
74 if (threadIdx.x % 2 == 0) {74 if (threadIdx.x % 2 == 0) {
75 coalesced_group active = coalesced_threads();75 coalesced_group active = coalesced_threads();
76- unsigned long long thread_num = active.meta_group_rank(); // 返回076+ unsigned int group_rank = active.meta_group_rank(); // 返回0
77 }77 }
78 ...78 ...
79 }79 }
Mdocs/zh/api/SIMT-API/cooperative_groups/coalesced_group/meta_group_size.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long meta_group_size() const33+unsigned int meta_group_size() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -58,7 +58,7 @@ unsigned long long meta_group_size() const
58 ...58 ...
59 if (threadIdx.x % 2 == 0) {59 if (threadIdx.x % 2 == 0) {
60 coalesced_group active = coalesced_threads();60 coalesced_group active = coalesced_threads();
61- unsigned long long thread_num = active.meta_group_size(); // 返回161+ unsigned int group_size = active.meta_group_size(); // 返回1
62 }62 }
63 ...63 ...
64 }64 }
@@ -73,7 +73,7 @@ unsigned long long meta_group_size() const
73 ...73 ...
74 if (threadIdx.x % 2 == 0) {74 if (threadIdx.x % 2 == 0) {
75 coalesced_group active = coalesced_threads();75 coalesced_group active = coalesced_threads();
76- unsigned long long thread_num = active.meta_group_size(); // 返回176+ unsigned int group_size = active.meta_group_size(); // 返回1
77 }77 }
78 ...78 ...
79 }79 }
Mdocs/zh/api/SIMT-API/cooperative_groups/coalesced_group/num_threads.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long num_threads() const33+unsigned int num_threads() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -58,7 +58,7 @@ unsigned long long num_threads() const
58 ...58 ...
59 if (threadIdx.x % 2 == 0) {59 if (threadIdx.x % 2 == 0) {
60 coalesced_group active = coalesced_threads();60 coalesced_group active = coalesced_threads();
61- unsigned long long thread_num = active.num_threads(); // 返回1661+ unsigned int thread_num = active.num_threads(); // 返回16
62 }62 }
63 ...63 ...
64 }64 }
@@ -73,7 +73,7 @@ unsigned long long num_threads() const
73 ...73 ...
74 if (threadIdx.x % 2 == 0) {74 if (threadIdx.x % 2 == 0) {
75 coalesced_group active = coalesced_threads();75 coalesced_group active = coalesced_threads();
76- unsigned long long thread_num = active.num_threads(); // 返回1676+ unsigned int thread_num = active.num_threads(); // 返回16
77 }77 }
78 ...78 ...
79 }79 }
Mdocs/zh/api/SIMT-API/cooperative_groups/coalesced_group/size.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long size() const33+unsigned int size() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -56,7 +56,7 @@ unsigned long long size() const
56 ...56 ...
57 if (threadIdx.x % 2 == 0) {57 if (threadIdx.x % 2 == 0) {
58 coalesced_group active = coalesced_threads();58 coalesced_group active = coalesced_threads();
59- unsigned long long thread_num = active.size(); // 返回1659+ unsigned int thread_num = active.size(); // 返回16
60 }60 }
61 ...61 ...
62 }62 }
@@ -71,7 +71,7 @@ unsigned long long size() const
71 ...71 ...
72 if (threadIdx.x % 2 == 0) {72 if (threadIdx.x % 2 == 0) {
73 coalesced_group active = coalesced_threads();73 coalesced_group active = coalesced_threads();
74- unsigned long long thread_num = active.size(); // 返回1674+ unsigned int thread_num = active.size(); // 返回16
75 }75 }
76 ...76 ...
77 }77 }
Mdocs/zh/api/SIMT-API/cooperative_groups/coalesced_group/thread_rank.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long thread_rank() const33+unsigned int thread_rank() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -61,7 +61,7 @@ unsigned long long thread_rank() const
61 ...61 ...
62 if (threadIdx.x % 2 == 0) {62 if (threadIdx.x % 2 == 0) {
63 coalesced_group active = coalesced_threads();63 coalesced_group active = coalesced_threads();
64- unsigned long long rank = active.thread_rank();64+ unsigned int rank = active.thread_rank();
65 }65 }
66 ...66 ...
67 }67 }
@@ -76,7 +76,7 @@ unsigned long long thread_rank() const
76 ...76 ...
77 if (threadIdx.x % 2 == 0) {77 if (threadIdx.x % 2 == 0) {
78 coalesced_group active = coalesced_threads();78 coalesced_group active = coalesced_threads();
79- unsigned long long rank = active.thread_rank();79+ unsigned int rank = active.thread_rank();
80 }80 }
81 ...81 ...
82 }82 }
Mdocs/zh/api/SIMT-API/cooperative_groups/thread_block_tile/all.md+2-2
@@ -60,7 +60,7 @@ int all(int predicate) const
60 ...60 ...
61 thread_block block = this_thread_block();61 thread_block block = this_thread_block();
62 auto tile4 = tiled_partition<4>(block);62 auto tile4 = tiled_partition<4>(block);
63- uint32_t result = tile4.all(1); // 返回163+ int result = tile4.all(1); // 返回1
64 ...64 ...
65 }65 }
66 ```66 ```
@@ -74,7 +74,7 @@ int all(int predicate) const
74 ...74 ...
75 thread_block block = this_thread_block();75 thread_block block = this_thread_block();
76 auto tile4 = tiled_partition<4>(block);76 auto tile4 = tiled_partition<4>(block);
77- uint32_t result = tile4.all(1); // 返回177+ int result = tile4.all(1); // 返回1
78 ...78 ...
79 }79 }
80 ```80 ```
Mdocs/zh/api/SIMT-API/cooperative_groups/thread_block_tile/any.md+2-2
@@ -60,7 +60,7 @@ int any(int predicate) const
60 ...60 ...
61 thread_block block = this_thread_block();61 thread_block block = this_thread_block();
62 auto tile4 = tiled_partition<4>(block);62 auto tile4 = tiled_partition<4>(block);
63- uint32_t result = tile4.any(1); // 返回163+ int result = tile4.any(1); // 返回1
64 ...64 ...
65 }65 }
66 ```66 ```
@@ -74,7 +74,7 @@ int any(int predicate) const
74 ...74 ...
75 thread_block block = this_thread_block();75 thread_block block = this_thread_block();
76 auto tile4 = tiled_partition<4>(block);76 auto tile4 = tiled_partition<4>(block);
77- uint32_t result = tile4.any(1); // 返回177+ int result = tile4.any(1); // 返回1
78 ...78 ...
79 }79 }
80 ```80 ```
Mdocs/zh/api/SIMT-API/cooperative_groups/thread_block_tile/meta_group_rank.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long meta_group_rank() const33+unsigned int meta_group_rank() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -56,7 +56,7 @@ unsigned long long meta_group_rank() const
56 ...56 ...
57 thread_block block = this_thread_block();57 thread_block block = this_thread_block();
58 auto tile4 = tiled_partition<4>(block);58 auto tile4 = tiled_partition<4>(block);
59- unsigned long long group_rank = tile4.meta_group_rank();59+ unsigned int group_rank = tile4.meta_group_rank();
60 ...60 ...
61 }61 }
62 ```62 ```
@@ -70,7 +70,7 @@ unsigned long long meta_group_rank() const
70 ...70 ...
71 thread_block block = this_thread_block();71 thread_block block = this_thread_block();
72 auto tile4 = tiled_partition<4>(block);72 auto tile4 = tiled_partition<4>(block);
73- unsigned long long group_rank = tile4.meta_group_rank();73+ unsigned int group_rank = tile4.meta_group_rank();
74 ...74 ...
75 }75 }
76 ```76 ```
Mdocs/zh/api/SIMT-API/cooperative_groups/thread_block_tile/meta_group_size.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long meta_group_size() const33+unsigned int meta_group_size() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -58,7 +58,7 @@ unsigned long long meta_group_size() const
58 ...58 ...
59 thread_block block = this_thread_block();59 thread_block block = this_thread_block();
60 auto tile4 = tiled_partition<4>(block);60 auto tile4 = tiled_partition<4>(block);
61- unsigned long long group_size = tile4.meta_group_size(); // 返回25661+ unsigned int group_size = tile4.meta_group_size(); // 返回256
62 ...62 ...
63 }63 }
64 ```64 ```
@@ -72,7 +72,7 @@ unsigned long long meta_group_size() const
72 ...72 ...
73 thread_block block = this_thread_block();73 thread_block block = this_thread_block();
74 auto tile4 = tiled_partition<4>(block);74 auto tile4 = tiled_partition<4>(block);
75- unsigned long long group_size = tile4.meta_group_size(); // 返回25675+ unsigned int group_size = tile4.meta_group_size(); // 返回256
76 ...76 ...
77 }77 }
78 ```78 ```
Mdocs/zh/api/SIMT-API/cooperative_groups/thread_block_tile/num_threads.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long num_threads() const33+unsigned int num_threads() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -56,7 +56,7 @@ unsigned long long num_threads() const
56 ...56 ...
57 thread_block block = this_thread_block();57 thread_block block = this_thread_block();
58 auto tile4 = tiled_partition<4>(block);58 auto tile4 = tiled_partition<4>(block);
59- unsigned long long thread_num = tile4.num_threads(); // 返回459+ unsigned int thread_num = tile4.num_threads(); // 返回4
60 ...60 ...
61 }61 }
62 ```62 ```
@@ -70,7 +70,7 @@ unsigned long long num_threads() const
70 ...70 ...
71 thread_block block = this_thread_block();71 thread_block block = this_thread_block();
72 auto tile4 = tiled_partition<4>(block);72 auto tile4 = tiled_partition<4>(block);
73- unsigned long long thread_num = tile4.num_threads(); // 返回473+ unsigned int thread_num = tile4.num_threads(); // 返回4
74 ...74 ...
75 }75 }
76 ```76 ```
Mdocs/zh/api/SIMT-API/cooperative_groups/thread_block_tile/size.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long size() const33+unsigned int size() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -56,7 +56,7 @@ unsigned long long size() const
56 ...56 ...
57 thread_block block = this_thread_block();57 thread_block block = this_thread_block();
58 auto tile4 = tiled_partition<4>(block);58 auto tile4 = tiled_partition<4>(block);
59- unsigned long long thread_num = tile4.size(); // 返回459+ unsigned int thread_num = tile4.size(); // 返回4
60 ...60 ...
61 }61 }
62 ```62 ```
@@ -70,7 +70,7 @@ unsigned long long size() const
70 ...70 ...
71 thread_block block = this_thread_block();71 thread_block block = this_thread_block();
72 auto tile4 = tiled_partition<4>(block);72 auto tile4 = tiled_partition<4>(block);
73- unsigned long long thread_num = tile4.size(); // 返回473+ unsigned int thread_num = tile4.size(); // 返回4
74 ...74 ...
75 }75 }
76 ```76 ```
Mdocs/zh/api/SIMT-API/cooperative_groups/thread_block_tile/thread_block_tile_intro.md+5-5
@@ -9,10 +9,10 @@
9 9 
10```cpp10```cpp
11void sync() const;11void sync() const;
12-unsigned long long num_threads() const;12+unsigned int num_threads() const;
13-unsigned long long thread_rank() const;13+unsigned int thread_rank() const;
14-unsigned long long meta_group_size() const;14+unsigned int meta_group_size() const;
15-unsigned long long meta_group_rank() const;15+unsigned int meta_group_rank() const;
16template <typename T>16template <typename T>
17T shfl(T var, int src_rank) const;17T shfl(T var, int src_rank) const;
18template <typename T>18template <typename T>
@@ -24,7 +24,7 @@ T shfl_xor(T var, unsigned int lane_mask) const;
24int any(int predicate) const;24int any(int predicate) const;
25int all(int predicate) const;25int all(int predicate) const;
26unsigned int ballot(int predicate) const;26unsigned int ballot(int predicate) const;
27-unsigned long long size() const;27+unsigned int size() const;
28```28```
29 29 
30## 接口支持范围30## 接口支持范围
Mdocs/zh/api/SIMT-API/cooperative_groups/thread_block_tile/thread_rank.md+3-3
@@ -30,7 +30,7 @@
30## 函数原型30## 函数原型
31 31 
32```cpp32```cpp
33-unsigned long long thread_rank() const33+unsigned int thread_rank() const
34```34```
35 35 
36## 参数说明36## 参数说明
@@ -61,7 +61,7 @@ unsigned long long thread_rank() const
61 ...61 ...
62 thread_block block = this_thread_block();62 thread_block block = this_thread_block();
63 auto tile4 = tiled_partition<4>(block);63 auto tile4 = tiled_partition<4>(block);
64- unsigned long long rank = tile4.thread_rank();64+ unsigned int rank = tile4.thread_rank();
65 ...65 ...
66 }66 }
67 ```67 ```
@@ -75,7 +75,7 @@ unsigned long long thread_rank() const
75 ...75 ...
76 thread_block block = this_thread_block();76 thread_block block = this_thread_block();
77 auto tile4 = tiled_partition<4>(block);77 auto tile4 = tiled_partition<4>(block);
78- unsigned long long rank = tile4.thread_rank();78+ unsigned int rank = tile4.thread_rank();
79 ...79 ...
80 }80 }
81 ```81 ```
Mimpl/simt_api/cooperative_groups_impl.h+19-24
@@ -280,16 +280,13 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline tiled_group::tiled_group(unsigned int num_
280 280 
281__SIMT_DEVICE_FUNCTIONS_DECL__ inline void tiled_group::sync() const { asc_threadfence_block(); }281__SIMT_DEVICE_FUNCTIONS_DECL__ inline void tiled_group::sync() const { asc_threadfence_block(); }
282 282 
283-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long tiled_group::num_threads() const283+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int tiled_group::num_threads() const { return _tiled_info.num_threads; }
284-{
285- return static_cast<unsigned long long>(_tiled_info.num_threads);
286-}
287 284 
288-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long tiled_group::size() const { return num_threads(); }285+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int tiled_group::size() const { return num_threads(); }
289 286 
290-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long tiled_group::thread_rank() const287+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int tiled_group::thread_rank() const
291{288{
292- return static_cast<unsigned long long>(__popc(_tiled_info.mask & lanemask_lt()));289+ return __popc(_tiled_info.mask & lanemask_lt());
293}290}
294 291 
295__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block::thread_block()292__SIMT_DEVICE_FUNCTIONS_DECL__ inline thread_block::thread_block()
@@ -541,26 +538,26 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline void coalesced_group::sync() const { asc_t
541 538 
542__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int coalesced_group::get_mask() const { return _tiled_info.mask; }539__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int coalesced_group::get_mask() const { return _tiled_info.mask; }
543 540 
544-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long coalesced_group::num_threads() const541+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int coalesced_group::num_threads() const
545{542{
546- return static_cast<unsigned long long>(_tiled_info.num_threads);543+ return _tiled_info.num_threads;
547}544}
548 545 
549-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long coalesced_group::size() const { return num_threads(); }546+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int coalesced_group::size() const { return num_threads(); }
550 547 
551-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long coalesced_group::thread_rank() const548+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int coalesced_group::thread_rank() const
552{549{
553- return static_cast<unsigned long long>(__popc(_tiled_info.mask & lanemask_lt()));550+ return __popc(_tiled_info.mask & lanemask_lt());
554}551}
555 552 
556-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long coalesced_group::meta_group_size() const553+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int coalesced_group::meta_group_size() const
557{554{
558- return static_cast<unsigned long long>(_tiled_info.meta_group_size);555+ return _tiled_info.meta_group_size;
559}556}
560 557 
561-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long coalesced_group::meta_group_rank() const558+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int coalesced_group::meta_group_rank() const
562{559{
563- return static_cast<unsigned long long>(_tiled_info.meta_group_rank);560+ return _tiled_info.meta_group_rank;
564}561}
565 562 
566template <typename T>563template <typename T>
@@ -703,20 +700,19 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline void single_warp_thread_block_tile_base<Si
703}700}
704 701 
705template <unsigned int Size>702template <unsigned int Size>
706-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long single_warp_thread_block_tile_base<Size>::thread_rank()703+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int single_warp_thread_block_tile_base<Size>::thread_rank()
707{704{
708 return laneid() & (Size - 1);705 return laneid() & (Size - 1);
709}706}
710 707 
711template <unsigned int Size>708template <unsigned int Size>
712-__SIMT_DEVICE_FUNCTIONS_DECL__ constexpr inline unsigned long long709+__SIMT_DEVICE_FUNCTIONS_DECL__ constexpr inline unsigned int single_warp_thread_block_tile_base<Size>::num_threads()
713-single_warp_thread_block_tile_base<Size>::num_threads()
714{710{
715 return numThreads;711 return numThreads;
716}712}
717 713 
718template <unsigned int Size>714template <unsigned int Size>
719-__SIMT_DEVICE_FUNCTIONS_DECL__ constexpr inline unsigned long long single_warp_thread_block_tile_base<Size>::size()715+__SIMT_DEVICE_FUNCTIONS_DECL__ constexpr inline unsigned int single_warp_thread_block_tile_base<Size>::size()
720{716{
721 return numThreads;717 return numThreads;
722}718}
@@ -748,20 +744,19 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline void multi_warp_thread_block_tile_base<Siz
748}744}
749 745 
750template <unsigned int Size>746template <unsigned int Size>
751-__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long multi_warp_thread_block_tile_base<Size>::thread_rank()747+__SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int multi_warp_thread_block_tile_base<Size>::thread_rank()
752{748{
753 return thread_block::thread_rank() & (Size - 1);749 return thread_block::thread_rank() & (Size - 1);
754}750}
755 751 
756template <unsigned int Size>752template <unsigned int Size>
757-__SIMT_DEVICE_FUNCTIONS_DECL__ constexpr inline unsigned long long753+__SIMT_DEVICE_FUNCTIONS_DECL__ constexpr inline unsigned int multi_warp_thread_block_tile_base<Size>::num_threads()
758-multi_warp_thread_block_tile_base<Size>::num_threads()
759{754{
760 return numThreads;755 return numThreads;
761}756}
762 757 
763template <unsigned int Size>758template <unsigned int Size>
764-__SIMT_DEVICE_FUNCTIONS_DECL__ constexpr inline unsigned long long multi_warp_thread_block_tile_base<Size>::size()759+__SIMT_DEVICE_FUNCTIONS_DECL__ constexpr inline unsigned int multi_warp_thread_block_tile_base<Size>::size()
765{760{
766 return numThreads;761 return numThreads;
767}762}
Minclude/simt_api/cooperative_groups.h+14-14
@@ -146,9 +146,9 @@ class tiled_group : public thread_group {
146 146 
147public:147public:
148 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;148 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;
149- __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long num_threads() const;149+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int num_threads() const;
150- __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long size() const;150+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int size() const;
151- __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long thread_rank() const;151+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int thread_rank() const;
152 152 
153protected:153protected:
154 __SIMT_DEVICE_FUNCTIONS_DECL__ inline tiled_group(unsigned int num_threads);154 __SIMT_DEVICE_FUNCTIONS_DECL__ inline tiled_group(unsigned int num_threads);
@@ -257,11 +257,11 @@ class coalesced_group : public thread_group {
257 257 
258public:258public:
259 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;259 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;
260- __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long num_threads() const;260+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int num_threads() const;
261- __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long size() const;261+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int size() const;
262- __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long thread_rank() const;262+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int thread_rank() const;
263- __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long meta_group_size() const;263+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int meta_group_size() const;
264- __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned long long meta_group_rank() const;264+ __SIMT_DEVICE_FUNCTIONS_DECL__ inline unsigned int meta_group_rank() const;
265 265 
266 template <typename T>266 template <typename T>
267 __SIMT_DEVICE_FUNCTIONS_DECL__ inline T shfl(T var, int src_rank) const;267 __SIMT_DEVICE_FUNCTIONS_DECL__ inline T shfl(T var, int src_rank) const;
@@ -354,9 +354,9 @@ class single_warp_thread_block_tile_base {
354 354 
355public:355public:
356 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;356 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;
357- __SIMT_DEVICE_FUNCTIONS_DECL__ static inline unsigned long long thread_rank();357+ __SIMT_DEVICE_FUNCTIONS_DECL__ static inline unsigned int thread_rank();
358- __SIMT_DEVICE_FUNCTIONS_DECL__ static constexpr inline unsigned long long num_threads();358+ __SIMT_DEVICE_FUNCTIONS_DECL__ static constexpr inline unsigned int num_threads();
359- __SIMT_DEVICE_FUNCTIONS_DECL__ static constexpr inline unsigned long long size();359+ __SIMT_DEVICE_FUNCTIONS_DECL__ static constexpr inline unsigned int size();
360 360 
361 template <typename T>361 template <typename T>
362 __SIMT_DEVICE_FUNCTIONS_DECL__ inline T shfl(T var, int src_rank) const;362 __SIMT_DEVICE_FUNCTIONS_DECL__ inline T shfl(T var, int src_rank) const;
@@ -395,9 +395,9 @@ class multi_warp_thread_block_tile_base {
395 395 
396public:396public:
397 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;397 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void sync() const;
398- __SIMT_DEVICE_FUNCTIONS_DECL__ static inline unsigned long long thread_rank();398+ __SIMT_DEVICE_FUNCTIONS_DECL__ static inline unsigned int thread_rank();
399- __SIMT_DEVICE_FUNCTIONS_DECL__ static constexpr inline unsigned long long num_threads();399+ __SIMT_DEVICE_FUNCTIONS_DECL__ static constexpr inline unsigned int num_threads();
400- __SIMT_DEVICE_FUNCTIONS_DECL__ static constexpr inline unsigned long long size();400+ __SIMT_DEVICE_FUNCTIONS_DECL__ static constexpr inline unsigned int size();
401 401 
402 template <typename T>402 template <typename T>
403 __SIMT_DEVICE_FUNCTIONS_DECL__ inline T shfl(T var, int src_rank) const;403 __SIMT_DEVICE_FUNCTIONS_DECL__ inline T shfl(T var, int src_rank) const;
Mtests/api/simt_api/ascendc_case_ascend950pr_9599_simt/test_operator_simt_c_cooperative_groups.cpp+8-8
@@ -1166,7 +1166,7 @@ TEST_F(CooperativeGroupsTestsuite, ThreadBlockTileConversionTest)
1166 thread_block tb = this_thread_block();1166 thread_block tb = this_thread_block();
1167 auto tile4 = tiled_partition<4>(tb);1167 auto tile4 = tiled_partition<4>(tb);
1168 thread_block_tile<4, void> tile_void = tile4;1168 thread_block_tile<4, void> tile_void = tile4;
1169- EXPECT_EQ(tile_void.size(), 4ull);1169+ EXPECT_EQ(tile_void.size(), 4u);
1170}1170}
1171 1171 
1172TEST_F(CooperativeGroupsTestsuite, NestedThreadBlockTileConversionTest)1172TEST_F(CooperativeGroupsTestsuite, NestedThreadBlockTileConversionTest)
@@ -1320,16 +1320,16 @@ TEST_F(CooperativeGroupsTestsuite, ThreadBlockTileConstructorTest)
1320{1320{
1321 thread_block tb = this_thread_block();1321 thread_block tb = this_thread_block();
1322 thread_block_tile<4, thread_block> tile4(tb);1322 thread_block_tile<4, thread_block> tile4(tb);
1323- EXPECT_EQ(tile4.size(), 4ull);1323+ EXPECT_EQ(tile4.size(), 4u);
1324- EXPECT_EQ(tile4.num_threads(), 4ull);1324+ EXPECT_EQ(tile4.num_threads(), 4u);
1325}1325}
1326 1326 
1327TEST_F(CooperativeGroupsTestsuite, ThreadBlockTileImplVoidConstructorCopiesMetadataTest)1327TEST_F(CooperativeGroupsTestsuite, ThreadBlockTileImplVoidConstructorCopiesMetadataTest)
1328{1328{
1329 thread_block_tile_impl<4, void> tile4(2u, 8u);1329 thread_block_tile_impl<4, void> tile4(2u, 8u);
1330 EXPECT_EQ(tile4.get_type(), group_type::tiled_group_type);1330 EXPECT_EQ(tile4.get_type(), group_type::tiled_group_type);
1331- EXPECT_EQ(tile4.size(), 4ull);1331+ EXPECT_EQ(tile4.size(), 4u);
1332- EXPECT_EQ(tile4.num_threads(), 4ull);1332+ EXPECT_EQ(tile4.num_threads(), 4u);
1333 EXPECT_EQ(tile4.meta_group_rank(), 2u);1333 EXPECT_EQ(tile4.meta_group_rank(), 2u);
1334 EXPECT_EQ(tile4.meta_group_size(), 8u);1334 EXPECT_EQ(tile4.meta_group_size(), 8u);
1335}1335}
@@ -1338,9 +1338,9 @@ TEST_F(CooperativeGroupsTestsuite, ThreadBlockTileFromThreadBlockTest)
1338{1338{
1339 thread_block tb = this_thread_block();1339 thread_block tb = this_thread_block();
1340 auto tile8 = tiled_partition<8>(tb);1340 auto tile8 = tiled_partition<8>(tb);
1341- EXPECT_EQ(tile8.size(), 8ull);1341+ EXPECT_EQ(tile8.size(), 8u);
1342- EXPECT_EQ(tile8.num_threads(), 8ull);1342+ EXPECT_EQ(tile8.num_threads(), 8u);
1343- EXPECT_GE(tile8.thread_rank(), 0ull);1343+ EXPECT_GE(tile8.thread_rank(), 0u);
1344 EXPECT_LT(tile8.thread_rank(), tile8.size());1344 EXPECT_LT(tile8.thread_rank(), tile8.size());
1345}1345}
1346 1346