diff --git a/docs/libcudacxx/ptx/instructions/bfind.rst b/docs/libcudacxx/ptx/instructions/bfind.rst
new file mode 100644
index 00000000000..694cdef27fe
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/bfind.rst
@@ -0,0 +1,9 @@
+.. _libcudacxx-ptx-instructions-bfind:
+
+bfind
+=====
+
+-  PTX ISA:
+   `bfind <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-bfind>`__
+
+.. include:: generated/bfind.rst
diff --git a/docs/libcudacxx/ptx/instructions/exit.rst b/docs/libcudacxx/ptx/instructions/exit.rst
new file mode 100644
index 00000000000..899735f1660
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/exit.rst
@@ -0,0 +1,9 @@
+.. _libcudacxx-ptx-instructions-exit:
+
+exit
+====
+
+-  PTX ISA:
+   `exit <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-exit>`__
+
+.. include:: generated/exit.rst
diff --git a/docs/libcudacxx/ptx/instructions/trap.rst b/docs/libcudacxx/ptx/instructions/trap.rst
new file mode 100644
index 00000000000..54b641f357e
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/trap.rst
@@ -0,0 +1,9 @@
+.. _libcudacxx-ptx-instructions-trap:
+
+trap
+====
+
+-  PTX ISA:
+   `trap <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-trap>`__
+
+.. include:: generated/trap.rst
diff --git a/libcudacxx/include/cuda/__ptx/instructions/bfind.h b/libcudacxx/include/cuda/__ptx/instructions/bfind.h
new file mode 100644
index 00000000000..af8b5de98f8
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/bfind.h
@@ -0,0 +1,37 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_PTX_BFIND_H_
+#define _CUDA_PTX_BFIND_H_
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__ptx/ptx_dot_variants.h>
+#include <cuda/__ptx/ptx_helper_functions.h>
+#include <cuda/std/cstdint>
+
+#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX
+
+#include <cuda/__ptx/instructions/generated/bfind.h>
+
+_LIBCUDACXX_END_NAMESPACE_CUDA_PTX
+
+#endif // _CUDA_PTX_BFIND_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/exit.h b/libcudacxx/include/cuda/__ptx/instructions/exit.h
new file mode 100644
index 00000000000..6ae28b71e54
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/exit.h
@@ -0,0 +1,37 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_PTX_EXIT_H_
+#define _CUDA_PTX_EXIT_H_
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__ptx/ptx_dot_variants.h>
+#include <cuda/__ptx/ptx_helper_functions.h>
+#include <cuda/std/cstdint>
+
+#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX
+
+#include <cuda/__ptx/instructions/generated/exit.h>
+
+_LIBCUDACXX_END_NAMESPACE_CUDA_PTX
+
+#endif // _CUDA_PTX_EXIT_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/trap.h b/libcudacxx/include/cuda/__ptx/instructions/trap.h
new file mode 100644
index 00000000000..08ca0a0ec9e
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/trap.h
@@ -0,0 +1,37 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_PTX_TRAP_H_
+#define _CUDA_PTX_TRAP_H_
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__ptx/ptx_dot_variants.h>
+#include <cuda/__ptx/ptx_helper_functions.h>
+#include <cuda/std/cstdint>
+
+#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX
+
+#include <cuda/__ptx/instructions/generated/trap.h>
+
+_LIBCUDACXX_END_NAMESPACE_CUDA_PTX
+
+#endif // _CUDA_PTX_TRAP_H_
diff --git a/libcudacxx/include/cuda/ptx b/libcudacxx/include/cuda/ptx
index 971288b456c..5ef738936ae 100644
--- a/libcudacxx/include/cuda/ptx
+++ b/libcudacxx/include/cuda/ptx
@@ -70,6 +70,7 @@
  */
 
 #include <cuda/__ptx/instructions/barrier_cluster.h>
+#include <cuda/__ptx/instructions/bfind.h>
 #include <cuda/__ptx/instructions/clusterlaunchcontrol.h>
 #include <cuda/__ptx/instructions/cp_async_bulk.h>
 #include <cuda/__ptx/instructions/cp_async_bulk_commit_group.h>
@@ -78,6 +79,7 @@
 #include <cuda/__ptx/instructions/cp_async_mbarrier_arrive.h>
 #include <cuda/__ptx/instructions/cp_reduce_async_bulk.h>
 #include <cuda/__ptx/instructions/cp_reduce_async_bulk_tensor.h>
+#include <cuda/__ptx/instructions/exit.h>
 #include <cuda/__ptx/instructions/fence.h>
 #include <cuda/__ptx/instructions/get_sreg.h>
 #include <cuda/__ptx/instructions/getctarank.h>
@@ -103,5 +105,5 @@
 #include <cuda/__ptx/instructions/tcgen05_wait.h>
 #include <cuda/__ptx/instructions/tensormap_cp_fenceproxy.h>
 #include <cuda/__ptx/instructions/tensormap_replace.h>
-
+#include <cuda/__ptx/instructions/trap.h>
 #endif // _CUDA_PTX
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.bfind.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.bfind.compile.pass.cpp
new file mode 100644
index 00000000000..77f70363cec
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.bfind.compile.pass.cpp
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: libcpp-has-no-threads
+
+// <cuda/ptx>
+
+#include <cuda/ptx>
+#include <cuda/std/utility>
+
+#include "generated/bfind.h"
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.exit.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.exit.compile.pass.cpp
new file mode 100644
index 00000000000..568cd75b9fa
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.exit.compile.pass.cpp
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: libcpp-has-no-threads
+
+// <cuda/ptx>
+
+#include <cuda/ptx>
+#include <cuda/std/utility>
+
+#include "generated/exit.h"
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.trap.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.trap.compile.pass.cpp
new file mode 100644
index 00000000000..cd9032033f5
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.trap.compile.pass.cpp
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: libcpp-has-no-threads
+
+// <cuda/ptx>
+
+#include <cuda/ptx>
+#include <cuda/std/utility>
+
+#include "generated/trap.h"
+
+int main(int, char**)
+{
+  return 0;
+}