Deprecate GridBarrier and GridBarrierLifetime (#3258)

Fixes: #1389

Author: bernhardmgruber

cub/cub/grid/grid_barrier.cuh

@@ -50,8 +50,10 @@ CUB_NAMESPACE_BEGIN
 
 /**
  * \brief GridBarrier implements a software global barrier among thread blocks within a CUDA grid
+ *
+ * deprecated [Since 2.9.0]
  */
-class GridBarrier
+class CCCL_DEPRECATED_BECAUSE("Use the APIs from cooperative groups instead") GridBarrier
 {
 protected:
   using SyncFlag = unsigned int;
@@ -131,8 +133,11 @@ public:
  *
  * Uses RAII for lifetime, i.e., device resources are reclaimed when
  * the destructor is called.
+ *
+ * deprecated [Since 2.9.0]
  */
-class GridBarrierLifetime : public GridBarrier
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+class CCCL_DEPRECATED_BECAUSE("Use the APIs from cooperative groups instead") GridBarrierLifetime : public GridBarrier
 {
 protected:
   // Number of bytes backed by d_sync
@@ -211,5 +216,6 @@ public:
     return retval;
   }
 };
+_CCCL_SUPPRESS_DEPRECATED_POP
 
 CUB_NAMESPACE_END

cub/test/test_grid_barrier.cu

@@ -47,7 +47,9 @@ using namespace cub;
 /**
  * Kernel that iterates through the specified number of software global barriers
  */
-__global__ void Kernel(GridBarrier global_barrier, int iterations)
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+__global__ void Kernel(GridBarrier global_barrier, int iterations) //
+  _CCCL_SUPPRESS_DEPRECATED_POP
 {
   for (int i = 0; i < iterations; i++)
   {
@@ -126,7 +128,9 @@ int main(int argc, char** argv)
   fflush(stdout);
 
   // Init global barrier
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   GridBarrierLifetime global_barrier;
+  _CCCL_SUPPRESS_DEPRECATED_POP
   global_barrier.Setup(grid_size);
 
   // Time kernel