From 13b8f7c72b095a3acb6b3240d2d831ee87babc6f Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Fri, 10 Jan 2025 17:47:39 +0100
Subject: [PATCH] Deprecate GridBarrier and GridBarrierLifetime (#3258) (#3288)

Fixes: #1389
---
 cub/cub/grid/grid_barrier.cuh | 10 ++++++++--
 cub/test/test_grid_barrier.cu |  6 +++++-
 2 files changed, 13 insertions(+), 3 deletions(-)

diff --git a/cub/cub/grid/grid_barrier.cuh b/cub/cub/grid/grid_barrier.cuh
index 74ff85d6153..f2ae69fc091 100644
--- a/cub/cub/grid/grid_barrier.cuh
+++ b/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
diff --git a/cub/test/test_grid_barrier.cu b/cub/test/test_grid_barrier.cu
index 2f5ecfa3ebb..e763b48d1e2 100644
--- a/cub/test/test_grid_barrier.cu
+++ b/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