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

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

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