Revert "PTX: Add cuda::ptx::elect_sync (#1537)" (#1638)

* Revert "PTX: Add `cuda::ptx::elect_sync` (#1537)" This reverts commit 3d629a9. The exposed API is not good enough. * Update libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/ptx.h --------- Co-authored-by: Michael Schellenberger Costa <[email protected]>
NVIDIA · Apr 17, 2024 · 9a9bffd · 9a9bffd
1 parent b68256e
commit 9a9bffd
Show file tree

Hide file tree

Showing 5 changed files with 2 additions and 140 deletions.
diff --git a/libcudacxx/docs/ptx.md b/libcudacxx/docs/ptx.md
@@ -400,7 +400,7 @@ notes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#release
 | [`activemask`]        | No                      |
 | [`redux.sync`]        | No                      |
 | [`griddepcontrol`]    | No                      |
-| [`elect.sync`]        | CTK-FUTURE, CCCL v2.5.0 |
+| [`elect.sync`]        | No                      |
 
 [`bar, barrier`]: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar-barrier
 [`bar.warp.sync`]: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar-warp-sync
@@ -416,7 +416,7 @@ notes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#release
 [`activemask`]: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-activemask
 [`redux.sync`]: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-redux-sync
 [`griddepcontrol`]: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol
-[`elect.sync`]: ptx/instructions/elect.sync.md
+[`elect.sync`]: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-elect-sync
 
 ### [Parallel Synchronization and Communication Instructions: mbarrier](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier)
 

diff --git a/libcudacxx/docs/ptx/instructions/elect.sync.md b/libcudacxx/docs/ptx/instructions/elect.sync.md
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/ptx.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/ptx.h
@@ -76,7 +76,6 @@
 #include <cuda/std/detail/libcxx/include/__cuda/ptx/instructions/cp_async_bulk_wait_group.h>
 #include <cuda/std/detail/libcxx/include/__cuda/ptx/instructions/cp_reduce_async_bulk.h>
 #include <cuda/std/detail/libcxx/include/__cuda/ptx/instructions/cp_reduce_async_bulk_tensor.h>
-#include <cuda/std/detail/libcxx/include/__cuda/ptx/instructions/elect_sync.h>
 #include <cuda/std/detail/libcxx/include/__cuda/ptx/instructions/fence.h>
 #include <cuda/std/detail/libcxx/include/__cuda/ptx/instructions/get_sreg.h>
 #include <cuda/std/detail/libcxx/include/__cuda/ptx/instructions/getctarank.h>

diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/ptx/instructions/elect_sync.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/ptx/instructions/elect_sync.h
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.elect.sync.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.elect.sync.compile.pass.cpp