diff --git a/paddle/fluid/framework/dim.h b/paddle/fluid/framework/dim.h index 531ee8445b78f2ffdb4f1197fdee9c5cf0b79081..66214b265fdf9078aeda4efa37c7ad1f2bbef62b 100644 --- a/paddle/fluid/framework/dim.h +++ b/paddle/fluid/framework/dim.h @@ -20,7 +20,6 @@ #include #include "paddle/fluid/framework/array.h" -#include "paddle/fluid/platform/assert.h" #include "paddle/fluid/platform/enforce.h" #include "paddle/fluid/platform/hostdevice.h" diff --git a/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu b/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu index 49420a7a667309eaf8a97c480ab361cec4c29e3e..9e927ed6800d0522dd3b5f6e74990348408b39b6 100644 --- a/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu +++ b/paddle/fluid/inference/tensorrt/plugin/split_op_plugin.cu @@ -111,10 +111,9 @@ int SplitPlugin::enqueue(int batchSize, const void* const* inputs, float const* input_ptr = reinterpret_cast(inputs[0]); float* const* h_odatas = reinterpret_cast(outputs); float** output_ptrs = thrust::raw_pointer_cast(&d_output_ptrs_[0]); - PADDLE_ENFORCE(cudaMemcpyAsync(output_ptrs, h_odatas, - d_output_ptrs_.size() * sizeof(float*), - cudaMemcpyHostToDevice, - stream) == cudaSuccess); + PADDLE_ENFORCE_CUDA_SUCCESS(cudaMemcpyAsync( + output_ptrs, h_odatas, d_output_ptrs_.size() * sizeof(float*), + cudaMemcpyHostToDevice, stream)); int outer_rows = outer_rows_ * batchSize; diff --git a/paddle/fluid/operators/center_loss_op.cu b/paddle/fluid/operators/center_loss_op.cu index eb172fb1f1e82ae3da2ce7d5cf4a76eb5a43e0dc..10b65fa215adc51ecc5c4ff482803a4c8379a757 100644 --- a/paddle/fluid/operators/center_loss_op.cu +++ b/paddle/fluid/operators/center_loss_op.cu @@ -14,7 +14,6 @@ limitations under the License. */ #include #include "paddle/fluid/operators/center_loss_op.h" -#include "paddle/fluid/platform/assert.h" #include "paddle/fluid/platform/cuda_primitives.h" #include "paddle/fluid/platform/gpu_info.h" namespace paddle { @@ -31,8 +30,8 @@ __global__ void ComputeDifferent(T *centers_diff, const T *X, const T *centers, while (idy < K) { int64_t id = ids[idy]; - PADDLE_ASSERT_MSG(id >= 0, "received id:", id); - PADDLE_ASSERT_MSG(id < N, "received id:", id); + PADDLE_ENFORCE(id >= 0, "received id:", id); + PADDLE_ENFORCE(id < N, "received id:", id); T *out = centers_diff + idy * D; const T *x = X + idy * D; const T *cent = centers + id * D; @@ -53,8 +52,8 @@ __global__ void UpdateCenters(T *centers, T *centers_diff, const int64_t *ids, while (idy < K) { int count = 1; int64_t id = ids[idy]; - PADDLE_ASSERT_MSG(id >= 0, "received id:", id); - PADDLE_ASSERT_MSG(id < N, "received id:", id); + PADDLE_ENFORCE(id >= 0, "received id:", id); + PADDLE_ENFORCE(id < N, "received id:", id); for (int i = 0; i < K; i++) { if (ids[i] == id) { diff --git a/paddle/fluid/operators/cross_entropy_op.h b/paddle/fluid/operators/cross_entropy_op.h old mode 100755 new mode 100644 index 2f136784f6b4758175ac38dd003ed4f068dd4bcd..a95e328742075227eb65245117586f5292ae252e --- a/paddle/fluid/operators/cross_entropy_op.h +++ b/paddle/fluid/operators/cross_entropy_op.h @@ -155,11 +155,11 @@ struct HardLabelCrossEntropyForwardFunctor { HOSTDEVICE void operator()(int64_t idx) const { auto label = label_[idx]; if (label != ignore_index_) { - PADDLE_ASSERT_MSG(label >= 0 && label < feature_size_, - "Variable value (label) of " - "OP(fluid.layers.cross_entropy) expected >= 0 " - "and < %ld, but got %ld. Please check label value.", - feature_size_, label); + PADDLE_ENFORCE(label >= 0 && label < feature_size_, + "Variable value (label) of " + "OP(fluid.layers.cross_entropy) expected >= 0 " + "and < %ld, but got %ld. Please check label value.", + feature_size_, label); auto match_x = x_[idx * feature_size_ + label]; y_[idx] = -math::TolerableValue()(real_log(match_x)); match_x_[idx] = match_x; diff --git a/paddle/fluid/operators/lookup_table_op.cu b/paddle/fluid/operators/lookup_table_op.cu index e4add1c746a007909e62acef3194c221c4603341..f9e12e01c13d2a9da2668c80e8479b09a31f280b 100644 --- a/paddle/fluid/operators/lookup_table_op.cu +++ b/paddle/fluid/operators/lookup_table_op.cu @@ -15,7 +15,6 @@ limitations under the License. */ #include "paddle/fluid/framework/eigen.h" #include "paddle/fluid/framework/op_registry.h" #include "paddle/fluid/operators/lookup_table_op.h" -#include "paddle/fluid/platform/assert.h" #include "paddle/fluid/platform/cuda_primitives.h" #include "paddle/fluid/platform/float16.h" @@ -32,12 +31,12 @@ __global__ void LookupTable(T *output, const T *table, const int64_t *ids, while (idy < K) { int64_t id = ids[idy]; - PADDLE_ASSERT_MSG( + PADDLE_ENFORCE( id >= 0, "Variable value (input) of OP(fluid.layers.embedding) " "expected >= 0 and < %ld, but got %ld. Please check input value.", N, id); - PADDLE_ASSERT_MSG( + PADDLE_ENFORCE( id < N, "Variable value (input) of OP(fluid.layers.embedding) " "expected >= 0 and < %ld, but got %ld. Please check input value.", @@ -67,12 +66,12 @@ __global__ void LookupTableGrad(T *table, const T *output, const int64_t *ids, while (idy < K) { int64_t id = ids[idy]; - PADDLE_ASSERT_MSG( + PADDLE_ENFORCE( id >= 0, "Variable value (input) of OP(fluid.layers.embedding) " "expected >= 0 and < %ld, but got %ld. Please check input value.", N, id); - PADDLE_ASSERT_MSG( + PADDLE_ENFORCE( id < N, "Variable value (input) of OP(fluid.layers.embedding) " "expected >= 0 and < %ld, but got %ld. Please check input value.", diff --git a/paddle/fluid/operators/math/CMakeLists.txt b/paddle/fluid/operators/math/CMakeLists.txt index d4837696241b8c4e3cca4f2afe872c6be559853c..88e750214ef31216a7dd35d4859fa2266b47bf86 100644 --- a/paddle/fluid/operators/math/CMakeLists.txt +++ b/paddle/fluid/operators/math/CMakeLists.txt @@ -7,7 +7,7 @@ function(math_library TARGET) set(cc_srcs) set(cu_srcs) set(hip_srcs) - set(math_common_deps device_context framework_proto) + set(math_common_deps device_context framework_proto enforce) set(multiValueArgs DEPS) cmake_parse_arguments(math_library "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN}) diff --git a/paddle/fluid/operators/math/cross_entropy.cu b/paddle/fluid/operators/math/cross_entropy.cu index 59f4485aa92c8dbaf219369ae0e0406758462920..2d871c6e14b855c01b7783bd90103a1e49c71ac2 100644 --- a/paddle/fluid/operators/math/cross_entropy.cu +++ b/paddle/fluid/operators/math/cross_entropy.cu @@ -27,10 +27,10 @@ __global__ void CrossEntropyKernel(T* Y, const T* X, const int64_t* label, const int ignore_index) { for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += blockDim.x * gridDim.x) { - PADDLE_ASSERT_MSG(label[i] >= 0 && label[i] < D || label[i] == ignore_index, - "label[%d] expected >= 0 and < %ld, or == %ld, but got " - "%ld. Please check input value.", - i, D, ignore_index, label[i]); + PADDLE_ENFORCE(label[i] >= 0 && label[i] < D || label[i] == ignore_index, + "label[%d] expected >= 0 and < %ld, or == %ld, but got " + "%ld. Please check input value.", + i, D, ignore_index, label[i]); Y[i] = ignore_index == label[i] ? static_cast(0) : -math::TolerableValue()(real_log(X[i * D + label[i]])); diff --git a/paddle/fluid/operators/math/cross_entropy.h b/paddle/fluid/operators/math/cross_entropy.h index 23d2cf4fd9f532f9217b63e84d928fce5e8e0acb..db19818951d7c9f7e55f8acf2f2de7e3e3819694 100644 --- a/paddle/fluid/operators/math/cross_entropy.h +++ b/paddle/fluid/operators/math/cross_entropy.h @@ -25,8 +25,8 @@ namespace math { template struct TolerableValue { HOSTDEVICE T operator()(const T& x) const { - PADDLE_ASSERT_MSG(std::is_floating_point::value, - "TolerableValue should be float in cross_entropy."); + PADDLE_ENFORCE(std::is_floating_point::value, + "TolerableValue should be float in cross_entropy."); const T kApproInf = 1e20; if (x == INFINITY) return kApproInf; diff --git a/paddle/fluid/operators/math/unpooling.cu b/paddle/fluid/operators/math/unpooling.cu index de6ee7c7cd6e8305af9386bb3d30d19c9846b690..d78e3385efb29cbba540d50433bf0fe35cedd448 100644 --- a/paddle/fluid/operators/math/unpooling.cu +++ b/paddle/fluid/operators/math/unpooling.cu @@ -37,10 +37,10 @@ __global__ void KernelUnpool2dMax(const int nthreads, const T* input_data, int cidx = boffset / in_c_stride; int out_offset = bidx * out_n_stride + cidx * out_c_stride; int out_index = indices_data[i]; - PADDLE_ASSERT_MSG(out_index < out_c_stride, - "out_index < out_c_stride. Expected %ld < %ld, but got " - "%ld >= %ld. Please check input value.", - out_index, out_c_stride, out_index, out_c_stride); + PADDLE_ENFORCE(out_index < out_c_stride, + "out_index < out_c_stride. Expected %ld < %ld, but got " + "%ld >= %ld. Please check input value.", + out_index, out_c_stride, out_index, out_c_stride); output_data[out_offset + out_index] = input_data[i]; } } @@ -62,10 +62,10 @@ __global__ void KernelUnpool2dMaxGrad( int cidx = boffset / in_c_stride; int out_offset = bidx * out_n_stride + cidx * out_c_stride; int out_index = indices_data[i]; - PADDLE_ASSERT_MSG(out_index < out_c_stride, - "out_index < out_c_stride. Expected %ld < %ld, but got " - "%ld >= %ld. Please check input value.", - out_index, out_c_stride, out_index, out_c_stride); + PADDLE_ENFORCE(out_index < out_c_stride, + "out_index < out_c_stride. Expected %ld < %ld, but got " + "%ld >= %ld. Please check input value.", + out_index, out_c_stride, out_index, out_c_stride); input_grad[i] = output_grad[out_offset + out_index]; } } diff --git a/paddle/fluid/operators/modified_huber_loss_op.h b/paddle/fluid/operators/modified_huber_loss_op.h index d7dbf791a7ee13d87836bb6b0292a44eafa982d9..d6dd5539af018c171c7ae00b945f7cfc858b7903 100644 --- a/paddle/fluid/operators/modified_huber_loss_op.h +++ b/paddle/fluid/operators/modified_huber_loss_op.h @@ -29,10 +29,10 @@ using EigenVector = framework::EigenVector; template struct CheckLabelValue { HOSTDEVICE T operator()(const T& val) const { - PADDLE_ASSERT_MSG(val == static_cast(0) || val == static_cast(1), - "LabelValue of modified_huber_loss_op expected to be 0 " - "or 1, but got %ld. Please check input value.", - val); + PADDLE_ENFORCE(val == static_cast(0) || val == static_cast(1), + "LabelValue of modified_huber_loss_op expected to be 0 " + "or 1, but got %ld. Please check input value.", + val); } }; diff --git a/paddle/fluid/operators/random_crop_op.h b/paddle/fluid/operators/random_crop_op.h index e1457eccb5b4d15941ad135e8a20a89ddddd26d8..ae58358cbb202c229cb8a96e20e4c48d926c5bf3 100644 --- a/paddle/fluid/operators/random_crop_op.h +++ b/paddle/fluid/operators/random_crop_op.h @@ -60,16 +60,16 @@ HOSTDEVICE inline void StridedMemcpy(const T* x, const size_t* x_dims, T* out, size_t offset_i = offsets[i]; if (i == rank - 1) { - PADDLE_ASSERT_MSG(x_stride == 1, - "When i:%d == rank:%d - 1, x_stride of random_crop_op " - "expected to be 1, but got %ld. Please check input " - "value.", - i, rank, x_stride); - PADDLE_ASSERT_MSG(out_stride == 1, - "When i:%d == rank:%d - 1, out_stride of random_crop_op " - "expected to be 1, but got %ld. Please check input " - "value.", - i, rank, out_stride); + PADDLE_ENFORCE(x_stride == 1, + "When i:%d == rank:%d - 1, x_stride of random_crop_op " + "expected to be 1, but got %ld. Please check input " + "value.", + i, rank, x_stride); + PADDLE_ENFORCE(out_stride == 1, + "When i:%d == rank:%d - 1, out_stride of random_crop_op " + "expected to be 1, but got %ld. Please check input " + "value.", + i, rank, out_stride); x += offset_i; for (size_t j = 0; j < out_dim_i; ++j) { *out++ = *x++; diff --git a/paddle/fluid/operators/sample_logits_op.h b/paddle/fluid/operators/sample_logits_op.h index 6da533da718ff1e70acf83329ab881fc72f5e441..18ef6c9d3fe62a913413cf8c84e23b7c6accfc5c 100644 --- a/paddle/fluid/operators/sample_logits_op.h +++ b/paddle/fluid/operators/sample_logits_op.h @@ -34,8 +34,8 @@ using EigenMatrix = framework::EigenMatrix; template struct TolerableValue { HOSTDEVICE T operator()(const T& x) const { - PADDLE_ASSERT_MSG(std::is_floating_point::value, - "TolerableValue should be float in sample_logits_op."); + PADDLE_ENFORCE(std::is_floating_point::value, + "TolerableValue should be float in sample_logits_op."); const T kApproInf = 1e20; if (x == INFINITY) return kApproInf; if (x == -INFINITY) return -kApproInf; diff --git a/paddle/fluid/platform/assert.h b/paddle/fluid/platform/assert.h deleted file mode 100644 index 2883bd5ed34834692cb0b637da372cb8e343d9bf..0000000000000000000000000000000000000000 --- a/paddle/fluid/platform/assert.h +++ /dev/null @@ -1,39 +0,0 @@ -/* Copyright (c) 2018 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 - -#define STRINGIFY(x) #x -#define TOSTRING(x) STRINGIFY(x) - -// For cuda, the assertions can affect performance and it is therefore -// recommended to disable them in production code -// https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#assertion -#if defined(__CUDA_ARCH__) -#include -#define EXIT() asm("trap;") -#else -#include -#define EXIT() throw std::runtime_error("Exception encounter.") -#endif - -// NOTE: PADDLE_ASSERT is mainly used in CUDA Kernel or HOSTDEVICE function. -#define PADDLE_ASSERT_MSG(_IS_NOT_ERROR, __FORMAT, ...) \ - do { \ - if (!(_IS_NOT_ERROR)) { \ - printf("Exception: %s:%d Assertion `%s` failed. " __FORMAT "\n", \ - __FILE__, __LINE__, TOSTRING(_IS_NOT_ERROR), ##__VA_ARGS__); \ - EXIT(); \ - } \ - } while (0) diff --git a/paddle/fluid/platform/enforce.h b/paddle/fluid/platform/enforce.h index 1e0cf0a4055a4d03e5ffbe8ff40baab6f3ff773a..6b1add67334adbd8fbf5347d5c21511329b95b7c 100644 --- a/paddle/fluid/platform/enforce.h +++ b/paddle/fluid/platform/enforce.h @@ -289,6 +289,19 @@ DEFINE_CUDA_STATUS_TYPE(ncclResult_t, ncclSuccess); ::paddle::string::Sprintf(__VA_ARGS__), __FILE__, __LINE__); \ } while (0) +#if defined(__CUDA_ARCH__) +// For cuda, the assertions can affect performance and it is therefore +// recommended to disable them in production code +// https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#assertion +#define PADDLE_ENFORCE(_IS_NOT_ERROR, __FORMAT, ...) \ + do { \ + if (!(_IS_NOT_ERROR)) { \ + printf("Exception: %s:%d Assertion `%s` failed. " __FORMAT "\n", \ + __FILE__, __LINE__, #_IS_NOT_ERROR, ##__VA_ARGS__); \ + asm("trap;"); \ + } \ + } while (0) +#else #define PADDLE_ENFORCE(COND, ...) \ do { \ auto __cond__ = (COND); \ @@ -302,6 +315,7 @@ DEFINE_CUDA_STATUS_TYPE(ncclResult_t, ncclSuccess); } \ } \ } while (0) +#endif #ifdef PADDLE_WITH_CUDA #define PADDLE_ENFORCE_CUDA_SUCCESS(COND, ...) \