codekingpro/portable-devtools
114k
1/*
2 * Copyright 1993-2021 NVIDIA Corporation. All rights reserved.
3 *
4 * NOTICE TO LICENSEE:
5 *
6 * This source code and/or documentation ("Licensed Deliverables") are
7 * subject to NVIDIA intellectual property rights under U.S. and
8 * international Copyright laws.
9 *
10 * These Licensed Deliverables contained herein is PROPRIETARY and
11 * CONFIDENTIAL to NVIDIA and is being provided under the terms and
12 * conditions of a form of NVIDIA software license agreement by and
13 * between NVIDIA and Licensee ("License Agreement") or electronically
14 * accepted by Licensee. Notwithstanding any terms or conditions to
15 * the contrary in the License Agreement, reproduction or disclosure
16 * of the Licensed Deliverables to any third party without the express
17 * written consent of NVIDIA is prohibited.
18 *
19 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
20 * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
21 * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
22 * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
23 * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
24 * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
25 * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
26 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
27 * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
28 * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
29 * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
30 * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
31 * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
32 * OF THESE LICENSED DELIVERABLES.
33 *
34 * U.S. Government End Users. These Licensed Deliverables are a
35 * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
36 * 1995), consisting of "commercial computer software" and "commercial
37 * computer software documentation" as such terms are used in 48
38 * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
39 * only as a commercial end item. Consistent with 48 C.F.R.12.212 and
40 * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
41 * U.S. Government End Users acquire the Licensed Deliverables with
42 * only those rights set forth herein.
43 *
44 * Any use of the Licensed Deliverables in individual and commercial
45 * software must include, in the user documentation and internal
46 * comments to the code, the above Disclaimer and U.S. Government End
47 * Users Notice.
48 */
49
50#ifndef _COOPERATIVE_GROUPS_H_
51#define _COOPERATIVE_GROUPS_H_
52
53#if defined(__cplusplus) && defined(__CUDACC__)
54
55#include "cooperative_groups/details/info.h"
56#include "cooperative_groups/details/driver_abi.h"
57#include "cooperative_groups/details/helpers.h"
58#include "cooperative_groups/details/memory.h"
59
60#if defined(_CG_HAS_STL_ATOMICS)
61#include <cuda/atomic>
62#define _CG_THREAD_SCOPE(scope) _CG_STATIC_CONST_DECL cuda::thread_scope thread_scope = scope;
63#else
64#define _CG_THREAD_SCOPE(scope)
65#endif
66
67_CG_BEGIN_NAMESPACE
68
69namespace details {
70 _CG_CONST_DECL unsigned int coalesced_group_id = 1;
71 _CG_CONST_DECL unsigned int multi_grid_group_id = 2;
72 _CG_CONST_DECL unsigned int grid_group_id = 3;
73 _CG_CONST_DECL unsigned int thread_block_id = 4;
74 _CG_CONST_DECL unsigned int multi_tile_group_id = 5;
75 _CG_CONST_DECL unsigned int cluster_group_id = 6;
76}
77
78/**
79 * class thread_group;
80 *
81 * Generic thread group type, into which all groups are convertible.
82 * It acts as a container for all storage necessary for the derived groups,
83 * and will dispatch the API calls to the correct derived group. This means
84 * that all derived groups must implement the same interface as thread_group.
85 */
86class thread_group
87{
88protected:
89 struct group_data {
90 unsigned int _unused : 1;
91 unsigned int type : 7, : 0;
92 };
93
94 struct gg_data {
95 details::grid_workspace *gridWs;
96 };
97
98#if defined(_CG_CPP11_FEATURES) && defined(_CG_ABI_EXPERIMENTAL)
99 struct mg_data {
100 unsigned long long _unused : 1;
101 unsigned long long type : 7;
102 unsigned long long handle : 56;
103 const details::multi_grid::multi_grid_functions *functions;
104 };
105#endif
106
107 struct tg_data {
108 unsigned int is_tiled : 1;
109 unsigned int type : 7;
110 unsigned int size : 24;
111 // packed to 4b
112 unsigned int metaGroupSize : 16;
113 unsigned int metaGroupRank : 16;
114 // packed to 8b
115 unsigned int mask;
116 // packed to 12b
117 unsigned int _res;
118 };
119
120 friend _CG_QUALIFIER thread_group tiled_partition(const thread_group& parent, unsigned int tilesz);
121 friend class thread_block;
122
123 union __align__(8) {
124 group_data group;
125 tg_data coalesced;
126 gg_data grid;
127#if defined(_CG_CPP11_FEATURES) && defined(_CG_ABI_EXPERIMENTAL)
128 mg_data multi_grid;
129#endif
130 } _data;
131
132 _CG_QUALIFIER thread_group operator=(const thread_group& src);
133
134 _CG_QUALIFIER thread_group(unsigned int type) {
135 _data.group.type = type;
136 _data.group._unused = false;
137 }
138
139#ifdef _CG_CPP11_FEATURES
140 static_assert(sizeof(tg_data) <= 16, "Failed size check");
141 static_assert(sizeof(gg_data) <= 16, "Failed size check");
142# ifdef _CG_ABI_EXPERIMENTAL
143 static_assert(sizeof(mg_data) <= 16, "Failed size check");
144# endif
145#endif
146
147public:
148 _CG_THREAD_SCOPE(cuda::thread_scope::thread_scope_device)
149
150 _CG_QUALIFIER unsigned long long size() const;
151 _CG_QUALIFIER unsigned long long num_threads() const;
152 _CG_QUALIFIER unsigned long long thread_rank() const;
153 _CG_QUALIFIER void sync() const;
154 _CG_QUALIFIER unsigned int get_type() const {
155 return _data.group.type;
156 }
157
158};
159
160template <unsigned int TyId>
161struct thread_group_base : public thread_group {
162 _CG_QUALIFIER thread_group_base() : thread_group(TyId) {}
163 _CG_STATIC_CONST_DECL unsigned int id = TyId;
164};
165
166#if defined(_CG_HAS_MULTI_GRID_GROUP)
167
168/**
169 * class multi_grid_group;
170 *
171 * Threads within this this group are guaranteed to be co-resident on the
172 * same system, on multiple devices within the same launched kernels.
173 * To use this group, the kernel must have been launched with
174 * cuLaunchCooperativeKernelMultiDevice (or the CUDA Runtime equivalent),
175 * and the device must support it (queryable device attribute).
176 *
177 * Constructed via this_multi_grid();
178 */
179
180
181# if defined(_CG_CPP11_FEATURES) && defined(_CG_ABI_EXPERIMENTAL)
182class multi_grid_group;
183
184// Multi grid group requires these functions to be templated to prevent ptxas from trying to use CG syscalls
185template <typename = void>
186__device__ _CG_DEPRECATED multi_grid_group this_multi_grid();
187
188class multi_grid_group : public thread_group_base<details::multi_grid_group_id>
189{
190private:
191 template <typename = void>
192 _CG_QUALIFIER multi_grid_group() {
193 _data.multi_grid.functions = details::multi_grid::load_grid_intrinsics();
194 _data.multi_grid.handle = _data.multi_grid.functions->get_intrinsic_handle();
195 }
196
197 friend multi_grid_group this_multi_grid<void>();
198
199public:
200 _CG_THREAD_SCOPE(cuda::thread_scope::thread_scope_system)
201
202 _CG_QUALIFIER bool is_valid() const {
203 return (_data.multi_grid.handle != 0);
204 }
205
206 _CG_QUALIFIER void sync() const {
207 if (!is_valid()) {
208 _CG_ABORT();
209 }
210 _data.multi_grid.functions->sync(_data.multi_grid.handle);
211 }
212
213 _CG_QUALIFIER unsigned long long num_threads() const {
214 _CG_ASSERT(is_valid());
215 return _data.multi_grid.functions->size(_data.multi_grid.handle);
216 }
217
218 _CG_QUALIFIER unsigned long long size() const {
219 return num_threads();
220 }
221
222 _CG_QUALIFIER unsigned long long thread_rank() const {
223 _CG_ASSERT(is_valid());
224 return _data.multi_grid.functions->thread_rank(_data.multi_grid.handle);
225 }
226
227 _CG_QUALIFIER unsigned int grid_rank() const {
228 _CG_ASSERT(is_valid());
229 return (_data.multi_grid.functions->grid_rank(_data.multi_grid.handle));
230 }
231
232 _CG_QUALIFIER unsigned int num_grids() const {
233 _CG_ASSERT(is_valid());
234 return (_data.multi_grid.functions->num_grids(_data.multi_grid.handle));
235 }
236};
237# else
238class multi_grid_group
239{
240private:
241 unsigned long long _handle;
242 unsigned int _size;
243 unsigned int _rank;
244
245 friend _CG_QUALIFIER multi_grid_group this_multi_grid();
246
247 _CG_QUALIFIER multi_grid_group() {
248 _handle = details::multi_grid::get_intrinsic_handle();
249 _size = details::multi_grid::size(_handle);
250 _rank = details::multi_grid::thread_rank(_handle);
251 }
252
253public:
254 _CG_THREAD_SCOPE(cuda::thread_scope::thread_scope_system)
255
256 _CG_QUALIFIER _CG_DEPRECATED bool is_valid() const {
257 return (_handle != 0);
258 }
259
260 _CG_QUALIFIER _CG_DEPRECATED void sync() const {
261 if (!is_valid()) {
262 _CG_ABORT();
263 }
264 details::multi_grid::sync(_handle);
265 }
266
267 _CG_QUALIFIER _CG_DEPRECATED unsigned long long num_threads() const {
268 _CG_ASSERT(is_valid());
269 return _size;
270 }
271
272 _CG_QUALIFIER _CG_DEPRECATED unsigned long long size() const {
273 return num_threads();
274 }
275
276 _CG_QUALIFIER _CG_DEPRECATED unsigned long long thread_rank() const {
277 _CG_ASSERT(is_valid());
278 return _rank;
279 }
280
281 _CG_QUALIFIER _CG_DEPRECATED unsigned int grid_rank() const {
282 _CG_ASSERT(is_valid());
283 return (details::multi_grid::grid_rank(_handle));
284 }
285
286 _CG_QUALIFIER _CG_DEPRECATED unsigned int num_grids() const {
287 _CG_ASSERT(is_valid());
288 return (details::multi_grid::num_grids(_handle));
289 }
290};
291# endif
292
293/**
294 * multi_grid_group this_multi_grid()
295 *
296 * Constructs a multi_grid_group
297 */
298# if defined(_CG_CPP11_FEATURES) && defined(_CG_ABI_EXPERIMENTAL)
299template <typename>
300__device__
301#else
302_CG_QUALIFIER
303# endif
304_CG_DEPRECATED
305multi_grid_group this_multi_grid()
306{
307 return multi_grid_group();
308}
309#endif
310
311/**
312 * class grid_group;
313 *
314 * Threads within this this group are guaranteed to be co-resident on the
315 * same device within the same launched kernel. To use this group, the kernel
316 * must have been launched with cuLaunchCooperativeKernel (or the CUDA Runtime equivalent),
317 * and the device must support it (queryable device attribute).
318 *
319 * Constructed via this_grid();
320 */
321class grid_group : public thread_group_base<details::grid_group_id>
322{
323 _CG_STATIC_CONST_DECL unsigned int _group_id = details::grid_group_id;
324 friend _CG_QUALIFIER grid_group this_grid();
325
326private:
327 _CG_QUALIFIER grid_group(details::grid_workspace *gridWs) {
328 _data.grid.gridWs = gridWs;
329 }
330
331 public:
332 _CG_THREAD_SCOPE(cuda::thread_scope::thread_scope_device)
333
334 _CG_QUALIFIER bool is_valid() const {
335 return (_data.grid.gridWs != NULL);
336 }
337
338 _CG_QUALIFIER void sync() const {
339 if (!is_valid()) {
340 _CG_ABORT();
341 }
342 details::grid::sync(&_data.grid.gridWs->barrier);
343 }
344
345#if defined(_CG_CPP11_FEATURES)
346 using arrival_token = unsigned int;
347
348 _CG_QUALIFIER arrival_token barrier_arrive() const {
349 if (!is_valid()) {
350 _CG_ABORT();
351 }
352 return details::grid::barrier_arrive(&_data.grid.gridWs->barrier);
353 }
354
355 _CG_QUALIFIER void barrier_wait(arrival_token&& token) const {
356 details::grid::barrier_wait(token, &_data.grid.gridWs->barrier);
357 }
358#endif
359
360 _CG_STATIC_QUALIFIER unsigned long long size() {
361 return details::grid::size();
362 }
363
364 _CG_STATIC_QUALIFIER dim3 group_dim() {
365 return details::grid::grid_dim();
366 }
367
368 _CG_STATIC_QUALIFIER dim3 dim_threads() {
369 return details::grid::dim_threads();
370 }
371
372 _CG_STATIC_QUALIFIER unsigned long long num_threads() {
373 return details::grid::num_threads();
374 }
375
376 _CG_STATIC_QUALIFIER dim3 thread_index() {
377 return details::grid::thread_index();
378 }
379
380 _CG_STATIC_QUALIFIER unsigned long long thread_rank() {
381 return details::grid::thread_rank();
382 }
383
384 _CG_STATIC_QUALIFIER dim3 dim_blocks() {
385 return details::grid::dim_blocks();
386 }
387
388 _CG_STATIC_QUALIFIER unsigned long long num_blocks() {
389 return details::grid::num_blocks();
390 }
391
392 _CG_STATIC_QUALIFIER dim3 block_index() {
393 return details::grid::block_index();
394 }
395
396 _CG_STATIC_QUALIFIER unsigned long long block_rank() {
397 return details::grid::block_rank();
398 }
399
400# if defined(_CG_HAS_CLUSTER_GROUP)
401 _CG_STATIC_QUALIFIER dim3 dim_clusters() {
402 return details::grid::dim_clusters();
403 }
404
405 _CG_STATIC_QUALIFIER unsigned long long num_clusters() {
406 return details::grid::num_clusters();
407 }
408
409 _CG_STATIC_QUALIFIER dim3 cluster_index() {
410 return details::grid::cluster_index();
411 }
412
413 _CG_STATIC_QUALIFIER unsigned long long cluster_rank() {
414 return details::grid::cluster_rank();
415 }
416# endif
417};
418
419_CG_QUALIFIER grid_group this_grid() {
420 // Load a workspace from the driver
421 grid_group gg(details::get_grid_workspace());
422#ifdef _CG_DEBUG
423 // *all* threads must be available to synchronize
424 gg.sync();
425#endif // _CG_DEBUG
426 return gg;
427}
428
429#if defined(_CG_HAS_CLUSTER_GROUP)
430/**
431 * class cluster_group
432 *
433 * Every GPU kernel is executed by a grid of thread blocks. A grid can be evenly
434 * divided along all dimensions to form groups of blocks, each group of which is
435 * a block cluster. Clustered grids are subject to various restrictions and
436 * limitations. Primarily, a cluster consists of at most 8 blocks by default
437 * (although the user is allowed to opt-in to non-standard sizes,) and clustered
438 * grids are subject to additional occupancy limitations due to per-cluster
439 * hardware resource consumption. In exchange, a block cluster is guaranteed to
440 * be a cooperative group, with access to all cooperative group capabilities, as
441 * well as cluster specific capabilities and accelerations. A cluster_group
442 * represents a block cluster.
443 *
444 * Constructed via this_cluster_group();
445 */
446class cluster_group : public thread_group_base<details::cluster_group_id>
447{
448 // Friends
449 friend _CG_QUALIFIER cluster_group this_cluster();
450
451 // Disable constructor
452 _CG_QUALIFIER cluster_group()
453 {
454 }
455
456 public:
457 //_CG_THREAD_SCOPE(cuda::thread_scope::thread_scope_cluster)
458
459 using arrival_token = struct {};
460
461 // Functionality exposed by the group
462 _CG_STATIC_QUALIFIER void sync()
463 {
464 return details::cluster::sync();
465 }
466
467 _CG_STATIC_QUALIFIER arrival_token barrier_arrive()
468 {
469 details::cluster::barrier_arrive();
470 return arrival_token();
471 }
472
473 _CG_STATIC_QUALIFIER void barrier_wait()
474 {
475 return details::cluster::barrier_wait();
476 }
477
478 _CG_STATIC_QUALIFIER void barrier_wait(arrival_token&&)
479 {
480 return details::cluster::barrier_wait();
481 }
482
483 _CG_STATIC_QUALIFIER unsigned int query_shared_rank(const void *addr)
484 {
485 return details::cluster::query_shared_rank(addr);
486 }
487
488 template <typename T>
489 _CG_STATIC_QUALIFIER T* map_shared_rank(T *addr, int rank)
490 {
491 return details::cluster::map_shared_rank(addr, rank);
492 }
493
494 _CG_STATIC_QUALIFIER dim3 block_index()
495 {
496 return details::cluster::block_index();
497 }
498
499 _CG_STATIC_QUALIFIER unsigned int block_rank()
500 {
501 return details::cluster::block_rank();
502 }
503
504 _CG_STATIC_QUALIFIER dim3 thread_index()
505 {
506 return details::cluster::thread_index();
507 }
508
509 _CG_STATIC_QUALIFIER unsigned int thread_rank()
510 {
511 return details::cluster::thread_rank();
512 }
513
514 _CG_STATIC_QUALIFIER dim3 dim_blocks()
515 {
516 return details::cluster::dim_blocks();
517 }
518
519 _CG_STATIC_QUALIFIER unsigned int num_blocks()
520 {
521 return details::cluster::num_blocks();
522 }
523
524 _CG_STATIC_QUALIFIER dim3 dim_threads()
525 {
526 return details::cluster::dim_threads();
527 }
528
529 _CG_STATIC_QUALIFIER unsigned int num_threads()
530 {
531 return details::cluster::num_threads();
532 }
533
534 // Legacy aliases
535 _CG_STATIC_QUALIFIER unsigned int size()
536 {
537 return num_threads();
538 }
539};
540
541/*
542 * cluster_group this_cluster()
543 *
544 * Constructs a cluster_group
545 */
546_CG_QUALIFIER cluster_group this_cluster()
547{
548 cluster_group cg;
549#ifdef _CG_DEBUG
550 cg.sync();
551#endif
552 return cg;
553}
554#endif
555
556#if defined(_CG_CPP11_FEATURES)
557class thread_block;
558template <unsigned int MaxBlockSize>
559_CG_QUALIFIER thread_block this_thread_block(block_tile_memory<MaxBlockSize>& scratch);
560#endif
561
562/**
563 * class thread_block
564 *
565 * Every GPU kernel is executed by a grid of thread blocks, and threads within
566 * each block are guaranteed to reside on the same streaming multiprocessor.
567 * A thread_block represents a thread block whose dimensions are not known until runtime.
568 *
569 * Constructed via this_thread_block();
570 */
571class thread_block : public thread_group_base<details::thread_block_id>
572{
573 // Friends
574 friend _CG_QUALIFIER thread_block this_thread_block();
575 friend _CG_QUALIFIER thread_group tiled_partition(const thread_group& parent, unsigned int tilesz);
576 friend _CG_QUALIFIER thread_group tiled_partition(const thread_block& parent, unsigned int tilesz);
577
578#if defined(_CG_CPP11_FEATURES)
579 template <unsigned int MaxBlockSize>
580 friend _CG_QUALIFIER thread_block this_thread_block(block_tile_memory<MaxBlockSize>& scratch);
581 template <unsigned int Size>
582 friend class __static_size_multi_warp_tile_base;
583
584 details::multi_warp_scratch* const tile_memory;
585
586 template <unsigned int MaxBlockSize>
587 _CG_QUALIFIER thread_block(block_tile_memory<MaxBlockSize>& scratch) :
588 tile_memory(details::get_scratch_ptr(&scratch)) {
589#ifdef _CG_DEBUG
590 if (num_threads() > MaxBlockSize) {
591 details::abort();
592 }
593#endif
594#if !defined(_CG_HAS_RESERVED_SHARED)
595 tile_memory->init_barriers(thread_rank());
596 sync();
597#endif
598 }
599#endif
600
601 // Disable constructor
602 _CG_QUALIFIER thread_block()
603#if defined(_CG_CPP11_FEATURES)
604 : tile_memory(details::get_scratch_ptr(NULL))
605#endif
606 { }
607
608 // Internal Use
609 _CG_QUALIFIER thread_group _get_tiled_threads(unsigned int tilesz) const {
610 const bool pow2_tilesz = ((tilesz & (tilesz - 1)) == 0);
611
612 // Invalid, immediately fail
613 if (tilesz == 0 || (tilesz > 32) || !pow2_tilesz) {
614 details::abort();
615 return (thread_block());
616 }
617
618 unsigned int mask;
619 unsigned int base_offset = thread_rank() & (~(tilesz - 1));
620 unsigned int masklength = min((unsigned int)size() - base_offset, tilesz);
621
622 mask = (unsigned int)(-1) >> (32 - masklength);
623 mask <<= (details::laneid() & ~(tilesz - 1));
624 thread_group tile = thread_group(details::coalesced_group_id);
625 tile._data.coalesced.mask = mask;
626 tile._data.coalesced.size = __popc(mask);
627 tile._data.coalesced.metaGroupSize = (details::cta::size() + tilesz - 1) / tilesz;
628 tile._data.coalesced.metaGroupRank = details::cta::thread_rank() / tilesz;
629 tile._data.coalesced.is_tiled = true;
630 return (tile);
631 }
632
633 public:
634 _CG_STATIC_CONST_DECL unsigned int _group_id = details::thread_block_id;
635 _CG_THREAD_SCOPE(cuda::thread_scope::thread_scope_block)
636
637 _CG_STATIC_QUALIFIER void sync() {
638 details::cta::sync();
639 }
640
641#if defined(_CG_CPP11_FEATURES)
642 struct arrival_token {};
643
644 _CG_QUALIFIER arrival_token barrier_arrive() const {
645 return arrival_token();
646 }
647
648 _CG_QUALIFIER void barrier_wait(arrival_token&&) const {
649 details::cta::sync();
650 }
651#endif
652
653 _CG_STATIC_QUALIFIER unsigned int size() {
654 return details::cta::size();
655 }
656
657 _CG_STATIC_QUALIFIER unsigned int thread_rank() {
658 return details::cta::thread_rank();
659 }
660
661 // Additional functionality exposed by the group
662 _CG_STATIC_QUALIFIER dim3 group_index() {
663 return details::cta::group_index();
664 }
665
666 _CG_STATIC_QUALIFIER dim3 thread_index() {
667 return details::cta::thread_index();
668 }
669
670 _CG_STATIC_QUALIFIER dim3 group_dim() {
671 return details::cta::block_dim();
672 }
673
674 _CG_STATIC_QUALIFIER dim3 dim_threads() {
675 return details::cta::dim_threads();
676 }
677
678 _CG_STATIC_QUALIFIER unsigned int num_threads() {
679 return details::cta::num_threads();
680 }
681
682};
683
684/**
685 * thread_block this_thread_block()
686 *
687 * Constructs a thread_block group
688 */
689_CG_QUALIFIER thread_block this_thread_block()
690{
691 return (thread_block());
692}
693
694#if defined(_CG_CPP11_FEATURES)
695template <unsigned int MaxBlockSize>
696_CG_QUALIFIER thread_block this_thread_block(block_tile_memory<MaxBlockSize>& scratch) {
697 return (thread_block(scratch));
698}
699#endif
700
701/**
702 * class coalesced_group
703 *
704 * A group representing the current set of converged threads in a warp.
705 * The size of the group is not guaranteed and it may return a group of
706 * only one thread (itself).
707 *
708 * This group exposes warp-synchronous builtins.
709 * Constructed via coalesced_threads();
710 */
711class coalesced_group : public thread_group_base<details::coalesced_group_id>
712{
713private:
714 friend _CG_QUALIFIER coalesced_group coalesced_threads();
715 friend _CG_QUALIFIER thread_group tiled_partition(const thread_group& parent, unsigned int tilesz);
716 friend _CG_QUALIFIER coalesced_group tiled_partition(const coalesced_group& parent, unsigned int tilesz);
717 friend class details::_coalesced_group_data_access;
718
719 _CG_QUALIFIER unsigned int _packLanes(unsigned laneMask) const {
720 unsigned int member_pack = 0;
721 unsigned int member_rank = 0;
722 for (int bit_idx = 0; bit_idx < 32; bit_idx++) {
723 unsigned int lane_bit = _data.coalesced.mask & (1 << bit_idx);
724 if (lane_bit) {
725 if (laneMask & lane_bit)
726 member_pack |= 1 << member_rank;
727 member_rank++;
728 }
729 }
730 return (member_pack);
731 }
732
733 // Internal Use
734 _CG_QUALIFIER coalesced_group _get_tiled_threads(unsigned int tilesz) const {
735 const bool pow2_tilesz = ((tilesz & (tilesz - 1)) == 0);
736
737 // Invalid, immediately fail
738 if (tilesz == 0 || (tilesz > 32) || !pow2_tilesz) {
739 details::abort();
740 return (coalesced_group(0));
741 }
742 if (size() <= tilesz) {
743 return (*this);
744 }
745
746 if ((_data.coalesced.is_tiled == true) && pow2_tilesz) {
747 unsigned int base_offset = (thread_rank() & (~(tilesz - 1)));
748 unsigned int masklength = min((unsigned int)size() - base_offset, tilesz);
749 unsigned int mask = (unsigned int)(-1) >> (32 - masklength);
750
751 mask <<= (details::laneid() & ~(tilesz - 1));
752 coalesced_group coalesced_tile = coalesced_group(mask);
753 coalesced_tile._data.coalesced.metaGroupSize = size() / tilesz;
754 coalesced_tile._data.coalesced.metaGroupRank = thread_rank() / tilesz;
755 coalesced_tile._data.coalesced.is_tiled = true;
756 return (coalesced_tile);
757 }
758 else if ((_data.coalesced.is_tiled == false) && pow2_tilesz) {
759 unsigned int mask = 0;
760 unsigned int member_rank = 0;
761 int seen_lanes = (thread_rank() / tilesz) * tilesz;
762 for (unsigned int bit_idx = 0; bit_idx < 32; bit_idx++) {
763 unsigned int lane_bit = _data.coalesced.mask & (1 << bit_idx);
764 if (lane_bit) {
765 if (seen_lanes <= 0 && member_rank < tilesz) {
766 mask |= lane_bit;
767 member_rank++;
768 }
769 seen_lanes--;
770 }
771 }
772 coalesced_group coalesced_tile = coalesced_group(mask);
773 // Override parent with the size of this group
774 coalesced_tile._data.coalesced.metaGroupSize = (size() + tilesz - 1) / tilesz;
775 coalesced_tile._data.coalesced.metaGroupRank = thread_rank() / tilesz;
776 return coalesced_tile;
777 }
778 else {
779 // None in _CG_VERSION 1000
780 details::abort();
781 }
782
783 return (coalesced_group(0));
784 }
785
786 protected:
787 _CG_QUALIFIER coalesced_group(unsigned int mask) {
788 _data.coalesced.mask = mask;
789 _data.coalesced.size = __popc(mask);
790 _data.coalesced.metaGroupRank = 0;
791 _data.coalesced.metaGroupSize = 1;
792 _data.coalesced.is_tiled = false;
793 }
794
795 _CG_QUALIFIER unsigned int get_mask() const {
796 return (_data.coalesced.mask);
797 }
798
799 public:
800 _CG_STATIC_CONST_DECL unsigned int _group_id = details::coalesced_group_id;
801 _CG_THREAD_SCOPE(cuda::thread_scope::thread_scope_block)
802
803 _CG_QUALIFIER unsigned int num_threads() const {
804 return _data.coalesced.size;
805 }
806
807 _CG_QUALIFIER unsigned int size() const {
808 return num_threads();
809 }
810
811 _CG_QUALIFIER unsigned int thread_rank() const {
812 return (__popc(_data.coalesced.mask & details::lanemask32_lt()));
813 }
814
815 // Rank of this group in the upper level of the hierarchy
816 _CG_QUALIFIER unsigned int meta_group_rank() const {
817 return _data.coalesced.metaGroupRank;
818 }
819
820 // Total num partitions created out of all CTAs when the group was created
821 _CG_QUALIFIER unsigned int meta_group_size() const {
822 return _data.coalesced.metaGroupSize;
823 }
824
825 _CG_QUALIFIER void sync() const {
826 __syncwarp(_data.coalesced.mask);
827 }
828
829#ifdef _CG_CPP11_FEATURES
830 template <typename TyElem, typename TyRet = details::remove_qual<TyElem>>
831 _CG_QUALIFIER TyRet shfl(TyElem&& elem, int srcRank) const {
832 unsigned int lane = (srcRank == 0) ? __ffs(_data.coalesced.mask) - 1 :
833 (size() == 32) ? srcRank : __fns(_data.coalesced.mask, 0, (srcRank + 1));
834
835 return details::tile::shuffle_dispatch<TyElem>::shfl(
836 _CG_STL_NAMESPACE::forward<TyElem>(elem), _data.coalesced.mask, lane, 32);
837 }
838
839 template <typename TyElem, typename TyRet = details::remove_qual<TyElem>>
840 _CG_QUALIFIER TyRet shfl_down(TyElem&& elem, unsigned int delta) const {
841 if (size() == 32) {
842 return details::tile::shuffle_dispatch<TyElem>::shfl_down(
843 _CG_STL_NAMESPACE::forward<TyElem>(elem), 0xFFFFFFFF, delta, 32);
844 }
845
846 unsigned int lane = __fns(_data.coalesced.mask, details::laneid(), delta + 1);
847
848 if (lane >= 32)
849 lane = details::laneid();
850
851 return details::tile::shuffle_dispatch<TyElem>::shfl(
852 _CG_STL_NAMESPACE::forward<TyElem>(elem), _data.coalesced.mask, lane, 32);
853 }
854
855 template <typename TyElem, typename TyRet = details::remove_qual<TyElem>>
856 _CG_QUALIFIER TyRet shfl_up(TyElem&& elem, int delta) const {
857 if (size() == 32) {
858 return details::tile::shuffle_dispatch<TyElem>::shfl_up(
859 _CG_STL_NAMESPACE::forward<TyElem>(elem), 0xFFFFFFFF, delta, 32);
860 }
861
862 unsigned lane = __fns(_data.coalesced.mask, details::laneid(), -(delta + 1));
863 if (lane >= 32)
864 lane = details::laneid();
865
866 return details::tile::shuffle_dispatch<TyElem>::shfl(
867 _CG_STL_NAMESPACE::forward<TyElem>(elem), _data.coalesced.mask, lane, 32);
868 }
869#else
870 template <typename TyIntegral>
871 _CG_QUALIFIER TyIntegral shfl(TyIntegral var, unsigned int src_rank) const {
872 details::assert_if_not_arithmetic<TyIntegral>();
873 unsigned int lane = (src_rank == 0) ? __ffs(_data.coalesced.mask) - 1 :
874 (size() == 32) ? src_rank : __fns(_data.coalesced.mask, 0, (src_rank + 1));
875 return (__shfl_sync(_data.coalesced.mask, var, lane, 32));
876 }
877
878 template <typename TyIntegral>
879 _CG_QUALIFIER TyIntegral shfl_up(TyIntegral var, int delta) const {
880 details::assert_if_not_arithmetic<TyIntegral>();
881 if (size() == 32) {
882 return (__shfl_up_sync(0xFFFFFFFF, var, delta, 32));
883 }
884 unsigned lane = __fns(_data.coalesced.mask, details::laneid(), -(delta + 1));
885 if (lane >= 32) lane = details::laneid();
886 return (__shfl_sync(_data.coalesced.mask, var, lane, 32));
887 }
888
889 template <typename TyIntegral>
890 _CG_QUALIFIER TyIntegral shfl_down(TyIntegral var, int delta) const {
891 details::assert_if_not_arithmetic<TyIntegral>();
892 if (size() == 32) {
893 return (__shfl_down_sync(0xFFFFFFFF, var, delta, 32));
894 }
895 unsigned int lane = __fns(_data.coalesced.mask, details::laneid(), delta + 1);
896 if (lane >= 32) lane = details::laneid();
897 return (__shfl_sync(_data.coalesced.mask, var, lane, 32));
898 }
899#endif
900
901 _CG_QUALIFIER int any(int predicate) const {
902 return (__ballot_sync(_data.coalesced.mask, predicate) != 0);
903 }
904 _CG_QUALIFIER int all(int predicate) const {
905 return (__ballot_sync(_data.coalesced.mask, predicate) == _data.coalesced.mask);
906 }
907 _CG_QUALIFIER unsigned int ballot(int predicate) const {
908 if (size() == 32) {
909 return (__ballot_sync(0xFFFFFFFF, predicate));
910 }
911 unsigned int lane_ballot = __ballot_sync(_data.coalesced.mask, predicate);
912 return (_packLanes(lane_ballot));
913 }
914
915#ifdef _CG_HAS_MATCH_COLLECTIVE
916
917 template <typename TyIntegral>
918 _CG_QUALIFIER unsigned int match_any(TyIntegral val) const {
919 details::assert_if_not_arithmetic<TyIntegral>();
920 if (size() == 32) {
921 return (__match_any_sync(0xFFFFFFFF, val));
922 }
923 unsigned int lane_match = __match_any_sync(_data.coalesced.mask, val);
924 return (_packLanes(lane_match));
925 }
926
927 template <typename TyIntegral>
928 _CG_QUALIFIER unsigned int match_all(TyIntegral val, int &pred) const {
929 details::assert_if_not_arithmetic<TyIntegral>();
930 if (size() == 32) {
931 return (__match_all_sync(0xFFFFFFFF, val, &pred));
932 }
933 unsigned int lane_match = __match_all_sync(_data.coalesced.mask, val, &pred);
934 return (_packLanes(lane_match));
935 }
936
937#endif /* !_CG_HAS_MATCH_COLLECTIVE */
938
939};
940
941_CG_QUALIFIER coalesced_group coalesced_threads()
942{
943 return (coalesced_group(__activemask()));
944}
945
946namespace details {
947 template <unsigned int Size> struct verify_thread_block_tile_size;
948 template <> struct verify_thread_block_tile_size<32> { typedef void OK; };
949 template <> struct verify_thread_block_tile_size<16> { typedef void OK; };
950 template <> struct verify_thread_block_tile_size<8> { typedef void OK; };
951 template <> struct verify_thread_block_tile_size<4> { typedef void OK; };
952 template <> struct verify_thread_block_tile_size<2> { typedef void OK; };
953 template <> struct verify_thread_block_tile_size<1> { typedef void OK; };
954
955#ifdef _CG_CPP11_FEATURES
956 template <unsigned int Size>
957 using _is_power_of_2 = _CG_STL_NAMESPACE::integral_constant<bool, (Size & (Size - 1)) == 0>;
958
959 template <unsigned int Size>
960 using _is_single_warp = _CG_STL_NAMESPACE::integral_constant<bool, Size <= 32>;
961 template <unsigned int Size>
962 using _is_multi_warp =
963 _CG_STL_NAMESPACE::integral_constant<bool, (Size > 32) && (Size <= 1024)>;
964
965 template <unsigned int Size>
966 using _is_valid_single_warp_tile =
967 _CG_STL_NAMESPACE::integral_constant<bool, _is_power_of_2<Size>::value && _is_single_warp<Size>::value>;
968 template <unsigned int Size>
969 using _is_valid_multi_warp_tile =
970 _CG_STL_NAMESPACE::integral_constant<bool, _is_power_of_2<Size>::value && _is_multi_warp<Size>::value>;
971#else
972 template <unsigned int Size>
973 struct _is_multi_warp {
974 static const bool value = false;
975 };
976#endif
977}
978
979template <unsigned int Size>
980class __static_size_tile_base
981{
982protected:
983 _CG_STATIC_CONST_DECL unsigned int numThreads = Size;
984
985public:
986 _CG_THREAD_SCOPE(cuda::thread_scope::thread_scope_block)
987
988 // Rank of thread within tile
989 _CG_STATIC_QUALIFIER unsigned int thread_rank() {
990 return (details::cta::thread_rank() & (numThreads - 1));
991 }
992
993 // Number of threads within tile
994 _CG_STATIC_CONSTEXPR_QUALIFIER unsigned int num_threads() {
995 return numThreads;
996 }
997
998 _CG_STATIC_CONSTEXPR_QUALIFIER unsigned int size() {
999 return num_threads();
1000 }
1001};
1002
1003template <unsigned int Size>
1004class __static_size_thread_block_tile_base : public __static_size_tile_base<Size>
1005{
1006 friend class details::_coalesced_group_data_access;
1007 typedef details::tile::tile_helpers<Size> th;
1008
1009#ifdef _CG_CPP11_FEATURES
1010 static_assert(details::_is_valid_single_warp_tile<Size>::value, "Size must be one of 1/2/4/8/16/32");
1011#else
1012 typedef typename details::verify_thread_block_tile_size<Size>::OK valid;
1013#endif
1014 using __static_size_tile_base<Size>::numThreads;
1015 _CG_STATIC_CONST_DECL unsigned int fullMask = 0xFFFFFFFF;
1016
1017 protected:
1018 _CG_STATIC_QUALIFIER unsigned int build_mask() {
1019 unsigned int mask = fullMask;
1020 if (numThreads != 32) {
1021 // [0,31] representing the current active thread in the warp
1022 unsigned int laneId = details::laneid();
1023 // shift mask according to the partition it belongs to
1024 mask = th::tileMask << (laneId & ~(th::laneMask));
1025 }
1026 return (mask);
1027 }
1028
1029public:
1030 _CG_STATIC_CONST_DECL unsigned int _group_id = details::coalesced_group_id;
1031
1032 _CG_STATIC_QUALIFIER void sync() {
1033 __syncwarp(build_mask());
1034 }
1035
1036#ifdef _CG_CPP11_FEATURES
1037 // PTX supported collectives
1038 template <typename TyElem, typename TyRet = details::remove_qual<TyElem>>
1039 _CG_QUALIFIER TyRet shfl(TyElem&& elem, int srcRank) const {
1040 return details::tile::shuffle_dispatch<TyElem>::shfl(
1041 _CG_STL_NAMESPACE::forward<TyElem>(elem), build_mask(), srcRank, numThreads);
1042 }
1043
1044 template <typename TyElem, typename TyRet = details::remove_qual<TyElem>>
1045 _CG_QUALIFIER TyRet shfl_down(TyElem&& elem, unsigned int delta) const {
1046 return details::tile::shuffle_dispatch<TyElem>::shfl_down(
1047 _CG_STL_NAMESPACE::forward<TyElem>(elem), build_mask(), delta, numThreads);
1048 }
1049
1050 template <typename TyElem, typename TyRet = details::remove_qual<TyElem>>
1051 _CG_QUALIFIER TyRet shfl_up(TyElem&& elem, unsigned int delta) const {
1052 return details::tile::shuffle_dispatch<TyElem>::shfl_up(
1053 _CG_STL_NAMESPACE::forward<TyElem>(elem), build_mask(), delta, numThreads);
1054 }
1055
1056 template <typename TyElem, typename TyRet = details::remove_qual<TyElem>>
1057 _CG_QUALIFIER TyRet shfl_xor(TyElem&& elem, unsigned int laneMask) const {
1058 return details::tile::shuffle_dispatch<TyElem>::shfl_xor(
1059 _CG_STL_NAMESPACE::forward<TyElem>(elem), build_mask(), laneMask, numThreads);
1060 }
1061#else
1062 template <typename TyIntegral>
1063 _CG_QUALIFIER TyIntegral shfl(TyIntegral var, int srcRank) const {
1064 details::assert_if_not_arithmetic<TyIntegral>();
1065 return (__shfl_sync(build_mask(), var, srcRank, numThreads));
1066 }
1067
1068 template <typename TyIntegral>
1069 _CG_QUALIFIER TyIntegral shfl_down(TyIntegral var, unsigned int delta) const {
1070 details::assert_if_not_arithmetic<TyIntegral>();
1071 return (__shfl_down_sync(build_mask(), var, delta, numThreads));
1072 }
1073
1074 template <typename TyIntegral>
1075 _CG_QUALIFIER TyIntegral shfl_up(TyIntegral var, unsigned int delta) const {
1076 details::assert_if_not_arithmetic<TyIntegral>();
1077 return (__shfl_up_sync(build_mask(), var, delta, numThreads));
1078 }
1079
1080 template <typename TyIntegral>
1081 _CG_QUALIFIER TyIntegral shfl_xor(TyIntegral var, unsigned int laneMask) const {
1082 details::assert_if_not_arithmetic<TyIntegral>();
1083 return (__shfl_xor_sync(build_mask(), var, laneMask, numThreads));
1084 }
1085#endif //_CG_CPP11_FEATURES
1086
1087 _CG_QUALIFIER int any(int predicate) const {
1088 unsigned int lane_ballot = __ballot_sync(build_mask(), predicate);
1089 return (lane_ballot != 0);
1090 }
1091 _CG_QUALIFIER int all(int predicate) const {
1092 unsigned int lane_ballot = __ballot_sync(build_mask(), predicate);
1093 return (lane_ballot == build_mask());
1094 }
1095 _CG_QUALIFIER unsigned int ballot(int predicate) const {
1096 unsigned int lane_ballot = __ballot_sync(build_mask(), predicate);
1097 return (lane_ballot >> (details::laneid() & (~(th::laneMask))));
1098 }
1099
1100#ifdef _CG_HAS_MATCH_COLLECTIVE
1101 template <typename TyIntegral>
1102 _CG_QUALIFIER unsigned int match_any(TyIntegral val) const {
1103 details::assert_if_not_arithmetic<TyIntegral>();
1104 unsigned int lane_match = __match_any_sync(build_mask(), val);
1105 return (lane_match >> (details::laneid() & (~(th::laneMask))));
1106 }
1107
1108 template <typename TyIntegral>
1109 _CG_QUALIFIER unsigned int match_all(TyIntegral val, int &pred) const {
1110 details::assert_if_not_arithmetic<TyIntegral>();
1111 unsigned int lane_match = __match_all_sync(build_mask(), val, &pred);
1112 return (lane_match >> (details::laneid() & (~(th::laneMask))));
1113 }
1114#endif
1115
1116};
1117
1118template <unsigned int Size, typename ParentT>
1119class __static_parent_thread_block_tile_base
1120{
1121public:
1122 // Rank of this group in the upper level of the hierarchy
1123 _CG_STATIC_QUALIFIER unsigned int meta_group_rank() {
1124 return ParentT::thread_rank() / Size;
1125 }
1126
1127 // Total num partitions created out of all CTAs when the group was created
1128 _CG_STATIC_QUALIFIER unsigned int meta_group_size() {
1129 return (ParentT::size() + Size - 1) / Size;
1130 }
1131};
1132
1133/**
1134 * class thread_block_tile<unsigned int Size, ParentT = void>
1135 *
1136 * Statically-sized group type, representing one tile of a thread block.
1137 * The only specializations currently supported are those with native
1138 * hardware support (1/2/4/8/16/32)
1139 *
1140 * This group exposes warp-synchronous builtins.
1141 * Can only be constructed via tiled_partition<Size>(ParentT&)
1142 */
1143
1144template <unsigned int Size, typename ParentT = void>
1145class __single_warp_thread_block_tile :
1146 public __static_size_thread_block_tile_base<Size>,
1147 public __static_parent_thread_block_tile_base<Size, ParentT>
1148{
1149 typedef __static_parent_thread_block_tile_base<Size, ParentT> staticParentBaseT;
1150 friend class details::_coalesced_group_data_access;
1151
1152protected:
1153 _CG_QUALIFIER __single_warp_thread_block_tile() { };
1154 _CG_QUALIFIER __single_warp_thread_block_tile(unsigned int, unsigned int) { };
1155
1156 _CG_STATIC_QUALIFIER unsigned int get_mask() {
1157 return __static_size_thread_block_tile_base<Size>::build_mask();
1158 }
1159};
1160
1161template <unsigned int Size>
1162class __single_warp_thread_block_tile<Size, void> :
1163 public __static_size_thread_block_tile_base<Size>,
1164 public thread_group_base<details::coalesced_group_id>
1165{
1166 _CG_STATIC_CONST_DECL unsigned int numThreads = Size;
1167
1168 template <unsigned int, typename ParentT> friend class __single_warp_thread_block_tile;
1169 friend class details::_coalesced_group_data_access;
1170
1171 typedef __static_size_thread_block_tile_base<numThreads> staticSizeBaseT;
1172
1173protected:
1174 _CG_QUALIFIER __single_warp_thread_block_tile(unsigned int meta_group_rank = 0, unsigned int meta_group_size = 1) {
1175 _data.coalesced.mask = staticSizeBaseT::build_mask();
1176 _data.coalesced.size = numThreads;
1177 _data.coalesced.metaGroupRank = meta_group_rank;
1178 _data.coalesced.metaGroupSize = meta_group_size;
1179 _data.coalesced.is_tiled = true;
1180 }
1181
1182 _CG_QUALIFIER unsigned int get_mask() const {
1183 return (_data.coalesced.mask);
1184 }
1185
1186public:
1187 using staticSizeBaseT::sync;
1188 using staticSizeBaseT::size;
1189 using staticSizeBaseT::num_threads;
1190 using staticSizeBaseT::thread_rank;
1191
1192 _CG_QUALIFIER unsigned int meta_group_rank() const {
1193 return _data.coalesced.metaGroupRank;
1194 }
1195
1196 _CG_QUALIFIER unsigned int meta_group_size() const {
1197 return _data.coalesced.metaGroupSize;
1198 }
1199};
1200