[Bug Fix] Support isfinite/isnan/isinf for float16/bfloat16 on CUDA/HIP (PaddlePaddle#75933)

- 在 isfinite_kernel_impl.h 的 GPU 侧 `Isfinite/Isnan/Isinf` 核函数里，把 “通用浮点” 模板拆成两支：一支只接受标准 `float/double`，另一支专门匹配 `phi::float16` 和 `phi::bfloat16`。这避免了 `std::is_floating_point` 对这两种自定义半精度类型返回 `false` 而导致完全没有匹配内核的情况，从而补齐了半精度在 CUDA/HIP 上的 `isfinite/isnan/isinf` 支持。
- 由于有了独立分支，调用的仍是对应的 `isfinite/isnan/isinf` 设备实现，逻辑保持一致，但现在 `float16/bfloat16` 会正确走到实际内核里，不再出现链接缺符号或运行时报 “未注册该数据类型” 的问题。
- 去掉三个模板 `IsfiniteKernel/IsinfKernel/IsnanKernel` 的 `PADDLE_API` 修饰，避免在头文件模板定义上做符号导出，引起重复导出或 Windows 下的装饰冲突。

Author: youge325

paddle/phi/kernels/impl/isfinite_kernel_impl.h

@@ -301,7 +301,23 @@ __global__ void IsfiniteCUDAKernel(
     const T* in_data,
     IndexType num,
     bool* out_data,
-    typename std::enable_if<std::is_floating_point<T>::value>::type* = 0) {
+    typename std::enable_if<std::is_floating_point<T>::value &&
+                            !std::is_same<T, phi::bfloat16>::value &&
+                            !std::is_same<T, phi::float16>::value>::type* = 0) {
+  IndexType idx = threadIdx.x + blockIdx.x * blockDim.x;
+  for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) {
+    const T& a = in_data[i];
+    out_data[i] = isfinite(a);
+  }
+}
+
+template <typename T, typename IndexType>
+__global__ void IsfiniteCUDAKernel(
+    const T* in_data,
+    IndexType num,
+    bool* out_data,
+    typename std::enable_if<std::is_same<T, phi::bfloat16>::value ||
+                            std::is_same<T, phi::float16>::value>::type* = 0) {
   IndexType idx = threadIdx.x + blockIdx.x * blockDim.x;
   for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) {
     const T& a = in_data[i];
@@ -340,7 +356,23 @@ __global__ void IsnanCUDAKernel(
     const T* in_data,
     IndexType num,
     bool* out_data,
-    typename std::enable_if<std::is_floating_point<T>::value>::type* = 0) {
+    typename std::enable_if<std::is_floating_point<T>::value &&
+                            !std::is_same<T, phi::bfloat16>::value &&
+                            !std::is_same<T, phi::float16>::value>::type* = 0) {
+  IndexType idx = threadIdx.x + blockIdx.x * blockDim.x;
+  for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) {
+    const T& a = in_data[i];
+    out_data[i] = isnan(a);
+  }
+}
+
+template <typename T, typename IndexType>
+__global__ void IsnanCUDAKernel(
+    const T* in_data,
+    IndexType num,
+    bool* out_data,
+    typename std::enable_if<std::is_same<T, phi::bfloat16>::value ||
+                            std::is_same<T, phi::float16>::value>::type* = 0) {
   IndexType idx = threadIdx.x + blockIdx.x * blockDim.x;
   for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) {
     const T& a = in_data[i];
@@ -379,7 +411,23 @@ __global__ void IsinfCUDAKernel(
     const T* in_data,
     IndexType num,
     bool* out_data,
-    typename std::enable_if<std::is_floating_point<T>::value>::type* = 0) {
+    typename std::enable_if<std::is_floating_point<T>::value &&
+                            !std::is_same<T, phi::bfloat16>::value &&
+                            !std::is_same<T, phi::float16>::value>::type* = 0) {
+  IndexType idx = threadIdx.x + blockIdx.x * blockDim.x;
+  for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) {
+    const T& a = in_data[i];
+    out_data[i] = isinf(a);
+  }
+}
+
+template <typename T, typename IndexType>
+__global__ void IsinfCUDAKernel(
+    const T* in_data,
+    IndexType num,
+    bool* out_data,
+    typename std::enable_if<std::is_same<T, phi::bfloat16>::value ||
+                            std::is_same<T, phi::float16>::value>::type* = 0) {
   IndexType idx = threadIdx.x + blockIdx.x * blockDim.x;
   for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) {
     const T& a = in_data[i];
@@ -477,9 +525,9 @@ struct IsinfFunctor<phi::GPUContext, T> {
 #endif
 
 template <typename T, typename Context>
-PADDLE_API void IsfiniteKernel(const Context& dev_ctx,
-                               const DenseTensor& x,
-                               DenseTensor* out) {
+void IsfiniteKernel(const Context& dev_ctx,
+                    const DenseTensor& x,
+                    DenseTensor* out) {
   if (out && out->numel() == 0) {
     dev_ctx.template Alloc<bool>(out);
     return;