Fixed decorators on max min clamp for CUDA 13.0

[morousg]

No description provided.

[Copilot]

Pull request overview

This PR updates the CUDA-side constexpr_cmath polyfills used when <cuda/std/algorithm> is unavailable (CUDA < 13.3), aiming to fix CUDA 13.0 compilation by adjusting the host/device decorators on max, min, and clamp.

Changes:

• Replaced ad-hoc constexpr __host__ __device__ decorators with the project macro FK_HOST_DEVICE_CNST for max, min, and clamp.
• Changed the polyfill signatures for max/min/clamp from const T& parameters/return to pass/return by value (T).

Suppressed comments (2)

include/fused_kernel/core/constexpr_libs/constexpr_cmath.h:64

• Same as max: the CUDA < 13.3 polyfill base::min returns by value while CUDA >= 13.3 aliases cuda::std::min (and non-CUDA aliases std::min) returning const T&. Keeping signatures consistent across branches avoids subtle compilation differences between CUDA versions.

template <typename T>
FK_HOST_DEVICE_CNST T min(const T a, const T b) {
    return (b < a) ? b : a;
}

include/fused_kernel/core/constexpr_libs/constexpr_cmath.h:69

• The CUDA < 13.3 polyfill base::clamp returns T by value, but CUDA >= 13.3 aliases cuda::std::clamp (and non-CUDA aliases std::clamp) returning const T&. This can lead to different types/behavior depending on CUDA version. Consider matching the standard signature here and only changing the decorators.

template <typename T>
FK_HOST_DEVICE_CNST T clamp(const T v, const T lo, const T hi) {
    return (v < lo) ? lo : ((hi < v) ? hi : v);
}

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.