diff --git a/libcudacxx/include/cuda/std/__numeric/gcd_lcm.h b/libcudacxx/include/cuda/std/__numeric/gcd_lcm.h
new file mode 100644
index 00000000000..f4a23fd73b6
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__numeric/gcd_lcm.h
@@ -0,0 +1,112 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBCUDACXX___NUMERIC_GCD_LCM_H
+#define _LIBCUDACXX___NUMERIC_GCD_LCM_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__type_traits/common_type.h>
+#include <cuda/std/__type_traits/is_integral.h>
+#include <cuda/std/__type_traits/is_same.h>
+#include <cuda/std/__type_traits/is_signed.h>
+#include <cuda/std/__type_traits/make_unsigned.h>
+#include <cuda/std/limits>
+
+// comes last
+#include <cuda/std/detail/__pragma_push>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+template <typename _Result, typename _Source, bool _IsSigned = _LIBCUDACXX_TRAIT(is_signed, _Source)>
+struct __ct_abs;
+
+template <typename _Result, typename _Source>
+struct __ct_abs<_Result, _Source, true>
+{
+  _CCCL_CONSTEXPR_CXX14 _LIBCUDACXX_INLINE_VISIBILITY _Result operator()(_Source __t) const noexcept
+  {
+    if (__t >= 0)
+    {
+      return __t;
+    }
+    _CCCL_DIAG_PUSH
+    _CCCL_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
+    if (__t == (numeric_limits<_Source>::min)())
+    {
+      return -static_cast<_Result>(__t);
+    }
+    _CCCL_DIAG_POP
+    return -__t;
+  }
+};
+
+template <typename _Result, typename _Source>
+struct __ct_abs<_Result, _Source, false>
+{
+  _CCCL_CONSTEXPR_CXX14 _LIBCUDACXX_INLINE_VISIBILITY _Result operator()(_Source __t) const noexcept
+  {
+    return __t;
+  }
+};
+
+template <class _Tp>
+_CCCL_CONSTEXPR_CXX14 _LIBCUDACXX_INLINE_VISIBILITY _Tp __gcd(_Tp __m, _Tp __n)
+{
+  static_assert((!_LIBCUDACXX_TRAIT(is_signed, _Tp)), "");
+  return __n == 0 ? __m : _CUDA_VSTD::__gcd<_Tp>(__n, __m % __n);
+}
+
+template <class _Tp, class _Up>
+_CCCL_CONSTEXPR_CXX14 _LIBCUDACXX_INLINE_VISIBILITY __common_type_t<_Tp, _Up> gcd(_Tp __m, _Up __n)
+{
+  static_assert((_LIBCUDACXX_TRAIT(is_integral, _Tp) && _LIBCUDACXX_TRAIT(is_integral, _Up)),
+                "Arguments to gcd must be integer types");
+  static_assert((!_LIBCUDACXX_TRAIT(is_same, __remove_cv_t<_Tp>, bool)), "First argument to gcd cannot be bool");
+  static_assert((!_LIBCUDACXX_TRAIT(is_same, __remove_cv_t<_Up>, bool)), "Second argument to gcd cannot be bool");
+  using _Rp = __common_type_t<_Tp, _Up>;
+  using _Wp = __make_unsigned_t<_Rp>;
+  return static_cast<_Rp>(
+    _CUDA_VSTD::__gcd(static_cast<_Wp>(__ct_abs<_Rp, _Tp>()(__m)), static_cast<_Wp>(__ct_abs<_Rp, _Up>()(__n))));
+}
+
+template <class _Tp, class _Up>
+_CCCL_CONSTEXPR_CXX14 _LIBCUDACXX_INLINE_VISIBILITY __common_type_t<_Tp, _Up> lcm(_Tp __m, _Up __n)
+{
+  static_assert((_LIBCUDACXX_TRAIT(is_integral, _Tp) && _LIBCUDACXX_TRAIT(is_integral, _Up)),
+                "Arguments to lcm must be integer types");
+  static_assert((!_LIBCUDACXX_TRAIT(is_same, __remove_cv_t<_Tp>, bool)), "First argument to lcm cannot be bool");
+  static_assert((!_LIBCUDACXX_TRAIT(is_same, __remove_cv_t<_Up>, bool)), "Second argument to lcm cannot be bool");
+  if (__m == 0 || __n == 0)
+  {
+    return 0;
+  }
+
+  using _Rp  = __common_type_t<_Tp, _Up>;
+  _Rp __val1 = __ct_abs<_Rp, _Tp>()(__m) / _CUDA_VSTD::gcd(__m, __n);
+  _Rp __val2 = __ct_abs<_Rp, _Up>()(__n);
+  _LIBCUDACXX_ASSERT((numeric_limits<_Rp>::max() / __val1 > __val2), "Overflow in lcm");
+  return __val1 * __val2;
+}
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#include <cuda/std/detail/__pragma_pop>
+
+#endif // _LIBCUDACXX___NUMERIC_GCD_LCM_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/numeric b/libcudacxx/include/cuda/std/detail/libcxx/include/numeric
index b57ec87c358..e128bf0a05f 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/numeric
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/numeric
@@ -159,6 +159,7 @@ floating_point  midpoint(floating_point a, floating_point b);    // C++20
 #include <cuda/std/__numeric/accumulate.h>
 #include <cuda/std/__numeric/adjacent_difference.h>
 #include <cuda/std/__numeric/exclusive_scan.h>
+#include <cuda/std/__numeric/gcd_lcm.h>
 #include <cuda/std/__numeric/inclusive_scan.h>
 #include <cuda/std/__numeric/inner_product.h>
 #include <cuda/std/__numeric/iota.h>
@@ -178,75 +179,6 @@ floating_point  midpoint(floating_point a, floating_point b);    // C++20
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
 #ifndef __cuda_std__
-#  if _CCCL_STD_VER > 2014
-template <typename _Result, typename _Source, bool _IsSigned = is_signed<_Source>::value>
-struct __ct_abs;
-
-template <typename _Result, typename _Source>
-struct __ct_abs<_Result, _Source, true>
-{
-  constexpr _LIBCUDACXX_INLINE_VISIBILITY _Result operator()(_Source __t) const noexcept
-  {
-    if (__t >= 0)
-    {
-      return __t;
-    }
-    if (__t == numeric_limits<_Source>::min())
-    {
-      return -static_cast<_Result>(__t);
-    }
-    return -__t;
-  }
-};
-
-template <typename _Result, typename _Source>
-struct __ct_abs<_Result, _Source, false>
-{
-  constexpr _LIBCUDACXX_INLINE_VISIBILITY _Result operator()(_Source __t) const noexcept
-  {
-    return __t;
-  }
-};
-
-template <class _Tp>
-constexpr _LIBCUDACXX_HIDDEN _Tp __gcd(_Tp __m, _Tp __n)
-{
-  static_assert((!is_signed<_Tp>::value), "");
-  return __n == 0 ? __m : _CUDA_VSTD::__gcd<_Tp>(__n, __m % __n);
-}
-
-template <class _Tp, class _Up>
-constexpr _LIBCUDACXX_INLINE_VISIBILITY common_type_t<_Tp, _Up> gcd(_Tp __m, _Up __n)
-{
-  static_assert((is_integral<_Tp>::value && is_integral<_Up>::value), "Arguments to gcd must be integer types");
-  static_assert((!is_same<__remove_cv_t<_Tp>, bool>::value), "First argument to gcd cannot be bool");
-  static_assert((!is_same<__remove_cv_t<_Up>, bool>::value), "Second argument to gcd cannot be bool");
-  using _Rp = common_type_t<_Tp, _Up>;
-  using _Wp = make_unsigned_t<_Rp>;
-  return static_cast<_Rp>(
-    _CUDA_VSTD::__gcd(static_cast<_Wp>(__ct_abs<_Rp, _Tp>()(__m)), static_cast<_Wp>(__ct_abs<_Rp, _Up>()(__n))));
-}
-
-template <class _Tp, class _Up>
-constexpr _LIBCUDACXX_INLINE_VISIBILITY common_type_t<_Tp, _Up> lcm(_Tp __m, _Up __n)
-{
-  static_assert((is_integral<_Tp>::value && is_integral<_Up>::value), "Arguments to lcm must be integer types");
-  static_assert((!is_same<__remove_cv_t<_Tp>, bool>::value), "First argument to lcm cannot be bool");
-  static_assert((!is_same<__remove_cv_t<_Up>, bool>::value), "Second argument to lcm cannot be bool");
-  if (__m == 0 || __n == 0)
-  {
-    return 0;
-  }
-
-  using _Rp  = common_type_t<_Tp, _Up>;
-  _Rp __val1 = __ct_abs<_Rp, _Tp>()(__m) / _CUDA_VSTD::gcd(__m, __n);
-  _Rp __val2 = __ct_abs<_Rp, _Up>()(__n);
-  _LIBCUDACXX_ASSERT((numeric_limits<_Rp>::max() / __val1 > __val2), "Overflow in lcm");
-  return __val1 * __val2;
-}
-
-#  endif /* _CCCL_STD_VER > 2014 */
-
 #  if _CCCL_STD_VER > 2017
 template <class _Tp>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr enable_if_t<
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool1.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool1.compile.fail.cpp
new file mode 100644
index 00000000000..7010950f0ba
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool1.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> gcd(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::gcd(false, 4);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool2.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool2.compile.fail.cpp
new file mode 100644
index 00000000000..1cc6bc10104
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool2.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> gcd(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::gcd(2, true);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool3.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool3.compile.fail.cpp
new file mode 100644
index 00000000000..98cbf36ef09
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool3.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> gcd(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::gcd<volatile bool, int>(false, 4);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool4.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool4.compile.fail.cpp
new file mode 100644
index 00000000000..a6e1240e794
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.bool4.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> gcd(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::gcd<int, const bool>(2, true);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.not_integral1.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.not_integral1.compile.fail.cpp
new file mode 100644
index 00000000000..72d99c08f70
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.not_integral1.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> gcd(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::gcd(2.0, 4);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.not_integral2.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.not_integral2.compile.fail.cpp
new file mode 100644
index 00000000000..73b68c4c11f
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.not_integral2.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> gcd(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::gcd(4, 6.0);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.pass.cpp
new file mode 100644
index 00000000000..c144e2cd488
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.gcd/gcd.pass.cpp
@@ -0,0 +1,127 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> gcd(_M __m, _N __n)
+
+#include <cuda/std/__numeric_>
+#include <cuda/std/cassert>
+#include <cuda/std/climits>
+#include <cuda/std/cstdint>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+struct TestCase
+{
+  int x;
+  int y;
+  int expect;
+};
+
+template <typename Input1, typename Input2, typename Output>
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool test0(int in1, int in2, int out)
+{
+  auto value1 = static_cast<Input1>(in1);
+  auto value2 = static_cast<Input2>(in2);
+  static_assert(cuda::std::is_same<Output, decltype(cuda::std::gcd(value1, value2))>::value, "");
+  static_assert(cuda::std::is_same<Output, decltype(cuda::std::gcd(value2, value1))>::value, "");
+  assert(static_cast<Output>(out) == cuda::std::gcd(value1, value2));
+  return true;
+}
+
+template <typename Input1, typename Input2 = Input1>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test()
+{
+  using S1                   = cuda::std::__make_signed_t<Input1>;
+  using S2                   = cuda::std::__make_signed_t<Input2>;
+  using U1                   = cuda::std::__make_signed_t<Input1>;
+  using U2                   = cuda::std::__make_signed_t<Input2>;
+  bool accumulate            = true;
+  constexpr TestCase Cases[] = {
+    {0, 0, 0}, {1, 0, 1}, {0, 1, 1}, {1, 1, 1}, {2, 3, 1}, {2, 4, 2}, {36, 17, 1}, {36, 18, 18}};
+  for (auto TC : Cases)
+  {
+    { // Test with two signed types
+      using Output = cuda::std::__common_type_t<S1, S2>;
+      accumulate &= test0<S1, S2, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<S1, S2, Output>(-TC.x, TC.y, TC.expect);
+      accumulate &= test0<S1, S2, Output>(TC.x, -TC.y, TC.expect);
+      accumulate &= test0<S1, S2, Output>(-TC.x, -TC.y, TC.expect);
+      accumulate &= test0<S2, S1, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<S2, S1, Output>(-TC.x, TC.y, TC.expect);
+      accumulate &= test0<S2, S1, Output>(TC.x, -TC.y, TC.expect);
+      accumulate &= test0<S2, S1, Output>(-TC.x, -TC.y, TC.expect);
+    }
+    { // test with two unsigned types
+      using Output = cuda::std::__common_type_t<U1, U2>;
+      accumulate &= test0<U1, U2, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<U2, U1, Output>(TC.x, TC.y, TC.expect);
+    }
+    { // Test with mixed signs
+      using Output = cuda::std::__common_type_t<S1, U2>;
+      accumulate &= test0<S1, U2, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<U2, S1, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<S1, U2, Output>(-TC.x, TC.y, TC.expect);
+      accumulate &= test0<U2, S1, Output>(TC.x, -TC.y, TC.expect);
+    }
+    { // Test with mixed signs
+      using Output = cuda::std::__common_type_t<S2, U1>;
+      accumulate &= test0<S2, U1, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<U1, S2, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<S2, U1, Output>(-TC.x, TC.y, TC.expect);
+      accumulate &= test0<U1, S2, Output>(TC.x, -TC.y, TC.expect);
+    }
+  }
+  assert(accumulate);
+}
+
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool test()
+{
+  test<signed char>();
+  test<short>();
+  test<int>();
+  test<long>();
+  test<long long>();
+
+  test<cuda::std::int8_t>();
+  test<cuda::std::int16_t>();
+  test<cuda::std::int32_t>();
+  test<cuda::std::int64_t>();
+
+  test<signed char, int>();
+  test<int, signed char>();
+  test<short, int>();
+  test<int, short>();
+  test<int, long>();
+  test<long, int>();
+  test<int, long long>();
+  test<long long, int>();
+
+  return true;
+}
+
+int main(int argc, char**)
+{
+  test();
+#if TEST_STD_VER >= 2014
+  static_assert(test(), "");
+#endif // TEST_STD_VER >= 2014
+
+  //  LWG#2837
+  {
+    auto res = cuda::std::gcd(static_cast<cuda::std::int64_t>(1234), INT32_MIN);
+    static_assert(cuda::std::is_same<decltype(res), cuda::std::int64_t>::value, "");
+    assert(res == 2);
+  }
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool1.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool1.compile.fail.cpp
new file mode 100644
index 00000000000..b351d42d543
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool1.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> lcm(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::lcm(false, 4);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool2.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool2.compile.fail.cpp
new file mode 100644
index 00000000000..987fcefa940
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool2.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> lcm(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::lcm(2, true);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool3.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool3.compile.fail.cpp
new file mode 100644
index 00000000000..fd96ea058b3
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool3.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> lcm(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::lcm<volatile bool, int>(false, 4);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool4.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool4.compile.fail.cpp
new file mode 100644
index 00000000000..9c9c6b5f532
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.bool4.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> lcm(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::lcm<int, const bool>(2, true);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.not_integral1.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.not_integral1.compile.fail.cpp
new file mode 100644
index 00000000000..3d64f03bee1
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.not_integral1.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> lcm(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::lcm(2.0, 4);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.not_integral2.compile.fail.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.not_integral2.compile.fail.cpp
new file mode 100644
index 00000000000..e07de81b886
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.not_integral2.compile.fail.cpp
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr common_type_t<_M,_N> lcm(_M __m, _N __n)
+
+// Remarks: If either M or N is not an integer type,
+// or if either is (a possibly cv-qualified) bool, the program is ill-formed.
+
+#include <cuda/std/__numeric_>
+
+int main(int, char**)
+{
+  cuda::std::lcm(4, 6.0);
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.pass.cpp
new file mode 100644
index 00000000000..ce1b98d471a
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.lcm/lcm.pass.cpp
@@ -0,0 +1,127 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template<class _M, class _N>
+// constexpr __common_type_t<_M,_N> lcm(_M __m, _N __n)
+
+#include <cuda/std/__numeric_>
+#include <cuda/std/cassert>
+#include <cuda/std/climits>
+#include <cuda/std/cstdint>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+struct TestCases
+{
+  int x;
+  int y;
+  int expect;
+};
+
+template <typename Input1, typename Input2, typename Output>
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool test0(int in1, int in2, int out)
+{
+  auto value1 = static_cast<Input1>(in1);
+  auto value2 = static_cast<Input2>(in2);
+  static_assert(cuda::std::is_same<Output, decltype(cuda::std::lcm(value1, value2))>::value, "");
+  static_assert(cuda::std::is_same<Output, decltype(cuda::std::lcm(value2, value1))>::value, "");
+  assert(static_cast<Output>(out) == cuda::std::lcm(value1, value2));
+  return true;
+}
+
+template <typename Input1, typename Input2 = Input1>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test()
+{
+  using S1                    = cuda::std::__make_signed_t<Input1>;
+  using S2                    = cuda::std::__make_signed_t<Input2>;
+  using U1                    = cuda::std::__make_signed_t<Input1>;
+  using U2                    = cuda::std::__make_signed_t<Input2>;
+  bool accumulate             = true;
+  constexpr TestCases Cases[] = {
+    {0, 0, 0}, {1, 0, 0}, {0, 1, 0}, {1, 1, 1}, {2, 3, 6}, {2, 4, 4}, {3, 17, 51}, {36, 18, 36}};
+  for (auto TC : Cases)
+  {
+    { // Test with two signed types
+      using Output = cuda::std::__common_type_t<S1, S2>;
+      accumulate &= test0<S1, S2, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<S1, S2, Output>(-TC.x, TC.y, TC.expect);
+      accumulate &= test0<S1, S2, Output>(TC.x, -TC.y, TC.expect);
+      accumulate &= test0<S1, S2, Output>(-TC.x, -TC.y, TC.expect);
+      accumulate &= test0<S2, S1, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<S2, S1, Output>(-TC.x, TC.y, TC.expect);
+      accumulate &= test0<S2, S1, Output>(TC.x, -TC.y, TC.expect);
+      accumulate &= test0<S2, S1, Output>(-TC.x, -TC.y, TC.expect);
+    }
+    { // test with two unsigned types
+      using Output = cuda::std::__common_type_t<U1, U2>;
+      accumulate &= test0<U1, U2, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<U2, U1, Output>(TC.x, TC.y, TC.expect);
+    }
+    { // Test with mixed signs
+      using Output = cuda::std::__common_type_t<S1, U2>;
+      accumulate &= test0<S1, U2, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<U2, S1, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<S1, U2, Output>(-TC.x, TC.y, TC.expect);
+      accumulate &= test0<U2, S1, Output>(TC.x, -TC.y, TC.expect);
+    }
+    { // Test with mixed signs
+      using Output = cuda::std::__common_type_t<S2, U1>;
+      accumulate &= test0<S2, U1, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<U1, S2, Output>(TC.x, TC.y, TC.expect);
+      accumulate &= test0<S2, U1, Output>(-TC.x, TC.y, TC.expect);
+      accumulate &= test0<U1, S2, Output>(TC.x, -TC.y, TC.expect);
+    }
+  }
+  assert(accumulate);
+}
+
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool test()
+{
+  test<signed char>();
+  test<short>();
+  test<int>();
+  test<long>();
+  test<long long>();
+
+  test<cuda::std::int8_t>();
+  test<cuda::std::int16_t>();
+  test<cuda::std::int32_t>();
+  test<cuda::std::int64_t>();
+
+  test<signed char, int>();
+  test<int, signed char>();
+  test<short, int>();
+  test<int, short>();
+  test<int, long>();
+  test<long, int>();
+  test<int, long long>();
+  test<long long, int>();
+
+  return true;
+}
+
+int main(int argc, char**)
+{
+  test();
+#if TEST_STD_VER >= 2014
+  static_assert(test(), "");
+#endif // TEST_STD_VER >= 2014
+
+  //  LWG#2837
+  {
+    auto res1 = cuda::std::lcm(static_cast<cuda::std::int64_t>(1234), INT32_MIN);
+    TEST_IGNORE_NODISCARD cuda::std::lcm(INT_MIN, 2UL); // this used to trigger UBSAN
+    static_assert(cuda::std::is_same<decltype(res1), cuda::std::int64_t>::value, "");
+    assert(res1 == 1324997410816LL);
+  }
+
+  return 0;
+}