Add __int128 and __float128 detection macros

docs/cccl_development/macro.rst

@@ -184,18 +184,20 @@ Non-standard Types Support
 +------------------------------+-------------------------------------------------------------------------------------------------------------------------------+
 | ``_CCCL_HAS_NVBF16``         | ``__nv_bfloat16/__nv_bfloat162`` data types are supported and enabled.  Prefer over ``__CUDA_BF16_TYPES_EXIST__``             |
 +------------------------------+-------------------------------------------------------------------------------------------------------------------------------+
-| ``_CCCL_HAS_FP8()``          | ``__nv_fp8_e5m2/__nv_fp8_e4m3/__nv_fp8_e8m0`` data types are supported and enabled.  Prefer over ``__CUDA_FP8_TYPES_EXIST__`` |
+| ``_CCCL_HAS_NVFP8()``        | ``__nv_fp8_e5m2/__nv_fp8_e4m3/__nv_fp8_e8m0`` data types are supported and enabled.  Prefer over ``__CUDA_FP8_TYPES_EXIST__`` |
 +------------------------------+-------------------------------------------------------------------------------------------------------------------------------+
 
-+------------------------------+-------------------------------------------------------------------------+
-| ``_CCCL_DISABLE_INT128``     | Disable ``__int128/__uint128_t`` support                                |
-+------------------------------+-------------------------------------------------------------------------+
-| ``_CCCL_DISABLE_FLOAT128``   | Disable ``__float128`` support                                          |
-+------------------------------+-------------------------------------------------------------------------+
-| ``_LIBCUDACXX_HAS_NVFP16``   | ``__half/__half2`` host/device are supported  (CUDA 12.2+)              |
-+------------------------------+-------------------------------------------------------------------------+
-| ``_LIBCUDACXX_HAS_NVBF16``   | ``__nv_bfloat16/__nv_bfloat162`` host/device are supported (CUDA 12.2+) |
-+------------------------------+-------------------------------------------------------------------------+
++---------------------------------+-------------------------------------------------------------------------+
+| ``_CCCL_DISABLE_NVFP8_SUPPORT`` | Disable ``__nv_fp8_e5m2/__nv_fp8_e4m3/__nv_fp8_e8m0`` support           |
++---------------------------------+-------------------------------------------------------------------------+
+| ``_CCCL_DISABLE_INT128``        | Disable ``__int128/__uint128_t`` support                                |
++---------------------------------+-------------------------------------------------------------------------+
+| ``_CCCL_DISABLE_FLOAT128``      | Disable ``__float128`` support                                          |
++---------------------------------+-------------------------------------------------------------------------+
+| ``_LIBCUDACXX_HAS_NVFP16``      | ``__half/__half2`` host/device are supported  (CUDA 12.2+)              |
++---------------------------------+-------------------------------------------------------------------------+
+| ``_LIBCUDACXX_HAS_NVBF16``      | ``__nv_bfloat16/__nv_bfloat162`` host/device are supported (CUDA 12.2+) |
++---------------------------------+-------------------------------------------------------------------------+
 
 ----
 

libcudacxx/include/cuda/__cccl_config

@@ -21,7 +21,7 @@
 #include <cuda/std/__cccl/dialect.h> // IWYU pragma: export
 #include <cuda/std/__cccl/exceptions.h> // IWYU pragma: export
 #include <cuda/std/__cccl/execution_space.h> // IWYU pragma: export
-#include <cuda/std/__cccl/extended_floating_point.h> // IWYU pragma: export
+#include <cuda/std/__cccl/extended_data_types.h> // IWYU pragma: export
 #include <cuda/std/__cccl/os.h> // IWYU pragma: export
 #include <cuda/std/__cccl/preprocessor.h> // IWYU pragma: export
 #include <cuda/std/__cccl/ptx_isa.h> // IWYU pragma: export

libcudacxx/include/cuda/std/__cccl/extended_floating_point.h → libcudacxx/include/cuda/std/__cccl/extended_data_types.h

@@ -4,12 +4,12 @@
 // 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.
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef __CCCL_EXTENDED_FLOATING_POINT_H
-#define __CCCL_EXTENDED_FLOATING_POINT_H
+#ifndef __CCCL_EXTENDED_DATA_TYPES_H
+#define __CCCL_EXTENDED_DATA_TYPES_H
 
 #include <cuda/std/__cccl/compiler.h>
 #include <cuda/std/__cccl/system_header.h>
@@ -23,8 +23,21 @@
 #endif // no system header
 
 #include <cuda/std/__cccl/diagnostic.h>
+#include <cuda/std/__cccl/os.h>
 #include <cuda/std/__cccl/preprocessor.h>
 
+#if !defined(_CCCL_DISABLE_INT128)
+#  if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && _CCCL_OS(LINUX)
+#    define _CCCL_HAS_INT128() 1
+#  elif defined(__SIZEOF_INT128__) && _CCCL_OS(LINUX)
+#    define _CCCL_HAS_INT128() 1
+#  else
+#    define _CCCL_HAS_INT128() 0
+#  endif
+#else
+#  define _CCCL_HAS_INT128() 0
+#endif // !_CCCL_DISABLE_INT128
+
 #if !defined(_CCCL_HAS_NVFP16)
 #  if _CCCL_HAS_INCLUDE(<cuda_fp16.h>) && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \
                         && !defined(CCCL_DISABLE_FP16_SUPPORT)
@@ -44,9 +57,28 @@
 #    define _CCCL_HAS_NVFP8() 1
 #  else
 #    define _CCCL_HAS_NVFP8() 0
-#  endif // _CCCL_HAS_INCLUDE(<cuda_fp8.h>)
+#  endif // _CCCL_HAS_INCLUDE(<cuda_fp8.h>) && defined(_CCCL_HAS_NVFP16) && defined(_CCCL_HAS_NVBF16)
 #else
 #  define _CCCL_HAS_NVFP8() 0
 #endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT)
 
-#endif // __CCCL_EXTENDED_FLOATING_POINT_H
+#if !defined(_CCCL_DISABLE_FLOAT128)
+#  if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_FLOAT128__) && _CCCL_OS(LINUX)
+#    if !defined(__CUDA_ARCH__) || (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1000)
+#      define _CCCL_HAS_FLOAT128() 1
+#    else
+#      define _CCCL_HAS_FLOAT128() 0
+#    endif
+// NVC++ support float128 only in host code
+#  elif (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) && !_CCCL_CUDA_COMPILER(NVHPC)
+#    if !defined(__CUDA_ARCH__) || (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1000)
+#      define _CCCL_HAS_FLOAT128() 1
+#    else
+#      define _CCCL_HAS_FLOAT128() 0
+#    endif
+#  else
+#    define _CCCL_HAS_FLOAT128() 0
+#  endif
+#endif // !defined(_CCCL_DISABLE_FLOAT128)
+
+#endif // __CCCL_EXTENDED_DATA_TYPES_H

libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.fail.cpp → libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.bf16.fail.cpp

@@ -6,37 +6,19 @@
 // SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
 //
 //===----------------------------------------------------------------------===//
-#include <cuda/std/__cccl/extended_floating_point.h>
+#include <cuda/std/__cccl/extended_data_types.h>
 
 #include "test_macros.h"
 
-#if !_CCCL_HAS_NVFP8()
-#  include <cuda_fp8.h>
-#endif
-#if !defined(_CCCL_HAS_NVFP16)
-#  include <cuda_fp16.h>
-#endif
 #if !defined(_CCCL_HAS_NVBF16)
 #  include <cuda_bf16.h>
 #endif
 
 int main(int, char**)
 {
-#if !_CCCL_HAS_NVFP8()
-  auto x = __nv_fp8_e4m3(1.0f);
-  unused(x);
-#else
-  static_assert(false);
-#endif
-#if !defined(_CCCL_HAS_NVFP16)
-  auto y = __half(1.0f);
-  unused(y);
-#else
-  static_assert(false);
-#endif
 #if !defined(_CCCL_HAS_NVBF16)
-  auto z = __nv_bfloat16(1.0f);
-  unused(z);
+  auto x3 = __nv_bfloat16(1.0f);
+  unused(x3);
 #else
   static_assert(false);
 #endif

libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp

@@ -0,0 +1,23 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+#include <cuda/std/__cccl/extended_data_types.h>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+#if !_CCCL_HAS_FLOAT128()
+  __float128 x4 = __float128(3.14) + __float128(3.14);
+  unused(x4);
+#else
+  static_assert(false);
+#endif
+  return 0;
+}

libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp16.fail.cpp

@@ -0,0 +1,26 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+#include <cuda/std/__cccl/extended_data_types.h>
+
+#include "test_macros.h"
+
+#if !defined(_CCCL_HAS_NVFP16)
+#  include <cuda_fp16.h>
+#endif
+
+int main(int, char**)
+{
+#if !defined(_CCCL_HAS_NVFP16)
+  auto x2 = __half(1.0f);
+  unused(x2);
+#else
+  static_assert(false);
+#endif
+  return 0;
+}

libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp8.fail.cpp

@@ -0,0 +1,26 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+#include <cuda/std/__cccl/extended_data_types.h>
+
+#include "test_macros.h"
+
+#if !_CCCL_HAS_NVFP8()
+#  include <cuda_fp8.h>
+#endif
+
+int main(int, char**)
+{
+#if !_CCCL_HAS_NVFP8()
+  auto x1 = __nv_fp8_e4m3(1.0f);
+  unused(x1);
+#else
+  static_assert(false);
+#endif
+  return 0;
+}

libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.i128.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
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+#include <cuda/std/__cccl/extended_data_types.h>
+
+#include "test_macros.h"
+
+int main(int, char**)
+{
+#if !_CCCL_HAS_INT128()
+  __int128 x    = __int128(123456789123) + __int128(123456789123);
+  __uint128_t y = __uint128_t(123456789123) + __uint128_t(123456789123);
+  unused(x);
+  unused(y);
+#else
+  static_assert(false);
+#endif
+  return 0;
+}

libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.pass.cpp → libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp

@@ -6,7 +6,7 @@
 // SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
 //
 //===----------------------------------------------------------------------===//
-#include <cuda/std/__cccl/extended_floating_point.h>
+#include <cuda/std/__cccl/extended_data_types.h>
 
 #include "test_macros.h"
 
@@ -22,17 +22,27 @@
 
 int main(int, char**)
 {
+#if _CCCL_HAS_INT128()
+  auto x1 = __int128(123456789123) + __int128(123456789123);
+  auto y1 = __uint128_t(123456789123) + __uint128_t(123456789123);
+  unused(x1);
+  unused(y1);
+#endif
 #if _CCCL_HAS_NVFP8()
-  auto x = __nv_fp8_e4m3(1.0f);
-  unused(x);
+  auto x2 = __nv_fp8_e4m3(1.0f);
+  unused(x2);
 #endif
 #if defined(_CCCL_HAS_NVFP16)
-  auto y = __half(1.0f);
-  unused(y);
+  auto x3 = __half(1.0f);
+  unused(x3);
 #endif
 #if defined(_CCCL_HAS_NVBF16)
-  auto z = __nv_bfloat16(1.0f);
-  unused(z);
+  auto x4 = __nv_bfloat16(1.0f);
+  unused(x4);
+#endif
+#if _CCCL_HAS_FLOAT128()
+  __float128 x5 = __float128(3.14) + __float128(3.14);
+  unused(x5);
 #endif
   return 0;
 }