[Backport 2.8] Add extended data type macro identification (#3586)

libcudacxx/include/cuda/__cccl_config

@@ -20,7 +20,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/preprocessor.h> // IWYU pragma: export
 #include <cuda/std/__cccl/ptx_isa.h> // IWYU pragma: export
 #include <cuda/std/__cccl/rtti.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/include/cuda/std/__cccl/os.h

@@ -0,0 +1,48 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CCCL_OS_H
+#define __CCCL_OS_H
+
+// The header provides the following macros to determine the host architecture:
+//
+// _CCCL_OS(WINDOWS)
+// _CCCL_OS(LINUX)
+// _CCCL_OS(ANDROID)
+// _CCCL_OS(QNX)
+
+// Determine the host compiler and its version
+#if defined(_WIN32) || defined(_WIN64) /* _WIN64 for NVRTC */
+#  define _CCCL_OS_WINDOWS_() 1
+#else
+#  define _CCCL_OS_WINDOWS_() 0
+#endif
+
+#if defined(__linux__) || defined(__LP64__) /* __LP64__ for NVRTC */
+#  define _CCCL_OS_LINUX_() 1
+#else
+#  define _CCCL_OS_LINUX_() 0
+#endif
+
+#if defined(__ANDROID__)
+#  define _CCCL_OS_ANDROID_() 1
+#else
+#  define _CCCL_OS_ANDROID_() 0
+#endif
+
+#if defined(__QNX__) || defined(__QNXNTO__)
+#  define _CCCL_OS_QNX_() 1
+#else
+#  define _CCCL_OS_QNX_() 0
+#endif
+
+#define _CCCL_OS(...) _CCCL_OS_##__VA_ARGS__##_()
+
+#endif // __CCCL_OS_H

libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.bf16.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_NVBF16)
+#  include <cuda_bf16.h>
+#endif
+
+int main(int, char**)
+{
+#if !defined(_CCCL_HAS_NVBF16)
+  auto x3 = __nv_bfloat16(1.0f);
+  unused(x3);
+#else
+  static_assert(false);
+#endif
+  return 0;
+}

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_data_types.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) 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
+#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_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 x2 = __nv_fp8_e4m3(1.0f);
+  unused(x2);
+#endif
+#if defined(_CCCL_HAS_NVFP16)
+  auto x3 = __half(1.0f);
+  unused(x3);
+#endif
+#if defined(_CCCL_HAS_NVBF16)
+  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;
+}

libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.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"
 

libcudacxx/test/libcudacxx/libcxx/macros/os.compile.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) 2023 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/__cccl/os.h>
+
+#if !defined(__CUDACC_RTC__)
+#  if _CCCL_OS(WINDOWS)
+#    include <windows.h>
+#  endif
+
+#  if _CCCL_OS(LINUX)
+#    include <unistd.h>
+#  endif
+
+#  if _CCCL_OS(ANDROID)
+#    include <android/api-level.h>
+#  endif
+
+#  if _CCCL_OS(QNX)
+#    include <qnx.h>
+#  endif
+#endif
+
+int main(int, char**)
+{
+  static_assert(_CCCL_OS(WINDOWS) + _CCCL_OS(LINUX) == 1, "");
+#if _CCCL_OS(ANDROID) || _CCCL_OS(QNX)
+  static_assert(_CCCL_OS(LINUX) == 1, "");
+  static_assert(_CCCL_OS(ANDROID) + _CCCL_OS(QNX) == 1, "");
+#endif
+#if _CCCL_OS(LINUX)
+  static_assert(_CCCL_OS(WINDOWS) == 0, "");
+#endif
+#if _CCCL_OS(WINDOWS)
+  static_assert(_CCCL_OS(ANDROID) + _CCCL_OS(QNX) + _CCCL_OS(LINUX) == 0, "");
+#endif
+  return 0;
+}