From e8b2e537d40d78fe70cccc254c062015ccadb45b 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          |  81 +++
 .../thread.mutex.class/try_lock.pass.cpp      |  85 +++
 .../assign.compile.fail.cpp                   |  36 ++
 .../copy.compile.fail.cpp                     |  35 ++
 .../thread.timedmutex.class/default.pass.cpp  |  31 +
 .../thread.timedmutex.class/lock.pass.cpp     |  81 +++
 .../thread.timedmutex.class/try_lock.pass.cpp |  85 +++
 .../try_lock_for.pass.cpp                     | 124 ++++
 .../try_lock_until.pass.cpp                   | 124 ++++
 .../thread.once.callonce/call_once.pass.cpp   | 350 +++++++++++
 .../thread.once.callonce/race.pass.cpp        |  79 +++
 .../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     |  64 +++
 .../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          |  78 +++
 .../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    |  78 +++
 .../mutex_adopt_lock.pass.cpp                 |  48 ++
 .../mutex_defer_lock.pass.cpp                 |  43 ++
 .../mutex_duration.pass.cpp                   |  86 +++
 .../mutex_time_point.pass.cpp                 |  85 +++
 .../mutex_try_to_lock.pass.cpp                |  77 +++
 .../thread.lock.unique.locking/lock.pass.cpp  |  87 +++
 .../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          |  61 ++
 .../thread.mutex.class/try_lock.pass.cpp      |  65 +++
 .../assign.compile.fail.cpp                   |  28 +
 .../copy.compile.fail.cpp                     |  27 +
 .../thread.timedmutex.class/default.pass.cpp  |  29 +
 .../thread.timedmutex.class/lock.pass.cpp     |  61 ++
 .../thread.timedmutex.class/try_lock.pass.cpp |  65 +++
 .../try_lock_for.pass.cpp                     |  83 +++
 .../try_lock_until.pass.cpp                   |  83 +++
 .../thread.once.callonce/call_once.pass.cpp   | 292 ++++++++++
 .../thread.once.callonce/race.pass.cpp        |  62 ++
 .../assign.compile.fail.cpp                   |  29 +
 .../copy.compile.fail.cpp                     |  28 +
 .../thread.once.onceflag/default.pass.cpp     |  36 ++
 .../support/heterogeneous_thread_handler.h    |  79 +++
 .../test/support/nasty_containers.h           |  24 +-
 79 files changed, 5780 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..deae9ac262
--- /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,81 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::mutex<cuda::thread_scope_system> system_scope;
+STATIC_TEST_GLOBAL_VAR cuda::mutex<cuda::thread_scope_device> device_scope;
+STATIC_TEST_GLOBAL_VAR cuda::mutex<cuda::thread_scope_block>  block_scope;
+STATIC_TEST_GLOBAL_VAR cuda::mutex<cuda::thread_scope_thread> thread_scope;
+
+template<class T> __host__ __device__ T& get_mutex();
+template<> __host__ __device__ cuda::mutex<cuda::thread_scope_system>& get_mutex() { return system_scope; }
+template<> __host__ __device__ cuda::mutex<cuda::thread_scope_device>& get_mutex() { return device_scope; }
+template<> __host__ __device__ cuda::mutex<cuda::thread_scope_block>&  get_mutex() { return block_scope; }
+template<> __host__ __device__ cuda::mutex<cuda::thread_scope_thread>& get_mutex() { return thread_scope; }
+
+using Clock = cuda::std::chrono::system_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 T>
+__host__ __device__ void f()
+{
+    time_point t0 = Clock::now();
+    get_mutex<T>().lock();
+    time_point t1 = Clock::now();
+    get_mutex<T>().unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+template<class T>
+__host__ __device__ void test() {
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&T::lock, get_mutex<T>());
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<T>);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&T::unlock, get_mutex<T>());
+    handler.syncthreads();
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+
+    // test<cuda::mutex<cuda::thread_scope_system>>();
+    test<cuda::mutex<cuda::thread_scope_device>>();
+    test<cuda::mutex<cuda::thread_scope_block>>();
+    test<cuda::mutex<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/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..84f8914791
--- /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,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: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class mutex;
+
+// bool try_lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::mutex<cuda::thread_scope_system> system_scope;
+STATIC_TEST_GLOBAL_VAR cuda::mutex<cuda::thread_scope_device> device_scope;
+STATIC_TEST_GLOBAL_VAR cuda::mutex<cuda::thread_scope_block>  block_scope;
+STATIC_TEST_GLOBAL_VAR cuda::mutex<cuda::thread_scope_thread> thread_scope;
+
+template<class T> __host__ __device__ T& get_mutex();
+template<> __host__ __device__ cuda::mutex<cuda::thread_scope_system>& get_mutex() { return system_scope; }
+template<> __host__ __device__ cuda::mutex<cuda::thread_scope_device>& get_mutex() { return device_scope; }
+template<> __host__ __device__ cuda::mutex<cuda::thread_scope_block>&  get_mutex() { return block_scope; }
+template<> __host__ __device__ cuda::mutex<cuda::thread_scope_thread>& get_mutex() { return thread_scope; }
+
+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 T>
+__host__ __device__ void f()
+{
+    time_point t0 = Clock::now();
+    assert(!get_mutex<T>().try_lock());
+    assert(!get_mutex<T>().try_lock());
+    assert(!get_mutex<T>().try_lock());
+    while(!get_mutex<T>().try_lock())
+        ;
+    time_point t1 = Clock::now();
+    get_mutex<T>().unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(200));  // within 200ms
+}
+
+template<class T>
+__host__ __device__ void test() {
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&T::lock, get_mutex<T>());
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<T>);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&T::unlock, get_mutex<T>());
+    handler.syncthreads();
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+
+    // test<cuda::mutex<cuda::thread_scope_system>>();
+    test<cuda::mutex<cuda::thread_scope_device>>();
+    test<cuda::mutex<cuda::thread_scope_block>>();
+    test<cuda::mutex<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/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..ff50319024
--- /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,81 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_system> system_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_device> device_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_block>  block_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_thread> thread_scope;
+
+template<class T> __host__ __device__ T& get_mutex();
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_system>& get_mutex() { return system_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_device>& get_mutex() { return device_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_block>&  get_mutex() { return block_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_thread>& get_mutex() { return thread_scope; }
+
+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 T>
+__host__ __device__ void f()
+{
+    time_point t0 = Clock::now();
+    get_mutex<T>().lock();
+    time_point t1 = Clock::now();
+    get_mutex<T>().unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+template<class T>
+__host__ __device__ void test() {
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&T::lock, get_mutex<T>());
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<T>);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&T::unlock, get_mutex<T>());
+    handler.syncthreads();
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+
+    // test<cuda::timed_mutex<cuda::thread_scope_system>>();
+    test<cuda::timed_mutex<cuda::thread_scope_device>>();
+    test<cuda::timed_mutex<cuda::thread_scope_block>>();
+    test<cuda::timed_mutex<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/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..a09c6915be
--- /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,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: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// bool try_lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_system> system_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_device> device_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_block>  block_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_thread> thread_scope;
+
+template<class T> __host__ __device__ T& get_mutex();
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_system>& get_mutex() { return system_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_device>& get_mutex() { return device_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_block>&  get_mutex() { return block_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_thread>& get_mutex() { return thread_scope; }
+
+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 T>
+__host__ __device__ void f()
+{
+    time_point t0 = Clock::now();
+    assert(!get_mutex<T>().try_lock());
+    assert(!get_mutex<T>().try_lock());
+    assert(!get_mutex<T>().try_lock());
+    while(!get_mutex<T>().try_lock())
+        ;
+    time_point t1 = Clock::now();
+    get_mutex<T>().unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(200));  // within 200ms
+}
+
+template<class T>
+__host__ __device__ void test() {
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&T::lock, get_mutex<T>());
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<T>);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&T::unlock, get_mutex<T>());
+    handler.syncthreads();
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+
+    // test<cuda::timed_mutex<cuda::thread_scope_system>>();
+    test<cuda::timed_mutex<cuda::thread_scope_device>>();
+    test<cuda::timed_mutex<cuda::thread_scope_block>>();
+    test<cuda::timed_mutex<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/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..08a8d8439a
--- /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,124 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_system> system_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_device> device_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_block>  block_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_thread> thread_scope;
+
+template<class T> __host__ __device__ T& get_mutex();
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_system>& get_mutex() { return system_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_device>& get_mutex() { return device_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_block>&  get_mutex() { return block_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_thread>& get_mutex() { return thread_scope; }
+
+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 T>
+__host__ __device__ void f1()
+{
+    time_point t0 = Clock::now();
+    assert(get_mutex<T>().try_lock_for(ms(300)) == true);
+    time_point t1 = Clock::now();
+    get_mutex<T>().unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+template<class T>
+__host__ __device__ void f2()
+{
+    time_point t0 = Clock::now();
+    assert(get_mutex<T>().try_lock_for(ms(250)) == false);
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+template<class T>
+__host__ __device__ void test() {
+    heterogeneous_thread_handler handler;
+    {
+        handler.run_on_first_thread(&T::lock, get_mutex<T>());
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1<T>);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&T::unlock, get_mutex<T>());
+        handler.syncthreads();
+        handler.join_test_thread();
+    }
+    {
+        handler.run_on_first_thread(&T::lock, get_mutex<T>());
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f2<T>);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&T::unlock, get_mutex<T>());
+        handler.syncthreads();
+        handler.join_test_thread();
+    }
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+
+    // test<cuda::timed_mutex<cuda::thread_scope_system>>();
+    test<cuda::timed_mutex<cuda::thread_scope_device>>();
+    test<cuda::timed_mutex<cuda::thread_scope_block>>();
+    test<cuda::timed_mutex<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/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..6eed75af1f
--- /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,124 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_system> system_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_device> device_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_block>  block_scope;
+STATIC_TEST_GLOBAL_VAR cuda::timed_mutex<cuda::thread_scope_thread> thread_scope;
+
+template<class T> __host__ __device__ T& get_mutex();
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_system>& get_mutex() { return system_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_device>& get_mutex() { return device_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_block>&  get_mutex() { return block_scope; }
+template<> __host__ __device__ cuda::timed_mutex<cuda::thread_scope_thread>& get_mutex() { return thread_scope; }
+
+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 T>
+__host__ __device__ void f1()
+{
+    time_point t0 = Clock::now();
+    assert(get_mutex<T>().try_lock_until(Clock::now() + ms(300)) == true);
+    time_point t1 = Clock::now();
+    get_mutex<T>().unlock();
+    ns d = t1 - t0 - ms(300);
+    assert(d < ms(50));  // within 50ms
+}
+
+template<class T>
+__host__ __device__ void f2()
+{
+    time_point t0 = Clock::now();
+    assert(get_mutex<T>().try_lock_until(Clock::now() + ms(200)) == false);
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(200);
+    assert(d < ms(50));  // within 50ms
+}
+
+template<class T>
+__host__ __device__ void test() {
+    heterogeneous_thread_handler handler;
+    {
+        handler.run_on_first_thread(&T::lock, get_mutex<T>());
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1<T>);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&T::unlock, get_mutex<T>());
+        handler.syncthreads();
+        handler.join_test_thread();
+    }
+    {
+        handler.run_on_first_thread(&T::lock, get_mutex<T>());
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f2<T>);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&T::unlock, get_mutex<T>());
+        handler.syncthreads();
+        handler.join_test_thread();
+    }
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+
+    // test<cuda::timed_mutex<cuda::thread_scope_system>>();
+    test<cuda::timed_mutex<cuda::thread_scope_device>>();
+    test<cuda::timed_mutex<cuda::thread_scope_block>>();
+    test<cuda::timed_mutex<cuda::thread_scope_thread>>();
+
+    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..871a96eb7d
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp
@@ -0,0 +1,350 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+typedef cuda::std::chrono::milliseconds ms;
+
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_system> flg0_system;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_device> flg0_device;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_block>  flg0_block;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_thread> flg0_thread;
+
+template<cuda::thread_scope Sco> __host__ __device__ cuda::once_flag<Sco>& get_flag0();
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_system>& get_flag0<cuda::thread_scope_system>() { return flg0_system; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_device>& get_flag0<cuda::thread_scope_device>() { return flg0_device; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_block>&  get_flag0<cuda::thread_scope_block>()  { return flg0_block; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_thread>& get_flag0<cuda::thread_scope_thread>() { return flg0_thread; }
+
+STATIC_TEST_GLOBAL_VAR long init0_called[cuda::thread_scope_thread] = {0};
+template<cuda::thread_scope Sco>
+__host__ __device__
+void init0()
+{
+#ifdef __CUDA_ARCH__
+    __libcpp_thread_sleep_for(ms(250));
+#else
+    std::this_thread::sleep_for(std::chrono::milliseconds(250));
+#endif
+    ++init0_called[Sco];
+}
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void f0()
+{
+    cuda::std::call_once(get_flag0<Sco>(), init0<Sco>);
+}
+
+#ifndef TEST_HAS_NO_EXCEPTIONS
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_system> flg3_system;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_device> flg3_device;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_block>  flg3_block;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_thread> flg3_thread;
+
+template<cuda::thread_scope Sco> __host__ __device__ cuda::once_flag<Sco>& get_flag3();
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_system>& get_flag3<cuda::thread_scope_system>() { return flg3_system; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_device>& get_flag3<cuda::thread_scope_device>() { return flg3_device; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_block>&  get_flag3<cuda::thread_scope_block>()  { return flg3_block; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_thread>& get_flag3<cuda::thread_scope_thread>() { return flg3_thread; }
+
+STATIC_TEST_GLOBAL_VAR int init3_called[cuda::thread_scope_thread] = {0};
+STATIC_TEST_GLOBAL_VAR int init3_completed[cuda::thread_scope_thread] = {0};
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void init3()
+{
+    ++init3_called[Sco];
+#ifdef __CUDA_ARCH__
+    __libcpp_thread_sleep_for(ms(250));
+#else
+    std::this_thread::sleep_for(std::chrono::milliseconds(250));
+#endif
+    if (init3_called[Sco] == 1)
+#ifdef __CUDA_ARCH__
+        _LIBCUDACXX_UNREACHABLE();
+#else
+        TEST_THROW(1);
+#endif
+    ++init3_completed[Sco];
+}
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void f3()
+{
+    try
+    {
+        cuda::std::call_once(get_flag3<Sco>(), init3<Sco>);
+    }
+    catch (...)
+    {
+    }
+}
+#endif // TEST_HAS_NO_EXCEPTIONS
+
+#if TEST_STD_VER >= 11
+
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_system> flg1_system;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_device> flg1_device;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_block>  flg1_block;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_thread> flg1_thread;
+
+template<cuda::thread_scope Sco> __host__ __device__ cuda::once_flag<Sco>& get_flag1();
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_system>& get_flag1<cuda::thread_scope_system>() { return flg1_system; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_device>& get_flag1<cuda::thread_scope_device>() { return flg1_device; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_block>&  get_flag1<cuda::thread_scope_block>()  { return flg1_block; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_thread>& get_flag1<cuda::thread_scope_thread>() { return flg1_thread; }
+
+STATIC_TEST_GLOBAL_VAR int init1_called[cuda::thread_scope_thread] = {0};
+template<cuda::thread_scope Sco>
+struct init1
+{
+    __host__ __device__ void operator()(int i) {init1_called[Sco] += i;}
+};
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void f1()
+{
+    cuda::std::call_once(get_flag1<Sco>(), init1<Sco>(), 1);
+}
+
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_system> flg2_system;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_device> flg2_device;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_block>  flg2_block;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_thread> flg2_thread;
+
+template<cuda::thread_scope Sco> __host__ __device__ cuda::once_flag<Sco>& get_flag2();
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_system>& get_flag2<cuda::thread_scope_system>() { return flg2_system; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_device>& get_flag2<cuda::thread_scope_device>() { return flg2_device; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_block>&  get_flag2<cuda::thread_scope_block>()  { return flg2_block; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_thread>& get_flag2<cuda::thread_scope_thread>() { return flg2_thread; }
+
+STATIC_TEST_GLOBAL_VAR int init2_called[cuda::thread_scope_thread] = {0};
+template<cuda::thread_scope Sco>
+struct init2
+{
+    __host__ __device__ void operator()(int i, int j) const {init2_called[Sco] += i + j;}
+};
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void f2()
+{
+    cuda::std::call_once(get_flag2<Sco>(), init2<Sco>(), 2, 3);
+    cuda::std::call_once(get_flag2<Sco>(), init2<Sco>(), 4, 5);
+}
+
+#endif // TEST_STD_VER >= 11
+
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_system> flg41_system;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_device> flg41_device;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_block>  flg41_block;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_thread> flg41_thread;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_system> flg42_system;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_device> flg42_device;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_block>  flg42_block;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_thread> flg42_thread;
+
+template<cuda::thread_scope Sco> __host__ __device__ cuda::once_flag<Sco>& get_flag41();
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_system>& get_flag41<cuda::thread_scope_system>() { return flg41_system; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_device>& get_flag41<cuda::thread_scope_device>() { return flg41_device; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_block>&  get_flag41<cuda::thread_scope_block>()  { return flg41_block; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_thread>& get_flag41<cuda::thread_scope_thread>() { return flg41_thread; }
+
+template<cuda::thread_scope Sco> __host__ __device__ cuda::once_flag<Sco>& get_flag42();
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_system>& get_flag42<cuda::thread_scope_system>() { return flg42_system; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_device>& get_flag42<cuda::thread_scope_device>() { return flg42_device; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_block>&  get_flag42<cuda::thread_scope_block>()  { return flg42_block; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_thread>& get_flag42<cuda::thread_scope_thread>() { return flg42_thread; }
+
+STATIC_TEST_GLOBAL_VAR int init41_called[cuda::thread_scope_thread] = {0};
+STATIC_TEST_GLOBAL_VAR int init42_called[cuda::thread_scope_thread] = {0};
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void init42();
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void init41()
+{
+#ifdef __CUDA_ARCH__
+    __libcpp_thread_sleep_for(ms(250));
+#else
+    std::this_thread::sleep_for(std::chrono::milliseconds(250));
+#endif
+    ++init41_called[Sco];
+}
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void init42()
+{
+#ifdef __CUDA_ARCH__
+    __libcpp_thread_sleep_for(ms(250));
+#else
+    std::this_thread::sleep_for(std::chrono::milliseconds(250));
+#endif
+    ++init42_called[Sco];
+}
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void f41()
+{
+    cuda::std::call_once(get_flag41<Sco>(), init41<Sco>);
+    cuda::std::call_once(get_flag42<Sco>(), init42<Sco>);
+}
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void f42()
+{
+    cuda::std::call_once(get_flag42<Sco>(), init42<Sco>);
+    cuda::std::call_once(get_flag41<Sco>(), init41<Sco>);
+}
+
+#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
+{
+    int lv_called, rv_called;
+
+    __host__ __device__ RefQual() : lv_called(0), rv_called(0) {}
+
+    __host__ __device__ void operator()() & { ++lv_called; }
+    __host__ __device__ void operator()() && { ++rv_called; }
+};
+
+#endif // TEST_STD_VER >= 11
+
+template<cuda::thread_scope Sco>
+__host__ __device__ void test() {
+    heterogeneous_thread_handler handler;
+    // check basic functionality
+    {
+        handler.run_on_first_thread(f0<Sco>);
+        handler.run_on_second_thread(f0<Sco>);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(init0_called[Sco] == 1);
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    // check basic exception safety
+    {
+        handler.run_on_first_thread(f3<Sco>);
+        handler.run_on_second_thread(f3<Sco>);
+        handler.join_test_thread();
+        handler.syncthreads();
+
+        assert(init3_called[Sco] == 2);
+        assert(init3_completed[Sco] == 1);
+    }
+#endif
+    // check deadlock avoidance
+    {
+        handler.run_on_first_thread(f41<Sco>);
+        handler.run_on_second_thread(f42<Sco>);
+        handler.join_test_thread();
+        handler.syncthreads();
+
+        assert(init41_called[Sco] == 1);
+        assert(init42_called[Sco] == 1);
+    }
+#if TEST_STD_VER >= 11
+    // check functors with 1 arg
+    {
+        handler.run_on_first_thread(f1<Sco>);
+        handler.run_on_second_thread(f1<Sco>);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(init1_called[Sco] == 1);
+    }
+    // check functors with 2 args
+    {
+        handler.run_on_first_thread(f2<Sco>);
+        handler.run_on_second_thread(f2<Sco>);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(init2_called[Sco] == 5);
+    }
+    {
+        cuda::once_flag<Sco> f;
+        cuda::call_once(f, MoveOnly(), MoveOnly());
+    }
+    // check LWG2442: call_once() shouldn't DECAY_COPY()
+    {
+        cuda::once_flag<Sco> f;
+        int i = 0;
+        cuda::call_once(f, NonCopyable(), i);
+    }
+    // reference qualifiers on functions are a C++11 extension
+    {
+        cuda::once_flag<Sco> f1, f2;
+        RefQual rq;
+        cuda::call_once(f1, rq);
+        assert(rq.lv_called == 1);
+        cuda::call_once(f2, cuda::std::move(rq));
+        assert(rq.rv_called == 1);
+    }
+#endif // TEST_STD_VER >= 11
+    handler.syncthreads();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+
+    // 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.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..190b1893ce
--- /dev/null
+++ b/.upstream-tests/test/cuda/thread/thread.mutex/thread.once/thread.once.callonce/race.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
+
+// <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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_system> system_scope;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_device> device_scope;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_block>  block_scope;
+STATIC_TEST_GLOBAL_VAR cuda::once_flag<cuda::thread_scope_thread> thread_scope;
+STATIC_TEST_GLOBAL_VAR long global[cuda::thread_scope_thread] = {0};
+
+template<cuda::thread_scope Sco> __host__ __device__ cuda::once_flag<Sco>& get_flag();
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_system>& get_flag<cuda::thread_scope_system>() { return system_scope; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_device>& get_flag<cuda::thread_scope_device>() { return device_scope; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_block>&  get_flag<cuda::thread_scope_block>() { return block_scope; }
+template<> __host__ __device__ cuda::once_flag<cuda::thread_scope_thread>& get_flag<cuda::thread_scope_thread>() { return thread_scope; }
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void init0()
+{
+    ++global[Sco];
+}
+
+template<cuda::thread_scope Sco>
+__host__ __device__
+void f0()
+{
+    cuda::call_once(get_flag<Sco>(), init0<Sco>);
+    assert(global[Sco] == 1);
+}
+
+template<cuda::thread_scope Sco>
+__host__ __device__ void test() {
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(f0<Sco>);
+    handler.run_on_second_thread(f0<Sco>);
+    handler.syncthreads();
+    handler.join_test_thread();
+    assert(global[Sco] == 1);
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    // 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/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..fca167e9ad
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/adopt_lock.pass.cpp
@@ -0,0 +1,64 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::mutex m;
+
+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;
+
+__host__ __device__
+void f()
+{
+    time_point t0 = Clock::now();
+    time_point t1;
+    {
+    cuda::std::lock_guard<cuda::std::mutex> lg(m, cuda::std::adopt_lock);
+    t1 = Clock::now();
+    }
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(200));  // within 200ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+    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..b10799f642
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.lock/thread.lock.guard/mutex.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 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::mutex m;
+
+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;
+
+__host__ __device__
+void f()
+{
+    time_point t0 = Clock::now();
+    time_point t1;
+    {
+    cuda::std::lock_guard<cuda::std::mutex> lg(m);
+    t1 = Clock::now();
+    }
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(200));  // within 200ms
+}
+
+__host__ __device__
+void check_ctad()
+{
+#if TEST_STD_VER >= 17
+    cuda::std::lock_guard lg(m);
+    static_assert((cuda::std::is_same<decltype(lg), cuda::std::lock_guard<cuda::std::mutex>>::value), "" );
+    unused(lg);
+#endif
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+
+    handler.run_on_first_thread(check_ctad);
+    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..eac432f199
--- /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,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
+
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::mutex m;
+
+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;
+
+__host__ __device__ void f()
+{
+    time_point t0 = Clock::now();
+    time_point t1;
+    {
+    cuda::std::unique_lock<cuda::std::mutex> ul(m);
+    t1 = Clock::now();
+    }
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+__host__ __device__
+void check_ctad()
+{
+#if TEST_STD_VER >= 17
+    cuda::std::unique_lock ul(m);
+    static_assert((cuda::std::is_same<decltype(ul), cuda::std::unique_lock<cuda::std::mutex>>::value), "" );
+    unused(ul);
+#endif
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+
+    handler.run_on_first_thread(check_ctad);
+    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..08b5a45078
--- /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,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;
+
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::timed_mutex m;
+
+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;
+
+__host__ __device__ void f1()
+{
+    time_point t0 = Clock::now();
+    cuda::std::unique_lock<cuda::std::timed_mutex> lk(m, ms(300));
+    assert(lk.owns_lock() == true);
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(300);
+    assert(d < ms(50));  // within 50ms
+}
+
+__host__ __device__ void f2()
+{
+    time_point t0 = Clock::now();
+    cuda::std::unique_lock<cuda::std::timed_mutex> lk(m, ms(200));
+    assert(lk.owns_lock() == false);
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(200);
+    assert(d < ms(50));  // within 50ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+
+    heterogeneous_thread_handler handler;
+    {
+      handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+      handler.syncthreads();
+
+      handler.run_on_second_thread(f1);
+      handler.sleep_first_thread(ms(250));
+      handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+      handler.syncthreads();
+      handler.join_test_thread();
+    }
+    {
+      handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+      handler.syncthreads();
+
+      handler.run_on_second_thread(f2);
+      handler.sleep_first_thread(ms(250));
+      handler.syncthreads(); // This sync should not be necessary, however we are running into timing issues otherwise
+      handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+      handler.syncthreads();
+      handler.join_test_thread();
+    }
+    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..03d49c58fd
--- /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,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: 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::timed_mutex m;
+
+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;
+
+__host__ __device__ void f1()
+{
+    time_point t0 = Clock::now();
+    cuda::std::unique_lock<cuda::std::timed_mutex> lk(m, Clock::now() + ms(300));
+    assert(lk.owns_lock() == true);
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(300);
+    assert(d < ms(50));  // within 50ms
+}
+
+__host__ __device__ void f2()
+{
+    time_point t0 = Clock::now();
+    cuda::std::unique_lock<cuda::std::timed_mutex> lk(m, Clock::now() + ms(200));
+    assert(lk.owns_lock() == false);
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(200);
+    assert(d < ms(50));  // within 50ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    {
+      handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+      handler.syncthreads();
+
+      handler.run_on_second_thread(f1);
+      handler.sleep_first_thread(ms(250));
+      handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+      handler.syncthreads();
+      handler.join_test_thread();
+    }
+    {
+      handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+      handler.syncthreads();
+
+      handler.run_on_second_thread(f2);
+      handler.sleep_first_thread(ms(250));
+      handler.syncthreads(); // This sync should not be necessary, however we are running into timing issues otherwise
+      handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+      handler.syncthreads();
+      handler.join_test_thread();
+    }
+    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..8f31e56295
--- /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,77 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::mutex m;
+
+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;
+
+__host__ __device__ void f()
+{
+    time_point t0 = Clock::now();
+    {
+        cuda::std::unique_lock<cuda::std::mutex> lk(m, cuda::std::try_to_lock);
+        assert(lk.owns_lock() == false);
+    }
+    {
+        cuda::std::unique_lock<cuda::std::mutex> lk(m, cuda::std::try_to_lock);
+        assert(lk.owns_lock() == false);
+    }
+    {
+        cuda::std::unique_lock<cuda::std::mutex> lk(m, cuda::std::try_to_lock);
+        assert(lk.owns_lock() == false);
+    }
+    while (true)
+    {
+        cuda::std::unique_lock<cuda::std::mutex> lk(m, cuda::std::try_to_lock);
+        if (lk.owns_lock())
+            break;
+    }
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(200));  // within 200ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+    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..8b930d0320
--- /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,87 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::mutex m;
+
+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;
+
+__host__ __device__ void f()
+{
+    cuda::std::unique_lock<cuda::std::mutex> lk(m, cuda::std::defer_lock);
+    time_point t0 = Clock::now();
+    lk.lock();
+    time_point t1 = Clock::now();
+    assert(lk.owns_lock() == true);
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(25));  // within 25ms
+#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
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+
+  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..7743421c6f
--- /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)
+        using Clock = cuda::std::chrono::system_clock;
+#else
+        using Clock = cuda::std::chrono::steady_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..e118f123b7
--- /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,61 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::mutex m;
+
+using Clock = cuda::std::chrono::system_clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+__host__ __device__ void f()
+{
+    time_point t0 = Clock::now();
+    m.lock();
+    time_point t1 = Clock::now();
+    m.unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+
+    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..077c58c2d1
--- /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,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
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class mutex;
+
+// bool try_lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+
+#include "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::mutex m;
+
+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;
+
+__host__ __device__ void f()
+{
+    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();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(200));  // within 200ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+
+    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..bdd004e6b5
--- /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,61 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::timed_mutex m;
+
+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;
+
+__host__ __device__ void f()
+{
+    time_point t0 = Clock::now();
+    m.lock();
+    time_point t1 = Clock::now();
+    m.unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+
+  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..17a5fb8aaf
--- /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,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
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class timed_mutex;
+
+// bool try_lock();
+
+#include<cuda/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::timed_mutex m;
+
+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;
+
+__host__ __device__ void f()
+{
+    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();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(200));  // within 200ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+    handler.syncthreads();
+    handler.join_test_thread();
+    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..78e59974c9
--- /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,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
+
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::timed_mutex m;
+
+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;
+
+__host__ __device__ void f1()
+{
+    time_point t0 = Clock::now();
+    assert(m.try_lock_for(ms(300)) == true);
+    time_point t1 = Clock::now();
+    m.unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+__host__ __device__ void f2()
+{
+    time_point t0 = Clock::now();
+    assert(m.try_lock_for(ms(250)) == false);
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    {
+        handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+        handler.syncthreads();
+        handler.join_test_thread();
+    }
+    {
+        handler.run_on_first_thread(&cuda::std::mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&cuda::std::mutex::unlock, m);
+        handler.syncthreads();
+        handler.join_test_thread();
+    }
+
+  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..22b2940a24
--- /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,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
+
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR cuda::std::timed_mutex m;
+
+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;
+
+__host__ __device__ void f1()
+{
+    time_point t0 = Clock::now();
+    assert(m.try_lock_until(Clock::now() + ms(300)) == true);
+    time_point t1 = Clock::now();
+    m.unlock();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+__host__ __device__ void f2()
+{
+    time_point t0 = Clock::now();
+    assert(m.try_lock_until(Clock::now() + ms(250)) == false);
+    time_point t1 = Clock::now();
+    ns d = t1 - t0 - ms(250);
+    assert(d < ms(50));  // within 50ms
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    {
+        handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+        handler.syncthreads();
+        handler.join_test_thread();
+    }
+    {
+        handler.run_on_first_thread(&cuda::std::timed_mutex::lock, m);
+        handler.syncthreads();
+
+        handler.run_on_second_thread(f1);
+        handler.sleep_first_thread(ms(250));
+        handler.run_on_first_thread(&cuda::std::timed_mutex::unlock, m);
+        handler.syncthreads();
+        handler.join_test_thread();
+    }
+
+    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..b520d1b30b
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/call_once.pass.cpp
@@ -0,0 +1,292 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/std/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+typedef cuda::std::chrono::milliseconds ms;
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::once_flag flg0;
+STATIC_TEST_GLOBAL_VAR
+int init0_called = 0;
+
+__host__ __device__
+void init0()
+{
+#ifdef __CUDA_ARCH__
+    __libcpp_thread_sleep_for(ms(250));
+#else
+    std::this_thread::sleep_for(std::chrono::milliseconds(250));
+#endif
+    ++init0_called;
+}
+
+__host__ __device__
+void f0()
+{
+    cuda::std::call_once(flg0, init0);
+}
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::once_flag flg3;
+
+STATIC_TEST_GLOBAL_VAR
+int init3_called = 0;
+STATIC_TEST_GLOBAL_VAR
+int init3_completed = 0;
+
+__host__ __device__
+void init3()
+{
+    ++init3_called;
+#ifdef __CUDA_ARCH__
+    __libcpp_thread_sleep_for(ms(250));
+#else
+    std::this_thread::sleep_for(std::chrono::milliseconds(250));
+#endif
+    if (init3_called == 1)
+#ifdef __CUDA_ARCH__
+        _LIBCUDACXX_UNREACHABLE();
+#else
+        TEST_THROW(1);
+#endif
+    ++init3_completed;
+}
+
+__host__ __device__
+void f3()
+{
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    try
+    {
+        cuda::std::call_once(flg3, init3);
+    }
+    catch (...)
+    {
+    }
+#endif
+    unused(flg3);
+}
+
+#if TEST_STD_VER >= 11
+
+STATIC_TEST_GLOBAL_VAR
+int init1_called = 0;
+
+struct init1
+{
+    __host__ __device__ void operator()(int i) {init1_called += i;}
+};
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::once_flag flg1;
+
+__host__ __device__
+void f1()
+{
+    cuda::std::call_once(flg1, init1(), 1);
+}
+
+STATIC_TEST_GLOBAL_VAR
+int init2_called = 0;
+
+struct init2
+{
+    __host__ __device__ void operator()(int i, int j) const {init2_called += i + j;}
+};
+
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::once_flag flg2;
+
+__host__ __device__
+void f2()
+{
+    cuda::std::call_once(flg2, init2(), 2, 3);
+    cuda::std::call_once(flg2, init2(), 4, 5);
+}
+
+#endif // TEST_STD_VER >= 11
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::once_flag flg41;
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::once_flag flg42;
+
+STATIC_TEST_GLOBAL_VAR
+int init41_called = 0;
+
+STATIC_TEST_GLOBAL_VAR
+int init42_called = 0;
+
+__host__ __device__
+void init42();
+
+__host__ __device__
+void init41()
+{
+#ifdef __CUDA_ARCH__
+    __libcpp_thread_sleep_for(ms(250));
+#else
+    std::this_thread::sleep_for(std::chrono::milliseconds(250));
+#endif
+    ++init41_called;
+}
+
+__host__ __device__
+void init42()
+{
+#ifdef __CUDA_ARCH__
+    __libcpp_thread_sleep_for(ms(250));
+#else
+    std::this_thread::sleep_for(std::chrono::milliseconds(250));
+#endif
+    ++init42_called;
+}
+
+__host__ __device__
+void f41()
+{
+    cuda::std::call_once(flg41, init41);
+    cuda::std::call_once(flg42, init42);
+}
+
+__host__ __device__
+void f42()
+{
+    cuda::std::call_once(flg42, init42);
+    cuda::std::call_once(flg41, init41);
+}
+
+#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
+{
+    int lv_called, rv_called;
+
+    __host__ __device__ RefQual() : lv_called(0), rv_called(0) {}
+
+    __host__ __device__ void operator()() & { ++lv_called; }
+    __host__ __device__ void operator()() && { ++rv_called; }
+};
+
+#endif // TEST_STD_VER >= 11
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    // check basic functionality
+    {
+        handler.run_on_first_thread(f0);
+        handler.run_on_second_thread(f0);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(init0_called == 1);
+    }
+#ifndef TEST_HAS_NO_EXCEPTIONS
+    // check basic exception safety
+    {
+        handler.run_on_first_thread(f3);
+        handler.run_on_second_thread(f3);
+        handler.join_test_thread();
+        handler.syncthreads();
+
+        assert(init3_called == 2);
+        assert(init3_completed == 1);
+    }
+#endif
+    // check deadlock avoidance
+    {
+        handler.run_on_first_thread(f41);
+        handler.run_on_second_thread(f42);
+        handler.join_test_thread();
+        handler.syncthreads();
+
+        assert(init41_called == 1);
+        assert(init42_called == 1);
+    }
+#if TEST_STD_VER >= 11
+    // check functors with 1 arg
+    {
+        handler.run_on_first_thread(f1);
+        handler.run_on_second_thread(f1);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(init1_called == 1);
+    }
+    // check functors with 2 args
+    {
+        handler.run_on_first_thread(f2);
+        handler.run_on_second_thread(f2);
+        handler.join_test_thread();
+        handler.syncthreads();
+        assert(init2_called == 5);
+    }
+    {
+        cuda::std::once_flag f;
+        cuda::std::call_once(f, MoveOnly(), MoveOnly());
+    }
+    // check LWG2442: call_once() shouldn't DECAY_COPY()
+    {
+        cuda::std::once_flag f;
+        int i = 0;
+        cuda::std::call_once(f, NonCopyable(), i);
+    }
+// reference qualifiers on functions are a C++11 extension
+    {
+        cuda::std::once_flag f1, f2;
+        RefQual rq;
+        cuda::std::call_once(f1, rq);
+        assert(rq.lv_called == 1);
+        cuda::std::call_once(f2, cuda::std::move(rq));
+        assert(rq.rv_called == 1);
+    }
+#endif // TEST_STD_VER >= 11
+
+  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..c5b1262856
--- /dev/null
+++ b/.upstream-tests/test/std/thread/thread.mutex/thread.once/thread.once.callonce/race.pass.cpp
@@ -0,0 +1,62 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 "heterogeneous_thread_handler.h"
+#include "test_macros.h"
+
+STATIC_TEST_GLOBAL_VAR
+cuda::std::once_flag flg0;
+STATIC_TEST_GLOBAL_VAR
+long global = 0;
+
+__host__ __device__
+void init0()
+{
+    ++global;
+}
+
+__host__ __device__
+void f0()
+{
+    cuda::std::call_once(flg0, init0);
+    assert(global == 1);
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+#endif
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(f0);
+    handler.run_on_second_thread(f0);
+    handler.syncthreads();
+    handler.join_test_thread();
+
+    assert(global == 1);
+    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..b76ec246b9
--- /dev/null
+++ b/.upstream-tests/test/support/heterogeneous_thread_handler.h
@@ -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.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef TEST_SUPPORT_HETEROGENEOUS_TEST_THREAD_H
+#define TEST_SUPPORT_HETEROGENEOUS_TEST_THREAD_H
+
+#include <cuda/std/chrono>
+
+#ifndef __CUDA_ARCH__
+#include <chrono>
+#include <thread>
+
+struct heterogeneous_thread_handler {
+    std::thread t_;
+
+    template <class F, class ...Args>
+    void run_on_first_thread(F&& f, Args&& ...args) {
+        cuda::std::__invoke(cuda::std::forward<F>(f), cuda::std::forward<Args>(args)...);
+    }
+
+    template <class F, class ...Args>
+    void run_on_second_thread(F&& f, Args&& ...args) {
+        t_ = std::thread(std::forward<F>(f), std::forward<Args>(args)...);
+    }
+
+    void sleep_first_thread(cuda::std::chrono::nanoseconds dur) {
+        std::this_thread::sleep_for(std::chrono::nanoseconds{dur.count()});
+    }
+
+    void syncthreads() const {}
+
+    void join_test_thread() {
+        t_.join();
+    }
+};
+
+#else
+
+struct heterogeneous_thread_handler {
+    template <class F, class ...Args>
+    __device__
+    void run_on_first_thread(F&& f, Args&& ...args) {
+        if (threadIdx.x == 0) {
+            cuda::std::__invoke(cuda::std::forward<F>(f), cuda::std::forward<Args>(args)...);
+        }
+    }
+
+    template <class F, class ...Args>
+    __device__
+    void run_on_second_thread(F&& f, Args&& ...args) {
+        if (threadIdx.x == 1) {
+            cuda::std::__invoke(cuda::std::forward<F>(f), cuda::std::forward<Args>(args)...);
+        }
+    }
+
+    __device__
+    void sleep_first_thread(cuda::std::chrono::nanoseconds dur) {
+        if (threadIdx.x == 0) {
+            __libcpp_thread_sleep_for(dur);
+        }
+    }
+
+    __device__
+    void syncthreads() const {
+        __syncthreads();
+    }
+
+    __device__
+    void join_test_thread() {}
+};
+#endif // __CUDA_ARCH__
+
+#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