// 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 "glog/logging.h" #include "paddle/phi/backends/cpu/cpu_context.h" #include "paddle/phi/backends/gpu/gpu_context.h" #include "paddle/phi/common/amp_type_traits.h" #include "paddle/phi/common/data_type.h" #include "paddle/phi/common/place.h" #include "paddle/phi/core/dense_tensor.h" #include "paddle/phi/common/memory_utils.h" namespace phi { template struct AccuracyCheckFunctor { void operator()(const Context& dev_ctx, const DenseTensor& in, const DenseTensor& other, const std::string& fn_name, const float rtol, const float atol, bool equal_nan, DenseTensor* output); }; template struct AccuracyCheckFunctor { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, const DenseTensor& other, const std::string& fn_name, const double rtol, const double atol, bool equal_nan, DenseTensor* output) { auto* in_a = in.data(); auto* in_b = other.data(); auto* out_data = dev_ctx.template Alloc(output); auto num = in.numel(); // *out_data = true; for (int i = 0; i < num; i++) { out_data[i] = true; } bool val; int res_index = -1; for (int i = 0; i < num; i++) { const double a = in_a[i], b = in_b[i]; if (std::isnan(a) || std::isnan(b)) { val = equal_nan && std::isnan(a) == std::isnan(b); } else { double left = (a > b ? a - b : b - a); double right = atol + (b > 0 ? rtol * b : (-rtol) * b); double diff = (left > right ? left - right : right - left); val = a == b || left <= right || diff <= 1e-10; } // *out_data &= val; out_data[i] = val; if (!val) { VLOG(2) << "Accuracy check failed between" << a << " and " << b << " at index= " << i; res_index = i; break; } } PADDLE_ENFORCE_EQ(val, true, common::errors::PreconditionNotMet( "Accuracy check failed, kernel name %s, res index %d", fn_name, res_index)); } }; template struct AccuracyCheckFunctor> { void operator()(const CPUContext& dev_ctx, const DenseTensor& in, const DenseTensor& other, const std::string& fn_name, const double rtol, const double atol, bool equal_nan, DenseTensor* output) { auto* in_a = in.data>(); auto* in_b = other.data>(); auto* out_data = dev_ctx.template Alloc(output); auto num = in.numel(); // *out_data = true; for (int i = 0; i < num; i++) { out_data[i] = true; } bool val = false; int res_index = -1; for (int i = 0; i < num; i++) { const dtype::complex a = in_a[i], b = in_b[i]; if (std::isnan(a) || std::isnan(b)) { val = equal_nan && std::isnan(a) == std::isnan(b); } else { T left = abs(a - b); T right = atol + rtol * abs(b); T diff = abs(left - right); val = a == b || left <= right || diff <= 1e-10; // *out_data &= val; out_data[i] = val; if (!val) { res_index = i; break; } } } PADDLE_ENFORCE_EQ(val, true, common::errors::PreconditionNotMet( "Accuracy check failed, kernel name %s, res index %d", fn_name, res_index)); } }; #if defined(__NVCC__) || defined(__HIPCC__) template __global__ void AccuracyCheckCUDAKernel(const T* in_data, const T* other_data, const double rtol, const double atol, bool equal_nan, int64_t num, bool* out_data) { int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; bool val; using MPType = typename MPTypeTrait::Type; for (int64_t i = idx; i < num; i += blockDim.x * gridDim.x) { const double a = static_cast(in_data[i]); const double b = static_cast(other_data[i]); if (isnan(a) || isnan(b)) { val = equal_nan && isnan(a) == isnan(b); } else { double left = (a > b ? a - b : b - a); double right = atol + (b > 0 ? rtol * b : (-rtol) * b); double diff = (left > right ? left - right : right - left); val = a == b || left <= right || diff <= 1e-10; } out_data[i] = val; if (!val) { *out_data = false; break; } } } template <> __global__ void AccuracyCheckCUDAKernel(const complex64* in_data, const complex64* other_data, const double rtol, const double atol, bool equal_nan, int64_t num, bool* out_data) { int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; bool val; for (int64_t i = idx; i < num; i += blockDim.x * gridDim.x) { const complex64 a = in_data[i]; const complex64 b = other_data[i]; if (isnan(a) || isnan(b)) { val = equal_nan && isnan(a) == isnan(b); } else { float left = abs(a - b); float right = atol + rtol * abs(b); float diff = abs(left - right); val = a == b || left <= right || diff <= 1e-10; } out_data[i] = val; if (!val) { *out_data = false; break; } } } template <> __global__ void AccuracyCheckCUDAKernel( const complex128* in_data, const complex128* other_data, const double rtol, const double atol, bool equal_nan, int64_t num, bool* out_data) { int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; bool val; for (int64_t i = idx; i < num; i += blockDim.x * gridDim.x) { const complex128 a = in_data[i]; const complex128 b = other_data[i]; if (isnan(a) || isnan(b)) { val = equal_nan && isnan(a) == isnan(b); } else { double left = abs(a - b); double right = atol + rtol * abs(b); double diff = abs(left - right); val = a == b || left <= right || diff <= 1e-10; } out_data[i] = val; if (!val) { *out_data = false; break; } } } template struct AccuracyCheckFunctor { void operator()(const GPUContext& dev_ctx, const DenseTensor& in, const DenseTensor& other, const std::string& fn_name, const double rtol, const double atol, bool equal_nan, DenseTensor* output) { int64_t num = in.numel(); const T* in_data = in.data(); const T* other_data = other.data(); bool* out_data = dev_ctx.template Alloc(output); int block = 1024; int64_t grid = (block - 1 + num) / block; grid = (grid > block) ? block : grid; #ifdef PADDLE_WITH_HIP hipMemset(out_data, true, num * sizeof(bool)); #else cudaMemset(out_data, true, num * sizeof(bool)); #endif AccuracyCheckCUDAKernel<<>>( in_data, other_data, rtol, atol, equal_nan, num, out_data); DenseTensor out_cpu; Copy(dev_ctx, *output, CPUPlace(), true, &out_cpu); auto data_ptr = out_cpu.data(); PADDLE_ENFORCE_EQ(*data_ptr, true, common::errors::PreconditionNotMet( "Accuracy check failed, kernel name %s", fn_name)); } }; #endif template void AccuracyCheckKernel(const Context& dev_ctx, const DenseTensor& x, const DenseTensor& y, const std::string& fn_name, const double rtol, const double atol, bool equal_nan, DenseTensor* out) { AccuracyCheckFunctor()( dev_ctx, x, y, fn_name, rtol, atol, equal_nan, out); } } // namespace phi