From 8bcd1faffcbe17f1879a18b04bab1bbf5a0eadd2 Mon Sep 17 00:00:00 2001 From: fengjiayi Date: Fri, 14 Jul 2017 18:12:14 +0800 Subject: [PATCH 1/5] refactor product(DDim ddim) --- paddle/framework/ddim.cc | 15 +++++++++------ paddle/framework/ddim_test.cc | 3 +++ 2 files changed, 12 insertions(+), 6 deletions(-) diff --git a/paddle/framework/ddim.cc b/paddle/framework/ddim.cc index 3f949a6595..a1ae079f4a 100644 --- a/paddle/framework/ddim.cc +++ b/paddle/framework/ddim.cc @@ -178,13 +178,16 @@ std::vector vectorize(const DDim& ddim) { return result; } -ssize_t product(const DDim& ddim) { - ssize_t result = 1; - std::vector v = vectorize(ddim); - for (auto i : v) { - result *= i; +struct ProductVisitor : public boost::static_visitor { + template + ssize_t operator()(const Dim& dim) { + return product(dim); } - return result; +}; + +ssize_t product(const DDim& ddim) { + ProductVisitor visitor; + return boost::apply_visitor(visitor, ddim); } ///\cond HIDDEN diff --git a/paddle/framework/ddim_test.cc b/paddle/framework/ddim_test.cc index 36eef02370..8ce7886f8a 100644 --- a/paddle/framework/ddim_test.cc +++ b/paddle/framework/ddim_test.cc @@ -52,6 +52,9 @@ TEST(DDim, Equality) { // product of a DDim EXPECT_EQ(paddle::framework::product(vddim), 45); + EXPECT_EQ( + paddle::framework::product(paddle::framework::make_ddim({3, 2, 5, 3})), + 90); } TEST(DDim, Print) { From ee90c2d22b2ff72fcc9983583bdb78962cb5ef72 Mon Sep 17 00:00:00 2001 From: fengjiayi Date: Sat, 15 Jul 2017 12:44:01 +0800 Subject: [PATCH 2/5] add slice_dim draft --- paddle/framework/dim.h | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/paddle/framework/dim.h b/paddle/framework/dim.h index 883fdc55eb..8dc1bab06d 100644 --- a/paddle/framework/dim.h +++ b/paddle/framework/dim.h @@ -401,5 +401,20 @@ HOSTDEVICE Dim linear_to_dimension(int linear_index, Dim extents) { return result; } +template +Dim slice(const Dim& dim, int begin, int end) { + PADDLE_ENFORCE(begin < end, + "Begin index must be less than end index in Dim slice."); + PADDLE_ENFORCE(begin >= 0 && end <= S && end - begin == D, + "Index error occurs in Dim slice."); + if (begin > 0) { + return slice(dim.tail, begin - 1, end - 1); + } + if (D == 1) { + return Dim<1>(dim.head); + } + return Dim(dim.head, slice(dim.tail, 0, end - 1)); +} + } // namespace framework } // namespace paddle From 85c4f488f32d5ca0dfc420f4f7b29fb5a969d946 Mon Sep 17 00:00:00 2001 From: fengjiayi Date: Sun, 16 Jul 2017 15:17:31 +0800 Subject: [PATCH 3/5] Refactor DDim's product() and add slice_ddim() 1. Refactor DDim's product() to make it more efficiently. 2. Add slice_ddim(). --- paddle/framework/ddim.cc | 41 +++++++++ paddle/framework/ddim.h | 2 + paddle/framework/ddim_test.cc | 17 ++++ paddle/framework/dim.h | 15 ---- paddle/framework/dim_test.cu | 163 +++++++++++++++++----------------- 5 files changed, 142 insertions(+), 96 deletions(-) diff --git a/paddle/framework/ddim.cc b/paddle/framework/ddim.cc index a1ae079f4a..c898b6e322 100644 --- a/paddle/framework/ddim.cc +++ b/paddle/framework/ddim.cc @@ -1,4 +1,5 @@ #include "paddle/framework/ddim.h" +#include "paddle/framework/enforce.h" namespace paddle { namespace framework { @@ -190,6 +191,46 @@ ssize_t product(const DDim& ddim) { return boost::apply_visitor(visitor, ddim); } +struct SliceVectorizeVisitor : public boost::static_visitor<> { + std::vector& vector; + int begin; + int end; + + SliceVectorizeVisitor(std::vector& v, int b, int e) + : vector(v), begin(b), end(e) { + PADDLE_ENFORCE(begin < end, + "Begin index must be less than end index in ddim slice."); + PADDLE_ENFORCE(begin >= 0, + "Begin index can't be less than zero in ddim slice."); + } + + template + void operator()(const Dim& dim) { + if (begin == 0) { + vector.push_back(dim.head); + } else { + --begin; + } + --end; + if (end > 0) { + this->operator()(dim.tail); + } + } + + void operator()(const Dim<1>& dim) { + PADDLE_ENFORCE(end == 1, "End index in ddim slice is out of bound."); + vector.push_back(dim.head); + } +}; + +DDim slice_ddim(const DDim& dim, int begin, int end) { + std::vector vec; + vec.reserve(end - begin); + SliceVectorizeVisitor visitor(vec, begin, end); + boost::apply_visitor(visitor, dim); + return make_ddim(vec); +} + ///\cond HIDDEN struct ArityVisitor : boost::static_visitor { diff --git a/paddle/framework/ddim.h b/paddle/framework/ddim.h index 223c4180be..675f8680f6 100644 --- a/paddle/framework/ddim.h +++ b/paddle/framework/ddim.h @@ -81,6 +81,8 @@ std::vector vectorize(const DDim& ddim); ssize_t product(const DDim& ddim); +DDim slice_ddim(const DDim& dim, int begin, int end); + /** * \brief What is the length of this dimension? * diff --git a/paddle/framework/ddim_test.cc b/paddle/framework/ddim_test.cc index 8ce7886f8a..408905b00b 100644 --- a/paddle/framework/ddim_test.cc +++ b/paddle/framework/ddim_test.cc @@ -55,6 +55,23 @@ TEST(DDim, Equality) { EXPECT_EQ( paddle::framework::product(paddle::framework::make_ddim({3, 2, 5, 3})), 90); + + // slice a DDim + paddle::framework::DDim ddim2 = + paddle::framework::make_ddim({1, 2, 3, 4, 5, 6}); + paddle::framework ::DDim ss = paddle::framework::slice_ddim(ddim2, 2, 5); + EXPECT_EQ(arity(ss), 3); + EXPECT_EQ(ss[0], 3); + EXPECT_EQ(ss[1], 4); + EXPECT_EQ(ss[2], 5); + paddle::framework ::DDim ss2 = paddle::framework::slice_ddim(ddim2, 0, 6); + EXPECT_EQ(arity(ss2), 6); + EXPECT_EQ(ss2[0], 1); + EXPECT_EQ(ss2[1], 2); + EXPECT_EQ(ss2[2], 3); + EXPECT_EQ(ss2[3], 4); + EXPECT_EQ(ss2[4], 5); + EXPECT_EQ(ss2[5], 6); } TEST(DDim, Print) { diff --git a/paddle/framework/dim.h b/paddle/framework/dim.h index 8dc1bab06d..883fdc55eb 100644 --- a/paddle/framework/dim.h +++ b/paddle/framework/dim.h @@ -401,20 +401,5 @@ HOSTDEVICE Dim linear_to_dimension(int linear_index, Dim extents) { return result; } -template -Dim slice(const Dim& dim, int begin, int end) { - PADDLE_ENFORCE(begin < end, - "Begin index must be less than end index in Dim slice."); - PADDLE_ENFORCE(begin >= 0 && end <= S && end - begin == D, - "Index error occurs in Dim slice."); - if (begin > 0) { - return slice(dim.tail, begin - 1, end - 1); - } - if (D == 1) { - return Dim<1>(dim.head); - } - return Dim(dim.head, slice(dim.tail, 0, end - 1)); -} - } // namespace framework } // namespace paddle diff --git a/paddle/framework/dim_test.cu b/paddle/framework/dim_test.cu index 0521741519..3898d0a447 100644 --- a/paddle/framework/dim_test.cu +++ b/paddle/framework/dim_test.cu @@ -1,100 +1,101 @@ #include #include -#include "paddle/framework/dim.h" #include "gtest/gtest.h" +#include "paddle/framework/dim.h" __global__ void test(paddle::framework::Dim<2>* o) { - o[0] = paddle::framework::make_dim(5, 6); + o[0] = paddle::framework::make_dim(5, 6); } __global__ void dyn_idx_gpu(int* o) { - auto d = paddle::framework::make_dim(5, 6); - o[0] = d[1]; + auto d = paddle::framework::make_dim(5, 6); + o[0] = d[1]; } 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); - - // 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); - - // 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); - - // dynamic access - paddle::framework::get(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[1], 11); - - // dynamic access on GPU - thrust::device_vector r(1); - dyn_idx_gpu<<<1,1>>>(thrust::raw_pointer_cast(r.data())); - int 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); + // 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); + + // 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); + + // 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); + + // dynamic access + paddle::framework::get(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[1], 11); + + // dynamic access on GPU + thrust::device_vector r(1); + dyn_idx_gpu<<<1, 1>>>(thrust::raw_pointer_cast(r.data())); + int 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) { - auto a = paddle::framework::make_dim(3, 4); - 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); - EXPECT_TRUE(a == c); + auto a = paddle::framework::make_dim(3, 4); + 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); + EXPECT_TRUE(a == c); } TEST(Dim, Print) { - { - std::stringstream ss; - auto a = paddle::framework::make_dim(2, 3); - ss << a; - EXPECT_EQ(ss.str(), "2, 3"); - } - { - std::stringstream ss; - ss << paddle::framework::make_dim(8); - EXPECT_EQ(ss.str(), "8"); - } + { + std::stringstream ss; + auto a = paddle::framework::make_dim(2, 3); + ss << a; + EXPECT_EQ(ss.str(), "2, 3"); + } + { + std::stringstream ss; + ss << paddle::framework::make_dim(8); + EXPECT_EQ(ss.str(), "8"); + } } From 7de10ee7e2c9febbe81b22ad83e3e4584ba4a988 Mon Sep 17 00:00:00 2001 From: fengjiayi Date: Sun, 16 Jul 2017 16:35:44 +0800 Subject: [PATCH 4/5] formate code --- paddle/framework/ddim_test.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/paddle/framework/ddim_test.cc b/paddle/framework/ddim_test.cc index 408905b00b..6a099f2aeb 100644 --- a/paddle/framework/ddim_test.cc +++ b/paddle/framework/ddim_test.cc @@ -59,12 +59,12 @@ TEST(DDim, Equality) { // slice a DDim paddle::framework::DDim ddim2 = paddle::framework::make_ddim({1, 2, 3, 4, 5, 6}); - paddle::framework ::DDim ss = paddle::framework::slice_ddim(ddim2, 2, 5); + paddle::framework::DDim ss = paddle::framework::slice_ddim(ddim2, 2, 5); EXPECT_EQ(arity(ss), 3); EXPECT_EQ(ss[0], 3); EXPECT_EQ(ss[1], 4); EXPECT_EQ(ss[2], 5); - paddle::framework ::DDim ss2 = paddle::framework::slice_ddim(ddim2, 0, 6); + paddle::framework::DDim ss2 = paddle::framework::slice_ddim(ddim2, 0, 6); EXPECT_EQ(arity(ss2), 6); EXPECT_EQ(ss2[0], 1); EXPECT_EQ(ss2[1], 2); From 778a1a9f8780405a32814d3f95accccb4304dc87 Mon Sep 17 00:00:00 2001 From: fengjiayi Date: Mon, 17 Jul 2017 12:03:46 +0800 Subject: [PATCH 5/5] add a comment for --- paddle/framework/ddim.h | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/paddle/framework/ddim.h b/paddle/framework/ddim.h index 675f8680f6..df26d73d48 100644 --- a/paddle/framework/ddim.h +++ b/paddle/framework/ddim.h @@ -81,6 +81,13 @@ std::vector vectorize(const DDim& ddim); ssize_t product(const DDim& ddim); +/** + * \brief Slice a ddim + * + * Slice dim with [begin, end). + * e.g. DDim d = make_ddim({1,2,3,4,5}); + * slice_ddim(d, 1, 3); ====> {2,3} + */ DDim slice_ddim(const DDim& dim, int begin, int end); /**