From c71df9e05c8d49cc61cf7b6918aaccedccc8834f Mon Sep 17 00:00:00 2001 From: Ben van Werkhoven Date: Fri, 25 Sep 2026 20:11:29 +0200 Subject: [PATCH 1/2] Remove deduction guides for alias templates to support C++20 The deduction guides for vec, vec_ptr, kscalar, kvec and kconstant were enabled only for `__cpp_deduction_guides >= 201907L`, i.e. C++20. Compilers reject guides that name an alias template, so kernel_float.h could not be included in C++20 builds. From C++20, the compiler derives deduction guides for alias templates from the class template's own guides, so these can be removed. Adds tests/deduction.cu with compile-time checks of class template argument deduction. --- include/kernel_float/memory.h | 8 ------ include/kernel_float/prelude.h | 12 --------- include/kernel_float/vector.h | 6 ----- single_include/kernel_float.h | 30 ++-------------------- tests/deduction.cu | 47 ++++++++++++++++++++++++++++++++++ 5 files changed, 49 insertions(+), 54 deletions(-) create mode 100644 tests/deduction.cu diff --git a/include/kernel_float/memory.h b/include/kernel_float/memory.h index 5a9f3ee..e3aaaa1 100644 --- a/include/kernel_float/memory.h +++ b/include/kernel_float/memory.h @@ -1090,14 +1090,6 @@ vector_ptr(T*) -> vector_ptr>; template vector_ptr(const T*) -> vector_ptr>; - -#if __cpp_deduction_guides >= 201907L -template -vec_ptr(T*) -> vec_ptr; - -template -vec_ptr(const T*) -> vec_ptr; -#endif #endif } // namespace kernel_float diff --git a/include/kernel_float/prelude.h b/include/kernel_float/prelude.h index a723820..711aac0 100644 --- a/include/kernel_float/prelude.h +++ b/include/kernel_float/prelude.h @@ -108,18 +108,6 @@ static constexpr kconstant operator""_c(unsigned long long int v) return static_cast(v); } -// Deduction guides for aliases are only supported from C++20 -#if defined(__cpp_deduction_guides) && __cpp_deduction_guides >= 201907L -template -kscalar(T&&) -> kscalar>; - -template -kvec(Args&&...) -> kvec, sizeof...(Args)>; - -template -kconstant(T&&) -> kconstant>; -#endif - } // namespace prelude } // namespace kernel_float diff --git a/include/kernel_float/vector.h b/include/kernel_float/vector.h index 222db2e..9da8c38 100644 --- a/include/kernel_float/vector.h +++ b/include/kernel_float/vector.h @@ -380,12 +380,6 @@ KERNEL_FLOAT_INLINE vec, sizeof...(Args)> make_vec(Args&&... // Deduction guide for `vector` template vector(Args&&... args) -> vector, extent>; - -// Deduction guides for aliases are only supported from C++20 -#if __cpp_deduction_guides >= 201907L -template -vec(Args&&... args) -> vec, sizeof...(Args)>; -#endif #endif } // namespace kernel_float diff --git a/single_include/kernel_float.h b/single_include/kernel_float.h index 1dc1d97..8dbdc40 100644 --- a/single_include/kernel_float.h +++ b/single_include/kernel_float.h @@ -16,8 +16,8 @@ //================================================================================ // this file has been auto-generated, do not modify its contents! -// date: 2026-08-28 13:42:54.102829 -// git hash: 14a94fbc188ca3285e7351e5bb79184ce89c75d7 +// date: 2026-09-25 19:45:19.408615 +// git hash: 5226c9232b45c090be46c41e17034c089ffbf8d0 //================================================================================ #ifndef KERNEL_FLOAT_MACROS_H @@ -3710,14 +3710,6 @@ vector_ptr(T*) -> vector_ptr>; template vector_ptr(const T*) -> vector_ptr>; - -#if __cpp_deduction_guides >= 201907L -template -vec_ptr(T*) -> vec_ptr; - -template -vec_ptr(const T*) -> vec_ptr; -#endif #endif } // namespace kernel_float @@ -4571,12 +4563,6 @@ KERNEL_FLOAT_INLINE vec, sizeof...(Args)> make_vec(Args&&... // Deduction guide for `vector` template vector(Args&&... args) -> vector, extent>; - -// Deduction guides for aliases are only supported from C++20 -#if __cpp_deduction_guides >= 201907L -template -vec(Args&&... args) -> vec, sizeof...(Args)>; -#endif #endif } // namespace kernel_float @@ -5788,18 +5774,6 @@ static constexpr kconstant operator""_c(unsigned long long int v) return static_cast(v); } -// Deduction guides for aliases are only supported from C++20 -#if defined(__cpp_deduction_guides) && __cpp_deduction_guides >= 201907L -template -kscalar(T&&) -> kscalar>; - -template -kvec(Args&&...) -> kvec, sizeof...(Args)>; - -template -kconstant(T&&) -> kconstant>; -#endif - } // namespace prelude } // namespace kernel_float diff --git a/tests/deduction.cu b/tests/deduction.cu new file mode 100644 index 0000000..f606294 --- /dev/null +++ b/tests/deduction.cu @@ -0,0 +1,47 @@ +#include + +#include "common.h" +#include "kernel_float/prelude.h" + +// Class template argument deduction. These checks are compile-time only: if this file compiles, +// they pass. The file is valid as both C++17 and C++20. + +namespace deduction_tests { +using namespace kernel_float::prelude; + +struct packed3 { + float x, y, z; // sizeof(packed3) != alignof(packed3) +}; + +__host__ __device__ void class_templates(float* fp, const packed3* cp) { + kf::vector v(1, 2.0f); + static_assert(std::is_same>::value, ""); + + kf::vector_ptr p(fp); + static_assert( + std::is_same>>::value, + ""); + + kf::vector_ptr q(cp); + using expected_q = kf::vector_ptr>; + static_assert(std::is_same::value, ""); + + kf::constant c(2.0); + static_assert(std::is_same>::value, ""); +} + +// From C++20, the compiler derives deduction guides for alias templates from those of the class +// template they name. +#if __cplusplus >= 202002L +__host__ __device__ void alias_templates() { + kf::scalar s(1.5); + static_assert(std::is_same>::value, ""); + + kscalar ks(1.5f); + static_assert(std::is_same>::value, ""); + + kconstant kc(2.0); + static_assert(std::is_same>::value, ""); +} +#endif +} // namespace deduction_tests From 1d03298f5f5514d2b906ad87c0ef769c4653b0ef Mon Sep 17 00:00:00 2001 From: stijn Date: Tue, 29 Sep 2026 19:38:03 +0200 Subject: [PATCH 2/2] Fix HIP build of `tests/deduction.cu` --- tests/deduction.cu | 32 ++++++++++++++++++-------------- 1 file changed, 18 insertions(+), 14 deletions(-) diff --git a/tests/deduction.cu b/tests/deduction.cu index f606294..2e904f2 100644 --- a/tests/deduction.cu +++ b/tests/deduction.cu @@ -3,8 +3,10 @@ #include "common.h" #include "kernel_float/prelude.h" -// Class template argument deduction. These checks are compile-time only: if this file compiles, -// they pass. The file is valid as both C++17 and C++20. +// Class template argument deduction. These tests run on host only: a bug in clang (HIP) +// treats deduction guides as host functions, so CTAD cannot be used in device code. +// See https://github.com/llvm/llvm-project/issues/146646 (fixed in LLVM 22 by +// https://github.com/llvm/llvm-project/pull/170481). namespace deduction_tests { using namespace kernel_float::prelude; @@ -13,35 +15,37 @@ struct packed3 { float x, y, z; // sizeof(packed3) != alignof(packed3) }; -__host__ __device__ void class_templates(float* fp, const packed3* cp) { +TEST_CASE("deduction guides") { + float* fp = nullptr; + const packed3* cp = nullptr; + kf::vector v(1, 2.0f); - static_assert(std::is_same>::value, ""); + STATIC_REQUIRE(std::is_same_v>); kf::vector_ptr p(fp); - static_assert( - std::is_same>>::value, - ""); + STATIC_REQUIRE(std::is_same_v>>); kf::vector_ptr q(cp); - using expected_q = kf::vector_ptr>; - static_assert(std::is_same::value, ""); + STATIC_REQUIRE( + std::is_same_v>>); kf::constant c(2.0); - static_assert(std::is_same>::value, ""); + STATIC_REQUIRE(std::is_same_v>); + CHECK(c.get() == 2.0); } // From C++20, the compiler derives deduction guides for alias templates from those of the class // template they name. #if __cplusplus >= 202002L -__host__ __device__ void alias_templates() { +TEST_CASE("deduction guides for alias templates") { kf::scalar s(1.5); - static_assert(std::is_same>::value, ""); + STATIC_REQUIRE(std::is_same_v>); kscalar ks(1.5f); - static_assert(std::is_same>::value, ""); + STATIC_REQUIRE(std::is_same_v>); kconstant kc(2.0); - static_assert(std::is_same>::value, ""); + STATIC_REQUIRE(std::is_same_v>); } #endif } // namespace deduction_tests