// Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // // http://www.apache.org/licenses/LICENSE-2.0 // // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // See the License for the specific language governing permissions and // limitations under the License. #pragma once #include #include #include "paddle/phi/backends/cpu/cpu_context.h" #include "paddle/phi/backends/gpu/gpu_context.h" #include "paddle/phi/common/data_type.h" #include "paddle/phi/common/place.h" #include "paddle/phi/core/dense_tensor.h" #include "paddle/phi/kernels/funcs/isfinite_functor.h" #include "paddle/phi/kernels/isfinite_kernel.h" // check if vanilla float/double template struct is_float_or_double : std::integral_constant::value || std::is_same::value> {}; // check ifspecial float type, e.g. float16/bfloat16 template struct is_other_float : std::integral_constant::value && !is_float_or_double::value> {}; // check if complex type template struct is_complex64_or_complex128 : std::integral_constant::value || std::is_same::value> {}; namespace phi { /* Codes for isfinite/isinf/isnan as constructed as below: 1. A general template, 2. partial specialization for regular floating-point numbers(float/double), 3. partial specialization for special floating-point numbers(float16/bfloat16 and other special float), 4. partial specialization for non-floating-point (integer) types. 5. partial specialization for complex types. */ /* IsfiniteFunctor */ template struct IsfiniteFunctor { void operator()(const Context& dev_ctx, const DenseTensor& in, DenseTensor* output); }; template struct IsfiniteFunctor< CPUContext, T, typename std::enable_if::value && !is_complex64_or_complex128::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { out_data[i] = true; } } }; template struct IsfiniteFunctor< CPUContext, T, typename std::enable_if::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = std::isfinite(a); } } }; template struct IsfiniteFunctor< CPUContext, T, typename std::enable_if::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = dtype::isfinite(a); } } }; template struct IsfiniteFunctor< CPUContext, T, typename std::enable_if::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = std::isfinite(a.real) && std::isfinite(a.imag); } } }; /* IsnanFunctor */ template struct IsnanFunctor { void operator()(const Context& dev_ctx, const DenseTensor& in, DenseTensor* output); }; template struct IsnanFunctor< CPUContext, T, typename std::enable_if::value && !is_complex64_or_complex128::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { out_data[i] = false; } } }; template struct IsnanFunctor< CPUContext, T, typename std::enable_if::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = std::isnan(a); } } }; template struct IsnanFunctor::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = dtype::isnan(a); } } }; template struct IsnanFunctor< CPUContext, T, typename std::enable_if::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = std::isnan(a.real) || std::isnan(a.imag); } } }; /* IsinfFunctor */ template struct IsinfFunctor { void operator()(const Context& dev_ctx, const DenseTensor& in, DenseTensor* output); }; template struct IsinfFunctor< CPUContext, T, typename std::enable_if::value && !is_complex64_or_complex128::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* out_data = dev_ctx.template Alloc(output); auto num = in.numel(); for (int64_t i = 0; i < num; i++) { out_data[i] = false; } } }; template struct IsinfFunctor< CPUContext, T, typename std::enable_if::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = std::isinf(a); } } }; template struct IsinfFunctor::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = dtype::isinf(a); } } }; template struct IsinfFunctor< CPUContext, T, typename std::enable_if::value>::type> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { auto* in_a = in.data(); auto* out_data = dev_ctx.template Alloc(output); int64_t num = in.numel(); for (int64_t i = 0; i < num; i++) { const T& a = in_a[i]; out_data[i] = std::isinf(a.real) || std::isinf(a.imag); } } }; #if defined(__NVCC__) || defined(__HIPCC__) /* IsfiniteFunctor */ template __global__ void IsfiniteCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value && !std::is_same::value && !std::is_same::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(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 __global__ void IsfiniteCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value || std::is_same::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(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 __global__ void IsfiniteCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(blockDim.x); for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) { out_data[i] = true; } } template __global__ void IsfiniteCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(blockDim.x); for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) { const T& a = in_data[i]; out_data[i] = isfinite(a.real) && isfinite(a.imag); } } /* IsnanFunctor */ template __global__ void IsnanCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value && !std::is_same::value && !std::is_same::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(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 __global__ void IsnanCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value || std::is_same::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(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 __global__ void IsnanCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(blockDim.x); for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) { out_data[i] = false; } } template __global__ void IsnanCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(blockDim.x); for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) { const T& a = in_data[i]; out_data[i] = isnan(a.real) || isnan(a.imag); } } /* IsinfFunctor */ template __global__ void IsinfCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value && !std::is_same::value && !std::is_same::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(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 __global__ void IsinfCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value || std::is_same::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(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 __global__ void IsinfCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(blockDim.x); for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) { out_data[i] = false; } } template __global__ void IsinfCUDAKernel( const T* in_data, IndexType num, bool* out_data, typename std::enable_if::value>::type* = 0) { IndexType idx = static_cast(threadIdx.x) + static_cast(blockIdx.x) * static_cast(blockDim.x); for (IndexType i = idx; i < num; i += blockDim.x * gridDim.x) { const T& a = in_data[i]; out_data[i] = isinf(a.real) || isinf(a.imag); } } template struct IsfiniteFunctor { void operator()(const GPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { int64_t num = in.numel(); const T* in_data = in.data(); bool* out_data = dev_ctx.template Alloc(output); int64_t block = 1024; int64_t grid = (block - 1 + num) / block; grid = (grid > block) ? block : grid; if (num + block * grid + 1 > std::numeric_limits::max()) { IsfiniteCUDAKernel <<>>(in_data, num, out_data); } else { IsfiniteCUDAKernel <<>>(in_data, num, out_data); } } }; template struct IsnanFunctor { void operator()(const GPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { int64_t num = in.numel(); const T* in_data = in.data(); bool* out_data = dev_ctx.template Alloc(output); int64_t block = 1024; int64_t grid = (block - 1 + num) / block; grid = (grid > block) ? block : grid; if (num + block * grid + 1 > std::numeric_limits::max()) { IsnanCUDAKernel <<>>(in_data, num, out_data); } else { IsnanCUDAKernel <<>>(in_data, num, out_data); } } }; template struct IsinfFunctor { void operator()(const GPUContext& dev_ctx, const DenseTensor& in, DenseTensor* output) { int64_t num = in.numel(); const T* in_data = in.data(); bool* out_data = dev_ctx.template Alloc(output); int64_t block = 1024; int64_t grid = (block - 1 + num) / block; grid = (grid > block) ? block : grid; if (num + block * grid + 1 > std::numeric_limits::max()) { IsinfCUDAKernel <<>>(in_data, num, out_data); } else { IsinfCUDAKernel <<>>(in_data, num, out_data); } } }; #endif template void IsfiniteKernel(const Context& dev_ctx, const DenseTensor& x, DenseTensor* out) { if (out && out->numel() == 0) { dev_ctx.template Alloc(out); return; } IsfiniteFunctor()(dev_ctx, x, out); } template void IsinfKernel(const Context& dev_ctx, const DenseTensor& x, DenseTensor* out) { if (out && out->numel() == 0) { dev_ctx.template Alloc(out); return; } IsinfFunctor()(dev_ctx, x, out); } template void IsnanKernel(const Context& dev_ctx, const DenseTensor& x, DenseTensor* out) { if (out && out->numel() == 0) { dev_ctx.template Alloc(out); return; } IsnanFunctor()(dev_ctx, x, out); } } // namespace phi