diff --git a/paddle/fluid/framework/ddim.cc b/paddle/fluid/framework/ddim.cc index bbc9982d9db4cd5bec872b44d2385afccd77ffd3..d4e469247f24b3091f94df1abb365b4933f3dcc7 100644 --- a/paddle/fluid/framework/ddim.cc +++ b/paddle/fluid/framework/ddim.cc @@ -48,52 +48,6 @@ bool DDim::operator==(const DDim& d) const { bool DDim::operator!=(const DDim& d) const { return !(*this == d); } -struct DDimPlusVisitor { - explicit DDimPlusVisitor(const int64_t* d1, const int64_t* d2) - : d1_(d1), d2_(d2) {} - - template - inline void operator()(Dim& self) const { - UnrollAdd::Run(d1_, d2_, self.GetMutable()); - } - - const int64_t* d1_; - const int64_t* d2_; -}; - -DDim DDim::operator+(const DDim& d) const { - PADDLE_ENFORCE(size() == d.size()); - DDim ret; - ret.rank_ = rank_; - ret.apply_visitor(DDimPlusVisitor(Get(), d.Get())); - return ret; -} - -struct DDimMulVisitor { - explicit DDimMulVisitor(const int64_t* d1, const int64_t* d2) - : d1_(d1), d2_(d2) {} - - template - inline void operator()(Dim& self) const { - UnrollMul::Run(d1_, d2_, self.GetMutable()); - } - - const int64_t* d1_; - const int64_t* d2_; -}; - -DDim DDim::operator*(const DDim& d) const { - PADDLE_ENFORCE(size() == d.size()); - DDim ret; - ret.rank_ = rank_; - ret.apply_visitor(DDimMulVisitor(Get(), d.Get())); - return ret; -} - -int64_t get(const DDim& ddim, int idx) { return ddim[idx]; } - -void set(DDim& ddim, int idx, int value) { ddim[idx] = value; } // NOLINT - std::vector vectorize(const DDim& ddim) { std::vector result(DDim::kMaxRank); dynamic_dim_assign(ddim.Get(), result.data(), ddim.size()); diff --git a/paddle/fluid/framework/ddim.h b/paddle/fluid/framework/ddim.h index 7d2e296b6c1a99180acc105eb73754233cfa15f4..bfe3e55a73da42b0865d4e02b59ec6be90ea0096 100644 --- a/paddle/fluid/framework/ddim.h +++ b/paddle/fluid/framework/ddim.h @@ -117,10 +117,6 @@ class DDim { bool operator!=(const DDim& d) const; - DDim operator+(const DDim& d) const; - - DDim operator*(const DDim& d) const; - inline const int64_t* Get() const { return dim_.Get(); } inline int64_t* GetMutable() { return dim_.GetMutable(); } @@ -174,9 +170,6 @@ DDim make_ddim(const std::vector& dims); */ DDim make_ddim(std::initializer_list dims); -int64_t get(const DDim& dim, int idx); -void set(DDim& dim, int idx, int val); // NOLINT - std::vector vectorize(const DDim& ddim); std::vector vectorize2int(const DDim& ddim); diff --git a/paddle/fluid/framework/ddim_test.cc b/paddle/fluid/framework/ddim_test.cc index c1eb3210fdc5ae1b0ddb9736184e152fdd10a242..b7b42fa019f6b2243719ebdb1628ea83c220abf9 100644 --- a/paddle/fluid/framework/ddim_test.cc +++ b/paddle/fluid/framework/ddim_test.cc @@ -34,8 +34,8 @@ TEST(DDim, Equality) { // mutate a DDim ddim[1] = 2; EXPECT_EQ(ddim[1], 2); - paddle::framework::set(ddim, 0, 6); - EXPECT_EQ(paddle::framework::get(ddim, 0), 6); + ddim[0] = 6; + EXPECT_EQ(ddim[0], 6); // vectorize a DDim std::vector res_vec = paddle::framework::vectorize(vddim); @@ -48,18 +48,6 @@ TEST(DDim, Equality) { EXPECT_EQ(res_vec[1], 2); EXPECT_EQ(res_vec[2], 1); - // add two DDims - paddle::framework::DDim ddim_sum = ddim + vddim; - EXPECT_EQ(ddim_sum[0], 15); - EXPECT_EQ(ddim_sum[1], 3); - EXPECT_EQ(ddim_sum[2], 10); - - // multiply two DDims - paddle::framework::DDim ddim_mul = ddim * vddim; - EXPECT_EQ(ddim_mul[0], 54); - EXPECT_EQ(ddim_mul[1], 2); - EXPECT_EQ(ddim_mul[2], 25); - // arity of a DDim EXPECT_EQ(paddle::framework::arity(ddim), 3); EXPECT_EQ(ddim.size(), 3); diff --git a/paddle/fluid/framework/details/CMakeLists.txt b/paddle/fluid/framework/details/CMakeLists.txt index 4eba8177c56b818d5890504c7a5c69e3e317d559..30f61ad1a4c5bf80b560cbfa1d05414bfcc4988d 100644 --- a/paddle/fluid/framework/details/CMakeLists.txt +++ b/paddle/fluid/framework/details/CMakeLists.txt @@ -94,5 +94,4 @@ cc_library(build_strategy SRCS build_strategy.cc DEPS fuse_relu_depthwise_conv_pass memory_optimize_pass lock_free_optimize_pass alloc_continuous_space_for_grad_pass fuse_all_reduce_op_pass backward_optimizer_op_deps_pass - fuse_adam_op_pass fuse_sgd_op_pass fuse_momentum_op_pass - record_skip_memory_opt_vars_pass) + fuse_adam_op_pass fuse_sgd_op_pass fuse_momentum_op_pass record_skip_memory_opt_vars_pass) diff --git a/paddle/fluid/framework/dim.h b/paddle/fluid/framework/dim.h index 88aee8379d835ce88b6b348aca99eb4a35bbeb5c..531ee8445b78f2ffdb4f1197fdee9c5cf0b79081 100644 --- a/paddle/fluid/framework/dim.h +++ b/paddle/fluid/framework/dim.h @@ -45,10 +45,6 @@ class Dim : public Array { HOSTDEVICE explicit Dim(int64_t head, Args... args) : BaseClass(head, args...) {} - /** Construct a Dim from a linear index and size. Uses Fortran order - * indexing. */ - HOSTDEVICE Dim(int64_t idx, const Dim& size); - /** Construct a Dim with each dimension set to the given index */ HOSTDEVICE explicit Dim(int64_t idx) { this->Fill(idx); } @@ -57,181 +53,12 @@ class Dim : public Array { HOST std::string to_string() const; }; -namespace detail { -template -struct FortranOrderIndexingConstructorFunctor { - HOSTDEVICE inline static void Run(const int64_t* in, int64_t* idx, - int64_t* out) { - out[kStart] = (*idx) % in[kStart]; - (*idx) /= in[kStart]; - FortranOrderIndexingConstructorFunctor::Run(in, idx, - out); - } -}; - -template -struct FortranOrderIndexingConstructorFunctor { - HOSTDEVICE inline static void Run(const int64_t* in, int64_t* idx, - int64_t* out) {} -}; -} // namespace detail - -template -HOSTDEVICE Dim::Dim(int64_t idx, const Dim& size) { - detail::FortranOrderIndexingConstructorFunctor<0, D, D == 0>::Run( - size.Get(), &idx, this->GetMutable()); -} - -template -HOSTDEVICE inline int64_t get(const Dim& dim) { - return dim[idx]; -} - -template -HOSTDEVICE inline int64_t& get(Dim& dim) { // NOLINT - return dim[idx]; -} - -template -HOSTDEVICE inline int64_t get(const Dim& dim, int idx) { - return dim[idx]; -} - -template -HOSTDEVICE inline int64_t& get(Dim& dim, int idx) { // NOLINT - return dim[idx]; -} - -// Dot product of two dims -template -HOSTDEVICE inline int64_t linearize(const Dim& a, const Dim& b) { - return UnrollProduct::Run(a.Get(), b.Get()); -} - // Product of a Dim template HOSTDEVICE inline int64_t product(const Dim& a) { return UnrollProduct::Run(a.Get()); } -// Is 0 <= idx_i < size_i for all i? -namespace detail { -template -struct ContainedFunctor { - HOSTDEVICE static inline bool Run(const int64_t* idx, const int64_t* size) { - return (idx[kStart] >= 0 && idx[kStart] < size[kStart]) && - ContainedFunctor::Run(idx, - size); - } -}; - -template -struct ContainedFunctor { - HOSTDEVICE static constexpr inline bool Run(const int64_t* idx, - const int64_t* size) { - return true; - } -}; -} // namespace detail - -template -HOSTDEVICE inline bool contained(const Dim& idx, const Dim& size) { - return detail::ContainedFunctor<0, D, D == 0>::Run(idx.Get(), size.Get()); -} - -/** - * \brief Compute exclusive prefix-multiply of a Dim. - */ -namespace detail { -template -struct ExPrefixMulFunctor { - HOSTDEVICE static inline void Run(const int64_t* in, int64_t* out) { - kStart == 0 ? out[kStart] = 1 : out[kStart] = - out[kStart - 1] * in[kStart - 1]; - detail::ExPrefixMulFunctor::Run(in, - out); - } -}; - -template -struct ExPrefixMulFunctor { - HOSTDEVICE static inline void Run(const int64_t* in, int64_t* out) {} -}; -} // namespace detail - -template -HOSTDEVICE inline Dim ex_prefix_mul(const Dim& src) { - Dim ret; - detail::ExPrefixMulFunctor<0, D, D == 0>::Run(src.Get(), ret.GetMutable()); - return ret; -} - -/** - * Add two dimensions together - */ -template -HOSTDEVICE inline Dim dim_plus(const Dim& a, const Dim& b) { - Dim ret; - UnrollAdd::Run(a.Get(), b.Get(), ret.GetMutable()); - return ret; -} - -template -HOSTDEVICE inline Dim operator+(const Dim& lhs, const Dim& rhs) { - return dim_plus(lhs, rhs); -} - -/** - * Multiply two dimensions together - */ -template -HOSTDEVICE inline Dim dim_mult(const Dim& a, const Dim& b) { - Dim ret; - UnrollMul::Run(a.Get(), b.Get(), ret.GetMutable()); - return ret; -} - -template -HOSTDEVICE Dim operator*(const Dim& lhs, const Dim& rhs) { - return dim_mult(lhs, rhs); -} - -/** - * \brief Normalize strides to ensure any dimension with extent 1 - * has stride 0. - * - * \param size Dim object containing the size of an array - * \param stride Dim object containing stride of an array - * \return Dim object the same size as \p size with normalized strides - * - */ -namespace detail { -template -struct NormalizeStridesFunctor { - HOSTDEVICE static void Run(const int64_t* size, const int64_t* stride, - int64_t* ret) { - ret[kStart] = (size[kStart] == 1 ? 0 : stride[kStart]); - NormalizeStridesFunctor::Run( - size, stride, ret); - } -}; - -template -struct NormalizeStridesFunctor { - HOSTDEVICE static void Run(const int64_t* size, const int64_t* stride, - int64_t* ret) {} -}; -} // namespace detail - -template -HOSTDEVICE Dim normalize_strides(const Dim& size, const Dim& stride) { - Dim ret; - detail::NormalizeStridesFunctor<0, D, D == 0>::Run(size.Get(), stride.Get(), - ret.GetMutable()); - return ret; -} - /** * Helper function to create a Dim * @@ -265,20 +92,6 @@ HOST std::string Dim::to_string() const { return stream.str(); } -template -HOSTDEVICE Dim linear_to_dimension(int linear_index, const Dim& extents) { - Dim result; - - for (int i = 0; i < D - 1; ++i) { - result[i] = linear_index % extents[i]; - linear_index /= extents[i]; - } - - result[D - 1] = linear_index; - - return result; -} - template inline void static_dim_assign(const T1* in, T2* out) { UnrollAssign::Run(in, out); diff --git a/paddle/fluid/framework/dim_test.cu b/paddle/fluid/framework/dim_test.cu index 0f384d12e6f04a96add03ccdbe7b11156c6cc67c..7add6d140c7e0942fca22df0c118e0f15460fb07 100644 --- a/paddle/fluid/framework/dim_test.cu +++ b/paddle/fluid/framework/dim_test.cu @@ -29,34 +29,28 @@ __global__ void dyn_idx_gpu(int64_t* o) { TEST(Dim, Equality) { // construct a Dim on the CPU auto a = paddle::framework::make_dim(3, 4); - EXPECT_EQ(paddle::framework::get<0>(a), 3); - EXPECT_EQ(paddle::framework::get<1>(a), 4); + EXPECT_EQ(a[0], 3); + EXPECT_EQ(a[1], 4); // construct a Dim on the GPU thrust::device_vector> t(2); test<<<1, 1>>>(thrust::raw_pointer_cast(t.data())); a = t[0]; - EXPECT_EQ(paddle::framework::get<0>(a), 5); - EXPECT_EQ(paddle::framework::get<1>(a), 6); - - // linearization - auto b = paddle::framework::make_dim(7, 8); - EXPECT_EQ(paddle::framework::linearize(a, b), 83); + EXPECT_EQ(a[0], 5); + EXPECT_EQ(a[1], 6); // product EXPECT_EQ(paddle::framework::product(a), 30); // mutate a Dim - paddle::framework::get<1>(b) = 10; - EXPECT_EQ(paddle::framework::get<0>(b), 7); - EXPECT_EQ(paddle::framework::get<1>(b), 10); + auto b = paddle::framework::make_dim(7, 8); + b[1] = 10; + EXPECT_EQ(b[0], 7); + EXPECT_EQ(b[1], 10); - // dynamic access - paddle::framework::get(b, 0) = 8; + b[0] = 8; b[1] = 11; - EXPECT_EQ(paddle::framework::get<0>(b), 8); - EXPECT_EQ(paddle::framework::get<1>(b), 11); - EXPECT_EQ(paddle::framework::get(b, 0), 8); + EXPECT_EQ(b[0], 8); EXPECT_EQ(b[1], 11); // dynamic access on GPU @@ -64,24 +58,6 @@ TEST(Dim, Equality) { dyn_idx_gpu<<<1, 1>>>(thrust::raw_pointer_cast(r.data())); int64_t res = r[0]; EXPECT_EQ(res, 6); - - // ex_prefix_mul - paddle::framework::Dim<3> c = - paddle::framework::ex_prefix_mul(paddle::framework::Dim<3>(3, 4, 5)); - EXPECT_EQ(paddle::framework::get<0>(c), 1); - EXPECT_EQ(paddle::framework::get<1>(c), 3); - EXPECT_EQ(paddle::framework::get<2>(c), 12); - - // generate from an index - auto size = paddle::framework::make_dim(4, 5, 2); - c = paddle::framework::Dim<3>(14, size); - EXPECT_EQ(paddle::framework::get<0>(c), 2); - EXPECT_EQ(paddle::framework::get<1>(c), 3); - EXPECT_EQ(paddle::framework::get<2>(c), 0); - c = paddle::framework::Dim<3>(25, size); - EXPECT_EQ(paddle::framework::get<0>(c), 1); - EXPECT_EQ(paddle::framework::get<1>(c), 1); - EXPECT_EQ(paddle::framework::get<2>(c), 1); } TEST(Dim, Bool) { @@ -89,10 +65,6 @@ TEST(Dim, Bool) { auto b = paddle::framework::make_dim(5, 6); auto c = paddle::framework::make_dim(3, 4); - // in_bounds check - EXPECT_TRUE(paddle::framework::contained(a, b)); - EXPECT_FALSE(paddle::framework::contained(b, a)); - // comparison EXPECT_TRUE(a == a); EXPECT_FALSE(a == b); diff --git a/paddle/fluid/framework/unroll_array_ops.h b/paddle/fluid/framework/unroll_array_ops.h index 731da74eff4d22da6730e589a1af919514f1c4b7..ab17641080551166f246c69d42ed85df3892a3e0 100644 --- a/paddle/fluid/framework/unroll_array_ops.h +++ b/paddle/fluid/framework/unroll_array_ops.h @@ -94,36 +94,6 @@ struct UnrollCompare { } }; -template -struct UnrollAdd { - template - HOSTDEVICE inline static void Run(const T *d1, const T *d2, T *d3) { - d3[kStart] = d1[kStart] + d2[kStart]; - UnrollAdd::Run(d1, d2, d3); - } -}; - -template -struct UnrollAdd { - template - HOSTDEVICE inline static void Run(const T *d1, const T *d2, T *d3) {} -}; - -template -struct UnrollMul { - template - HOSTDEVICE inline static void Run(const T *d1, const T *d2, T *d3) { - d3[kStart] = d1[kStart] * d2[kStart]; - UnrollMul::Run(d1, d2, d3); - } -}; - -template -struct UnrollMul { - template - HOSTDEVICE inline static void Run(const T *d1, const T *d2, T *d3) {} -}; - template struct UnrollProduct { template @@ -131,12 +101,6 @@ struct UnrollProduct { return d[kStart] * UnrollProduct::Run(d); } - - template - HOSTDEVICE inline static T Run(const T *d1, const T *d2) { - return d1[kStart] * d2[kStart] + - UnrollProduct::Run(d1, d2); - } }; template @@ -145,11 +109,6 @@ struct UnrollProduct { HOSTDEVICE inline constexpr static T Run(const T *d) { return 1; } - - template - HOSTDEVICE inline constexpr static T Run(const T *d1, const T *d2) { - return 0; - } }; } // namespace detail @@ -166,12 +125,6 @@ using UnrollVarArgsAssign = detail::UnrollVarArgsAssign; template using UnrollCompare = detail::UnrollCompare<0, N, N == 0>; -template -using UnrollAdd = detail::UnrollAdd<0, N, N == 0>; - -template -using UnrollMul = detail::UnrollMul<0, N, N == 0>; - template using UnrollProduct = detail::UnrollProduct<0, N, N == 0>; diff --git a/paddle/fluid/framework/unroll_array_ops_test.cc b/paddle/fluid/framework/unroll_array_ops_test.cc index 51433c83c801765d8df10590abdd319ba60e4873..be811478eec17e0986ae7579ff323d94dea3155a 100644 --- a/paddle/fluid/framework/unroll_array_ops_test.cc +++ b/paddle/fluid/framework/unroll_array_ops_test.cc @@ -74,34 +74,9 @@ TEST(unroll_ops, compare) { EXPECT_FALSE(UnrollCompare<1>::Run(a, b)); } -TEST(unroll_ops, add) { - int a[] = {2, 3, 4}; - int b[] = {5, 10, 102}; - int c[] = {0, 0, 0}; - UnrollAdd<2>::Run(a, b, c); - EXPECT_EQ(a[0] + b[0], c[0]); - EXPECT_EQ(a[1] + b[1], c[1]); - EXPECT_EQ(c[2], 0); -} - -TEST(unroll_ops, mul) { - int a[] = {2, 3, 4}; - int b[] = {5, 10, 102}; - int c[] = {0, 0, 0}; - UnrollMul<2>::Run(a, b, c); - EXPECT_EQ(a[0] * b[0], c[0]); - EXPECT_EQ(a[1] * b[1], c[1]); - EXPECT_EQ(c[2], 0); -} - TEST(unroll_ops, product) { int a[] = {2, 3, 4}; - int b[] = {5, 10, 102}; - EXPECT_EQ(UnrollProduct<3>::Run(a), a[0] * a[1] * a[2]); - - EXPECT_EQ(UnrollProduct<3>::Run(a, b), - a[0] * b[0] + a[1] * b[1] + a[2] * b[2]); } } // namespace framework diff --git a/paddle/fluid/operators/hierarchical_sigmoid_op.h b/paddle/fluid/operators/hierarchical_sigmoid_op.h index 7cfe0aabcb7f3ce86ccc3a9a1c54b3b60d384aa1..a0af514509d87ce64ea4abab687a0f03607f7fc1 100644 --- a/paddle/fluid/operators/hierarchical_sigmoid_op.h +++ b/paddle/fluid/operators/hierarchical_sigmoid_op.h @@ -248,8 +248,7 @@ class HierarchicalSigmoidGradOpKernel : public framework::OpKernel { w_grad->set_height(w.dims()[0]); auto* w_grad_value = w_grad->mutable_value(); framework::DDim temp_dim(w.dims()); - set(temp_dim, 0, real_rows.size()); - + temp_dim[0] = real_rows.size(); w_grad_value->mutable_data(temp_dim, ctx.GetPlace()); zero(dev_ctx, w_grad_value, static_cast(0.0)); bit_code->MulGradWeight(pre_out_grad, w_grad, in); diff --git a/paddle/fluid/platform/place.cc b/paddle/fluid/platform/place.cc index 60b2d83f15746eab0a4d29c7965c064690b6d46d..50a401bb160ce71aca4cbddb321849c328f8f686 100644 --- a/paddle/fluid/platform/place.cc +++ b/paddle/fluid/platform/place.cc @@ -40,15 +40,6 @@ class PlacePrinter : public boost::static_visitor<> { } // namespace detail -static Place the_default_place; - -void set_place(const Place &place) { the_default_place = place; } -const Place &get_place() { return the_default_place; } - -const CUDAPlace default_gpu() { return CUDAPlace(0); } -const CPUPlace default_cpu() { return CPUPlace(); } -const CUDAPinnedPlace default_cuda_pinned() { return CUDAPinnedPlace(); } - bool is_gpu_place(const Place &p) { return boost::apply_visitor(IsCUDAPlace(), p); } diff --git a/paddle/fluid/platform/place.h b/paddle/fluid/platform/place.h index a095d4929ec2130b4af48d32bf016d9fe108b418..daa70e943ded1909785458833349e4f8091847d1 100644 --- a/paddle/fluid/platform/place.h +++ b/paddle/fluid/platform/place.h @@ -80,13 +80,6 @@ typedef boost::variant Place; using PlaceList = std::vector; -void set_place(const Place &); -const Place &get_place(); - -const CUDAPlace default_gpu(); -const CPUPlace default_cpu(); -const CUDAPinnedPlace default_cuda_pinned(); - bool is_gpu_place(const Place &); bool is_cpu_place(const Place &); bool is_cuda_pinned_place(const Place &); diff --git a/paddle/fluid/platform/place_test.cc b/paddle/fluid/platform/place_test.cc index 6a919c56252243ddfbd653b768bcea3bdea77831..e4c1d3def90f191194b46bb9ea27dd27d69dcb8b 100644 --- a/paddle/fluid/platform/place_test.cc +++ b/paddle/fluid/platform/place_test.cc @@ -30,16 +30,6 @@ TEST(Place, Equality) { EXPECT_FALSE(paddle::platform::places_are_same_class(g0, cpu)); } -TEST(Place, Default) { - EXPECT_TRUE(paddle::platform::is_gpu_place(paddle::platform::get_place())); - EXPECT_TRUE(paddle::platform::is_gpu_place(paddle::platform::default_gpu())); - EXPECT_TRUE(paddle::platform::is_cpu_place(paddle::platform::default_cpu())); - - EXPECT_FALSE(paddle::platform::is_cpu_place(paddle::platform::get_place())); - paddle::platform::set_place(paddle::platform::CPUPlace()); - EXPECT_TRUE(paddle::platform::is_cpu_place(paddle::platform::get_place())); -} - TEST(Place, Print) { { std::stringstream ss;