From c3843532d331b5ca74505772010254d3acb489b6 Mon Sep 17 00:00:00 2001
From: Michael Schellenberger Costa <miscco@nvidia.com>
Date: Wed, 22 Mar 2023 15:36:23 +0100
Subject: [PATCH] Add some tests

---
 .../assign.compile.fail.cpp                   |  36 ++
 .../thread.mutex.class/copy.compile.fail.cpp  |  35 ++
 .../thread.mutex.class/default.pass.cpp       |  32 ++
 .../thread.mutex.class/lock.pass.cpp          |  86 +++
 .../thread.mutex.class/try_lock.pass.cpp      |  90 +++
 .../assign.compile.fail.cpp                   |  36 ++
 .../copy.compile.fail.cpp                     |  35 ++
 .../thread.timedmutex.class/default.pass.cpp  |  31 +
 .../thread.timedmutex.class/lock.pass.cpp     |  86 +++
 .../thread.timedmutex.class/try_lock.pass.cpp |  90 +++
 .../try_lock_for.pass.cpp                     | 130 +++++
 .../try_lock_until.pass.cpp                   | 130 +++++
 .../thread.once.callonce/call_once.pass.cpp   | 313 ++++++++++
 .../thread.once.callonce/race.pass.cpp        |  91 +++
 .../assign.compile.fail.cpp                   |  37 ++
 .../copy.compile.fail.cpp                     |  36 ++
 .../thread.once.onceflag/default.pass.cpp     |  60 ++
 .../thread.lock.algorithm/lock.pass.cpp       | 543 ++++++++++++++++++
 .../thread.lock.algorithm/try_lock.pass.cpp   | 529 +++++++++++++++++
 .../thread.lock.guard/adopt_lock.pass.cpp     | 106 ++++
 .../thread.lock.guard/assign.compile.fail.cpp |  30 +
 .../thread.lock.guard/copy.compile.fail.cpp   |  28 +
 .../thread.lock.guard/implicit_ctad.pass.cpp  |  33 ++
 .../thread.lock.guard/mutex.fail.cpp          |  28 +
 .../thread.lock.guard/mutex.pass.cpp          | 117 ++++
 .../thread.lock.guard/types.pass.cpp          |  34 ++
 .../thread.lock.scoped/adopt_lock.pass.cpp    |  75 +++
 .../thread.lock.scoped/assign.fail.cpp        |  54 ++
 .../thread.lock.scoped/copy.fail.cpp          |  50 ++
 .../thread.lock.scoped/implicit_ctad.pass.cpp |  46 ++
 .../thread.lock.scoped/mutex.fail.cpp         |  57 ++
 .../thread.lock.scoped/mutex.pass.cpp         | 163 ++++++
 .../thread.lock.scoped/types.pass.cpp         |  85 +++
 .../thread.lock.unique/implicit_ctad.pass.cpp |  33 ++
 .../copy_assign.compile.fail.cpp              |  38 ++
 .../copy_ctor.compile.fail.cpp                |  36 ++
 .../thread.lock.unique.cons/default.pass.cpp  |  31 +
 .../move_assign.pass.cpp                      |  53 ++
 .../move_ctor.pass.cpp                        |  49 ++
 .../thread.lock.unique.cons/mutex.pass.cpp    |  94 +++
 .../mutex_adopt_lock.pass.cpp                 |  48 ++
 .../mutex_defer_lock.pass.cpp                 |  43 ++
 .../mutex_duration.pass.cpp                   | 102 ++++
 .../mutex_time_point.pass.cpp                 | 100 ++++
 .../mutex_try_to_lock.pass.cpp                |  92 +++
 .../thread.lock.unique.locking/lock.pass.cpp  | 101 ++++
 .../try_lock.pass.cpp                         |  75 +++
 .../try_lock_for.pass.cpp                     |  78 +++
 .../try_lock_until.pass.cpp                   |  83 +++
 .../unlock.pass.cpp                           |  65 +++
 .../member_swap.pass.cpp                      |  43 ++
 .../nonmember_swap.pass.cpp                   |  44 ++
 .../thread.lock.unique.mod/release.pass.cpp   |  48 ++
 .../thread.lock.unique.obs/mutex.pass.cpp     |  36 ++
 .../thread.lock.unique.obs/op_bool.pass.cpp   |  41 ++
 .../thread.lock.unique.obs/owns_lock.pass.cpp |  36 ++
 .../thread.lock.unique/types.pass.cpp         |  34 ++
 .../thread.mutex/thread.lock/types.fail.cpp   |  32 ++
 .../thread.mutex/thread.lock/types.pass.cpp   |  38 ++
 .../nothing_to_do.pass.cpp                    |  14 +
 .../assign.compile.fail.cpp                   |  28 +
 .../thread.mutex.class/copy.compile.fail.cpp  |  27 +
 .../thread.mutex.class/default.pass.cpp       |  30 +
 .../thread.mutex.class/lock.pass.cpp          |  75 +++
 .../thread.mutex.class/try_lock.pass.cpp      |  79 +++
 .../assign.compile.fail.cpp                   |  28 +
 .../copy.compile.fail.cpp                     |  27 +
 .../thread.timedmutex.class/default.pass.cpp  |  29 +
 .../thread.timedmutex.class/lock.pass.cpp     |  75 +++
 .../thread.timedmutex.class/try_lock.pass.cpp |  80 +++
 .../try_lock_for.pass.cpp                     |  99 ++++
 .../try_lock_until.pass.cpp                   |  99 ++++
 .../thread.once.callonce/call_once.pass.cpp   | 303 ++++++++++
 .../thread.once.callonce/race.pass.cpp        |  80 +++
 .../assign.compile.fail.cpp                   |  29 +
 .../copy.compile.fail.cpp                     |  28 +
 .../thread.once.onceflag/default.pass.cpp     |  36 ++
 .../support/heterogeneous_thread_handler.h    |  91 +++
 .../test/support/nasty_containers.h           |  24 +-
 79 files changed, 6074 insertions(+), 12 deletions(-)
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/assign.compile.fail.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/copy.compile.fail.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/default.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/assign.compile.fail.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/copy.compile.fail.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/default.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/lock.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_for.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_until.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/assign.compile.fail.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/copy.compile.fail.cpp
 create mode 100644 .upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/default.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock.algorithm/lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock.algorithm/try_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/adopt_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/assign.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/copy.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/implicit_ctad.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/types.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/adopt_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/assign.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/copy.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/implicit_ctad.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/mutex.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/mutex.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/types.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/implicit_ctad.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/copy_assign.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/copy_ctor.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/default.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/move_assign.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/move_ctor.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_adopt_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_defer_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_duration.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_time_point.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_try_to_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock_for.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock_until.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/unlock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/member_swap.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/nonmember_swap.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/release.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/mutex.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/op_bool.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/owns_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/types.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/types.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.lock/types.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.general/nothing_to_do.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/assign.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/copy.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/default.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/assign.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/copy.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/default.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_for.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_until.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/assign.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/copy.compile.fail.cpp
 create mode 100644 .upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/default.pass.cpp
 create mode 100644 .upstream-tests/test/support/heterogeneous_thread_handler.h

diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/assign.compile.fail.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/assign.compile.fail.cpp
new file mode 100644
index 0000000000..72c8cfc9d8
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/assign.compile.fail.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class mutex;
+
+// mutex& operator=(const mutex&) = delete;
+
+#include<cuda/mutex>
+
+template<int scope>
+void test() {
+    cuda::mutex<scope> m0;
+    cuda::mutex<scope> m1;
+    m1 = m0;
+}
+
+int main(int, char**)
+{
+  test<cuda::thread_scope_system>();
+  test<cuda::thread_scope_device>();
+  test<cuda::thread_scope_block>();
+  test<cuda::thread_scope_thread>();
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/copy.compile.fail.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/copy.compile.fail.cpp
new file mode 100644
index 0000000000..e6284eaa49
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/copy.compile.fail.cpp
@@ -0,0 +1,35 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class mutex;
+
+// mutex(const mutex&) = delete;
+
+#include<cuda/mutex>
+
+template<int scope>
+void test() {
+    cuda::mutex<scope> m0;
+    cuda::mutex<scope> m1{m0};
+}
+
+int main(int, char**)
+{
+  test<cuda::thread_scope_system>();
+  test<cuda::thread_scope_device>();
+  test<cuda::thread_scope_block>();
+  test<cuda::thread_scope_thread>();
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/default.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/default.pass.cpp
new file mode 100644
index 0000000000..b12a450468
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/default.pass.cpp
@@ -0,0 +1,32 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class mutex;
+
+// mutex();
+
+#include<cuda/mutex>
+#include<cuda/std/type_traits>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::mutex<cuda::thread_scope_system>>::value, "");
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::mutex<cuda::thread_scope_device>>::value, "");
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::mutex<cuda::thread_scope_block>>::value, "");
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::mutex<cuda::thread_scope_thread>>::value, "");
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp
new file mode 100644
index 0000000000..f02cd5b84c
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp
@@ -0,0 +1,86 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class mutex;
+
+// void lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    m->lock();
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::mutex<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::mutex<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::mutex<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp
new file mode 100644
index 0000000000..2aad657911
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp
@@ -0,0 +1,90 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class mutex;
+
+// bool try_lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    while(!m->try_lock())
+        ;
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::mutex<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::mutex<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::mutex<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/assign.compile.fail.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/assign.compile.fail.cpp
new file mode 100644
index 0000000000..65ea3bfe88
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/assign.compile.fail.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class timed_mutex;
+
+// timed_mutex& operator=(const timed_mutex&) = delete;
+
+#include<cuda/mutex>
+
+template<int scope>
+void test() {
+    cuda::timed_mutex<scope> m0;
+    cuda::timed_mutex<scope> m1;
+    m1 = m0;
+}
+
+int main(int, char**)
+{
+  test<cuda::thread_scope_system>();
+  test<cuda::thread_scope_device>();
+  test<cuda::thread_scope_block>();
+  test<cuda::thread_scope_thread>();
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/copy.compile.fail.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/copy.compile.fail.cpp
new file mode 100644
index 0000000000..2d0af92a88
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/copy.compile.fail.cpp
@@ -0,0 +1,35 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class timed_mutex;
+
+// timed_mutex(const timed_mutex&) = delete;
+
+#include<cuda/mutex>
+
+template<int scope>
+void test() {
+    cuda::timed_mutex<scope> m0;
+    cuda::timed_mutex<scope> m1{m0};
+}
+
+int main(int, char**)
+{
+  test<cuda::thread_scope_system>();
+  test<cuda::thread_scope_device>();
+  test<cuda::thread_scope_block>();
+  test<cuda::thread_scope_thread>();
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/default.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/default.pass.cpp
new file mode 100644
index 0000000000..cca4e1207f
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/default.pass.cpp
@@ -0,0 +1,31 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class timed_mutex;
+
+// timed_mutex();
+
+#include<cuda/mutex>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::timed_mutex<cuda::thread_scope_system>>::value, "");
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::timed_mutex<cuda::thread_scope_device>>::value, "");
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::timed_mutex<cuda::thread_scope_block>>::value, "");
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::timed_mutex<cuda::thread_scope_thread>>::value, "");
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/lock.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/lock.pass.cpp
new file mode 100644
index 0000000000..5a0dcd474f
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/lock.pass.cpp
@@ -0,0 +1,86 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// void lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    m->lock();
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::mutex<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::mutex<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::mutex<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock.pass.cpp
new file mode 100644
index 0000000000..29c68dd3bb
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock.pass.cpp
@@ -0,0 +1,90 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// bool try_lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    while(!m->try_lock())
+        ;
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::mutex<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::mutex<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::mutex<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_for.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_for.pass.cpp
new file mode 100644
index 0000000000..1a1d9c62d8
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_for.pass.cpp
@@ -0,0 +1,130 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+// UNSUPPORTED: true
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// template <class Rep, class Period>
+//     bool try_lock_for(const chrono::duration<Rep, Period>& rel_time);
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// bool try_lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f1(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(m->try_lock_for(ms(300)) == true);
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex>
+__host__ __device__ void f2(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(m->try_lock_for(ms(250)) == false);
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    {
+        handler.run_on_first_thread(&Mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1<Mutex>, m);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&Mutex::unlock, m);
+        handler.join_test_thread();
+    }
+
+    handler.syncthreads();
+
+    {
+        handler.run_on_first_thread(&Mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f2<Mutex>, m);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&Mutex::unlock, m);
+        handler.join_test_thread();
+    }
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::mutex<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::mutex<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::mutex<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_until.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_until.pass.cpp
new file mode 100644
index 0000000000..6048b6614b
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_until.pass.cpp
@@ -0,0 +1,130 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+// UNSUPPORTED: true
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// template <class Clock, class Duration>
+//     bool try_lock_until(const chrono::time_point<Clock, Duration>& abs_time);
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// bool try_lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f1(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(m->try_lock_until(Clock::now() + ms(300)) == true);
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex>
+__host__ __device__ void f2(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(m->try_lock_until(Clock::now() + ms(200)) == false);
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    {
+        handler.run_on_first_thread(&Mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1<Mutex>, m);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&Mutex::unlock, m);
+        handler.join_test_thread();
+    }
+
+    handler.syncthreads();
+
+    {
+        handler.run_on_first_thread(&Mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f2<Mutex>, m);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&Mutex::unlock, m);
+        handler.join_test_thread();
+    }
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::mutex<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::mutex<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::mutex<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp
new file mode 100644
index 0000000000..f8fb38bcd3
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp
@@ -0,0 +1,313 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// template<class Callable, class ...Args>
+//   void call_once(once_flag& flag, Callable&& func, Args&&... args);
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::milliseconds ms;
+
+__host__ __device__
+void init0(int& init0_called)
+{
+    test_sleep_thread(ms(250));
+    ++init0_called;
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f0(OnceFlag* flg0, int* init0_called)
+{
+    cuda::std::call_once(*flg0, init0, *init0_called);
+}
+
+#ifndef TEST_HAS_NO_EXCEPTIONS
+__host__ __device__
+void init3(int& init3_called, int& init3_completed)
+{
+    ++init3_called;
+    test_sleep_thread(ms(250));
+    if (init3_called == 1)
+#ifdef __CUDA_ARCH__
+        _LIBCUDACXX_UNREACHABLE();
+#else
+        TEST_THROW(1);
+#endif
+    ++init3_completed;
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f3(OnceFlag* flg3, int* init3_called, int* init3_completed)
+{
+    try
+    {
+        cuda::std::call_once(*flg3, init3, *init3_called, *init3_completed);
+    }
+    catch (...)
+    {
+    }
+}
+#endif // TEST_HAS_NO_EXCEPTIONS
+
+#if TEST_STD_VER >= 11
+struct init1
+{
+    int& init1_called;
+    __host__ __device__ void operator()(int i) { init1_called += i; }
+};
+
+template<class OnceFlag>
+__host__ __device__
+void f1(OnceFlag* flg1, int* init1_called)
+{
+    cuda::std::call_once(*flg1, init1{*init1_called}, 1);
+}
+
+struct init2
+{
+    int& init2_called;
+    __host__ __device__ void operator()(int i, int j) const {init2_called += i + j;}
+};
+
+template<class OnceFlag>
+__host__ __device__
+void f2(OnceFlag* flg2, int* init2_called)
+{
+    cuda::std::call_once(*flg2, init2{*init2_called}, 2, 3);
+    cuda::std::call_once(*flg2, init2{*init2_called}, 4, 5);
+}
+
+#endif // TEST_STD_VER >= 11
+__host__ __device__
+void init41(int& init41_called)
+{
+    test_sleep_thread(ms(250));
+    ++init41_called;
+}
+
+__host__ __device__
+void init42(int& init42_called)
+{
+    test_sleep_thread(ms(250));
+    ++init42_called;
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f41(OnceFlag* flg41, OnceFlag* flg42, int* init41_called, int* init42_called)
+{
+    cuda::std::call_once(*flg41, init41, *init41_called);
+    cuda::std::call_once(*flg42, init42, *init42_called);
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f42(OnceFlag* flg41, OnceFlag* flg42, int* init41_called, int* init42_called)
+{
+    cuda::std::call_once(*flg41, init42, *init42_called);
+    cuda::std::call_once(*flg42, init41, *init41_called);
+}
+
+#if TEST_STD_VER >= 11
+
+class MoveOnly
+{
+    __host__ __device__ MoveOnly(const MoveOnly&);
+public:
+    __host__ __device__ MoveOnly() {}
+    __host__ __device__ MoveOnly(MoveOnly&&) {}
+
+    __host__ __device__ void operator()(MoveOnly&&) {}
+};
+
+class NonCopyable
+{
+    __host__ __device__ NonCopyable(const NonCopyable&);
+public:
+    __host__ __device__ NonCopyable() {}
+
+    __host__ __device__ void operator()(int&) {}
+};
+
+// reference qualifiers on functions are a C++11 extension
+struct RefQual
+{
+    __host__ __device__ void operator()( int& lv_called) & { ++lv_called; }
+    __host__ __device__ void operator()( int& rv_called) && { ++rv_called; }
+};
+
+#endif // TEST_STD_VER >= 11
+
+template<class OnceFlag,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<OnceFlag, Initializer> sel;
+    Selector<int, Initializer> sel_count;
+    heterogeneous_thread_handler handler;
+
+    // check basic functionality
+    {
+        SHARED OnceFlag* flg0;
+        flg0 = sel.construct();
+
+        SHARED int* init0_called;
+        init0_called = sel_count.construct();
+
+        handler.run_on_first_thread(f0<OnceFlag>, flg0, init0_called);
+        handler.run_on_second_thread(f0<OnceFlag>, flg0, init0_called);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(*init0_called == 1);
+    }
+
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    // check basic exception safety
+    {
+        SHARED OnceFlag* flg0;
+        flg0 = sel.construct();
+
+        SHARED int* init3_called;
+        init3_called = sel_count.construct();
+        SHARED int* init3_completed;
+        init3_completed = sel_count.construct();
+
+        handler.run_on_first_thread( f3<OnceFlag>, flg3, init3_called, init3_completed);
+        handler.run_on_second_thread(f3<OnceFlag>, flg3, init3_called, init3_completed);
+        handler.join_test_thread();
+        handler.syncthreads();
+
+        assert(*init3_called == 2);
+        assert(*init3_completed == 1);
+    }
+#endif
+    // check deadlock avoidance
+    {
+        SHARED OnceFlag* flg41;
+        flg41 = sel.construct();
+        SHARED OnceFlag* flg42;
+        flg42 = sel.construct();
+
+        SHARED int* init41_called;
+        init41_called = sel_count.construct();
+        SHARED int* init42_called;
+        init42_called = sel_count.construct();
+
+        handler.run_on_first_thread( f41<OnceFlag>, flg41, flg42, init41_called, init42_called);
+        handler.run_on_second_thread(f42<OnceFlag>, flg41, flg42, init41_called, init42_called);
+        handler.join_test_thread();
+        handler.syncthreads();
+
+        assert(*init41_called == 1);
+        assert(*init42_called == 1);
+    }
+#if TEST_STD_VER >= 11
+    // check functors with 1 arg
+    {
+        SHARED OnceFlag* flg1;
+        flg1 = sel.construct();
+
+        SHARED int* init1_called;
+        init1_called = sel_count.construct();
+
+        handler.run_on_first_thread( f1<OnceFlag>, flg1, init1_called);
+        handler.run_on_second_thread(f1<OnceFlag>, flg1, init1_called);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(*init1_called == 1);
+    }
+    // check functors with 2 args
+    {
+        SHARED OnceFlag* flg2;
+        flg2 = sel.construct();
+
+        SHARED int* init2_called;
+        init2_called = sel_count.construct();
+
+        handler.run_on_first_thread( f2<OnceFlag>, flg2, init2_called);
+        handler.run_on_second_thread(f2<OnceFlag>, flg2, init2_called);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(*init2_called == 5);
+    }
+
+    {
+        SHARED OnceFlag* f;
+        f = sel.construct();
+
+        cuda::call_once(*f, MoveOnly(), MoveOnly());
+    }
+    // check LWG2442: call_once() shouldn't DECAY_COPY()
+    {
+        SHARED OnceFlag* f;
+        f = sel.construct();
+        int i = 0;
+
+        cuda::call_once(*f, NonCopyable(), i);
+    }
+    // reference qualifiers on functions are a C++11 extension
+    {
+        SHARED OnceFlag* f1, *f2;
+        f1 = sel.construct();
+        f2 = sel.construct();
+
+        SHARED int* lv_called;
+        lv_called = sel_count.construct();
+
+        SHARED int* rv_called;
+        rv_called = sel_count.construct();
+
+        RefQual rq;
+        cuda::call_once(*f1, rq, *lv_called);
+        assert(*lv_called == 1);
+        cuda::call_once(*f2, cuda::std::move(rq), *rv_called);
+        assert(*rv_called == 1);
+    }
+#endif // TEST_STD_VER >= 11
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::once_flag<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::once_flag<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::once_flag<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp
new file mode 100644
index 0000000000..4da2b0380a
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp
@@ -0,0 +1,91 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// template<class Callable, class ...Args>
+//   void call_once(once_flag& flag, Callable&& func, Args&&... args);
+
+// This test is supposed to be run with ThreadSanitizer and verifies that
+// call_once properly synchronizes user state, a data race that was fixed
+// in r280621.
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::milliseconds ms;
+
+__host__ __device__
+void init0(int& global)
+{
+    ++global;
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f0(OnceFlag* flg, int* global)
+{
+    cuda::call_once(*flg, init0, *global);
+    assert(*global == 1);
+}
+
+template<class OnceFlag,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<OnceFlag, Initializer> sel;
+    SHARED OnceFlag* flg;
+    flg = sel.construct();
+
+    Selector<int, Initializer> sel_count;
+    SHARED int* global;
+    global = sel_count.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(f0<OnceFlag>, flg, global);
+    handler.run_on_second_thread(f0<OnceFlag>, flg, global);
+    handler.syncthreads();
+    handler.join_test_thread();
+    assert(*global == 1);
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::once_flag<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::once_flag<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::once_flag<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::once_flag<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/assign.compile.fail.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/assign.compile.fail.cpp
new file mode 100644
index 0000000000..ab55f1db89
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/assign.compile.fail.cpp
@@ -0,0 +1,37 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// once_flag& operator=(const once_flag&) = delete;
+
+#include<cuda/mutex>
+
+template<int scope>
+void test() {
+    cuda::once_flag<scope> f0;
+    cuda::once_flag<scope> f1;
+    f1 = f0;
+}
+
+int main(int, char**)
+{
+  test<cuda::thread_scope_system>();
+  test<cuda::thread_scope_device>();
+  test<cuda::thread_scope_block>();
+  test<cuda::thread_scope_thread>();
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/copy.compile.fail.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/copy.compile.fail.cpp
new file mode 100644
index 0000000000..c1c0f85ff4
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/copy.compile.fail.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// once_flag(const once_flag&) = delete;
+
+#include<cuda/mutex>
+
+template<int scope>
+void test() {
+    cuda::once_flag<scope> f0;
+    cuda::once_flag<scope> f1{f0};
+}
+
+int main(int, char**)
+{
+  test<cuda::thread_scope_system>();
+  test<cuda::thread_scope_device>();
+  test<cuda::thread_scope_block>();
+  test<cuda::thread_scope_thread>();
+
+  return 0;
+}
diff --git a/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/default.pass.cpp b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/default.pass.cpp
new file mode 100644
index 0000000000..31ebaf0c4d
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.onceflag/default.pass.cpp
@@ -0,0 +1,60 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// constexpr once_flag() noexcept;
+
+#include<cuda/mutex>
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    {
+    cuda::once_flag<cuda::thread_scope_system> f;
+    unused(f);
+    }
+    {
+    cuda::once_flag<cuda::thread_scope_device> f;
+    unused(f);
+    }
+    {
+    cuda::once_flag<cuda::thread_scope_block> f;
+    unused(f);
+    }
+    {
+    cuda::once_flag<cuda::thread_scope_thread> f;
+    unused(f);
+    }
+#if TEST_STD_VER >= 11
+    {
+    constexpr cuda::once_flag<cuda::thread_scope_system> f;
+    unused(f);
+    }
+    {
+    constexpr cuda::once_flag<cuda::thread_scope_device> f;
+    unused(f);
+    }
+    {
+    constexpr cuda::once_flag<cuda::thread_scope_block> f;
+    unused(f);
+    }
+    {
+    constexpr cuda::once_flag<cuda::thread_scope_thread> f;
+    unused(f);
+    }
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock.algorithm/lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock.algorithm/lock.pass.cpp
new file mode 100644
index 0000000000..126b289ec5
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock.algorithm/lock.pass.cpp
@@ -0,0 +1,543 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// This test hangs forever when built against libstdc++ (Oct 2016).
+// UNSUPPORTED: stdlib=libstdc++
+
+// This test isn't quite standards-conforming: it's testing our specific
+// algorithm, where when lx.try_lock() fails we start the next attempt
+// with an unconditional lx.lock(). Thus our algorithm can handle a list
+// of mutexes where at-most-one of them is of the evil type `class L1`,
+// but will loop forever if two or more of them are `class L1`.
+
+// <mutex>
+
+// template <class L1, class L2, class... L3>
+//   void lock(L1&, L2&, L3&...);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+class L0
+{
+    bool locked_;
+
+public:
+    __host__ __device__ L0() : locked_(false) {}
+
+    __host__ __device__ void lock()
+    {
+        locked_ = true;
+    }
+
+    __host__ __device__ bool try_lock()
+    {
+        locked_ = true;
+        return locked_;
+    }
+
+    __host__ __device__ void unlock() {locked_ = false;}
+
+    __host__ __device__ bool locked() const {return locked_;}
+};
+
+class L1
+{
+    bool locked_;
+
+public:
+    __host__ __device__ L1() : locked_(false) {}
+
+    __host__ __device__ void lock()
+    {
+        locked_ = true;
+    }
+
+    __host__ __device__ bool try_lock()
+    {
+        locked_ = false;
+        return locked_;
+    }
+
+    __host__ __device__ void unlock() {locked_ = false;}
+
+    __host__ __device__ bool locked() const {return locked_;}
+};
+
+class L2
+{
+    bool locked_;
+
+public:
+    __host__ __device__ L2() : locked_(false) {}
+
+    __host__ __device__ void lock()
+    {
+        TEST_THROW(1);
+    }
+
+    __host__ __device__ bool try_lock()
+    {
+        TEST_THROW(1);
+        return locked_;
+    }
+
+    __host__ __device__ void unlock() {locked_ = false;}
+
+    __host__ __device__ bool locked() const {return locked_;}
+};
+
+__host__ __device__
+void with_one_or_two_locks() {
+    {
+        L0 l0;
+        L0 l1;
+        cuda::std::lock(l0, l1);
+        assert(l0.locked());
+        assert(l1.locked());
+    }
+    {
+        L0 l0;
+        L1 l1;
+        cuda::std::lock(l0, l1);
+        assert(l0.locked());
+        assert(l1.locked());
+    }
+    {
+        L1 l0;
+        L0 l1;
+        cuda::std::lock(l0, l1);
+        assert(l0.locked());
+        assert(l1.locked());
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    {
+        L0 l0;
+        L2 l1;
+        try
+        {
+            cuda::std::lock(l0, l1);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+        }
+    }
+    {
+        L2 l0;
+        L0 l1;
+        try
+        {
+            cuda::std::lock(l0, l1);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+        }
+    }
+    {
+        L1 l0;
+        L2 l1;
+        try
+        {
+            cuda::std::lock(l0, l1);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+        }
+    }
+    {
+        L2 l0;
+        L1 l1;
+        try
+        {
+            cuda::std::lock(l0, l1);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+        }
+    }
+    {
+        L2 l0;
+        L2 l1;
+        try
+        {
+            cuda::std::lock(l0, l1);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+        }
+    }
+#endif
+}
+
+__host__ __device__
+void with_three_locks() {
+    {
+        L0 l0;
+        L0 l1;
+        L0 l2;
+        cuda::std::lock(l0, l1, l2);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    {
+        L2 l0;
+        L2 l1;
+        L2 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+#endif
+    {
+        L0 l0;
+        L0 l1;
+        L1 l2;
+        cuda::std::lock(l0, l1, l2);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+    }
+    {
+        L0 l0;
+        L1 l1;
+        L0 l2;
+        cuda::std::lock(l0, l1, l2);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+    }
+    {
+        L1 l0;
+        L0 l1;
+        L0 l2;
+        cuda::std::lock(l0, l1, l2);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    {
+        L0 l0;
+        L0 l1;
+        L2 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L0 l0;
+        L2 l1;
+        L0 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L0 l1;
+        L0 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L2 l1;
+        L0 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L0 l1;
+        L2 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L0 l0;
+        L2 l1;
+        L2 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L2 l1;
+        L1 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L1 l1;
+        L2 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L1 l0;
+        L2 l1;
+        L2 l2;
+        try
+        {
+            cuda::std::lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+#endif // TEST_HAS_NO_EXCEPTIONS
+}
+
+__host__ __device__
+void with_four_locks() {
+{
+        L0 l0;
+        L0 l1;
+        L0 l2;
+        L0 l3;
+        cuda::std::lock(l0, l1, l2, l3);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+        assert(l3.locked());
+    }
+    {
+        L0 l0;
+        L0 l1;
+        L0 l2;
+        L1 l3;
+        cuda::std::lock(l0, l1, l2, l3);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+        assert(l3.locked());
+    }
+    {
+        L0 l0;
+        L0 l1;
+        L1 l2;
+        L0 l3;
+        cuda::std::lock(l0, l1, l2, l3);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+        assert(l3.locked());
+    }
+    {
+        L0 l0;
+        L1 l1;
+        L0 l2;
+        L0 l3;
+        cuda::std::lock(l0, l1, l2, l3);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+        assert(l3.locked());
+    }
+    {
+        L1 l0;
+        L0 l1;
+        L0 l2;
+        L0 l3;
+        cuda::std::lock(l0, l1, l2, l3);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+        assert(l3.locked());
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    {
+        L0 l0;
+        L0 l1;
+        L0 l2;
+        L2 l3;
+        try
+        {
+            cuda::std::lock(l0, l1, l2, l3);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+            assert(!l3.locked());
+        }
+    }
+    {
+        L0 l0;
+        L0 l1;
+        L2 l2;
+        L0 l3;
+        try
+        {
+            cuda::std::lock(l0, l1, l2, l3);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+            assert(!l3.locked());
+        }
+    }
+    {
+        L0 l0;
+        L2 l1;
+        L0 l2;
+        L0 l3;
+        try
+        {
+            cuda::std::lock(l0, l1, l2, l3);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+            assert(!l3.locked());
+        }
+    }
+    {
+        L2 l0;
+        L0 l1;
+        L0 l2;
+        L0 l3;
+        try
+        {
+            cuda::std::lock(l0, l1, l2, l3);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+            assert(!l3.locked());
+        }
+    }
+#endif // TEST_HAS_NO_EXCEPTIONS
+}
+
+int main(int, char**)
+{
+    with_one_or_two_locks();
+    with_three_locks();
+#ifndef __CUDA_ARCH__ // explodes stack space
+    with_four_locks();
+#endif
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock.algorithm/try_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock.algorithm/try_lock.pass.cpp
new file mode 100644
index 0000000000..b0e473109f
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock.algorithm/try_lock.pass.cpp
@@ -0,0 +1,529 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class L1, class L2, class... L3>
+//   int try_lock(L1&, L2&, L3&...);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+class L0
+{
+    bool locked_;
+
+public:
+    __host__ __device__ L0() : locked_(false) {}
+
+    __host__ __device__ bool try_lock()
+    {
+        locked_ = true;
+        return locked_;
+    }
+
+    __host__ __device__ void unlock() {locked_ = false;}
+
+    __host__ __device__ bool locked() const {return locked_;}
+};
+
+class L1
+{
+    bool locked_;
+
+public:
+    __host__ __device__ L1() : locked_(false) {}
+
+    __host__ __device__ bool try_lock()
+    {
+        locked_ = false;
+        return locked_;
+    }
+
+    __host__ __device__ void unlock() {locked_ = false;}
+
+    __host__ __device__ bool locked() const {return locked_;}
+};
+
+class L2
+{
+    bool locked_;
+
+public:
+    __host__ __device__ L2() : locked_(false) {}
+
+    __host__ __device__ bool try_lock()
+    {
+        TEST_THROW(1);
+        return locked_;
+    }
+
+    __host__ __device__ void unlock() {locked_ = false;}
+
+    __host__ __device__ bool locked() const {return locked_;}
+};
+
+int main(int, char**)
+{
+    {
+        L0 l0;
+        L0 l1;
+        assert(cuda::std::try_lock(l0, l1) == -1);
+        assert(l0.locked());
+        assert(l1.locked());
+    }
+    {
+        L0 l0;
+        L1 l1;
+        assert(cuda::std::try_lock(l0, l1) == 1);
+        assert(!l0.locked());
+        assert(!l1.locked());
+    }
+    {
+        L1 l0;
+        L0 l1;
+        assert(cuda::std::try_lock(l0, l1) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    {
+        L0 l0;
+        L2 l1;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+        }
+    }
+    {
+        L2 l0;
+        L0 l1;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+        }
+    }
+#endif
+#if TEST_STD_VER >= 11
+    {
+        L0 l0;
+        L0 l1;
+        L0 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == -1);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+    }
+    {
+        L1 l0;
+        L1 l1;
+        L1 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    {
+        L2 l0;
+        L2 l1;
+        L2 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L0 l0;
+        L1 l1;
+        L2 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 1);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+#endif
+    {
+        L0 l0;
+        L0 l1;
+        L1 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 2);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L0 l0;
+        L1 l1;
+        L0 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 1);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L1 l0;
+        L0 l1;
+        L0 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    {
+        L0 l0;
+        L0 l1;
+        L2 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L0 l0;
+        L2 l1;
+        L0 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L0 l1;
+        L0 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+#endif
+    {
+        L1 l0;
+        L1 l1;
+        L0 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L1 l0;
+        L0 l1;
+        L1 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L0 l0;
+        L1 l1;
+        L1 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 1);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    {
+        L1 l0;
+        L1 l1;
+        L2 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L1 l0;
+        L2 l1;
+        L1 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L2 l0;
+        L1 l1;
+        L1 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L2 l1;
+        L0 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L0 l1;
+        L2 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L0 l0;
+        L2 l1;
+        L2 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L2 l1;
+        L1 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L1 l1;
+        L2 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L1 l0;
+        L2 l1;
+        L2 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L0 l0;
+        L2 l1;
+        L1 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L1 l0;
+        L0 l1;
+        L2 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L1 l0;
+        L2 l1;
+        L0 l2;
+        assert(cuda::std::try_lock(l0, l1, l2) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+    }
+    {
+        L2 l0;
+        L0 l1;
+        L1 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+    {
+        L2 l0;
+        L1 l1;
+        L0 l2;
+        try
+        {
+            (void)cuda::std::try_lock(l0, l1, l2);
+            assert(false);
+        }
+        catch (int)
+        {
+            assert(!l0.locked());
+            assert(!l1.locked());
+            assert(!l2.locked());
+        }
+    }
+#endif // TEST_HAS_NO_EXCEPTIONS
+    {
+        L0 l0;
+        L0 l1;
+        L0 l2;
+        L0 l3;
+        assert(cuda::std::try_lock(l0, l1, l2, l3) == -1);
+        assert(l0.locked());
+        assert(l1.locked());
+        assert(l2.locked());
+        assert(l3.locked());
+    }
+    {
+        L1 l0;
+        L0 l1;
+        L0 l2;
+        L0 l3;
+        assert(cuda::std::try_lock(l0, l1, l2, l3) == 0);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+        assert(!l3.locked());
+    }
+    {
+        L0 l0;
+        L1 l1;
+        L0 l2;
+        L0 l3;
+        assert(cuda::std::try_lock(l0, l1, l2, l3) == 1);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+        assert(!l3.locked());
+    }
+    {
+        L0 l0;
+        L0 l1;
+        L1 l2;
+        L0 l3;
+        assert(cuda::std::try_lock(l0, l1, l2, l3) == 2);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+        assert(!l3.locked());
+    }
+    {
+        L0 l0;
+        L0 l1;
+        L0 l2;
+        L1 l3;
+        assert(cuda::std::try_lock(l0, l1, l2, l3) == 3);
+        assert(!l0.locked());
+        assert(!l1.locked());
+        assert(!l2.locked());
+        assert(!l3.locked());
+    }
+#endif // TEST_STD_VER >= 11
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/adopt_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/adopt_lock.pass.cpp
new file mode 100644
index 0000000000..104d4604b5
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/adopt_lock.pass.cpp
@@ -0,0 +1,106 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+#define _LIBCUDACXX_INLINE_THREADING
+// <mutex>
+
+// template <class Mutex> class lock_guard;
+
+// lock_guard(mutex_type& m, adopt_lock_t);
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+#ifndef __CUDA_ARCH__
+template<class Mutex>
+__host__ void do_try_lock(Mutex* m) {
+  assert(m->try_lock() == false);
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ void test_host() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    {
+      m->lock();
+      cuda::std::lock_guard<Mutex> lg(*m, cuda::std::adopt_lock);
+      std::thread t(do_try_lock<Mutex>, m);
+      t.join();
+    }
+
+    m->lock();
+    m->unlock();
+}
+#endif
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__device__ void test_device() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    Selector<cuda::std::atomic_flag, Initializer> sel_flag;
+    SHARED cuda::std::atomic_flag* thread_ready;
+    thread_ready = sel_flag.construct(false);
+
+    if (threadIdx.x == 0) {
+      m->lock();
+    }
+    __syncthreads();
+
+    if (threadIdx.x == 0) {
+      cuda::std::lock_guard<Mutex> lg(*m, cuda::std::adopt_lock);
+      thread_ready->test_and_set(cuda::std::memory_order_relaxed);
+      thread_ready->notify_one();
+      thread_ready->wait(true, cuda::std::memory_order_relaxed);  // Waits until clear
+    } else {
+      thread_ready->wait(false, cuda::std::memory_order_relaxed); // Waits until test_and_set
+      assert(m->try_lock() == false);
+      thread_ready->clear(cuda::std::memory_order_relaxed);
+      thread_ready->notify_one();
+    }
+    __syncthreads();
+
+    m->lock();
+    m->unlock();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test_host<cuda::std::mutex, local_memory_selector>();
+#else
+    test_device<cuda::std::mutex, shared_memory_selector>();
+    test_device<cuda::std::mutex, global_memory_selector>();
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/assign.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/assign.compile.fail.cpp
new file mode 100644
index 0000000000..7d0940f07a
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/assign.compile.fail.cpp
@@ -0,0 +1,30 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class lock_guard;
+
+// lock_guard& operator=(lock_guard const&) = delete;
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::mutex m0;
+    cuda::std::mutex m1;
+    cuda::std::lock_guard<cuda::std::mutex> lg0(m0);
+    cuda::std::lock_guard<cuda::std::mutex> lg(m1);
+    lg = lg0;
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/copy.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/copy.compile.fail.cpp
new file mode 100644
index 0000000000..82776bb157
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/copy.compile.fail.cpp
@@ -0,0 +1,28 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class lock_guard;
+
+// lock_guard(lock_guard const&) = delete;
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::mutex m;
+    cuda::std::lock_guard<cuda::std::mutex> lg0(m);
+    cuda::std::lock_guard<cuda::std::mutex> lg(lg0);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/implicit_ctad.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/implicit_ctad.pass.cpp
new file mode 100644
index 0000000000..e4d7553866
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/implicit_ctad.pass.cpp
@@ -0,0 +1,33 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++98, c++03, c++11, c++14
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// lock_guard
+
+// Make sure that the implicitly-generated CTAD works.
+
+#include<cuda/std/mutex>
+
+#include "test_macros.h"
+
+int main(int, char**) {
+  cuda::std::mutex mutex;
+  {
+    cuda::std::lock_guard lock(mutex);
+    ASSERT_SAME_TYPE(decltype(lock), cuda::std::lock_guard<cuda::std::mutex>);
+  }
+
+  return 0;
+}
+
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.fail.cpp
new file mode 100644
index 0000000000..1352499513
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.fail.cpp
@@ -0,0 +1,28 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class lock_guard;
+
+// explicit lock_guard(mutex_type& m);
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::mutex m;
+    cuda::std::lock_guard<cuda::std::mutex> lg = m; // expected-error{{no viable conversion}}
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.pass.cpp
new file mode 100644
index 0000000000..3611a8ca09
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.pass.cpp
@@ -0,0 +1,117 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class lock_guard;
+
+// explicit lock_guard(mutex_type& m);
+
+// template<class _Mutex> lock_guard(lock_guard<_Mutex>)
+//     -> lock_guard<_Mutex>;  // C++17
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__
+void check_ctad(Mutex* m)
+{
+#if TEST_STD_VER >= 17
+    cuda::std::lock_guard lg(*m);
+    static_assert((cuda::std::is_same<decltype(lg), cuda::std::lock_guard<Mutex>>::value), "" );
+    unused(lg);
+#endif
+}
+
+#ifndef __CUDA_ARCH__
+template<class Mutex>
+__host__ void do_try_lock(Mutex* m) {
+  assert(m->try_lock() == false);
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ void test_host() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    {
+      cuda::std::lock_guard<Mutex> lg(*m);
+      std::thread t(do_try_lock<Mutex>, m);
+      t.join();
+    }
+
+    m->lock();
+    m->unlock();
+
+    check_ctad(m);
+}
+#endif
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__device__ void test_device() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    Selector<cuda::std::atomic_flag, Initializer> sel_flag;
+    SHARED cuda::std::atomic_flag* thread_ready;
+    thread_ready = sel_flag.construct(false);
+
+    if (threadIdx.x == 0) {
+      cuda::std::lock_guard<Mutex> lg(*m);
+      thread_ready->test_and_set(cuda::std::memory_order_relaxed);
+      thread_ready->notify_one();
+      thread_ready->wait(true, cuda::std::memory_order_relaxed);
+    } else {
+      thread_ready->wait(false, cuda::std::memory_order_relaxed);
+      assert(m->try_lock() == false);
+      thread_ready->clear(cuda::std::memory_order_relaxed);
+      thread_ready->notify_one();
+    }
+    __syncthreads();
+
+    m->lock();
+    m->unlock();
+
+    check_ctad(m);
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test_host<cuda::std::mutex, local_memory_selector>();
+#else
+    test_device<cuda::std::mutex, shared_memory_selector>();
+    test_device<cuda::std::mutex, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/types.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/types.pass.cpp
new file mode 100644
index 0000000000..67204b6449
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/types.pass.cpp
@@ -0,0 +1,34 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex>
+// class lock_guard
+// {
+// public:
+//     typedef Mutex mutex_type;
+//     ...
+// };
+
+#include<cuda/std/mutex>
+#include<cuda/std/type_traits>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    static_assert((cuda::std::is_same<cuda::std::lock_guard<cuda::std::mutex>::mutex_type,
+                   cuda::std::mutex>::value), "");
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/adopt_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/adopt_lock.pass.cpp
new file mode 100644
index 0000000000..ed0110b84e
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/adopt_lock.pass.cpp
@@ -0,0 +1,75 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++03, c++11, c++14
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class ...Mutex> class scoped_lock;
+
+// scoped_lock(adopt_lock_t, Mutex&...);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include "test_macros.h"
+
+struct TestMutex {
+    bool locked = false;
+    TestMutex() = default;
+
+    __host__ __device__ void lock() { assert(!locked); locked = true; }
+    __host__ __device__ bool try_lock() { if (locked) return false; locked = true; return true; }
+    __host__ __device__ void unlock() { assert(locked); locked = false; }
+
+    TestMutex(TestMutex const&) = delete;
+    TestMutex& operator=(TestMutex const&) = delete;
+};
+
+int main(int, char**)
+{
+    {
+        using LG = cuda::std::scoped_lock<>;
+        LG lg(cuda::std::adopt_lock);
+    }
+    {
+        TestMutex m1;
+        using LG = cuda::std::scoped_lock<TestMutex>;
+        m1.lock();
+        {
+            LG lg(cuda::std::adopt_lock, m1);
+            assert(m1.locked);
+        }
+        assert(!m1.locked);
+    }
+    {
+        TestMutex m1, m2;
+        using LG = cuda::std::scoped_lock<TestMutex, TestMutex>;
+        m1.lock(); m2.lock();
+        {
+            LG lg(cuda::std::adopt_lock, m1, m2);
+            assert(m1.locked && m2.locked);
+        }
+        assert(!m1.locked && !m2.locked);
+    }
+    {
+        TestMutex m1, m2, m3;
+        using LG = cuda::std::scoped_lock<TestMutex, TestMutex, TestMutex>;
+        m1.lock(); m2.lock(); m3.lock();
+        {
+            LG lg(cuda::std::adopt_lock, m1, m2, m3);
+            assert(m1.locked && m2.locked && m3.locked);
+        }
+        assert(!m1.locked && !m2.locked && !m3.locked);
+    }
+
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/assign.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/assign.fail.cpp
new file mode 100644
index 0000000000..473d6218fc
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/assign.fail.cpp
@@ -0,0 +1,54 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++03, c++11, c++14
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class ...Mutex> class scoped_lock;
+
+// scoped_lock& operator=(scoped_lock const&) = delete;
+
+#include<cuda/std/mutex>
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    using M = cuda::std::mutex;
+    M m0, m1, m2;
+    M om0, om1, om2;
+    {
+        using LG = cuda::std::scoped_lock<>;
+        LG lg1, lg2;
+        lg1 = lg2; // expected-error{{overload resolution selected deleted operator '='}}
+    }
+    {
+        using LG = cuda::std::scoped_lock<M>;
+        LG lg1(m0);
+        LG lg2(om0);
+        lg1 = lg2; // expected-error{{overload resolution selected deleted operator '='}}
+    }
+    {
+        using LG = cuda::std::scoped_lock<M, M>;
+        LG lg1(m0, m1);
+        LG lg2(om0, om1);
+        lg1 = lg2; // expected-error{{overload resolution selected deleted operator '='}}
+    }
+    {
+        using LG = cuda::std::scoped_lock<M, M, M>;
+        LG lg1(m0, m1, m2);
+        LG lg2(om0, om1, om2);
+        lg1 = lg2; // expected-error{{overload resolution selected deleted operator '='}}
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/copy.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/copy.fail.cpp
new file mode 100644
index 0000000000..0849ffd0aa
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/copy.fail.cpp
@@ -0,0 +1,50 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++03, c++11, c++14
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class ...Mutex> class scoped_lock;
+
+// scoped_lock(scoped_lock const&) = delete;
+
+#include<cuda/std/mutex>
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    using M = cuda::std::mutex;
+    M m0, m1, m2;
+    {
+        using LG = cuda::std::scoped_lock<>;
+        const LG Orig;
+        LG Copy(Orig); // expected-error{{call to deleted constructor of 'LG'}}
+    }
+    {
+        using LG = cuda::std::scoped_lock<M>;
+        const LG Orig(m0);
+        LG Copy(Orig); // expected-error{{call to deleted constructor of 'LG'}}
+    }
+    {
+        using LG = cuda::std::scoped_lock<M, M>;
+        const LG Orig(m0, m1);
+        LG Copy(Orig); // expected-error{{call to deleted constructor of 'LG'}}
+    }
+    {
+        using LG = cuda::std::scoped_lock<M, M, M>;
+        const LG Orig(m0, m1, m2);
+        LG Copy(Orig); // expected-error{{call to deleted constructor of 'LG'}}
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/implicit_ctad.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/implicit_ctad.pass.cpp
new file mode 100644
index 0000000000..2a7b84aaae
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/implicit_ctad.pass.cpp
@@ -0,0 +1,46 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++98, c++03, c++11, c++14
+// UNSUPPORTED: gcc-7
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// scoped_lock
+
+// Make sure that the implicitly-generated CTAD works.
+
+#include<cuda/std/mutex>
+
+#include "test_macros.h"
+
+int main(int, char**) {
+  cuda::std::mutex m1;
+  {
+    cuda::std::scoped_lock lock(m1);
+    ASSERT_SAME_TYPE(decltype(lock), cuda::std::scoped_lock<cuda::std::mutex>);
+  }
+#if 0 // No recursive mutex
+  cuda::std::recursive_mutex m2;
+  cuda::std::recursive_timed_mutex m3;
+  {
+    cuda::std::scoped_lock lock(m1, m2);
+    ASSERT_SAME_TYPE(decltype(lock), cuda::std::scoped_lock<cuda::std::mutex, cuda::std::recursive_mutex>);
+  }
+  {
+    cuda::std::scoped_lock lock(m1, m2, m3);
+    ASSERT_SAME_TYPE(decltype(lock), cuda::std::scoped_lock<cuda::std::mutex, cuda::std::recursive_mutex, cuda::std::recursive_timed_mutex>);
+  }
+#endif
+
+  return 0;
+}
+
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/mutex.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/mutex.fail.cpp
new file mode 100644
index 0000000000..a353a043cf
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/mutex.fail.cpp
@@ -0,0 +1,57 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++03, c++11, c++14
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class ...Mutex> class scoped_lock;
+
+// explicit scoped_lock(Mutex&...);
+
+#include<cuda/std/mutex>
+#include "test_macros.h"
+
+template <class LG>
+__host__ __device__ void test_conversion(LG) {}
+
+int main(int, char**)
+{
+    using M = cuda::std::mutex;
+    M m0, m1, m2;
+    M n0, n1, n2;
+    {
+        using LG = cuda::std::scoped_lock<>;
+        LG lg = {}; // expected-error{{chosen constructor is explicit in copy-initialization}}
+        test_conversion<LG>({}); // expected-error{{no matching function for call}}
+        ((void)lg);
+    }
+    {
+        using LG = cuda::std::scoped_lock<M>;
+        LG lg = {m0}; // expected-error{{chosen constructor is explicit in copy-initialization}}
+        test_conversion<LG>({n0}); // expected-error{{no matching function for call}}
+        ((void)lg);
+    }
+    {
+        using LG = cuda::std::scoped_lock<M, M>;
+        LG lg = {m0, m1}; // expected-error{{chosen constructor is explicit in copy-initialization}}
+        test_conversion<LG>({n0, n1}); // expected-error{{no matching function for call}}
+        ((void)lg);
+    }
+    {
+        using LG = cuda::std::scoped_lock<M, M, M>;
+        LG lg = {m0, m1, m2}; // expected-error{{chosen constructor is explicit in copy-initialization}}
+        test_conversion<LG>({n0, n1, n2}); // expected-error{{no matching function for call}}
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/mutex.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/mutex.pass.cpp
new file mode 100644
index 0000000000..b1099634b8
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/mutex.pass.cpp
@@ -0,0 +1,163 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++03, c++11, c++14
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class ...Mutex> class scoped_lock;
+
+// explicit scoped_lock(mutex_type& m);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include "test_macros.h"
+
+struct TestMutex {
+    bool locked = false;
+    TestMutex() = default;
+    __host__ __device__ ~TestMutex() { assert(!locked); }
+
+    __host__ __device__ void lock() { assert(!locked); locked = true; }
+    __host__ __device__ bool try_lock() { if (locked) return false; locked = true; return true; }
+    __host__ __device__ void unlock() { assert(locked); locked = false; }
+
+    TestMutex(TestMutex const&) = delete;
+    TestMutex& operator=(TestMutex const&) = delete;
+};
+
+#if !defined(TEST_HAS_NO_EXCEPTIONS)
+struct TestMutexThrows {
+    bool locked = false;
+    bool throws_on_lock = false;
+
+    __host__ __device__ TestMutexThrows() = default;
+    __host__ __device__ ~TestMutexThrows() { assert(!locked); }
+
+    __host__ __device__ void lock() {
+        assert(!locked);
+        if (throws_on_lock) {
+            throw 42;
+        }
+        locked = true;
+    }
+
+    __host__ __device__ bool try_lock() {
+        if (locked) return false;
+        lock();
+        return true;
+    }
+
+    __host__ __device__ void unlock() { assert(locked); locked = false; }
+
+    TestMutexThrows(TestMutexThrows const&) = delete;
+    TestMutexThrows& operator=(TestMutexThrows const&) = delete;
+};
+#endif // !defined(TEST_HAS_NO_EXCEPTIONS)
+
+int main(int, char**)
+{
+    {
+        using LG = cuda::std::scoped_lock<>;
+        LG lg;
+        unused(lg);
+    }
+    {
+        using LG = cuda::std::scoped_lock<TestMutex>;
+        TestMutex m1;
+        {
+            LG lg(m1);
+            assert(m1.locked);
+        }
+        assert(!m1.locked);
+    }
+    {
+        using LG = cuda::std::scoped_lock<TestMutex, TestMutex>;
+        TestMutex m1, m2;
+        {
+            LG lg(m1, m2);
+            assert(m1.locked && m2.locked);
+        }
+        assert(!m1.locked && !m2.locked);
+    }
+    {
+        using LG = cuda::std::scoped_lock<TestMutex, TestMutex, TestMutex>;
+        TestMutex m1, m2, m3;
+        {
+            LG lg(m1, m2, m3);
+            assert(m1.locked && m2.locked && m3.locked);
+        }
+        assert(!m1.locked && !m2.locked && !m3.locked);
+    }
+#if !defined(TEST_HAS_NO_EXCEPTIONS)
+    {
+        using MT = TestMutexThrows;
+        using LG = cuda::std::scoped_lock<MT>;
+        MT m1;
+        m1.throws_on_lock = true;
+        try {
+            LG lg(m1);
+            assert(false);
+        } catch (int) {}
+        assert(!m1.locked);
+    }
+    {
+        using MT = TestMutexThrows;
+        using LG = cuda::std::scoped_lock<MT, MT>;
+        MT m1, m2;
+        m1.throws_on_lock = true;
+        try {
+            LG lg(m1, m2);
+            assert(false);
+        } catch (int) {}
+        assert(!m1.locked && !m2.locked);
+    }
+    {
+        using MT = TestMutexThrows;
+        using LG = cuda::std::scoped_lock<MT, MT, MT>;
+        MT m1, m2, m3;
+        m2.throws_on_lock = true;
+        try {
+            LG lg(m1, m2, m3);
+            assert(false);
+        } catch (int) {}
+        assert(!m1.locked && !m2.locked && !m3.locked);
+    }
+#endif
+
+#if TEST_STD_VER >= 17
+    {
+    TestMutex m1, m2, m3;
+        {
+        cuda::std::scoped_lock sl{};
+        static_assert((cuda::std::is_same<decltype(sl), cuda::std::scoped_lock<>>::value), "" );
+        unused(sl);
+        }
+        {
+        cuda::std::scoped_lock sl{m1};
+        static_assert((cuda::std::is_same<decltype(sl), cuda::std::scoped_lock<decltype(m1)>>::value), "" );
+        unused(sl);
+        }
+        {
+        cuda::std::scoped_lock sl{m1, m2};
+        static_assert((cuda::std::is_same<decltype(sl), cuda::std::scoped_lock<decltype(m1), decltype(m2)>>::value), "" );
+        unused(sl);
+        }
+        {
+        cuda::std::scoped_lock sl{m1, m2, m3};
+        static_assert((cuda::std::is_same<decltype(sl), cuda::std::scoped_lock<decltype(m1), decltype(m2), decltype(m3)>>::value), "" );
+        unused(sl);
+        }
+    }
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/types.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/types.pass.cpp
new file mode 100644
index 0000000000..3019295978
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.scoped/types.pass.cpp
@@ -0,0 +1,85 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++03, c++11, c++14
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class... MutexTypes>
+// class scoped_lock
+// {
+// public:
+//     typedef Mutex mutex_type;  // Only if sizeof...(MutexTypes) == 1
+//     ...
+// };
+
+#include<cuda/std/mutex>
+#include<cuda/std/type_traits>
+#include "test_macros.h"
+
+struct NAT {};
+
+template <class LG>
+__host__ __device__ auto test_typedef(int) -> typename LG::mutex_type;
+
+template <class LG>
+__host__ __device__ auto test_typedef(...) -> NAT;
+
+template <class LG>
+__host__ __device__ constexpr bool has_mutex_type() {
+    return !cuda::std::is_same<decltype(test_typedef<LG>(0)), NAT>::value;
+}
+
+int main(int, char**)
+{
+    {
+        using T = cuda::std::scoped_lock<>;
+        static_assert(!has_mutex_type<T>(), "");
+    }
+    {
+        using M1 = cuda::std::mutex;
+        using T = cuda::std::scoped_lock<M1>;
+        static_assert(cuda::std::is_same<T::mutex_type, M1>::value, "");
+    }
+#if 0 // No recursive mutex
+    {
+        using M1 = cuda::std::recursive_mutex;
+        using T = cuda::std::scoped_lock<M1>;
+        static_assert(cuda::std::is_same<T::mutex_type, M1>::value, "");
+    }
+    {
+        using M1 = cuda::std::mutex;
+        using M2 = cuda::std::recursive_mutex;
+        using T = cuda::std::scoped_lock<M1, M2>;
+        static_assert(!has_mutex_type<T>(), "");
+    }
+    {
+        using M1 = cuda::std::mutex;
+        using M2 = cuda::std::recursive_mutex;
+        using T = cuda::std::scoped_lock<M1, M1, M2>;
+        static_assert(!has_mutex_type<T>(), "");
+    }
+#endif
+    {
+        using M1 = cuda::std::mutex;
+        using T = cuda::std::scoped_lock<M1, M1>;
+        static_assert(!has_mutex_type<T>(), "");
+    }
+#if 0 // No recursive mutex
+    {
+        using M1 = cuda::std::recursive_mutex;
+        using T = cuda::std::scoped_lock<M1, M1, M1>;
+        static_assert(!has_mutex_type<T>(), "");
+    }
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/implicit_ctad.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/implicit_ctad.pass.cpp
new file mode 100644
index 0000000000..a65275b87b
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/implicit_ctad.pass.cpp
@@ -0,0 +1,33 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: c++98, c++03, c++11, c++14
+// UNSUPPORTED: gcc-7
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// unique_lock
+
+// Make sure that the implicitly-generated CTAD works.
+
+#include<cuda/std/mutex>
+
+#include "test_macros.h"
+
+int main(int, char**) {
+  cuda::std::mutex mutex;
+  {
+    cuda::std::unique_lock lock(mutex);
+    ASSERT_SAME_TYPE(decltype(lock), cuda::std::unique_lock<cuda::std::mutex>);
+  }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/copy_assign.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/copy_assign.compile.fail.cpp
new file mode 100644
index 0000000000..b89063c263
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/copy_assign.compile.fail.cpp
@@ -0,0 +1,38 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// unique_lock& operator=(unique_lock const&) = delete;
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+int main(int, char**)
+{
+    {
+    typedef cuda::std::mutex M;
+    M m0;
+    M m1;
+    cuda::std::unique_lock<M> lk0(m0);
+    cuda::std::unique_lock<M> lk1(m1);
+    lk1 = lk0;
+    assert(lk1.mutex() == &m0);
+    assert(lk1.owns_lock() == true);
+    assert(lk0.mutex() == nullptr);
+    assert(lk0.owns_lock() == false);
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/copy_ctor.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/copy_ctor.compile.fail.cpp
new file mode 100644
index 0000000000..e77179514d
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/copy_ctor.compile.fail.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// unique_lock(unique_lock const&) = delete;
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+int main(int, char**)
+{
+    {
+    typedef cuda::std::mutex M;
+    M m;
+    cuda::std::unique_lock<M> lk0(m);
+    cuda::std::unique_lock<M> lk = lk0;
+    assert(lk.mutex() == &m);
+    assert(lk.owns_lock() == true);
+    assert(lk0.mutex() == nullptr);
+    assert(lk0.owns_lock() == false);
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/default.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/default.pass.cpp
new file mode 100644
index 0000000000..c673cdddd9
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/default.pass.cpp
@@ -0,0 +1,31 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// unique_lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<cuda::std::mutex> ul;
+    assert(!ul.owns_lock());
+    assert(ul.mutex() == nullptr);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/move_assign.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/move_assign.pass.cpp
new file mode 100644
index 0000000000..8ca41a304b
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/move_assign.pass.cpp
@@ -0,0 +1,53 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads, c++03
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// unique_lock& operator=(unique_lock&& u);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include "nasty_containers.h"
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    {
+    typedef cuda::std::mutex M;
+    M m0;
+    M m1;
+    cuda::std::unique_lock<M> lk0(m0);
+    cuda::std::unique_lock<M> lk1(m1);
+    lk1 = cuda::std::move(lk0);
+    assert(lk1.mutex() == cuda::std::addressof(m0));
+    assert(lk1.owns_lock() == true);
+    assert(lk0.mutex() == nullptr);
+    assert(lk0.owns_lock() == false);
+    }
+    {
+    typedef nasty_mutex M;
+    M m0;
+    M m1;
+    cuda::std::unique_lock<M> lk0(m0);
+    cuda::std::unique_lock<M> lk1(m1);
+    lk1 = cuda::std::move(lk0);
+    assert(lk1.mutex() == cuda::std::addressof(m0));
+    assert(lk1.owns_lock() == true);
+    assert(lk0.mutex() == nullptr);
+    assert(lk0.owns_lock() == false);
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/move_ctor.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/move_ctor.pass.cpp
new file mode 100644
index 0000000000..23b5235a36
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/move_ctor.pass.cpp
@@ -0,0 +1,49 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads, c++03
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// unique_lock(unique_lock&& u);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include "nasty_containers.h"
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    {
+    typedef cuda::std::mutex M;
+    M m;
+    cuda::std::unique_lock<M> lk0(m);
+    cuda::std::unique_lock<M> lk = cuda::std::move(lk0);
+    assert(lk.mutex() == cuda::std::addressof(m));
+    assert(lk.owns_lock() == true);
+    assert(lk0.mutex() == nullptr);
+    assert(lk0.owns_lock() == false);
+    }
+    {
+    typedef nasty_mutex M;
+    M m;
+    cuda::std::unique_lock<M> lk0(m);
+    cuda::std::unique_lock<M> lk = cuda::std::move(lk0);
+    assert(lk.mutex() == cuda::std::addressof(m));
+    assert(lk.owns_lock() == true);
+    assert(lk0.mutex() == nullptr);
+    assert(lk0.owns_lock() == false);
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex.pass.cpp
new file mode 100644
index 0000000000..979beda8c5
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex.pass.cpp
@@ -0,0 +1,94 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// explicit unique_lock(mutex_type& m);
+
+// template<class _Mutex> unique_lock(unique_lock<_Mutex>)
+//     -> unique_lock<_Mutex>;  // C++17
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    time_point t1;
+    {
+        cuda::std::unique_lock<Mutex> ul(*m);
+        t1 = Clock::now();
+    }
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex>
+__host__ __device__
+void check_ctad(Mutex* m)
+{
+#if TEST_STD_VER >= 17
+    cuda::std::unique_lock ul(*m);
+    static_assert((cuda::std::is_same<decltype(ul), cuda::std::unique_lock<Mutex>>::value), "" );
+    unused(ul);
+#endif
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+
+    handler.syncthreads();
+    handler.run_on_first_thread(check_ctad<Mutex>, m);
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::mutex, local_memory_selector>();
+#else
+    test<cuda::std::mutex, shared_memory_selector>();
+    test<cuda::std::mutex, global_memory_selector>();
+#endif
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_adopt_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_adopt_lock.pass.cpp
new file mode 100644
index 0000000000..3442878be3
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_adopt_lock.pass.cpp
@@ -0,0 +1,48 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// unique_lock(mutex_type& m, adopt_lock_t);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include "nasty_containers.h"
+
+#include "test_macros.h"
+
+__host__ __device__ void test() {
+    {
+    typedef cuda::std::mutex M;
+    M m;
+    m.lock();
+    cuda::std::unique_lock<M> lk(m, cuda::std::adopt_lock);
+    assert(lk.mutex() == cuda::std::addressof(m));
+    assert(lk.owns_lock() == true);
+    }
+    {
+    typedef nasty_mutex M;
+    M m;
+    m.lock();
+    cuda::std::unique_lock<M> lk(m, cuda::std::adopt_lock);
+    assert(lk.mutex() == cuda::std::addressof(m));
+    assert(lk.owns_lock() == true);
+    }
+}
+
+int main(int, char**)
+{
+  test();
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_defer_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_defer_lock.pass.cpp
new file mode 100644
index 0000000000..cfa8538705
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_defer_lock.pass.cpp
@@ -0,0 +1,43 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// unique_lock(mutex_type& m, defer_lock_t);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include "nasty_containers.h"
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    {
+    typedef cuda::std::mutex M;
+    M m;
+    cuda::std::unique_lock<M> lk(m, cuda::std::defer_lock);
+    assert(lk.mutex() == cuda::std::addressof(m));
+    assert(lk.owns_lock() == false);
+    }
+    {
+    typedef nasty_mutex M;
+    M m;
+    cuda::std::unique_lock<M> lk(m, cuda::std::defer_lock);
+    assert(lk.mutex() == cuda::std::addressof(m));
+    assert(lk.owns_lock() == false);
+    }
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_duration.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_duration.pass.cpp
new file mode 100644
index 0000000000..d32e38e30c
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_duration.pass.cpp
@@ -0,0 +1,102 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// template <class Rep, class Period>
+//   unique_lock(mutex_type& m, const chrono::duration<Rep, Period>& rel_time);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f1(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    cuda::std::unique_lock<Mutex> lk(*m, ms(500));
+    assert(lk.owns_lock() == true);
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex>
+__host__ __device__ void f2(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    cuda::std::unique_lock<Mutex> lk(*m, ms(200));
+    assert(lk.owns_lock() == false);
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(200));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    {
+      handler.run_on_first_thread(&Mutex::lock, m);
+      handler.syncthreads();
+
+      handler.run_on_second_thread(f1<Mutex>, m);
+      handler.sleep_first_thread(ms(250));
+      handler.run_on_first_thread(&Mutex::unlock, m);
+      handler.join_test_thread();
+    }
+
+    handler.syncthreads();
+
+    {
+      handler.run_on_first_thread(&Mutex::lock, m);
+      handler.syncthreads();
+
+      handler.run_on_second_thread(f2<Mutex>, m);
+      handler.sleep_first_thread(ms(250));
+      handler.run_on_first_thread(&Mutex::unlock, m);
+      handler.join_test_thread();
+    }
+
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::timed_mutex, local_memory_selector>();
+#else
+    test<cuda::std::timed_mutex, shared_memory_selector>();
+    test<cuda::std::timed_mutex, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_time_point.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_time_point.pass.cpp
new file mode 100644
index 0000000000..e751821f19
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_time_point.pass.cpp
@@ -0,0 +1,100 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// template <class Clock, class Duration>
+//   unique_lock(mutex_type& m, const chrono::time_point<Clock, Duration>& abs_time);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::high_resolution_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f1(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    cuda::std::unique_lock<Mutex> lk(*m, Clock::now() + ms(500));
+    assert(lk.owns_lock() == true);
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex>
+__host__ __device__ void f2(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    cuda::std::unique_lock<Mutex> lk(*m, Clock::now() + ms(200));
+    assert(lk.owns_lock() == false);
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(200));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    {
+      handler.run_on_first_thread(&Mutex::lock, m);
+      handler.syncthreads();
+
+      handler.run_on_second_thread(f1<Mutex>, m);
+      handler.sleep_first_thread(ms(250));
+      handler.run_on_first_thread(&Mutex::unlock, m);
+      handler.join_test_thread();
+    }
+
+    handler.syncthreads();
+
+    {
+      handler.run_on_first_thread(&Mutex::lock, m);
+      handler.syncthreads();
+
+      handler.run_on_second_thread(f2<Mutex>, m);
+      handler.sleep_first_thread(ms(250));
+      handler.run_on_first_thread(&Mutex::unlock, m);
+      handler.join_test_thread();
+    }
+}
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::timed_mutex, local_memory_selector>();
+#else
+    test<cuda::std::timed_mutex, shared_memory_selector>();
+    test<cuda::std::timed_mutex, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_try_to_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_try_to_lock.pass.cpp
new file mode 100644
index 0000000000..64e9941ade
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.cons/mutex_try_to_lock.pass.cpp
@@ -0,0 +1,92 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// unique_lock(mutex_type& m, try_to_lock_t);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    {
+        cuda::std::unique_lock<Mutex> lk(*m, cuda::std::try_to_lock);
+        assert(lk.owns_lock() == false);
+    }
+    {
+        cuda::std::unique_lock<Mutex> lk(*m, cuda::std::try_to_lock);
+        assert(lk.owns_lock() == false);
+    }
+    {
+        cuda::std::unique_lock<Mutex> lk(*m, cuda::std::try_to_lock);
+        assert(lk.owns_lock() == false);
+    }
+    while (true)
+    {
+        cuda::std::unique_lock<Mutex> lk(*m, cuda::std::try_to_lock);
+        if (lk.owns_lock())
+            break;
+    }
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+
+}
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::mutex, local_memory_selector>();
+#else
+    test<cuda::std::mutex, shared_memory_selector>();
+    test<cuda::std::mutex, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/lock.pass.cpp
new file mode 100644
index 0000000000..36689f1abd
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/lock.pass.cpp
@@ -0,0 +1,101 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// void lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    cuda::std::unique_lock<Mutex> lk(*m, cuda::std::defer_lock);
+    time_point t0 = Clock::now();
+    lk.lock();
+    time_point t1 = Clock::now();
+    assert(lk.owns_lock() == true);
+    assert(t1 - t0 >= ms(250));
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        lk.lock();
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EDEADLK);
+    }
+#endif
+    lk.unlock();
+    lk.release();
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        lk.lock();
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EPERM);
+    }
+#endif
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::mutex, local_memory_selector>();
+#else
+    test<cuda::std::mutex, shared_memory_selector>();
+    test<cuda::std::mutex, global_memory_selector>();
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock.pass.cpp
new file mode 100644
index 0000000000..dc9a883eb0
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock.pass.cpp
@@ -0,0 +1,75 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+//
+// FLAKY_TEST.
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// bool try_lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR bool try_lock_called = false;
+
+struct mutex
+{
+    __host__ __device__ bool try_lock()
+    {
+        try_lock_called = !try_lock_called;
+        return try_lock_called;
+    }
+    __host__ __device__ void unlock() {}
+};
+
+STATIC_TEST_GLOBAL_VAR mutex m;
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<mutex> lk(m, cuda::std::defer_lock);
+    assert(lk.try_lock() == true);
+    assert(try_lock_called == true);
+    assert(lk.owns_lock() == true);
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        TEST_IGNORE_NODISCARD lk.try_lock();
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EDEADLK);
+    }
+#endif
+    lk.unlock();
+    assert(lk.try_lock() == false);
+    assert(try_lock_called == false);
+    assert(lk.owns_lock() == false);
+    lk.release();
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        TEST_IGNORE_NODISCARD lk.try_lock();
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EPERM);
+    }
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock_for.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock_for.pass.cpp
new file mode 100644
index 0000000000..56dbbce8d2
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock_for.pass.cpp
@@ -0,0 +1,78 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// template <class Rep, class Period>
+//   bool try_lock_for(const chrono::duration<Rep, Period>& rel_time);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR bool try_lock_for_called = false;
+
+typedef cuda::std::chrono::milliseconds ms;
+
+struct mutex
+{
+    template <class Rep, class Period>
+    __host__ __device__ bool try_lock_for(const cuda::std::chrono::duration<Rep, Period>& rel_time)
+    {
+        assert(rel_time == ms(5));
+        try_lock_for_called = !try_lock_for_called;
+        return try_lock_for_called;
+    }
+    __host__ __device__ void unlock() {}
+};
+
+STATIC_TEST_GLOBAL_VAR mutex m;
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<mutex> lk(m, cuda::std::defer_lock);
+    assert(lk.try_lock_for(ms(5)) == true);
+    assert(try_lock_for_called == true);
+    assert(lk.owns_lock() == true);
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        TEST_IGNORE_NODISCARD lk.try_lock_for(ms(5));
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EDEADLK);
+    }
+#endif
+    lk.unlock();
+    assert(lk.try_lock_for(ms(5)) == false);
+    assert(try_lock_for_called == false);
+    assert(lk.owns_lock() == false);
+    lk.release();
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        TEST_IGNORE_NODISCARD lk.try_lock_for(ms(5));
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EPERM);
+    }
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock_until.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock_until.pass.cpp
new file mode 100644
index 0000000000..90130e3de1
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/try_lock_until.pass.cpp
@@ -0,0 +1,83 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// template <class Clock, class Duration>
+//   bool try_lock_until(const chrono::time_point<Clock, Duration>& abs_time);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR bool try_lock_until_called = false;
+
+struct mutex
+{
+    template <class Clock, class Duration>
+    __host__ __device__ bool try_lock_until(const cuda::std::chrono::time_point<Clock, Duration>& abs_time)
+    {
+        typedef cuda::std::chrono::milliseconds ms;
+        assert(Clock::now() - abs_time < ms(5));
+        try_lock_until_called = !try_lock_until_called;
+        return try_lock_until_called;
+    }
+    __host__ __device__ void unlock() {}
+};
+
+STATIC_TEST_GLOBAL_VAR mutex m;
+
+#if defined(_LIBCUDACXX_HAS_NO_MONOTONIC_CLOCK)
+typedef cuda::std::chrono::system_clock Clock;
+#else
+typedef cuda::std::chrono::steady_clock Clock;
+#endif
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<mutex> lk(m, cuda::std::defer_lock);
+    assert(lk.try_lock_until(Clock::now()) == true);
+    assert(try_lock_until_called == true);
+    assert(lk.owns_lock() == true);
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        TEST_IGNORE_NODISCARD lk.try_lock_until(Clock::now());
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EDEADLK);
+    }
+#endif
+    lk.unlock();
+    assert(lk.try_lock_until(Clock::now()) == false);
+    assert(try_lock_until_called == false);
+    assert(lk.owns_lock() == false);
+    lk.release();
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        TEST_IGNORE_NODISCARD lk.try_lock_until(Clock::now());
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EPERM);
+    }
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/unlock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/unlock.pass.cpp
new file mode 100644
index 0000000000..ee8704ff1a
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.locking/unlock.pass.cpp
@@ -0,0 +1,65 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// void unlock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR bool unlock_called = false;
+
+struct mutex
+{
+    __host__ __device__ void lock() {}
+    __host__ __device__ void unlock() {unlock_called = true;}
+};
+
+STATIC_TEST_GLOBAL_VAR mutex m;
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<mutex> lk(m);
+    lk.unlock();
+    assert(unlock_called == true);
+    assert(lk.owns_lock() == false);
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        lk.unlock();
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EPERM);
+    }
+#endif
+    lk.release();
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        lk.unlock();
+        assert(false);
+    }
+    catch (cuda::std::system_error& e)
+    {
+        assert(e.code().value() == EPERM);
+    }
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/member_swap.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/member_swap.pass.cpp
new file mode 100644
index 0000000000..ab9402db29
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/member_swap.pass.cpp
@@ -0,0 +1,43 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// void swap(unique_lock& u);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+struct mutex
+{
+    __host__ __device__ void lock() {}
+    __host__ __device__ void unlock() {}
+};
+
+STATIC_TEST_GLOBAL_VAR mutex m;
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<mutex> lk1(m);
+    cuda::std::unique_lock<mutex> lk2;
+    lk1.swap(lk2);
+    assert(lk1.mutex() == nullptr);
+    assert(lk1.owns_lock() == false);
+    assert(lk2.mutex() == &m);
+    assert(lk2.owns_lock() == true);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/nonmember_swap.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/nonmember_swap.pass.cpp
new file mode 100644
index 0000000000..7be3f15fee
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/nonmember_swap.pass.cpp
@@ -0,0 +1,44 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// template <class Mutex>
+//   void swap(unique_lock<Mutex>& x, unique_lock<Mutex>& y);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+struct mutex
+{
+    __host__ __device__ void lock() {}
+    __host__ __device__ void unlock() {}
+};
+
+STATIC_TEST_GLOBAL_VAR mutex m;
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<mutex> lk1(m);
+    cuda::std::unique_lock<mutex> lk2;
+    swap(lk1, lk2);
+    assert(lk1.mutex() == nullptr);
+    assert(lk1.owns_lock() == false);
+    assert(lk2.mutex() == &m);
+    assert(lk2.owns_lock() == true);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/release.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/release.pass.cpp
new file mode 100644
index 0000000000..88444921f1
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.mod/release.pass.cpp
@@ -0,0 +1,48 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// mutex_type* release() noexcept;
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR int lock_count = 0;
+STATIC_TEST_GLOBAL_VAR int unlock_count = 0;
+struct mutex
+{
+    __host__ __device__ void lock() {++lock_count;}
+    __host__ __device__ void unlock() {++unlock_count;}
+};
+
+STATIC_TEST_GLOBAL_VAR mutex m;
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<mutex> lk(m);
+    assert(lk.mutex() == &m);
+    assert(lk.owns_lock() == true);
+    assert(lock_count == 1);
+    assert(unlock_count == 0);
+    assert(lk.release() == &m);
+    assert(lk.mutex() == nullptr);
+    assert(lk.owns_lock() == false);
+    assert(lock_count == 1);
+    assert(unlock_count == 0);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/mutex.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/mutex.pass.cpp
new file mode 100644
index 0000000000..9dae3f3c20
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/mutex.pass.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// mutex_type *mutex() const;
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::mutex m;
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<cuda::std::mutex> lk0;
+    assert(lk0.mutex() == nullptr);
+    cuda::std::unique_lock<cuda::std::mutex> lk1(m);
+    assert(lk1.mutex() == &m);
+    lk1.unlock();
+    assert(lk1.mutex() == &m);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/op_bool.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/op_bool.pass.cpp
new file mode 100644
index 0000000000..7c8b7b8cf3
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/op_bool.pass.cpp
@@ -0,0 +1,41 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// explicit operator bool() const noexcept;
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/type_traits>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::mutex m;
+
+int main(int, char**)
+{
+    static_assert(cuda::std::is_constructible<bool, cuda::std::unique_lock<cuda::std::mutex> >::value, "");
+    static_assert(!cuda::std::is_convertible<cuda::std::unique_lock<cuda::std::mutex>, bool>::value, "");
+
+    cuda::std::unique_lock<cuda::std::mutex> lk0;
+    assert(static_cast<bool>(lk0) == false);
+    cuda::std::unique_lock<cuda::std::mutex> lk1(m);
+    assert(static_cast<bool>(lk1) == true);
+    lk1.unlock();
+    assert(static_cast<bool>(lk1) == false);
+    ASSERT_NOEXCEPT(static_cast<bool>(lk0));
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/owns_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/owns_lock.pass.cpp
new file mode 100644
index 0000000000..a69487e3a5
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/thread.lock.unique.obs/owns_lock.pass.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex> class unique_lock;
+
+// bool owns_lock() const;
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::mutex m;
+
+int main(int, char**)
+{
+    cuda::std::unique_lock<cuda::std::mutex> lk0;
+    assert(lk0.owns_lock() == false);
+    cuda::std::unique_lock<cuda::std::mutex> lk1(m);
+    assert(lk1.owns_lock() == true);
+    lk1.unlock();
+    assert(lk1.owns_lock() == false);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/types.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/types.pass.cpp
new file mode 100644
index 0000000000..be9f262dff
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.unique/types.pass.cpp
@@ -0,0 +1,34 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// template <class Mutex>
+// class unique_lock
+// {
+// public:
+//     typedef Mutex mutex_type;
+//     ...
+// };
+
+#include<cuda/std/mutex>
+#include<cuda/std/type_traits>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    static_assert((cuda::std::is_same<cuda::std::unique_lock<cuda::std::mutex>::mutex_type,
+                   cuda::std::mutex>::value), "");
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/types.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/types.fail.cpp
new file mode 100644
index 0000000000..31b9170cdb
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/types.fail.cpp
@@ -0,0 +1,32 @@
+//===----------------------------------------------------------------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is dual licensed under the MIT and the University of Illinois Open
+// Source Licenses. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct defer_lock_t { explicit defer_lock_t() = default; };
+// struct try_to_lock_t { explicit try_to_lock_t() = default; };
+// struct adopt_lock_t { explicit adopt_lock_t() = default; };
+
+// This test checks for LWG 2510.
+
+#include<cuda/std/mutex>
+
+
+cuda::std::defer_lock_t __host__ __device__ f1() { return {}; } // expected-error 1 {{chosen constructor is explicit in copy-initialization}}
+cuda::std::try_to_lock_t __host__ __device__ f2() { return {}; } // expected-error 1 {{chosen constructor is explicit in copy-initialization}}
+cuda::std::adopt_lock_t __host__ __device__ f3() { return {}; } // expected-error 1 {{chosen constructor is explicit in copy-initialization}}
+
+int main(int, char**) {
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.lock/types.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/types.pass.cpp
new file mode 100644
index 0000000000..8176711a2e
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/types.pass.cpp
@@ -0,0 +1,38 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct defer_lock_t { explicit defer_lock_t() = default; };
+// struct try_to_lock_t { explicit try_to_lock_t() = default; };
+// struct adopt_lock_t { explicit adopt_lock_t() = default; };
+//
+// constexpr defer_lock_t  defer_lock{};
+// constexpr try_to_lock_t try_to_lock{};
+// constexpr adopt_lock_t  adopt_lock{};
+
+#include<cuda/std/mutex>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    typedef cuda::std::defer_lock_t T1;
+    typedef cuda::std::try_to_lock_t T2;
+    typedef cuda::std::adopt_lock_t T3;
+
+    T1 t1 = cuda::std::defer_lock; unused(t1);
+    T2 t2 = cuda::std::try_to_lock; unused(t2);
+    T3 t3 = cuda::std::adopt_lock; unused(t3);
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.general/nothing_to_do.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.general/nothing_to_do.pass.cpp
new file mode 100644
index 0000000000..8224a0f7bf
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.general/nothing_to_do.pass.cpp
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+int main(int, char**)
+{
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/assign.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/assign.compile.fail.cpp
new file mode 100644
index 0000000000..cde71d1d40
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/assign.compile.fail.cpp
@@ -0,0 +1,28 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class mutex;
+
+// mutex& operator=(const mutex&) = delete;
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::mutex m0;
+    cuda::std::mutex m1;
+    m1 = m0;
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/copy.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/copy.compile.fail.cpp
new file mode 100644
index 0000000000..a673ef99d5
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/copy.compile.fail.cpp
@@ -0,0 +1,27 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class mutex;
+
+// mutex(const mutex&) = delete;
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::mutex m0;
+    cuda::std::mutex m1(m0);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/default.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/default.pass.cpp
new file mode 100644
index 0000000000..c54fed1c12
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/default.pass.cpp
@@ -0,0 +1,30 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class mutex;
+
+// mutex();
+
+#include<cuda/std/mutex>
+#include<cuda/std/type_traits>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+  static_assert(cuda::std::is_nothrow_default_constructible<cuda::std::mutex>::value, "");
+  cuda::std::mutex m;
+  unused(m);
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp
new file mode 100644
index 0000000000..88a9e3849d
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp
@@ -0,0 +1,75 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class mutex;
+
+// void lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    m->lock();
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::mutex, local_memory_selector>();
+#else
+    test<cuda::std::mutex, shared_memory_selector>();
+    test<cuda::std::mutex, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp
new file mode 100644
index 0000000000..02bb7f64c6
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp
@@ -0,0 +1,79 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class mutex;
+
+// bool try_lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    while(!m->try_lock())
+        ;
+    m->unlock();
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::mutex, local_memory_selector>();
+#else
+    test<cuda::std::mutex, shared_memory_selector>();
+    test<cuda::std::mutex, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/assign.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/assign.compile.fail.cpp
new file mode 100644
index 0000000000..1050b0820f
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/assign.compile.fail.cpp
@@ -0,0 +1,28 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class timed_mutex;
+
+// timed_mutex& operator=(const timed_mutex&) = delete;
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::timed_mutex m0;
+    cuda::std::timed_mutex m1;
+    m1 = m0;
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/copy.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/copy.compile.fail.cpp
new file mode 100644
index 0000000000..5923ccbd05
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/copy.compile.fail.cpp
@@ -0,0 +1,27 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class timed_mutex;
+
+// timed_mutex(const timed_mutex&) = delete;
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::timed_mutex m0;
+    cuda::std::timed_mutex m1(m0);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/default.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/default.pass.cpp
new file mode 100644
index 0000000000..3c421ede81
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/default.pass.cpp
@@ -0,0 +1,29 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// class timed_mutex;
+
+// timed_mutex();
+
+#include<cuda/std/mutex>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    cuda::std::timed_mutex m;
+    unused(m);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/lock.pass.cpp
new file mode 100644
index 0000000000..ca63da3f7a
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/lock.pass.cpp
@@ -0,0 +1,75 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// void lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    m->lock();
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::timed_mutex, local_memory_selector>();
+#else
+    test<cuda::std::timed_mutex, shared_memory_selector>();
+    test<cuda::std::timed_mutex, global_memory_selector>();
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock.pass.cpp
new file mode 100644
index 0000000000..5a927e3a47
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock.pass.cpp
@@ -0,0 +1,80 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// bool try_lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    while(!m->try_lock())
+        ;
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::timed_mutex, local_memory_selector>();
+#else
+    test<cuda::std::timed_mutex, shared_memory_selector>();
+    test<cuda::std::timed_mutex, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_for.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_for.pass.cpp
new file mode 100644
index 0000000000..605d4856a5
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_for.pass.cpp
@@ -0,0 +1,99 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// template <class Rep, class Period>
+//     bool try_lock_for(const chrono::duration<Rep, Period>& rel_time);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f1(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(m->try_lock_for(ms(500)) == true);
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex>
+__host__ __device__ void f2(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(m->try_lock_for(ms(200)) == false);
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(200));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    {
+        handler.run_on_first_thread(&Mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1<Mutex>, m);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&Mutex::unlock, m);
+        handler.join_test_thread();
+    }
+
+    handler.syncthreads();
+
+    {
+        handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f2<Mutex>, m);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+        handler.join_test_thread();
+    }
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::timed_mutex, local_memory_selector>();
+#else
+    test<cuda::std::timed_mutex, shared_memory_selector>();
+    test<cuda::std::timed_mutex, global_memory_selector>();
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_until.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_until.pass.cpp
new file mode 100644
index 0000000000..2d3f9dedc5
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.mutex.requirements/thread.timedmutex.requirements/thread.timedmutex.class/try_lock_until.pass.cpp
@@ -0,0 +1,99 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// template <class Clock, class Duration>
+//     bool try_lock_until(const chrono::time_point<Clock, Duration>& abs_time);
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f1(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(m->try_lock_until(Clock::now() + ms(500)) == true);
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex>
+__host__ __device__ void f2(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(m->try_lock_until(Clock::now() + ms(200)) == false);
+    time_point t1 = Clock::now();
+    assert(t1 - t0 >= ms(200));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    {
+        handler.run_on_first_thread(&Mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1<Mutex>, m);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&Mutex::unlock, m);
+        handler.join_test_thread();
+    }
+
+    handler.syncthreads();
+
+    {
+        handler.run_on_first_thread(&Mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f2<Mutex>, m);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&Mutex::unlock, m);
+        handler.join_test_thread();
+    }
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::timed_mutex, local_memory_selector>();
+#else
+    test<cuda::std::timed_mutex, shared_memory_selector>();
+    test<cuda::std::timed_mutex, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp
new file mode 100644
index 0000000000..9a429ec161
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp
@@ -0,0 +1,303 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// template<class Callable, class ...Args>
+//   void call_once(once_flag& flag, Callable&& func, Args&&... args);
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::milliseconds ms;
+
+__host__ __device__
+void init0(int& init0_called)
+{
+    test_sleep_thread(ms(250));
+    ++init0_called;
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f0(OnceFlag* flg0, int* init0_called)
+{
+    cuda::std::call_once(*flg0, init0, *init0_called);
+}
+
+#ifndef TEST_HAS_NO_EXCEPTIONS
+__host__ __device__
+void init3(int& init3_called, int& init3_completed)
+{
+    ++init3_called;
+    test_sleep_thread(ms(250));
+    if (init3_called == 1)
+#ifdef __CUDA_ARCH__
+        _LIBCUDACXX_UNREACHABLE();
+#else
+        TEST_THROW(1);
+#endif
+    ++init3_completed;
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f3(OnceFlag* flg3, int* init3_called, int* init3_completed)
+{
+    try
+    {
+        cuda::std::call_once(*flg3, init3, *init3_called, *init3_completed);
+    }
+    catch (...)
+    {
+    }
+}
+#endif // TEST_HAS_NO_EXCEPTIONS
+
+#if TEST_STD_VER >= 11
+struct init1
+{
+    int& init1_called;
+    __host__ __device__ void operator()(int i) { init1_called += i; }
+};
+
+template<class OnceFlag>
+__host__ __device__
+void f1(OnceFlag* flg1, int* init1_called)
+{
+    cuda::std::call_once(*flg1, init1{*init1_called}, 1);
+}
+
+struct init2
+{
+    int& init2_called;
+    __host__ __device__ void operator()(int i, int j) const {init2_called += i + j;}
+};
+
+template<class OnceFlag>
+__host__ __device__
+void f2(OnceFlag* flg2, int* init2_called)
+{
+    cuda::std::call_once(*flg2, init2{*init2_called}, 2, 3);
+    cuda::std::call_once(*flg2, init2{*init2_called}, 4, 5);
+}
+
+#endif // TEST_STD_VER >= 11
+__host__ __device__
+void init41(int& init41_called)
+{
+    test_sleep_thread(ms(250));
+    ++init41_called;
+}
+
+__host__ __device__
+void init42(int& init42_called)
+{
+    test_sleep_thread(ms(250));
+    ++init42_called;
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f41(OnceFlag* flg41, OnceFlag* flg42, int* init41_called, int* init42_called)
+{
+    cuda::std::call_once(*flg41, init41, *init41_called);
+    cuda::std::call_once(*flg42, init42, *init42_called);
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f42(OnceFlag* flg41, OnceFlag* flg42, int* init41_called, int* init42_called)
+{
+    cuda::std::call_once(*flg41, init42, *init42_called);
+    cuda::std::call_once(*flg42, init41, *init41_called);
+}
+
+#if TEST_STD_VER >= 11
+
+class MoveOnly
+{
+    __host__ __device__ MoveOnly(const MoveOnly&);
+public:
+    __host__ __device__ MoveOnly() {}
+    __host__ __device__ MoveOnly(MoveOnly&&) {}
+
+    __host__ __device__ void operator()(MoveOnly&&) {}
+};
+
+class NonCopyable
+{
+    __host__ __device__ NonCopyable(const NonCopyable&);
+public:
+    __host__ __device__ NonCopyable() {}
+
+    __host__ __device__ void operator()(int&) {}
+};
+
+// reference qualifiers on functions are a C++11 extension
+struct RefQual
+{
+    __host__ __device__ void operator()( int& lv_called) & { ++lv_called; }
+    __host__ __device__ void operator()( int& rv_called) && { ++rv_called; }
+};
+
+#endif // TEST_STD_VER >= 11
+
+template<class OnceFlag,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<OnceFlag, Initializer> sel;
+    Selector<int, Initializer> sel_count;
+    heterogeneous_thread_handler handler;
+
+    // check basic functionality
+    {
+        SHARED OnceFlag* flg0;
+        flg0 = sel.construct();
+
+        SHARED int* init0_called;
+        init0_called = sel_count.construct();
+
+        handler.run_on_first_thread(f0<OnceFlag>, flg0, init0_called);
+        handler.run_on_second_thread(f0<OnceFlag>, flg0, init0_called);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(*init0_called == 1);
+    }
+
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    // check basic exception safety
+    {
+        SHARED OnceFlag* flg0;
+        flg0 = sel.construct();
+
+        SHARED int* init3_called;
+        init3_called = sel_count.construct();
+        SHARED int* init3_completed;
+        init3_completed = sel_count.construct();
+
+        handler.run_on_first_thread( f3<OnceFlag>, flg3, init3_called, init3_completed);
+        handler.run_on_second_thread(f3<OnceFlag>, flg3, init3_called, init3_completed);
+        handler.join_test_thread();
+        handler.syncthreads();
+
+        assert(*init3_called == 2);
+        assert(*init3_completed == 1);
+    }
+#endif
+    // check deadlock avoidance
+    {
+        SHARED OnceFlag* flg41;
+        flg41 = sel.construct();
+        SHARED OnceFlag* flg42;
+        flg42 = sel.construct();
+
+        SHARED int* init41_called;
+        init41_called = sel_count.construct();
+        SHARED int* init42_called;
+        init42_called = sel_count.construct();
+
+        handler.run_on_first_thread( f41<OnceFlag>, flg41, flg42, init41_called, init42_called);
+        handler.run_on_second_thread(f42<OnceFlag>, flg41, flg42, init41_called, init42_called);
+        handler.join_test_thread();
+        handler.syncthreads();
+
+        assert(*init41_called == 1);
+        assert(*init42_called == 1);
+    }
+#if TEST_STD_VER >= 11
+    // check functors with 1 arg
+    {
+        SHARED OnceFlag* flg1;
+        flg1 = sel.construct();
+
+        SHARED int* init1_called;
+        init1_called = sel_count.construct();
+
+        handler.run_on_first_thread( f1<OnceFlag>, flg1, init1_called);
+        handler.run_on_second_thread(f1<OnceFlag>, flg1, init1_called);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(*init1_called == 1);
+    }
+    // check functors with 2 args
+    {
+        SHARED OnceFlag* flg2;
+        flg2 = sel.construct();
+
+        SHARED int* init2_called;
+        init2_called = sel_count.construct();
+
+        handler.run_on_first_thread( f2<OnceFlag>, flg2, init2_called);
+        handler.run_on_second_thread(f2<OnceFlag>, flg2, init2_called);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(*init2_called == 5);
+    }
+
+    {
+        SHARED OnceFlag* f;
+        f = sel.construct();
+
+        cuda::std::call_once(*f, MoveOnly(), MoveOnly());
+    }
+    // check LWG2442: call_once() shouldn't DECAY_COPY()
+    {
+        SHARED OnceFlag* f;
+        f = sel.construct();
+        int i = 0;
+
+        cuda::std::call_once(*f, NonCopyable(), i);
+    }
+    // reference qualifiers on functions are a C++11 extension
+    {
+        SHARED OnceFlag* f1, *f2;
+        f1 = sel.construct();
+        f2 = sel.construct();
+
+        SHARED int* lv_called;
+        lv_called = sel_count.construct();
+
+        SHARED int* rv_called;
+        rv_called = sel_count.construct();
+
+        RefQual rq;
+        cuda::std::call_once(*f1, rq, *lv_called);
+        assert(*lv_called == 1);
+        cuda::std::call_once(*f2, cuda::std::move(rq), *rv_called);
+        assert(*rv_called == 1);
+    }
+#endif // TEST_STD_VER >= 11
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::once_flag, local_memory_selector>();
+#else
+    test<cuda::std::once_flag, shared_memory_selector>();
+    test<cuda::std::once_flag, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp
new file mode 100644
index 0000000000..bb5a89ca96
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp
@@ -0,0 +1,80 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// template<class Callable, class ...Args>
+//   void call_once(once_flag& flag, Callable&& func, Args&&... args);
+
+// This test is supposed to be run with ThreadSanitizer and verifies that
+// call_once properly synchronizes user state, a data race that was fixed
+// in r280621.
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::milliseconds ms;
+
+__host__ __device__
+void init0(int& global)
+{
+    ++global;
+}
+
+template<class OnceFlag>
+__host__ __device__
+void f0(OnceFlag* flg, int* global)
+{
+    cuda::std::call_once(*flg, init0, *global);
+    assert(*global == 1);
+}
+
+template<class OnceFlag,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<OnceFlag, Initializer> sel;
+    SHARED OnceFlag* flg;
+    flg = sel.construct();
+
+    Selector<int, Initializer> sel_count;
+    SHARED int* global;
+    global = sel_count.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(f0<OnceFlag>, flg, global);
+    handler.run_on_second_thread(f0<OnceFlag>, flg, global);
+    handler.syncthreads();
+    handler.join_test_thread();
+    assert(*global == 1);
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::std::once_flag, local_memory_selector>();
+#else
+    test<cuda::std::once_flag, shared_memory_selector>();
+    test<cuda::std::once_flag, global_memory_selector>();
+#endif
+
+    return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/assign.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/assign.compile.fail.cpp
new file mode 100644
index 0000000000..e2b8ffa589
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/assign.compile.fail.cpp
@@ -0,0 +1,29 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// once_flag& operator=(const once_flag&) = delete;
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::once_flag f;
+    cuda::std::once_flag f2;
+    f2 = f;
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/copy.compile.fail.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/copy.compile.fail.cpp
new file mode 100644
index 0000000000..b92c198d86
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/copy.compile.fail.cpp
@@ -0,0 +1,28 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: nvrtc
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// once_flag(const once_flag&) = delete;
+
+#include<cuda/std/mutex>
+
+int main(int, char**)
+{
+    cuda::std::once_flag f;
+    cuda::std::once_flag f2(f);
+
+  return 0;
+}
diff --git a/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/default.pass.cpp b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/default.pass.cpp
new file mode 100644
index 0000000000..363f0d3bfa
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.onceflag/default.pass.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// <mutex>
+
+// struct once_flag;
+
+// constexpr once_flag() noexcept;
+
+#include<cuda/std/mutex>
+#include "test_macros.h"
+
+int main(int, char**)
+{
+    {
+    cuda::std::once_flag f;
+    unused(f);
+    }
+#if TEST_STD_VER >= 11
+    {
+    constexpr cuda::std::once_flag f;
+    unused(f);
+    }
+#endif
+
+  return 0;
+}
diff --git a/.upstream-tests/test/support/heterogeneous_thread_handler.h b/.upstream-tests/test/support/heterogeneous_thread_handler.h
new file mode 100644
index 0000000000..d3bddb333f
--- /dev/null
+++ b/.upstream-tests/test/support/heterogeneous_thread_handler.h
@@ -0,0 +1,91 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef TEST_SUPPORT_HETEROGENEOUS_TEST_THREAD_H
+#define TEST_SUPPORT_HETEROGENEOUS_TEST_THREAD_H
+
+#include <nv/target>
+
+#include <cuda/std/chrono>
+
+#ifndef __CUDACC_RTC__
+#include <chrono>
+#include <thread>
+#endif // __CUDACC_RTC__
+
+__host__ __device__
+void test_sleep_thread(cuda::std::chrono::milliseconds dur) {
+    constexpr cuda::std::chrono::nanoseconds max_sleep_duration{500000};
+    cuda::std::chrono::nanoseconds dur_mu{dur};
+    cuda::std::chrono::nanoseconds waited{0};
+    const auto start = cuda::std::chrono::high_resolution_clock::now();
+    while (dur_mu > waited) {
+        NV_IF_TARGET(
+            NV_IS_DEVICE,
+            __libcpp_thread_sleep_for(dur_mu - waited < max_sleep_duration ? dur_mu - waited : max_sleep_duration);,
+            std::this_thread::sleep_for(std::chrono::nanoseconds{(dur_mu - waited).count()});
+        )
+        waited = cuda::std::chrono::high_resolution_clock::now() - start;
+    }
+}
+
+struct heterogeneous_thread_handler {
+#ifndef __CUDACC_RTC__
+    union { std::thread t_; };
+#endif // __CUDACC_RTC__
+
+    __host__ __device__
+    heterogeneous_thread_handler() noexcept {}
+    __host__ __device__
+    ~heterogeneous_thread_handler() noexcept {}
+
+    template <class F, class ...Args>
+    __host__ __device__
+    void run_on_first_thread(F&& f, Args&& ...args) {
+        NV_IF_TARGET(
+            NV_IS_DEVICE,
+            if (threadIdx.x == 0) { cuda::std::__invoke(cuda::std::forward<F>(f), cuda::std::forward<Args>(args)...); },
+            cuda::std::__invoke(cuda::std::forward<F>(f), cuda::std::forward<Args>(args)...);
+        )
+    }
+
+    template <class F, class ...Args>
+    __host__ __device__
+    void run_on_second_thread(F&& f, Args&& ...args) {
+        NV_IF_TARGET(
+            NV_IS_DEVICE,
+            if (threadIdx.x == 1) { cuda::std::__invoke(cuda::std::forward<F>(f), cuda::std::forward<Args>(args)...); },
+            ::new((void*)std::addressof(t_)) std::thread(std::forward<F>(f), std::forward<Args>(args)...);
+        )
+    }
+
+    __host__ __device__
+    void sleep_first_thread(cuda::std::chrono::milliseconds dur) {
+        test_sleep_thread(dur);
+    }
+
+    __host__ __device__
+    void syncthreads() const {
+        NV_IF_TARGET(
+            NV_IS_DEVICE,
+            __syncthreads();
+        )
+    }
+
+    __host__ __device__
+    void join_test_thread() {
+        NV_IF_TARGET(
+            NV_IS_DEVICE,
+            (),
+            t_.join();
+        )
+    }
+};
+
+#endif // TEST_SUPPORT_HETEROGENEOUS_TEST_THREAD_H
diff --git a/.upstream-tests/test/support/nasty_containers.h b/.upstream-tests/test/support/nasty_containers.h
index bfc89170f1..f361288773 100644
--- a/.upstream-tests/test/support/nasty_containers.h
+++ b/.upstream-tests/test/support/nasty_containers.h
@@ -295,26 +295,26 @@ bool operator==(const nasty_list<T>& x, const nasty_list<T>& y) { return x.l_ ==
 class nasty_mutex
 {
 public:
-     nasty_mutex() TEST_NOEXCEPT {}
-     ~nasty_mutex() {}
+     __host__ __device__ nasty_mutex() TEST_NOEXCEPT {}
+     __host__ __device__ ~nasty_mutex() {}
 
-    nasty_mutex *operator& ()   { assert(false); return nullptr; }
+    __host__ __device__ nasty_mutex *operator& ()   { assert(false); return nullptr; }
     template <typename T>
-    void operator, (const T &) { assert(false); }
+    __host__ __device__ void operator, (const T &) { assert(false); }
 
 private:
-    nasty_mutex(const nasty_mutex&)            { assert(false); }
-    nasty_mutex& operator=(const nasty_mutex&) { assert(false); return *this; }
+    __host__ __device__ nasty_mutex(const nasty_mutex&)            { assert(false); }
+    __host__ __device__ nasty_mutex& operator=(const nasty_mutex&) { assert(false); return *this; }
 
 public:
-    void lock()               {}
-    bool try_lock() TEST_NOEXCEPT { return true; }
-    void unlock() TEST_NOEXCEPT   {}
+    __host__ __device__ void lock()               {}
+    __host__ __device__ bool try_lock() TEST_NOEXCEPT { return true; }
+    __host__ __device__ void unlock() TEST_NOEXCEPT   {}
 
     // Shared ownership
-    void lock_shared()     {}
-    bool try_lock_shared() { return true; }
-    void unlock_shared()   {}
+    __host__ __device__ void lock_shared()     {}
+    __host__ __device__ bool try_lock_shared() { return true; }
+    __host__ __device__ void unlock_shared()   {}
 };
 
 #endif