From 0808605a6dd3433d182dc67174002fd85f17343b Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 15 Jan 2025 15:29:30 -0800 Subject: [PATCH 01/29] adding int128 and float128 detection macros --- libcudacxx/include/cuda/__cccl_config | 1 + .../include/cuda/std/__cccl/128bit_types.h | 40 +++++++++++++++++++ .../std/cccl/128bit_types.compile.pass.cpp | 27 +++++++++++++ 3 files changed, 68 insertions(+) create mode 100644 libcudacxx/include/cuda/std/__cccl/128bit_types.h create mode 100644 libcudacxx/test/libcudacxx/std/cccl/128bit_types.compile.pass.cpp diff --git a/libcudacxx/include/cuda/__cccl_config b/libcudacxx/include/cuda/__cccl_config index fc28d63ace1e..08bf8a481813 100644 --- a/libcudacxx/include/cuda/__cccl_config +++ b/libcudacxx/include/cuda/__cccl_config @@ -11,6 +11,7 @@ #ifndef _CUDA__CCCL_CONFIG #define _CUDA__CCCL_CONFIG +#include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export diff --git a/libcudacxx/include/cuda/std/__cccl/128bit_types.h b/libcudacxx/include/cuda/std/__cccl/128bit_types.h new file mode 100644 index 000000000000..60ffcfaee9f9 --- /dev/null +++ b/libcudacxx/include/cuda/std/__cccl/128bit_types.h @@ -0,0 +1,40 @@ +//===----------------------------------------------------------------------===// +// +// 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 __CCCL_128BIT_TYPES_H +#define __CCCL_128BIT_TYPES_H + +#include +#include + +#ifndef _CCCL_DISABLE_INT128 +# if defined(__SIZEOF_INT128__) +# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) +# define _CCCL_HAS_INT128() 1 +# elif (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !_CCCL_OS(WINDOWS) +# define _CCCL_HAS_INT128() 1 +# else +# define _CCCL_HAS_INT128() 0 +# endif +# endif // defined(__SIZEOF_INT128__) +#else +# define _CCCL_HAS_INT128() 0 +#endif // #ifndef _CCCL_DISABLE_INT128 + +#ifndef _CCCL_DISABLE_FLOAT128 +# if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) \ + && (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !defined(__CUDA_ARCH__) +# define _CCCL_HAS_FLOAT128() 1 +# else +# define _CCCL_HAS_FLOAT128() 0 +# endif +#endif // #ifndef _CCCL_DISABLE_FLOAT128 + +#endif // __CCCL_128BIT_TYPES_H diff --git a/libcudacxx/test/libcudacxx/std/cccl/128bit_types.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/cccl/128bit_types.compile.pass.cpp new file mode 100644 index 000000000000..3a1a776c76b6 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/cccl/128bit_types.compile.pass.cpp @@ -0,0 +1,27 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#include "test_macros.h" + +int main(int, char**) +{ +#if _CCCL_HAS_INT128() + auto x = __int128(123456789123) + __int128(123456789123); + auto y = __uint128_t(123456789123) + __uint128_t(123456789123); + unused(x); + unused(y); +#endif +#if _CCCL_HAS_FLOAT128() + auto z = __float128(3.14) + __float128(3.14); + unused(z); +#endif + return 0; +} From 24e5f782cb48adef845b613cee9e9fed2b97ad65 Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 15 Jan 2025 15:32:05 -0800 Subject: [PATCH 02/29] document new macros --- docs/cccl_development/macro.rst | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/docs/cccl_development/macro.rst b/docs/cccl_development/macro.rst index 6bf1b0b67ab3..e37a8ee00a82 100644 --- a/docs/cccl_development/macro.rst +++ b/docs/cccl_development/macro.rst @@ -87,6 +87,14 @@ The following macros are used to check the target architecture. They comply with | ``_CCCL_ARCH(X86_64)`` | X86 64-bit | +-------------------------+-------------------------------------+ +**128-bit Types Macros**: + ++--------------------------+-------------------------------------+ +| ``_CCCL_HAS_INT128()`` | Support for 128-bit integer | ++--------------------------+-------------------------------------+ +| ``_CCCL_HAS_FLOAT128()`` | Support for 128-bit floating-point | ++--------------------------+-------------------------------------+ + ---- OS Macros From 12664e3597a077eec5f95a0ba1bd7edaafc3cf3c Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 15 Jan 2025 17:19:10 -0800 Subject: [PATCH 03/29] fix missing preprocessing paths --- libcudacxx/include/cuda/std/__cccl/128bit_types.h | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/128bit_types.h b/libcudacxx/include/cuda/std/__cccl/128bit_types.h index 60ffcfaee9f9..99a839d99e9f 100644 --- a/libcudacxx/include/cuda/std/__cccl/128bit_types.h +++ b/libcudacxx/include/cuda/std/__cccl/128bit_types.h @@ -15,15 +15,17 @@ #include #ifndef _CCCL_DISABLE_INT128 -# if defined(__SIZEOF_INT128__) -# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) +# if defined(__SIZEOF_INT128__) // defined(__SIZEOF_INT128__) vvvv +# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && !_CCCL_OS(WINDOWS) # define _CCCL_HAS_INT128() 1 # elif (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !_CCCL_OS(WINDOWS) # define _CCCL_HAS_INT128() 1 # else # define _CCCL_HAS_INT128() 0 # endif -# endif // defined(__SIZEOF_INT128__) +# else // defined(__SIZEOF_INT128__) ^^^^ / !defined(__SIZEOF_INT128__) vvvv +# define _CCCL_HAS_INT128() 0 +# endif // !defined(__SIZEOF_INT128__) ^^^^ #else # define _CCCL_HAS_INT128() 0 #endif // #ifndef _CCCL_DISABLE_INT128 @@ -35,6 +37,8 @@ # else # define _CCCL_HAS_FLOAT128() 0 # endif +#else +# define _CCCL_HAS_INT128() 0 #endif // #ifndef _CCCL_DISABLE_FLOAT128 #endif // __CCCL_128BIT_TYPES_H From ab8f299fa9a7447183faa6d4b638e7d271b8c449 Mon Sep 17 00:00:00 2001 From: fbusato Date: Thu, 16 Jan 2025 10:17:50 -0800 Subject: [PATCH 04/29] improve documentation --- docs/cccl_development/macro.rst | 59 +++++++++++++++++---------------- 1 file changed, 31 insertions(+), 28 deletions(-) diff --git a/docs/cccl_development/macro.rst b/docs/cccl_development/macro.rst index e37a8ee00a82..5ec75811cd1f 100644 --- a/docs/cccl_development/macro.rst +++ b/docs/cccl_development/macro.rst @@ -68,11 +68,11 @@ The ``_CCCL_CUDA_COMPILER`` function-like macro can also be used to check the ve **PTX macros**: -+-------------------------+-------------------------------------------------------------------------------------------------------------------+ -| ``_CCCL_PTX_ARCH`` | Alias of ``__CUDA_ARCH__`` with value equal to 0 if cuda compiler is not available | -+-------------------------+-------------------------------------------------------------------------------------------------------------------+ -| ``__cccl_ptx_isa`` | PTX ISA version available with the current CUDA compiler, e.g. PTX ISA 8.4 (``840``) is available from CUDA 12.4 | -+-------------------------+-------------------------------------------------------------------------------------------------------------------+ ++--------------------+-------------------------------------------------------------------------------------------------------------------+ +| ``_CCCL_PTX_ARCH`` | Alias of ``__CUDA_ARCH__`` with value equal to 0 if cuda compiler is not available | ++--------------------+-------------------------------------------------------------------------------------------------------------------+ +| ``__cccl_ptx_isa`` | PTX ISA version available with the current CUDA compiler, e.g. PTX ISA 8.4 (``840``) is available from CUDA 12.4 | ++--------------------+-------------------------------------------------------------------------------------------------------------------+ ---- @@ -87,14 +87,6 @@ The following macros are used to check the target architecture. They comply with | ``_CCCL_ARCH(X86_64)`` | X86 64-bit | +-------------------------+-------------------------------------+ -**128-bit Types Macros**: - -+--------------------------+-------------------------------------+ -| ``_CCCL_HAS_INT128()`` | Support for 128-bit integer | -+--------------------------+-------------------------------------+ -| ``_CCCL_HAS_FLOAT128()`` | Support for 128-bit floating-point | -+--------------------------+-------------------------------------+ - ---- OS Macros @@ -167,10 +159,8 @@ Usage example: ---- -CUDA Extension Macros ---------------------- - -**CUDA attributes**: +CUDA attributes +--------------- +------------------------------+----------------------------------------------------------+ | ``_CCCL_GRID_CONSTANT`` | Grid constant kernel parameter | @@ -178,19 +168,30 @@ CUDA Extension Macros | ``_CCCL_GLOBAL_CONSTANT`` | Host/device global scope constant (``inline constexpr``) | +------------------------------+----------------------------------------------------------+ -**Extended floating-point types**: +---- -+------------------------------+-----------------------------------------------------------------------------------------------------------------+ -| ``_CCCL_HAS_NVFP16`` | `__half/__half2` data types are supported and enabled. Prefer over ``__CUDA_FP16_TYPES_EXIST__`` | -+------------------------------+-----------------------------------------------------------------------------------------------------------------+ -| ``_CCCL_HAS_NVBF16`` | `__nv_bfloat16/__nv_bfloat162` data types are supported and enabled. Prefer over ``__CUDA_BF16_TYPES_EXIST__`` | -+------------------------------+-----------------------------------------------------------------------------------------------------------------+ +Non-standard Types Support +-------------------------- -+------------------------------+----------------------------------------------------------------+ -| ``_LIBCUDACXX_HAS_NVFP16`` | `__half/__half2` host/device support (CUDA 12.2) | -+------------------------------+----------------------------------------------------------------+ -| ``_LIBCUDACXX_HAS_NVBF16`` | `__nv_bfloat16/__nv_bfloat162` host/device support (CUDA 12.2) | -+------------------------------+----------------------------------------------------------------+ ++------------------------------+-------------------------------------------------------------------------------------------------------------------+ +| ``_CCCL_HAS_INT128()`` | ``__int128`` and ``__uint128_t`` for 128-bit integer are supported and enabled | ++------------------------------+-------------------------------------------------------------------------------------------------------------------+ +| ``_CCCL_HAS_FLOAT128()`` | ``__float128`` for 128-bit floating-point are supported and enabled | ++------------------------------+-------------------------------------------------------------------------------------------------------------------+ +| ``_CCCL_HAS_NVFP16`` | ``__half/__half2`` data types are supported and enabled. Prefer over ``__CUDA_FP16_TYPES_EXIST__`` | ++------------------------------+-------------------------------------------------------------------------------------------------------------------+ +| ``_CCCL_HAS_NVBF16`` | ``__nv_bfloat16/__nv_bfloat162`` data types are supported and enabled. Prefer over ``__CUDA_BF16_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+) | ++------------------------------+-------------------------------------------------------------------------+ ---- @@ -252,6 +253,8 @@ Usage example: **Portable attributes**: ++----------------------------------+------------------------------------------------------------------------------+ +| ``_CCCL_ASSUME()`` | Portable ``[[assume]]`` attribute (before C++23) | +----------------------------------+------------------------------------------------------------------------------+ | ``_CCCL_FALLTHROUGH()`` | Portable ``[[fallthrough]]`` attribute (before C++17) | +----------------------------------+------------------------------------------------------------------------------+ From e1b01896a998686638200c11d74ac14c721fe3c5 Mon Sep 17 00:00:00 2001 From: fbusato Date: Thu, 16 Jan 2025 10:29:49 -0800 Subject: [PATCH 05/29] move macro tests in libcxx directory --- .../{std/cccl => libcxx/macros}/128bit_types.compile.pass.cpp | 0 .../{std/cccl => libcxx/macros}/architecture.compile.pass.cpp | 0 .../libcudacxx/{std/cccl => libcxx/macros}/os.compile.pass.cpp | 0 3 files changed, 0 insertions(+), 0 deletions(-) rename libcudacxx/test/libcudacxx/{std/cccl => libcxx/macros}/128bit_types.compile.pass.cpp (100%) rename libcudacxx/test/libcudacxx/{std/cccl => libcxx/macros}/architecture.compile.pass.cpp (100%) rename libcudacxx/test/libcudacxx/{std/cccl => libcxx/macros}/os.compile.pass.cpp (100%) diff --git a/libcudacxx/test/libcudacxx/std/cccl/128bit_types.compile.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp similarity index 100% rename from libcudacxx/test/libcudacxx/std/cccl/128bit_types.compile.pass.cpp rename to libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp diff --git a/libcudacxx/test/libcudacxx/std/cccl/architecture.compile.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/architecture.compile.pass.cpp similarity index 100% rename from libcudacxx/test/libcudacxx/std/cccl/architecture.compile.pass.cpp rename to libcudacxx/test/libcudacxx/libcxx/macros/architecture.compile.pass.cpp diff --git a/libcudacxx/test/libcudacxx/std/cccl/os.compile.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/os.compile.pass.cpp similarity index 100% rename from libcudacxx/test/libcudacxx/std/cccl/os.compile.pass.cpp rename to libcudacxx/test/libcudacxx/libcxx/macros/os.compile.pass.cpp From 377c7f69af65208982d5048d6a1e9a4c666fa273 Mon Sep 17 00:00:00 2001 From: fbusato Date: Fri, 17 Jan 2025 11:06:03 -0800 Subject: [PATCH 06/29] add negative tests --- .../include/cuda/std/__cccl/128bit_types.h | 11 +++---- .../macros/128bit_types.compile.fail.cpp | 30 +++++++++++++++++++ 2 files changed, 36 insertions(+), 5 deletions(-) create mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp diff --git a/libcudacxx/include/cuda/std/__cccl/128bit_types.h b/libcudacxx/include/cuda/std/__cccl/128bit_types.h index 99a839d99e9f..95ebe4a37c2a 100644 --- a/libcudacxx/include/cuda/std/__cccl/128bit_types.h +++ b/libcudacxx/include/cuda/std/__cccl/128bit_types.h @@ -16,9 +16,9 @@ #ifndef _CCCL_DISABLE_INT128 # if defined(__SIZEOF_INT128__) // defined(__SIZEOF_INT128__) vvvv -# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && !_CCCL_OS(WINDOWS) +# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && _CCCL_OS(LINUX) # define _CCCL_HAS_INT128() 1 -# elif (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !_CCCL_OS(WINDOWS) +# elif (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && _CCCL_OS(LINUX) # define _CCCL_HAS_INT128() 1 # else # define _CCCL_HAS_INT128() 0 @@ -31,14 +31,15 @@ #endif // #ifndef _CCCL_DISABLE_INT128 #ifndef _CCCL_DISABLE_FLOAT128 -# if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) \ - && (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !defined(__CUDA_ARCH__) +# if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) \ + && (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !defined(__CUDA_ARCH__) \ + && !_CCCL_COMPILER(NVRTC) # define _CCCL_HAS_FLOAT128() 1 # else # define _CCCL_HAS_FLOAT128() 0 # endif #else -# define _CCCL_HAS_INT128() 0 +# define _CCCL_HAS_FLOAT128() 0 #endif // #ifndef _CCCL_DISABLE_FLOAT128 #endif // __CCCL_128BIT_TYPES_H diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp new file mode 100644 index 000000000000..d7dff72444f0 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp @@ -0,0 +1,30 @@ +//===----------------------------------------------------------------------===// +// +// 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. +// +//===----------------------------------------------------------------------===// +#define _CCCL_DISABLE_FLOAT128 +#include + +#include "test_macros.h" + +int main(int, char**) +{ +#if !_CCCL_HAS_INT128() + auto x = __int128(123456789123) + __int128(123456789123); + auto y = __uint128_t(123456789123) + __uint128_t(123456789123); + unused(x); + unused(y); +#endif +#if !_CCCL_HAS_FLOAT128() +# if defined(__CUDA_ARCH__) + static_assert(false); // nvcc raises a warning in device code but doesn't fail +# endif + auto z = __float128(3.14) + __float128(3.14); + unused(z); +#endif + return 0; +} From d767bb4545ca4ec46ffb961f0778dadbb7055457 Mon Sep 17 00:00:00 2001 From: fbusato Date: Fri, 17 Jan 2025 11:27:59 -0800 Subject: [PATCH 07/29] add static_assert to negative tests --- .../libcxx/macros/128bit_types.compile.fail.cpp | 8 ++++---- .../libcxx/macros/128bit_types.compile.pass.cpp | 1 - 2 files changed, 4 insertions(+), 5 deletions(-) diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp index d7dff72444f0..2cdb9af2eb43 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp @@ -6,7 +6,6 @@ // SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. // //===----------------------------------------------------------------------===// -#define _CCCL_DISABLE_FLOAT128 #include #include "test_macros.h" @@ -18,13 +17,14 @@ int main(int, char**) auto y = __uint128_t(123456789123) + __uint128_t(123456789123); unused(x); unused(y); +#else + static_assert(false); #endif #if !_CCCL_HAS_FLOAT128() -# if defined(__CUDA_ARCH__) - static_assert(false); // nvcc raises a warning in device code but doesn't fail -# endif auto z = __float128(3.14) + __float128(3.14); unused(z); +#else + static_assert(false); #endif return 0; } diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp index 3a1a776c76b6..7bfacb4d66f4 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp @@ -6,7 +6,6 @@ // SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. // //===----------------------------------------------------------------------===// - #include #include "test_macros.h" From d1ad8e801e46807aa0390549161f02d2d7eaadbf Mon Sep 17 00:00:00 2001 From: fbusato Date: Fri, 17 Jan 2025 14:34:20 -0800 Subject: [PATCH 08/29] fix nvrtc test case --- .../test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp index 2cdb9af2eb43..38b9edcfd7fe 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp @@ -20,7 +20,7 @@ int main(int, char**) #else static_assert(false); #endif -#if !_CCCL_HAS_FLOAT128() +#if !_CCCL_HAS_FLOAT128() && !_CCCL_COMPILER(NVRTC) auto z = __float128(3.14) + __float128(3.14); unused(z); #else From 6290337e66aa91b26d1c58abd7549b5dfb5e55f1 Mon Sep 17 00:00:00 2001 From: fbusato Date: Tue, 21 Jan 2025 12:32:57 -0800 Subject: [PATCH 09/29] rename to extend_data_types.h --- libcudacxx/include/cuda/__cccl_config | 2 +- ...floating_point.h => extended_data_types.h} | 37 +++++++++++++++++-- 2 files changed, 34 insertions(+), 5 deletions(-) rename libcudacxx/include/cuda/std/__cccl/{extended_floating_point.h => extended_data_types.h} (51%) diff --git a/libcudacxx/include/cuda/__cccl_config b/libcudacxx/include/cuda/__cccl_config index 08bf8a481813..456903394a4d 100644 --- a/libcudacxx/include/cuda/__cccl_config +++ b/libcudacxx/include/cuda/__cccl_config @@ -22,7 +22,7 @@ #include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export -#include // IWYU pragma: export +#include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export diff --git a/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h similarity index 51% rename from libcudacxx/include/cuda/std/__cccl/extended_floating_point.h rename to libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 5169ea4ad67a..dd10f9706ebf 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h +++ b/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 #include @@ -23,6 +23,7 @@ #endif // no system header #include +#include #include #if !defined(_CCCL_HAS_NVFP16) @@ -39,4 +40,32 @@ # endif #endif // !_CCCL_HAS_NVBF16 -#endif // __CCCL_EXTENDED_FLOATING_POINT_H +#if !defined(_CCCL_DISABLE_INT128) +# if defined(__SIZEOF_INT128__) // defined(__SIZEOF_INT128__) vvvv +# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && _CCCL_OS(LINUX) +# define _CCCL_HAS_INT128() 1 +# elif (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && _CCCL_OS(LINUX) +# define _CCCL_HAS_INT128() 1 +# else +# define _CCCL_HAS_INT128() 0 +# endif +# else // defined(__SIZEOF_INT128__) ^^^^ / !defined(__SIZEOF_INT128__) vvvv +# define _CCCL_HAS_INT128() 0 +# endif // !defined(__SIZEOF_INT128__) ^^^^ +#else +# define _CCCL_HAS_INT128() 0 +#endif // #if !defined(_CCCL_DISABLE_INT128) + +#if !defined(_CCCL_DISABLE_FLOAT128) +# if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) \ + && (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !defined(__CUDA_ARCH__) \ + && !_CCCL_COMPILER(NVRTC) +# define _CCCL_HAS_FLOAT128() 1 +# else +# define _CCCL_HAS_FLOAT128() 0 +# endif +#else +# define _CCCL_HAS_FLOAT128() 0 +#endif // #if !defined(_CCCL_DISABLE_FLOAT128) + +#endif // __CCCL_EXTENDED_DATA_TYPES_H From 46f4eb24af61fb514a774a201882ac334533124d Mon Sep 17 00:00:00 2001 From: fbusato Date: Tue, 21 Jan 2025 15:02:35 -0800 Subject: [PATCH 10/29] fix conflict --- .../cuda/std/__cccl/extended_floating_point.h | 52 +++++++++++++++++++ 1 file changed, 52 insertions(+) create mode 100644 libcudacxx/include/cuda/std/__cccl/extended_floating_point.h diff --git a/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h b/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h new file mode 100644 index 000000000000..fe83fbb03c18 --- /dev/null +++ b/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h @@ -0,0 +1,52 @@ +//===----------------------------------------------------------------------===// +// +// 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 __CCCL_EXTENDED_FLOATING_POINT_H +#define __CCCL_EXTENDED_FLOATING_POINT_H + +#include +#include + +#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 +#include + +#if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) +# if _CCCL_HAS_INCLUDE() +# define _CCCL_HAS_NVFP8() 1 +# else +# define _CCCL_HAS_NVFP8() 0 +# endif // _CCCL_CUDACC_AT_LEAST(11, 8) && _CCCL_HAS_INCLUDE() +#else +# define _CCCL_HAS_NVFP8() 0 +#endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT) + +#if !defined(_CCCL_HAS_NVFP16) +# if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ + && !defined(CCCL_DISABLE_FP16_SUPPORT) +# define _CCCL_HAS_NVFP16 1 +# endif +#endif // !_CCCL_HAS_NVFP16 + +#if !defined(_CCCL_HAS_NVBF16) +# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && !defined(CCCL_DISABLE_BF16_SUPPORT) \ + && !defined(CUB_DISABLE_BF16_SUPPORT) +# define _CCCL_HAS_NVBF16 1 +# endif +#endif // !_CCCL_HAS_NVBF16 + +#endif // __CCCL_EXTENDED_FLOATING_POINT_H From be5d7db2876de4acc422c6b70d2b34670d02376e Mon Sep 17 00:00:00 2001 From: fbusato Date: Tue, 21 Jan 2025 15:04:27 -0800 Subject: [PATCH 11/29] move _CCCL_HAS_NVFP8 to extended data types --- .../cuda/std/__cccl/extended_data_types.h | 38 +++++++++----- .../cuda/std/__cccl/extended_floating_point.h | 52 ------------------- 2 files changed, 24 insertions(+), 66 deletions(-) delete mode 100644 libcudacxx/include/cuda/std/__cccl/extended_floating_point.h diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index dd10f9706ebf..556492f08033 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -26,20 +26,6 @@ #include #include -#if !defined(_CCCL_HAS_NVFP16) -# if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ - && !defined(CCCL_DISABLE_FP16_SUPPORT) -# define _CCCL_HAS_NVFP16 1 -# endif -#endif // !_CCCL_HAS_NVFP16 - -#if !defined(_CCCL_HAS_NVBF16) -# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && !defined(CCCL_DISABLE_BF16_SUPPORT) \ - && !defined(CUB_DISABLE_BF16_SUPPORT) -# define _CCCL_HAS_NVBF16 1 -# endif -#endif // !_CCCL_HAS_NVBF16 - #if !defined(_CCCL_DISABLE_INT128) # if defined(__SIZEOF_INT128__) // defined(__SIZEOF_INT128__) vvvv # if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && _CCCL_OS(LINUX) @@ -56,6 +42,30 @@ # define _CCCL_HAS_INT128() 0 #endif // #if !defined(_CCCL_DISABLE_INT128) +#if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) +# if _CCCL_HAS_INCLUDE() +# define _CCCL_HAS_NVFP8() 1 +# else +# define _CCCL_HAS_NVFP8() 0 +# endif // _CCCL_CUDACC_AT_LEAST(11, 8) && _CCCL_HAS_INCLUDE() +#else +# define _CCCL_HAS_NVFP8() 0 +#endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT) + +#if !defined(_CCCL_HAS_NVFP16) +# if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ + && !defined(CCCL_DISABLE_FP16_SUPPORT) +# define _CCCL_HAS_NVFP16 1 +# endif +#endif // !_CCCL_HAS_NVFP16 + +#if !defined(_CCCL_HAS_NVBF16) +# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && !defined(CCCL_DISABLE_BF16_SUPPORT) \ + && !defined(CUB_DISABLE_BF16_SUPPORT) +# define _CCCL_HAS_NVBF16 1 +# endif +#endif // !_CCCL_HAS_NVBF16 + #if !defined(_CCCL_DISABLE_FLOAT128) # if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) \ && (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !defined(__CUDA_ARCH__) \ diff --git a/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h b/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h deleted file mode 100644 index fe83fbb03c18..000000000000 --- a/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h +++ /dev/null @@ -1,52 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 __CCCL_EXTENDED_FLOATING_POINT_H -#define __CCCL_EXTENDED_FLOATING_POINT_H - -#include -#include - -#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 -#include - -#if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) -# if _CCCL_HAS_INCLUDE() -# define _CCCL_HAS_NVFP8() 1 -# else -# define _CCCL_HAS_NVFP8() 0 -# endif // _CCCL_CUDACC_AT_LEAST(11, 8) && _CCCL_HAS_INCLUDE() -#else -# define _CCCL_HAS_NVFP8() 0 -#endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT) - -#if !defined(_CCCL_HAS_NVFP16) -# if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ - && !defined(CCCL_DISABLE_FP16_SUPPORT) -# define _CCCL_HAS_NVFP16 1 -# endif -#endif // !_CCCL_HAS_NVFP16 - -#if !defined(_CCCL_HAS_NVBF16) -# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && !defined(CCCL_DISABLE_BF16_SUPPORT) \ - && !defined(CUB_DISABLE_BF16_SUPPORT) -# define _CCCL_HAS_NVBF16 1 -# endif -#endif // !_CCCL_HAS_NVBF16 - -#endif // __CCCL_EXTENDED_FLOATING_POINT_H From ac4ecc512513045f85da90700d9a2e6006a45939 Mon Sep 17 00:00:00 2001 From: fbusato Date: Tue, 21 Jan 2025 15:36:27 -0800 Subject: [PATCH 12/29] unify extended data type tests --- libcudacxx/include/cuda/__cccl_config | 1 - .../include/cuda/std/__cccl/128bit_types.h | 45 ------------------- .../macros/128bit_types.compile.fail.cpp | 30 ------------- .../macros/128bit_types.compile.pass.cpp | 26 ----------- ....fail.cpp => extended_data_types.fail.cpp} | 28 +++++++++--- ....pass.cpp => extended_data_types.pass.cpp} | 24 +++++++--- 6 files changed, 38 insertions(+), 116 deletions(-) delete mode 100644 libcudacxx/include/cuda/std/__cccl/128bit_types.h delete mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp delete mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp rename libcudacxx/test/libcudacxx/libcxx/macros/{extended_floating_point.fail.cpp => extended_data_types.fail.cpp} (63%) rename libcudacxx/test/libcudacxx/libcxx/macros/{extended_floating_point.pass.cpp => extended_data_types.pass.cpp} (63%) diff --git a/libcudacxx/include/cuda/__cccl_config b/libcudacxx/include/cuda/__cccl_config index 456903394a4d..e7fdf26dcc51 100644 --- a/libcudacxx/include/cuda/__cccl_config +++ b/libcudacxx/include/cuda/__cccl_config @@ -11,7 +11,6 @@ #ifndef _CUDA__CCCL_CONFIG #define _CUDA__CCCL_CONFIG -#include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export diff --git a/libcudacxx/include/cuda/std/__cccl/128bit_types.h b/libcudacxx/include/cuda/std/__cccl/128bit_types.h deleted file mode 100644 index 95ebe4a37c2a..000000000000 --- a/libcudacxx/include/cuda/std/__cccl/128bit_types.h +++ /dev/null @@ -1,45 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 __CCCL_128BIT_TYPES_H -#define __CCCL_128BIT_TYPES_H - -#include -#include - -#ifndef _CCCL_DISABLE_INT128 -# if defined(__SIZEOF_INT128__) // defined(__SIZEOF_INT128__) vvvv -# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && _CCCL_OS(LINUX) -# define _CCCL_HAS_INT128() 1 -# elif (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && _CCCL_OS(LINUX) -# define _CCCL_HAS_INT128() 1 -# else -# define _CCCL_HAS_INT128() 0 -# endif -# else // defined(__SIZEOF_INT128__) ^^^^ / !defined(__SIZEOF_INT128__) vvvv -# define _CCCL_HAS_INT128() 0 -# endif // !defined(__SIZEOF_INT128__) ^^^^ -#else -# define _CCCL_HAS_INT128() 0 -#endif // #ifndef _CCCL_DISABLE_INT128 - -#ifndef _CCCL_DISABLE_FLOAT128 -# if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) \ - && (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !defined(__CUDA_ARCH__) \ - && !_CCCL_COMPILER(NVRTC) -# define _CCCL_HAS_FLOAT128() 1 -# else -# define _CCCL_HAS_FLOAT128() 0 -# endif -#else -# define _CCCL_HAS_FLOAT128() 0 -#endif // #ifndef _CCCL_DISABLE_FLOAT128 - -#endif // __CCCL_128BIT_TYPES_H diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp deleted file mode 100644 index 38b9edcfd7fe..000000000000 --- a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.fail.cpp +++ /dev/null @@ -1,30 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 - -#include "test_macros.h" - -int main(int, char**) -{ -#if !_CCCL_HAS_INT128() - auto x = __int128(123456789123) + __int128(123456789123); - auto y = __uint128_t(123456789123) + __uint128_t(123456789123); - unused(x); - unused(y); -#else - static_assert(false); -#endif -#if !_CCCL_HAS_FLOAT128() && !_CCCL_COMPILER(NVRTC) - auto z = __float128(3.14) + __float128(3.14); - unused(z); -#else - static_assert(false); -#endif - return 0; -} diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp deleted file mode 100644 index 7bfacb4d66f4..000000000000 --- a/libcudacxx/test/libcudacxx/libcxx/macros/128bit_types.compile.pass.cpp +++ /dev/null @@ -1,26 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 - -#include "test_macros.h" - -int main(int, char**) -{ -#if _CCCL_HAS_INT128() - auto x = __int128(123456789123) + __int128(123456789123); - auto y = __uint128_t(123456789123) + __uint128_t(123456789123); - unused(x); - unused(y); -#endif -#if _CCCL_HAS_FLOAT128() - auto z = __float128(3.14) + __float128(3.14); - unused(z); -#endif - return 0; -} diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fail.cpp similarity index 63% rename from libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.fail.cpp rename to libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fail.cpp index 3b6b457e6338..df712700c772 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fail.cpp @@ -6,7 +6,7 @@ // SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. // //===----------------------------------------------------------------------===// -#include +#include #include "test_macros.h" @@ -22,21 +22,35 @@ int main(int, char**) { -#if !_CCCL_HAS_NVFP8() - auto x = __nv_fp8_e4m3(1.0f); +#if !_CCCL_HAS_INT128() + auto x = __int128(123456789123) + __int128(123456789123); + auto y = __uint128_t(123456789123) + __uint128_t(123456789123); unused(x); + unused(y); +#else + static_assert(false); +#endif +#if !_CCCL_HAS_NVFP8() + auto x1 = __nv_fp8_e4m3(1.0f); + unused(x1); #else static_assert(false); #endif #if !defined(_CCCL_HAS_NVFP16) - auto y = __half(1.0f); - unused(y); + auto x2 = __half(1.0f); + unused(x2); #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 +#if !_CCCL_HAS_FLOAT128() && !_CCCL_COMPILER(NVRTC) + auto x4 = __float128(3.14) + __float128(3.14); + unused(x4); #else static_assert(false); #endif diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp similarity index 63% rename from libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.pass.cpp rename to libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp index fa1476611bcb..3acb73bed50e 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/extended_floating_point.pass.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp @@ -6,7 +6,7 @@ // SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. // //===----------------------------------------------------------------------===// -#include +#include #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() + auto x5 = __float128(3.14) + __float128(3.14); + unused(x5); #endif return 0; } From 9168c04f8defdddb5a07dde8b7f59a6d3e5577e7 Mon Sep 17 00:00:00 2001 From: fbusato Date: Tue, 21 Jan 2025 15:40:00 -0800 Subject: [PATCH 13/29] split failing test cases --- .../macros/extended_data_types.bf16.fail.cpp | 26 ++++++++++++++ .../macros/extended_data_types.fp128.fail.cpp | 22 ++++++++++++ .../macros/extended_data_types.fp16.fail.cpp | 26 ++++++++++++++ .../macros/extended_data_types.fp8.fail.cpp | 26 ++++++++++++++ ....cpp => extended_data_types.i128.fail.cpp} | 34 ------------------- 5 files changed, 100 insertions(+), 34 deletions(-) create mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.bf16.fail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp16.fail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp8.fail.cpp rename libcudacxx/test/libcudacxx/libcxx/macros/{extended_data_types.fail.cpp => extended_data_types.i128.fail.cpp} (53%) diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.bf16.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.bf16.fail.cpp new file mode 100644 index 000000000000..fef41309a677 --- /dev/null +++ b/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 + +#include "test_macros.h" + +#if !defined(_CCCL_HAS_NVBF16) +# include +#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; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp new file mode 100644 index 000000000000..d31edf2d34b4 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp @@ -0,0 +1,22 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#include "test_macros.h" + +int main(int, char**) +{ +#if !_CCCL_HAS_FLOAT128() && !_CCCL_COMPILER(NVRTC) + auto x4 = __float128(3.14) + __float128(3.14); + unused(x4); +#else + static_assert(false); +#endif + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp16.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp16.fail.cpp new file mode 100644 index 000000000000..d4925743db92 --- /dev/null +++ b/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 + +#include "test_macros.h" + +#if !defined(_CCCL_HAS_NVFP16) +# include +#endif + +int main(int, char**) +{ +#if !defined(_CCCL_HAS_NVFP16) + auto x2 = __half(1.0f); + unused(x2); +#else + static_assert(false); +#endif + return 0; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp8.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp8.fail.cpp new file mode 100644 index 000000000000..7418e0030b71 --- /dev/null +++ b/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 + +#include "test_macros.h" + +#if !_CCCL_HAS_NVFP8() +# include +#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; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.i128.fail.cpp similarity index 53% rename from libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fail.cpp rename to libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.i128.fail.cpp index df712700c772..ef2be6cc91ae 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.i128.fail.cpp @@ -10,16 +10,6 @@ #include "test_macros.h" -#if !_CCCL_HAS_NVFP8() -# include -#endif -#if !defined(_CCCL_HAS_NVFP16) -# include -#endif -#if !defined(_CCCL_HAS_NVBF16) -# include -#endif - int main(int, char**) { #if !_CCCL_HAS_INT128() @@ -29,30 +19,6 @@ int main(int, char**) unused(y); #else static_assert(false); -#endif -#if !_CCCL_HAS_NVFP8() - auto x1 = __nv_fp8_e4m3(1.0f); - unused(x1); -#else - static_assert(false); -#endif -#if !defined(_CCCL_HAS_NVFP16) - auto x2 = __half(1.0f); - unused(x2); -#else - static_assert(false); -#endif -#if !defined(_CCCL_HAS_NVBF16) - auto x3 = __nv_bfloat16(1.0f); - unused(x3); -#else - static_assert(false); -#endif -#if !_CCCL_HAS_FLOAT128() && !_CCCL_COMPILER(NVRTC) - auto x4 = __float128(3.14) + __float128(3.14); - unused(x4); -#else - static_assert(false); #endif return 0; } From a21ec9a8b7d22f400cfd4a39ddd81a47fb23979f Mon Sep 17 00:00:00 2001 From: fbusato Date: Tue, 21 Jan 2025 17:29:38 -0800 Subject: [PATCH 14/29] __SIZEOF_INT128__ not defined with NVRTC --- .../cuda/std/__cccl/extended_data_types.h | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 556492f08033..dcbdf672a6d9 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -27,17 +27,14 @@ #include #if !defined(_CCCL_DISABLE_INT128) -# if defined(__SIZEOF_INT128__) // defined(__SIZEOF_INT128__) vvvv -# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && _CCCL_OS(LINUX) -# define _CCCL_HAS_INT128() 1 -# elif (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && _CCCL_OS(LINUX) -# define _CCCL_HAS_INT128() 1 -# else -# define _CCCL_HAS_INT128() 0 -# endif -# else // defined(__SIZEOF_INT128__) ^^^^ / !defined(__SIZEOF_INT128__) vvvv +# if _CCCL_COMPILER(NVRTC) && defined(__CUDACC_RTC_INT128__) && _CCCL_OS(LINUX) +# define _CCCL_HAS_INT128() 1 +# elif defined(__SIZEOF_INT128__) && (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) \ + && _CCCL_OS(LINUX) +# define _CCCL_HAS_INT128() 1 +# else # define _CCCL_HAS_INT128() 0 -# endif // !defined(__SIZEOF_INT128__) ^^^^ +# endif #else # define _CCCL_HAS_INT128() 0 #endif // #if !defined(_CCCL_DISABLE_INT128) From 8a2c5a79d097b060d7609c4062bbc643e75bd104 Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 22 Jan 2025 12:38:02 -0800 Subject: [PATCH 15/29] workaround for nvrtc compiler bug --- .../libcxx/macros/extended_data_types.i128.fail.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.i128.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.i128.fail.cpp index ef2be6cc91ae..04dd7ea75254 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.i128.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.i128.fail.cpp @@ -13,8 +13,8 @@ int main(int, char**) { #if !_CCCL_HAS_INT128() - auto x = __int128(123456789123) + __int128(123456789123); - auto y = __uint128_t(123456789123) + __uint128_t(123456789123); + __int128 x = __int128(123456789123) + __int128(123456789123); + __uint128_t y = __uint128_t(123456789123) + __uint128_t(123456789123); unused(x); unused(y); #else From 76cd412ca0863b28f1915c28a8bf49bec4dbc597 Mon Sep 17 00:00:00 2001 From: Federico Busato <50413820+fbusato@users.noreply.github.com> Date: Thu, 23 Jan 2025 09:27:19 -0800 Subject: [PATCH 16/29] Update libcudacxx/include/cuda/std/__cccl/extended_data_types.h Co-authored-by: Michael Schellenberger Costa --- libcudacxx/include/cuda/std/__cccl/extended_data_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index dcbdf672a6d9..816ef2cc1416 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -73,6 +73,6 @@ # endif #else # define _CCCL_HAS_FLOAT128() 0 -#endif // #if !defined(_CCCL_DISABLE_FLOAT128) +#endif // !_CCCL_DISABLE_FLOAT128 #endif // __CCCL_EXTENDED_DATA_TYPES_H From bf0e1fa7cdb40cbc6e490095a271778ee97ac0b6 Mon Sep 17 00:00:00 2001 From: Federico Busato <50413820+fbusato@users.noreply.github.com> Date: Thu, 23 Jan 2025 09:27:31 -0800 Subject: [PATCH 17/29] Update libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp Co-authored-by: Michael Schellenberger Costa --- .../libcxx/macros/extended_data_types.fp128.fail.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp index d31edf2d34b4..0e1329531dda 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp @@ -1,6 +1,7 @@ //===----------------------------------------------------------------------===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// 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. From 6300bbd1703403d83d545553ea3eb2a626b51d26 Mon Sep 17 00:00:00 2001 From: Federico Busato <50413820+fbusato@users.noreply.github.com> Date: Thu, 23 Jan 2025 09:27:54 -0800 Subject: [PATCH 18/29] Update libcudacxx/include/cuda/std/__cccl/extended_data_types.h Co-authored-by: Michael Schellenberger Costa --- libcudacxx/include/cuda/std/__cccl/extended_data_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 816ef2cc1416..cc08bd28c33c 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -37,7 +37,7 @@ # endif #else # define _CCCL_HAS_INT128() 0 -#endif // #if !defined(_CCCL_DISABLE_INT128) +#endif // !_CCCL_DISABLE_INT128 #if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) # if _CCCL_HAS_INCLUDE() From 48fdc551e162e7333f3852be30fd3368ce7797fa Mon Sep 17 00:00:00 2001 From: Federico Busato <50413820+fbusato@users.noreply.github.com> Date: Thu, 23 Jan 2025 09:29:26 -0800 Subject: [PATCH 19/29] Update libcudacxx/include/cuda/std/__cccl/extended_data_types.h Co-authored-by: Michael Schellenberger Costa --- libcudacxx/include/cuda/std/__cccl/extended_data_types.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index cc08bd28c33c..5268b4166a1a 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -29,8 +29,7 @@ #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_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) \ - && _CCCL_OS(LINUX) +# elif defined(__SIZEOF_INT128__) && _CCCL_OS(LINUX) # define _CCCL_HAS_INT128() 1 # else # define _CCCL_HAS_INT128() 0 From e057cfa176a543c7b54d85ad0c923ceb7cc5f0fe Mon Sep 17 00:00:00 2001 From: fbusato Date: Thu, 23 Jan 2025 10:22:50 -0800 Subject: [PATCH 20/29] improve float128 detection --- .../include/cuda/std/__cccl/extended_data_types.h | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 5268b4166a1a..4c78991e3c8f 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -63,15 +63,15 @@ #endif // !_CCCL_HAS_NVBF16 #if !defined(_CCCL_DISABLE_FLOAT128) -# if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) \ - && (_CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(NVHPC)) && !defined(__CUDA_ARCH__) \ - && !_CCCL_COMPILER(NVRTC) -# define _CCCL_HAS_FLOAT128() 1 +# if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) && !_CCCL_COMPILER(NVRTC) +# 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 -#else -# define _CCCL_HAS_FLOAT128() 0 -#endif // !_CCCL_DISABLE_FLOAT128 +#endif // !defined(_CCCL_DISABLE_FLOAT128) #endif // __CCCL_EXTENDED_DATA_TYPES_H From 603f8b7d27eeccf7dd532afb8c2de2e0f98def3c Mon Sep 17 00:00:00 2001 From: fbusato Date: Thu, 23 Jan 2025 14:35:59 -0800 Subject: [PATCH 21/29] sync extended_floating_point.h --- .../cuda/std/__cccl/extended_floating_point.h | 52 +++++++++++++++++++ 1 file changed, 52 insertions(+) create mode 100644 libcudacxx/include/cuda/std/__cccl/extended_floating_point.h diff --git a/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h b/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h new file mode 100644 index 000000000000..dee553633d89 --- /dev/null +++ b/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h @@ -0,0 +1,52 @@ +//===----------------------------------------------------------------------===// +// +// 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 __CCCL_EXTENDED_FLOATING_POINT_H +#define __CCCL_EXTENDED_FLOATING_POINT_H + +#include +#include + +#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 +#include + +#if !defined(_CCCL_HAS_NVFP16) +# if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ + && !defined(CCCL_DISABLE_FP16_SUPPORT) +# define _CCCL_HAS_NVFP16 1 +# endif +#endif // !_CCCL_HAS_NVFP16 + +#if !defined(_CCCL_HAS_NVBF16) +# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && !defined(CCCL_DISABLE_BF16_SUPPORT) \ + && !defined(CUB_DISABLE_BF16_SUPPORT) +# define _CCCL_HAS_NVBF16 1 +# endif +#endif // !_CCCL_HAS_NVBF16 + +#if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) +# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && defined(_CCCL_HAS_NVBF16) +# define _CCCL_HAS_NVFP8() 1 +# else +# define _CCCL_HAS_NVFP8() 0 +# endif // _CCCL_HAS_INCLUDE() +#else +# define _CCCL_HAS_NVFP8() 0 +#endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT) + +#endif // __CCCL_EXTENDED_FLOATING_POINT_H From 6b1e89f3175822a3308b55d9d9acffe693a9831f Mon Sep 17 00:00:00 2001 From: fbusato Date: Thu, 23 Jan 2025 14:36:28 -0800 Subject: [PATCH 22/29] delete extended_floating_point.h --- .../cuda/std/__cccl/extended_floating_point.h | 52 ------------------- 1 file changed, 52 deletions(-) delete mode 100644 libcudacxx/include/cuda/std/__cccl/extended_floating_point.h diff --git a/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h b/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h deleted file mode 100644 index dee553633d89..000000000000 --- a/libcudacxx/include/cuda/std/__cccl/extended_floating_point.h +++ /dev/null @@ -1,52 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 __CCCL_EXTENDED_FLOATING_POINT_H -#define __CCCL_EXTENDED_FLOATING_POINT_H - -#include -#include - -#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 -#include - -#if !defined(_CCCL_HAS_NVFP16) -# if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ - && !defined(CCCL_DISABLE_FP16_SUPPORT) -# define _CCCL_HAS_NVFP16 1 -# endif -#endif // !_CCCL_HAS_NVFP16 - -#if !defined(_CCCL_HAS_NVBF16) -# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && !defined(CCCL_DISABLE_BF16_SUPPORT) \ - && !defined(CUB_DISABLE_BF16_SUPPORT) -# define _CCCL_HAS_NVBF16 1 -# endif -#endif // !_CCCL_HAS_NVBF16 - -#if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) -# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && defined(_CCCL_HAS_NVBF16) -# define _CCCL_HAS_NVFP8() 1 -# else -# define _CCCL_HAS_NVFP8() 0 -# endif // _CCCL_HAS_INCLUDE() -#else -# define _CCCL_HAS_NVFP8() 0 -#endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT) - -#endif // __CCCL_EXTENDED_FLOATING_POINT_H From 683725275bd42e87124881200d5375e6e9a41cbf Mon Sep 17 00:00:00 2001 From: fbusato Date: Fri, 24 Jan 2025 09:27:53 -0800 Subject: [PATCH 23/29] revert bf16, fp16 conditions for fp8 --- libcudacxx/include/cuda/std/__cccl/extended_data_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 4c78991e3c8f..25daab597a58 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -39,7 +39,7 @@ #endif // !_CCCL_DISABLE_INT128 #if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) -# if _CCCL_HAS_INCLUDE() +# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && defined(_CCCL_HAS_NVBF16) # define _CCCL_HAS_NVFP8() 1 # else # define _CCCL_HAS_NVFP8() 0 From 1d6c90524786cb9380d58ae6f55f2700776d84e3 Mon Sep 17 00:00:00 2001 From: fbusato Date: Fri, 24 Jan 2025 09:44:50 -0800 Subject: [PATCH 24/29] add nvrtc support for float128 --- .../cuda/std/__cccl/extended_data_types.h | 28 +++++++++++-------- .../macros/extended_data_types.fp128.fail.cpp | 4 +-- .../macros/extended_data_types.pass.cpp | 2 +- 3 files changed, 20 insertions(+), 14 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 25daab597a58..9ad127190386 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -38,16 +38,6 @@ # define _CCCL_HAS_INT128() 0 #endif // !_CCCL_DISABLE_INT128 -#if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) -# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && defined(_CCCL_HAS_NVBF16) -# define _CCCL_HAS_NVFP8() 1 -# else -# define _CCCL_HAS_NVFP8() 0 -# endif // _CCCL_CUDACC_AT_LEAST(11, 8) && _CCCL_HAS_INCLUDE() -#else -# define _CCCL_HAS_NVFP8() 0 -#endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT) - #if !defined(_CCCL_HAS_NVFP16) # if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ && !defined(CCCL_DISABLE_FP16_SUPPORT) @@ -62,8 +52,24 @@ # endif #endif // !_CCCL_HAS_NVBF16 +#if !defined(_CCCL_DISABLE_NVFP8_SUPPORT) +# if _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && defined(_CCCL_HAS_NVBF16) +# define _CCCL_HAS_NVFP8() 1 +# else +# define _CCCL_HAS_NVFP8() 0 +# endif // _CCCL_CUDACC_AT_LEAST(11, 8) && _CCCL_HAS_INCLUDE() +#else +# define _CCCL_HAS_NVFP8() 0 +#endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT) + #if !defined(_CCCL_DISABLE_FLOAT128) -# if (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) && !_CCCL_COMPILER(NVRTC) +# 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 +# elif (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) # if !defined(__CUDA_ARCH__) || (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1000) # define _CCCL_HAS_FLOAT128() 1 # else diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp index 0e1329531dda..5cebaa961680 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.fp128.fail.cpp @@ -13,8 +13,8 @@ int main(int, char**) { -#if !_CCCL_HAS_FLOAT128() && !_CCCL_COMPILER(NVRTC) - auto x4 = __float128(3.14) + __float128(3.14); +#if !_CCCL_HAS_FLOAT128() + __float128 x4 = __float128(3.14) + __float128(3.14); unused(x4); #else static_assert(false); diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp index 3acb73bed50e..3b230d8710b1 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp @@ -41,7 +41,7 @@ int main(int, char**) unused(x4); #endif #if _CCCL_HAS_FLOAT128() - auto x5 = __float128(3.14) + __float128(3.14); + __float128 x5 = __float128(3.14) + __float128(3.14); unused(x5); #endif return 0; From da06349e2f38a599315379c3cbb58b4ec6a7888d Mon Sep 17 00:00:00 2001 From: fbusato Date: Mon, 27 Jan 2025 11:18:21 -0800 Subject: [PATCH 25/29] Disable float128 with nvc++ --- libcudacxx/include/cuda/std/__cccl/extended_data_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 9ad127190386..123d8c9021ca 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -69,7 +69,7 @@ # else # define _CCCL_HAS_FLOAT128() 0 # endif -# elif (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) +# elif (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) && !_CCCL_COMPILER(NVHPC) # if !defined(__CUDA_ARCH__) || (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1000) # define _CCCL_HAS_FLOAT128() 1 # else From 724146417bc312ee8890b7d58491fec31219070c Mon Sep 17 00:00:00 2001 From: fbusato Date: Mon, 27 Jan 2025 11:52:55 -0800 Subject: [PATCH 26/29] fix nvc++ float128 detection on device code --- libcudacxx/include/cuda/std/__cccl/extended_data_types.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 123d8c9021ca..c30c8333e6cc 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -69,7 +69,8 @@ # else # define _CCCL_HAS_FLOAT128() 0 # endif -# elif (defined(__SIZEOF_FLOAT128__) || defined(__FLOAT128__)) && _CCCL_OS(LINUX) && !_CCCL_COMPILER(NVHPC) +// 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 From d77c768de1beff142b7b15fff03531cbd2c21063 Mon Sep 17 00:00:00 2001 From: fbusato Date: Tue, 28 Jan 2025 10:44:03 -0800 Subject: [PATCH 27/29] sync macro.rst --- docs/cccl_development/macro.rst | 51 +++++++++++++-------------------- 1 file changed, 20 insertions(+), 31 deletions(-) diff --git a/docs/cccl_development/macro.rst b/docs/cccl_development/macro.rst index 2016c2d1a534..30de6aa8b105 100644 --- a/docs/cccl_development/macro.rst +++ b/docs/cccl_development/macro.rst @@ -68,11 +68,11 @@ The ``_CCCL_CUDA_COMPILER`` function-like macro can also be used to check the ve **PTX macros**: -+--------------------+-------------------------------------------------------------------------------------------------------------------+ -| ``_CCCL_PTX_ARCH`` | Alias of ``__CUDA_ARCH__`` with value equal to 0 if cuda compiler is not available | -+--------------------+-------------------------------------------------------------------------------------------------------------------+ -| ``__cccl_ptx_isa`` | PTX ISA version available with the current CUDA compiler, e.g. PTX ISA 8.4 (``840``) is available from CUDA 12.4 | -+--------------------+-------------------------------------------------------------------------------------------------------------------+ ++-------------------------+-------------------------------------------------------------------------------------------------------------------+ +| ``_CCCL_PTX_ARCH`` | Alias of ``__CUDA_ARCH__`` with value equal to 0 if cuda compiler is not available | ++-------------------------+-------------------------------------------------------------------------------------------------------------------+ +| ``__cccl_ptx_isa`` | PTX ISA version available with the current CUDA compiler, e.g. PTX ISA 8.4 (``840``) is available from CUDA 12.4 | ++-------------------------+-------------------------------------------------------------------------------------------------------------------+ ---- @@ -159,8 +159,10 @@ Usage example: ---- -CUDA attributes ---------------- +CUDA Extension Macros +--------------------- + +**CUDA attributes**: +------------------------------+----------------------------------------------------------+ | ``_CCCL_GRID_CONSTANT`` | Grid constant kernel parameter | @@ -168,30 +170,19 @@ CUDA attributes | ``_CCCL_GLOBAL_CONSTANT`` | Host/device global scope constant (``inline constexpr``) | +------------------------------+----------------------------------------------------------+ ----- +**Extended floating-point types**: -Non-standard Types Support --------------------------- ++------------------------------+-----------------------------------------------------------------------------------------------------------------+ +| ``_CCCL_HAS_NVFP16`` | `__half/__half2` data types are supported and enabled. Prefer over ``__CUDA_FP16_TYPES_EXIST__`` | ++------------------------------+-----------------------------------------------------------------------------------------------------------------+ +| ``_CCCL_HAS_NVBF16`` | `__nv_bfloat16/__nv_bfloat162` data types are supported and enabled. Prefer over ``__CUDA_BF16_TYPES_EXIST__`` | ++------------------------------+-----------------------------------------------------------------------------------------------------------------+ -+------------------------------+-------------------------------------------------------------------------------------------------------------------+ -| ``_CCCL_HAS_INT128()`` | ``__int128`` and ``__uint128_t`` for 128-bit integer are supported and enabled | -+------------------------------+-------------------------------------------------------------------------------------------------------------------+ -| ``_CCCL_HAS_FLOAT128()`` | ``__float128`` for 128-bit floating-point are supported and enabled | -+------------------------------+-------------------------------------------------------------------------------------------------------------------+ -| ``_CCCL_HAS_NVFP16`` | ``__half/__half2`` data types are supported and enabled. Prefer over ``__CUDA_FP16_TYPES_EXIST__`` | -+------------------------------+-------------------------------------------------------------------------------------------------------------------+ -| ``_CCCL_HAS_NVBF16`` | ``__nv_bfloat16/__nv_bfloat162`` data types are supported and enabled. Prefer over ``__CUDA_BF16_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+) | -+------------------------------+-------------------------------------------------------------------------+ ++------------------------------+----------------------------------------------------------------+ +| ``_LIBCUDACXX_HAS_NVFP16`` | `__half/__half2` host/device support (CUDA 12.2) | ++------------------------------+----------------------------------------------------------------+ +| ``_LIBCUDACXX_HAS_NVBF16`` | `__nv_bfloat16/__nv_bfloat162` host/device support (CUDA 12.2) | ++------------------------------+----------------------------------------------------------------+ ---- @@ -253,8 +244,6 @@ Usage example: **Portable attributes**: -+----------------------------------+------------------------------------------------------------------------------+ -| ``_CCCL_ASSUME()`` | Portable ``[[assume]]`` attribute (before C++23) | +----------------------------------+------------------------------------------------------------------------------+ | ``_CCCL_FALLTHROUGH()`` | Portable ``[[fallthrough]]`` attribute (before C++17) | +----------------------------------+------------------------------------------------------------------------------+ From ebba451bc17c109d18ee761a08c3c8a9d2bc8ca8 Mon Sep 17 00:00:00 2001 From: fbusato Date: Tue, 28 Jan 2025 10:48:10 -0800 Subject: [PATCH 28/29] improve fp8 docs --- docs/cccl_development/macro.rst | 22 ++++++++++++---------- 1 file changed, 12 insertions(+), 10 deletions(-) diff --git a/docs/cccl_development/macro.rst b/docs/cccl_development/macro.rst index 5cc39f0f0d33..6b832da00008 100644 --- a/docs/cccl_development/macro.rst +++ b/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+) | ++---------------------------------+-------------------------------------------------------------------------+ ---- From 738bd6676bc5327d9df28811f83663fd774fbf13 Mon Sep 17 00:00:00 2001 From: Federico Busato <50413820+fbusato@users.noreply.github.com> Date: Wed, 29 Jan 2025 09:07:15 -0800 Subject: [PATCH 29/29] Update libcudacxx/include/cuda/std/__cccl/extended_data_types.h Co-authored-by: Bernhard Manfred Gruber --- libcudacxx/include/cuda/std/__cccl/extended_data_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index c30c8333e6cc..034ebc51b5f0 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -57,7 +57,7 @@ # define _CCCL_HAS_NVFP8() 1 # else # define _CCCL_HAS_NVFP8() 0 -# endif // _CCCL_CUDACC_AT_LEAST(11, 8) && _CCCL_HAS_INCLUDE() +# endif // _CCCL_HAS_INCLUDE() && defined(_CCCL_HAS_NVFP16) && defined(_CCCL_HAS_NVBF16) #else # define _CCCL_HAS_NVFP8() 0 #endif // !defined(_CCCL_DISABLE_NVFP8_SUPPORT)