Optimize for CG again

Author: bernhardmgruber

libcudacxx/include/cuda/__memcpy_async/cp_async_bulk_shared_global.h

@@ -37,8 +37,46 @@
 
 #    include <cuda/std/__cccl/prologue.h>
 
+namespace cooperative_groups
+{
+namespace __v1
+{
+class thread_block;
+
+template <unsigned int Size, typename ParentT>
+class thread_block_tile;
+} // namespace __v1
+using namespace __v1;
+} // namespace cooperative_groups
+
 _CCCL_BEGIN_NAMESPACE_CUDA
 
+template <typename Group>
+[[nodiscard]] _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int __thread_rank(const Group& __g)
+{
+  return __g.thread_rank();
+}
+
+// elect from the whole thread block
+[[nodiscard]] _CCCL_DEVICE _CCCL_FORCEINLINE bool
+__elect_from_group(const cooperative_groups::thread_block& __g) noexcept
+{
+  // Cannot call __g.thread_rank(), because we only forward declared the thread_block type
+  // cooperative groups (and we here) maps a multidimensional thread id into the thread rank the same way as warps do
+  const unsigned int tid             = threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;
+  const unsigned int warp_id         = tid / 32;
+  const unsigned int uniform_warp_id = __shfl_sync(0xFFFFFFFF, warp_id, 0); // broadcast from lane 0
+  return uniform_warp_id == 0 && ::cuda::ptx::elect_sync(0xFFFFFFFF); // elect a leader thread among warp 0
+}
+
+// elect from a single warp
+template <typename Parent>
+[[nodiscard]] _CCCL_DEVICE _CCCL_FORCEINLINE bool
+__elect_from_group(const cooperative_groups::thread_block_tile<32, Parent>&) noexcept
+{
+  return ::cuda::ptx::elect_sync(0xFFFFFFFF); // elect a leader thread among warp 0
+}
+
 template <typename _Group>
 [[nodiscard]] _CCCL_DEVICE _CCCL_FORCEINLINE bool __elect_from_group(const _Group& __g) noexcept
 {