// SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. All rights reserved. // SPDX-License-Identifier: BSD-3-Clause #pragma once #include #include #include #include #include #include #include /*********************************************************************************************************************** * CUB operator to identity **********************************************************************************************************************/ template inline constexpr T identity_v = cuda::identity_element(); template inline const T identity_v, T> = T{}; // e.g. short2, float2, complex<__half> etc. /*********************************************************************************************************************** * half_t specializations **********************************************************************************************************************/ template <> inline const half_t identity_v, half_t> = half_t{0.0f}; template <> inline const half_t identity_v, half_t> = half_t{1.0f}; template <> inline const half_t identity_v, half_t> = cuda::std::numeric_limits::max(); template <> inline const half_t identity_v, half_t> = cuda::std::numeric_limits::lowest(); /*********************************************************************************************************************** * bfloat16_t specializations **********************************************************************************************************************/ template <> inline const bfloat16_t identity_v, bfloat16_t> = bfloat16_t{0.0f}; template <> inline const bfloat16_t identity_v, bfloat16_t> = bfloat16_t{1.0f}; template <> inline const bfloat16_t identity_v, bfloat16_t> = cuda::std::numeric_limits::max(); template <> inline const bfloat16_t identity_v, bfloat16_t> = cuda::std::numeric_limits::lowest(); /*********************************************************************************************************************** * short2, ushort2, float2 specializations **********************************************************************************************************************/ template <> inline constexpr short2 identity_v, short2> = short2{cuda::std::numeric_limits::lowest(), cuda::std::numeric_limits::lowest()}; template <> inline constexpr ushort2 identity_v, ushort2> = ushort2{0, 0}; template <> inline constexpr float2 identity_v, float2> = float2{cuda::std::numeric_limits::lowest(), cuda::std::numeric_limits::lowest()}; template <> inline const __half2 identity_v, __half2> = __half2{cuda::std::numeric_limits<__half>::lowest(), cuda::std::numeric_limits<__half>::lowest()}; template <> inline const __nv_bfloat162 identity_v, __nv_bfloat162> = __nv_bfloat162{ cuda::std::numeric_limits<__nv_bfloat16>::lowest(), cuda::std::numeric_limits<__nv_bfloat16>::lowest()}; template <> inline constexpr short2 identity_v, short2> = short2{cuda::std::numeric_limits::max(), cuda::std::numeric_limits::max()}; template <> inline constexpr ushort2 identity_v, ushort2> = ushort2{cuda::std::numeric_limits::max(), cuda::std::numeric_limits::max()}; template <> inline const __half2 identity_v, __half2> = __half2{cuda::std::numeric_limits<__half>::max(), cuda::std::numeric_limits<__half>::max()}; template <> inline const __nv_bfloat162 identity_v, __nv_bfloat162> = __nv_bfloat162{cuda::std::numeric_limits<__nv_bfloat16>::max(), cuda::std::numeric_limits<__nv_bfloat16>::max()}; template