Use macro guard CUDA functions for back compatibility in grouped_topk_kernel.cu

[minosfuture]

Summary:
cuda::std::isfinite is not available with earlier CUDA versions. We guard it
with macros and extract a device function for is_finite.

Test Plan: build with 12.4 and 12.8

Reviewed By: houseroad

Differential Revision: D82918389

[facebook-github-bot]

@minosfuture has exported this pull request. If you are a Meta employee, you can view the originating diff in D82918389.

[gemini-code-assist]

Code Review

The pull request introduces a compatibility wrapper is_finite to handle differences in isfinite support across CUDA versions. This is a good approach to maintain backward compatibility. My feedback focuses on improving the implementation of this new function for better robustness and readability.

[gemini-code-assist]

The implementation of is_finite can be improved for robustness and conciseness.

1. The preprocessor check for the CUDA version can be made more robust and standard by using the __CUDACC_VER__ macro. This macro combines major, minor, and patch versions into a single integer (e.g., 120800 for 12.8.0), which makes the check cleaner and correctly handles patch versions if needed in the future. The current check only considers major and minor versions.
2. The function body can be simplified by directly returning the result of the isfinite calls within the #if/#else branches, removing the need for the intermediate res variable.

Here is a suggested improved version:

template <typename T>
__device__ inline bool is_finite(const T val) {
#if defined(__CUDACC_VER__) && __CUDACC_VER__ >= 120800
  return cuda::std::isfinite(val);
#else
  return isfinite(cuda_cast<float, T>(val));
#endif
}

[houseroad]

CUDACC_VER should be deprecated, rigth?

[minosfuture]

the training data of gemini needs some refreshing :D

[facebook-github-bot]

@minosfuture has exported this pull request. If you are a Meta employee, you can view the originating diff in D82918389.

[yewentao256]

LGTM, thanks for the work!