Add some tests

.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;
+}

.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;
+}

.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;
+}

.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/lock.pass.cpp

@@ -0,0 +1,86 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class mutex;
+
+// void lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    m->lock();
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::mutex<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::mutex<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::mutex<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}

.upstream-tests/test/cuda/thread/thread.mutex/thread.mutex.requirements/thread.mutex.requirements.mutex/thread.mutex.class/try_lock.pass.cpp

@@ -0,0 +1,90 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: no-threads
+// UNSUPPORTED: pre-sm-70
+
+// FLAKY_TEST.
+
+// <mutex>
+
+// class mutex;
+
+// bool try_lock();
+
+#include<cuda/mutex>
+#include<cuda/std/cassert>
+#include<cuda/std/chrono>
+
+#include "test_macros.h"
+
+#include "cuda_space_selector.h"
+#include "heterogeneous_thread_handler.h"
+
+typedef cuda::std::chrono::system_clock Clock;
+typedef Clock::time_point time_point;
+typedef Clock::duration duration;
+typedef cuda::std::chrono::milliseconds ms;
+typedef cuda::std::chrono::nanoseconds ns;
+
+template<class Mutex>
+__host__ __device__ void f(Mutex* m)
+{
+    time_point t0 = Clock::now();
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    assert(!m->try_lock());
+    while(!m->try_lock())
+        ;
+    time_point t1 = Clock::now();
+    m->unlock();
+    assert(t1 - t0 >= ms(250));
+}
+
+template<class Mutex,
+    template<typename, typename> typename Selector,
+    typename Initializer = constructor_initializer>
+__host__ __device__ void test() {
+    Selector<Mutex, Initializer> sel;
+    SHARED Mutex* m;
+    m = sel.construct();
+
+    heterogeneous_thread_handler handler;
+    handler.run_on_first_thread(&Mutex::lock, m);
+    handler.syncthreads();
+
+    handler.run_on_second_thread(f<Mutex>, m);
+    handler.sleep_first_thread(ms(250));
+    handler.run_on_first_thread(&Mutex::unlock, m);
+    handler.join_test_thread();
+}
+
+int main(int, char**)
+{
+#ifndef __CUDA_ARCH__
+    cuda_thread_count = 2;
+
+    test<cuda::mutex<cuda::thread_scope_system>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  local_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, local_memory_selector>();
+#else
+    test<cuda::mutex<cuda::thread_scope_system>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  shared_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, shared_memory_selector>();
+
+    test<cuda::mutex<cuda::thread_scope_system>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_device>, global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_block>,  global_memory_selector>();
+    test<cuda::mutex<cuda::thread_scope_thread>, global_memory_selector>();
+#endif
+
+    return 0;
+}

.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;
+}

.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;
+}

.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;
+}