// Copyright (c) 2022 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. #include #include "paddle/phi/backends/gpu/gpu_launch_config.h" #include "paddle/phi/core/kernel_registry.h" #include "paddle/phi/kernels/dist_kernel.h" #include "paddle/phi/kernels/elementwise_subtract_kernel.h" #include "paddle/phi/kernels/full_kernel.h" #include "paddle/phi/kernels/funcs/math_cuda_utils.h" #include "paddle/phi/kernels/gpu/reduce.h" #include "paddle/phi/kernels/legacy/reduce_max_kernel.h" #include "paddle/phi/kernels/p_norm_kernel.h" #include "paddle/phi/kernels/reduce_min_kernel.h" namespace phi { #define FULL_MASK 0xffffffff template struct ZeroOrderFunctor { public: HOSTDEVICE explicit inline ZeroOrderFunctor() {} HOSTDEVICE inline Ty operator()(const Tx& x, const Tx& y) const { return static_cast(x != y); } }; template struct OtherOrderFunctor { HOSTDEVICE explicit inline OtherOrderFunctor(const Ty& p_order) : p_order_(p_order) {} HOSTDEVICE inline Ty operator()(const Tx& x, const Tx& y) const { return static_cast( pow(abs(static_cast(x) - static_cast(y)), p_order_)); } private: Ty p_order_; }; template struct PowFunctor { HOSTDEVICE explicit inline PowFunctor(const Ty& p_order) : p_order_(p_order) {} HOSTDEVICE inline Tx operator()(const Tx x) const { return static_cast(pow(static_cast(x), p_order_)); } Ty p_order_; }; template // Tx is high precision, Tout is low/out precision struct PowFunctorHighPrecision { HOSTDEVICE explicit inline PowFunctorHighPrecision(const Ty& p_order) : p_order_(p_order) {} HOSTDEVICE inline Tx operator()(const Tx x) const { return static_cast(pow(static_cast(x), p_order_)); } Ty p_order_; }; template __global__ void ReduceSumWithSubtract( const T* x, const T* y, T* out, int64_t N, Functor func) { using MT = typename MPTypeTrait::Type; MT sum_val(0.0); CUDA_KERNEL_LOOP_TYPE(i, N, int64_t) { sum_val += func(x[i], y[i]); } sum_val = funcs::BlockReduceSum(sum_val, FULL_MASK); if (threadIdx.x == 0) { out[blockIdx.x] = static_cast(sum_val); } } template __global__ void ReduceMaxWithSubtract(const T* x, const T* y, T* out, int64_t N) { using MT = typename MPTypeTrait::Type; MT max_val = std::numeric_limits::min(); CUDA_KERNEL_LOOP_TYPE(i, N, int64_t) { max_val = max(max_val, abs(static_cast(x[i]) - static_cast(y[i]))); } max_val = funcs::BlockReduceMax(max_val, FULL_MASK); if (threadIdx.x == 0) { out[blockIdx.x] = static_cast(max_val); } } template __global__ void ReduceMinWithSubtract(const T* x, const T* y, T* out, int64_t N) { using MT = typename MPTypeTrait::Type; MT min_val = std::numeric_limits::max(); CUDA_KERNEL_LOOP_TYPE(i, N, int64_t) { min_val = min(min_val, abs(static_cast(x[i]) - static_cast(y[i]))); } min_val = funcs::BlockReduceMin(min_val, FULL_MASK); if (threadIdx.x == 0) { out[blockIdx.x] = static_cast(min_val); } } template void DistKernel(const Context& dev_ctx, const DenseTensor& x, const DenseTensor& y, float p, DenseTensor* out) { if (x.numel() == 0 || y.numel() == 0) { Full(dev_ctx, out->dims(), 0, out); return; } using MT = typename MPTypeTrait::Type; DenseTensor intermediate; const T* x_ptr = x.data(); const T* y_ptr = y.data(); T* o_ptr = dev_ctx.template Alloc(out); auto stream = dev_ctx.stream(); auto xdim = x.dims(); if (xdim == y.dims()) { // same shape int64_t n = x.numel(); auto config = backends::gpu::GetGpuLaunchConfig1D(dev_ctx, n); intermediate.Resize({config.block_per_grid.x}); T* i_ptr = dev_ctx.template Alloc(&intermediate); std::vector axis_dims = {static_cast(-1)}; std::vector reduce_axis = funcs::details::GetReduceDim(axis_dims, xdim.size(), true); if (p == 0) { ReduceSumWithSubtract <<>>( x_ptr, y_ptr, i_ptr, n, ZeroOrderFunctor()); funcs::ReduceKernel>( dev_ctx, intermediate, out, kps::IdentityFunctor(), reduce_axis); } else if (p == INFINITY) { ReduceMaxWithSubtract <<>>( x_ptr, y_ptr, i_ptr, n); MaxRawKernel( dev_ctx, intermediate, reduce_axis, true, true, out); } else if (p == -INFINITY) { ReduceMinWithSubtract <<>>( x_ptr, y_ptr, i_ptr, n); MinRawKernel( dev_ctx, intermediate, reduce_axis, true, true, out); } else { MT p_order = static_cast(p); ReduceSumWithSubtract <<>>( x_ptr, y_ptr, i_ptr, n, OtherOrderFunctor(p_order)); DenseTensor out_other; out_other.Resize(out->dims()); dev_ctx.template Alloc(&out_other); funcs::ReduceKernel>( dev_ctx, intermediate, &out_other, kps::IdentityFunctor(), reduce_axis); std::vector ins = {&out_other}; std::vector outs = {out}; MT p_order_ = static_cast(1.f / p_order); funcs::ElementwiseKernel( dev_ctx, ins, &outs, PowFunctorHighPrecision(p_order_)); } } else { auto t = Subtract(dev_ctx, x, y); PNormKernel(dev_ctx, t, p, -1, 1e-12, false, true, out); } } } // namespace phi PD_REGISTER_KERNEL(dist, GPU, ALL_LAYOUT, phi::DistKernel, float, double, phi::bfloat16, phi::float16) {}