From b2d2ea537b2e0d6b36e266fd5ae1ab2de7c6dddd Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 21 Sep 2022 02:44:23 +0800 Subject: [PATCH 01/55] add sparse reshape --- paddle/phi/api/yaml/sparse_backward.yaml | 11 ++ paddle/phi/api/yaml/sparse_ops.yaml | 11 ++ paddle/phi/infermeta/sparse/unary.cc | 8 + paddle/phi/infermeta/sparse/unary.h | 4 + .../kernels/sparse/cpu/reshape_grad_kernel.cc | 82 +++++++++ .../phi/kernels/sparse/cpu/reshape_kernel.cc | 140 +++++++++++++++ .../kernels/sparse/cpu/sparse_utils_kernel.cc | 6 +- .../kernels/sparse/gpu/reshape_grad_kernel.cu | 107 +++++++++++ .../phi/kernels/sparse/gpu/reshape_kernel.cu | 169 ++++++++++++++++++ .../kernels/sparse/gpu/sparse_utils_kernel.cu | 6 +- paddle/phi/kernels/sparse/unary_grad_kernel.h | 14 ++ paddle/phi/kernels/sparse/unary_kernel.h | 45 +++++ .../tests/unittests/test_sparse_reshape_op.py | 78 ++++++++ python/paddle/incubate/sparse/__init__.py | 2 + python/paddle/incubate/sparse/unary.py | 23 +++ 15 files changed, 702 insertions(+), 4 deletions(-) create mode 100644 paddle/phi/kernels/sparse/cpu/reshape_grad_kernel.cc create mode 100644 paddle/phi/kernels/sparse/cpu/reshape_kernel.cc create mode 100644 paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu create mode 100644 paddle/phi/kernels/sparse/gpu/reshape_kernel.cu create mode 100644 python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py diff --git a/paddle/phi/api/yaml/sparse_backward.yaml b/paddle/phi/api/yaml/sparse_backward.yaml index 41816898c3a50..4592f6f159a9a 100644 --- a/paddle/phi/api/yaml/sparse_backward.yaml +++ b/paddle/phi/api/yaml/sparse_backward.yaml @@ -404,3 +404,14 @@ func : fused_attention_csr_grad{dense, dense, dense, sparse_csr, dense -> dense, dense, dense} layout : softmax data_type: query + +- backward_op : reshape_grad + forward : reshape(Tensor x, int64_t[] new_shape) -> Tensor(out) + args : (Tensor x, Tensor out_grad) + output : Tensor(x_grad) + infer_meta : + func : UnchangedInferMeta + param : [x] + kernel : + func : reshape_coo_grad {sparse_coo, sparse_coo -> sparse_coo}, + reshape_csr_grad {sparse_csr, sparse_csr -> sparse_csr} \ No newline at end of file diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index 043c12615fb7f..33c4f6d7f5591 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -456,3 +456,14 @@ mv_csr{sparse_csr, dense -> dense} layout : x backward: mv_grad + +- op : reshape + args : (Tensor x, int64_t[] new_shape) + output : Tensor(out) + infer_meta : + func : sparse::ReshapeInferMeta + kernel : + func : reshape_coo{sparse_coo -> sparse_coo}, + reshape_csr{sparse_csr -> sparse_csr} + layout : x + backward : reshape_grad \ No newline at end of file diff --git a/paddle/phi/infermeta/sparse/unary.cc b/paddle/phi/infermeta/sparse/unary.cc index 45cb4f75e38c9..6037c38fd96ed 100644 --- a/paddle/phi/infermeta/sparse/unary.cc +++ b/paddle/phi/infermeta/sparse/unary.cc @@ -32,5 +32,13 @@ void ValuesInferMeta(const MetaTensor& x, MetaTensor* out) { out->set_layout(x.layout()); } +void ReshapeInferMeta(const MetaTensor& x, + const std::vector& new_shape, + MetaTensor* out) { + out->set_dtype(x.dtype()); + out->set_dims(phi::make_ddim(new_shape)); + out->set_layout(x.layout()); +} + } // namespace sparse } // namespace phi diff --git a/paddle/phi/infermeta/sparse/unary.h b/paddle/phi/infermeta/sparse/unary.h index 880e90b7ae697..360852c559fe4 100644 --- a/paddle/phi/infermeta/sparse/unary.h +++ b/paddle/phi/infermeta/sparse/unary.h @@ -24,5 +24,9 @@ void IndicesInferMeta(const MetaTensor& x, MetaTensor* out); void ValuesInferMeta(const MetaTensor& x, MetaTensor* out); +void ReshapeInferMeta(const MetaTensor& x, + const std::vector& new_shape, + MetaTensor* out); + } // namespace sparse } // namespace phi diff --git a/paddle/phi/kernels/sparse/cpu/reshape_grad_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_grad_kernel.cc new file mode 100644 index 0000000000000..57ae338af8fe9 --- /dev/null +++ b/paddle/phi/kernels/sparse/cpu/reshape_grad_kernel.cc @@ -0,0 +1,82 @@ +// 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 "paddle/phi/kernels/sparse/unary_grad_kernel.h" +#include "paddle/phi/kernels/sparse/unary_kernel.h" + +#include "paddle/phi/backends/cpu/cpu_context.h" +#include "paddle/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/sparse/empty_kernel.h" +#include "paddle/phi/kernels/sparse/impl/unary_grad_kernel_impl.h" + +namespace phi { +namespace sparse { + +template +void ReshapeCooGradKernel(const Context& dev_ctx, + const SparseCooTensor& x, + const SparseCooTensor& dout, + // const std::vector& perm, + SparseCooTensor* dx) { + EmptyLikeCooKernel(dev_ctx, x, dx); + std::vector x_shape(x.dims().size()); + for (int i=0; i(dev_ctx, dout, x_shape, dx); +} + +template +void ReshapeCsrGradKernel(const Context& dev_ctx, + const SparseCsrTensor& x, + const SparseCsrTensor& dout, + // const std::vector& perm, + SparseCsrTensor* dx) { + EmptyLikeCsrKernel(dev_ctx, x, dx); + std::vector x_shape(x.dims().size()); + for (int i=0; i(dev_ctx, dout, x_shape, dx); +} + + +} // namespace sparse +} // namespace phi + +PD_REGISTER_KERNEL(reshape_coo_grad, + CPU, + ALL_LAYOUT, + phi::sparse::ReshapeCooGradKernel, + float, + double, + int8_t, + uint8_t, + int16_t, + int, + int64_t, + bool) {} + +PD_REGISTER_KERNEL(reshape_csr_grad, + CPU, + ALL_LAYOUT, + phi::sparse::ReshapeCsrGradKernel, + float, + double, + int8_t, + uint8_t, + int16_t, + int, + int64_t, + bool) {} diff --git a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc new file mode 100644 index 0000000000000..0a51afc04137c --- /dev/null +++ b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc @@ -0,0 +1,140 @@ +// 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 "paddle/phi/kernels/sparse/unary_kernel.h" + +// #include "paddle/phi/core/ddim.cc" +#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" +// #include "paddle/phi/kernels/sparse/cpu/sparse_utils_kernel.cc" + +#include "paddle/phi/backends/cpu/cpu_context.h" +#include "paddle/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/funcs/eigen/common.h" +#include "paddle/phi/kernels/funcs/eigen/eigen_function.h" +#include "paddle/phi/kernels/sparse/empty_kernel.h" +#include "paddle/phi/kernels/sparse/impl/unary_grad_kernel_impl.h" +#include "paddle/phi/kernels/sparse/impl/unary_kernel_impl.h" + +namespace phi { +namespace sparse { + +template +void ReshapeCooKernel(const Context& dev_ctx, + const SparseCooTensor& x, + // const std::vector& sparse_part_permutation, + const std::vector& new_shape, + SparseCooTensor* out) { + /* + 目前只能针对 sparse part dims 部分进行reshape + */ + // create "out" sparse tensor + int64_t x_nnz = x.nnz(); + // DDim out_dims = x.dims().transpose(perm); + DDim out_dims = phi::make_ddim(new_shape); + ////get sparse part dimensions of x and out + std::vector x_sparse_part_dims; + std::vector out_sparse_part_dims; + for (int i = 0; i < x.sparse_dim(); ++i) { + x_sparse_part_dims.push_back(x.dims()[i]); + } + for (int i = 0; i < out_dims.size() - x.dense_dim(); ++i) { + out_sparse_part_dims.push_back(out_dims[i]); + } + // DenseTensor out_indices = EmptyLike(dev_ctx, x.indices()); + DenseTensor out_indices = Empty(dev_ctx, + {static_cast(out_sparse_part_dims.size()), x_nnz} + ); + DenseTensor out_values(x.values()); + out->SetMember(out_indices, out_values, out_dims, x.coalesced()); + + // compute values of indices + const DenseTensor& x_indices = x.indices(); + const auto* x_indices_data = x_indices.data(); + auto* out_indices_data = out_indices.data(); +// //i 表示 indices 的 行标 +// for (unsigned int i = 0; i < perm.size(); ++i) { +// // j 表示 indices 的 列标 +// for (int64_t j = 0; j < x_nnz; ++j) { +// // 修改 out indices 的索引为 (i, j)的元素值 +// /* Caution : 这是原来的计算逻辑，我认为是 错误的， +// 这里计算逻辑是： 原tensor的shape是 (10, 20, 30, 40, 50) +// 一个非零元素的索引为 (1, 2, 3, 4, 5) +// 进行transpose 后, tensor的shape 是 (30, 10, 50, 20, 40) +// 这里的计算逻辑就认为该非零元素的新索引就是 (3, 1, 5, 2, 4) +// 没错，这就是transpose的计算逻辑，transpose后元素在内存中的位置改变了 +// 你更改的逻辑其实是 reshape的计算逻辑，reshape后所有元素在内存中的位置均不变 +// */ +// out_indices_data[j + i * x_nnz] = x_indices_data[j + perm[i] * x_nnz]; +// } +// } + + // 我的更改后的计算逻辑如下： + const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); + const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); + int64_t location = 0; + + for (int64_t j = 0; j < x_nnz; ++j) { + location = 0; + for (int i = 0; i < x.sparse_dim(); ++i) { + location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; + } + for (size_t i = 0; i < out_sparse_part_dims.size(); ++i) { + out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; + location %= out_sparse_part_strides[i]; + } + } + +} + + +template +void ReshapeCsrKernel(const Context& dev_ctx, + const SparseCsrTensor& x, + const std::vector& new_shape, + SparseCsrTensor* out) { + /*将csr格式转化为coo格式后处理*/ +const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); +SparseCooTensor out_coo; +ReshapeCooKernel(dev_ctx, x_coo, new_shape, &out_coo); +CooToCsrKernel(dev_ctx, out_coo, out); +} + +} // namespace sparse +} // namespace phi + +PD_REGISTER_KERNEL(reshape_coo, + CPU, + ALL_LAYOUT, + phi::sparse::ReshapeCooKernel, + float, + double, + int8_t, + uint8_t, + int16_t, + int, + int64_t, + bool) {} + +PD_REGISTER_KERNEL(reshape_csr, + CPU, + ALL_LAYOUT, + phi::sparse::ReshapeCsrKernel, + float, + double, + int8_t, + uint8_t, + int16_t, + int, + int64_t, + bool) {} \ No newline at end of file diff --git a/paddle/phi/kernels/sparse/cpu/sparse_utils_kernel.cc b/paddle/phi/kernels/sparse/cpu/sparse_utils_kernel.cc index d0016099cd759..dcb4399aa2862 100644 --- a/paddle/phi/kernels/sparse/cpu/sparse_utils_kernel.cc +++ b/paddle/phi/kernels/sparse/cpu/sparse_utils_kernel.cc @@ -329,7 +329,8 @@ PD_REGISTER_KERNEL(csr_to_coo, int8_t, int16_t, int, - int64_t) {} + int64_t, + bool) {} PD_REGISTER_KERNEL(coo_to_csr, CPU, @@ -342,7 +343,8 @@ PD_REGISTER_KERNEL(coo_to_csr, int8_t, int16_t, int, - int64_t) {} + int64_t, + bool) {} PD_REGISTER_KERNEL(dense_to_csr, CPU, diff --git a/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu new file mode 100644 index 0000000000000..fc24c6f39aca8 --- /dev/null +++ b/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu @@ -0,0 +1,107 @@ +// 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 "paddle/phi/kernels/sparse/unary_grad_kernel.h" +#include "paddle/phi/kernels/sparse/unary_kernel.h" + +#include "paddle/phi/backends/gpu/gpu_context.h" +#include "paddle/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/sparse/empty_kernel.h" +#include "paddle/phi/kernels/sparse/impl/unary_grad_kernel_impl.h" + +namespace phi { +namespace sparse { + +// std::vector get_gpu_grad_perm(std::vector perm) { +// std::vector grad_perm(perm.size()); +// for (unsigned int i = 0; i < perm.size(); ++i) { +// grad_perm[perm[i]] = i; +// } +// return grad_perm; +// } + + +// just copy from paddle\phi\kernels\sparse\cpu\reshape_grad_kernel.cc +template +void ReshapeCooGradKernel(const Context& dev_ctx, + const SparseCooTensor& x, + const SparseCooTensor& dout, + // const std::vector& perm, + SparseCooTensor* dx) { + EmptyLikeCooKernel(dev_ctx, x, dx); + // std::vector grad_perm = get_gpu_grad_perm(perm); + // TransposeCooKernel(dev_ctx, dout, grad_perm, dx); + std::vector x_shape(x.dims().size()); + for (int i=0; i(dev_ctx, dout, x_shape, dx); +} + +// template +// void TransposeCsrGradKernel(const Context& dev_ctx, +// const SparseCsrTensor& x, +// const SparseCsrTensor& dout, +// const std::vector& perm, +// SparseCsrTensor* dx) { +// EmptyLikeCsrKernel(dev_ctx, x, dx); +// std::vector grad_perm = get_gpu_grad_perm(perm); +// TransposeCsrKernel(dev_ctx, dout, grad_perm, dx); +// } + +// just copy from paddle\phi\kernels\sparse\cpu\reshape_grad_kernel.cc +template +void ReshapeCsrGradKernel(const Context& dev_ctx, + const SparseCsrTensor& x, + const SparseCsrTensor& dout, + // const std::vector& perm, + SparseCsrTensor* dx) { + EmptyLikeCsrKernel(dev_ctx, x, dx); + std::vector x_shape(x.dims().size()); + for (int i=0; i(dev_ctx, dout, x_shape, dx); +} + +} // namespace sparse +} // namespace phi + +PD_REGISTER_KERNEL(reshape_coo_grad, + GPU, + ALL_LAYOUT, + phi::sparse::ReshapeCooGradKernel, + phi::dtype::float16, + float, + double, + int8_t, + uint8_t, + int16_t, + int, + int64_t, + bool) {} + +PD_REGISTER_KERNEL(reshape_csr_grad, + GPU, + ALL_LAYOUT, + phi::sparse::ReshapeCsrGradKernel, + phi::dtype::float16, + float, + double, + int8_t, + uint8_t, + int16_t, + int, + int64_t, + bool) {} diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu new file mode 100644 index 0000000000000..7cf8268be8ef1 --- /dev/null +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -0,0 +1,169 @@ +// 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 "paddle/phi/kernels/sparse/unary_kernel.h" + +#include "paddle/phi/backends/gpu/gpu_context.h" +#include "paddle/phi/core/kernel_registry.h" +#include "paddle/phi/kernels/funcs/elementwise_base.h" +#include "paddle/phi/kernels/sparse/empty_kernel.h" +#include "paddle/phi/kernels/sparse/impl/unary_kernel_impl.h" + +// #include "paddle/phi/kernels/sparse/gpu/sparse_utils_kernel.cu" +#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" + + +namespace phi { +namespace sparse { + +__global__ void ReshapeCooCudaKernel(const int64_t *x_indices_data, + const int& num_x_sparse_part_dims, + // const int *perm, + const int& num_out_sparse_part_dims, + const int64_t& x_nnz, + const int64_t* x_sparse_part_strides, + const int64_t* out_sparse_part_strides, + int64_t *out_indices_data) { + + // for (std::size_t i = 0; i < n_dim; ++i) { + // CUDA_KERNEL_LOOP_TYPE(j, x_nnz, int64_t) { + // out_indices_data[j + i * x_nnz] = x_indices_data[j + perm[i] * x_nnz]; + // } + // } + + CUDA_KERNEL_LOOP_TYPE(j, x_nnz, int64_t) { + int64_t location = 0; + for (int i = 0; i < num_x_sparse_part_dims; ++i) { + location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; + } + for (int i = 0; i < num_out_sparse_part_dims; ++i) { + out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; + location %= out_sparse_part_strides[i]; + } + } +} + + + +template +void ReshapeCooKernel(const Context &dev_ctx, + const SparseCooTensor &x, + const std::vector& new_shape, + SparseCooTensor *out) { + // create "out" sparse tensor + int64_t x_nnz = x.nnz(); + // DDim out_dims = x.dims().transpose(perm); + DDim out_dims = phi::make_ddim(new_shape); + + /////// get sparse part dimensions of x and out + std::vector x_sparse_part_dims; + std::vector out_sparse_part_dims; + for (int i = 0; i < x.sparse_dim(); ++i) { + x_sparse_part_dims.push_back(x.dims()[i]); + } + for (int i = 0; i < out_dims.size() - x.dense_dim(); ++i) { + out_sparse_part_dims.push_back(out_dims[i]); + } + // DenseTensor out_indices = EmptyLike(dev_ctx, x.indices()); + DenseTensor out_indices = Empty(dev_ctx, + {static_cast(out_sparse_part_dims.size()), x_nnz} + ); + + DenseTensor out_values(x.values()); + + out->SetMember(out_indices, out_values, out_dims, x.coalesced()); + + // compute values of indices + const DenseTensor& x_indices = x.indices(); + const auto *x_indices_data = x_indices.data(); + auto *out_indices_data = out_indices.data(); + + + + + +////////////// +// int *d_perm; +// #ifdef PADDLE_WITH_HIP +// hipMalloc(reinterpret_cast(&d_perm), sizeof(int) * perm.size()); +// hipMemcpy( +// d_perm, perm.data(), sizeof(int) * perm.size(), hipMemcpyHostToDevice); +// #else +// cudaMalloc(reinterpret_cast(&d_perm), sizeof(int) * perm.size()); +// cudaMemcpy( +// d_perm, perm.data(), sizeof(int) * perm.size(), cudaMemcpyHostToDevice); +// #endif +////////// + + + const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); + const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); + + auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, x_nnz, 1); + ReshapeCooCudaKernel<<>>( + // x_indices_data, d_perm, perm.size(), x_nnz, out_indices_data + x_indices_data, + x_sparse_part_dims.size(), out_sparse_part_dims.size(), x_nnz, + x_sparse_part_strides.Get(), out_sparse_part_strides.Get(), + out_indices_data + ); +} + + +// just copy from paddle\phi\kernels\sparse\cpu\reshape_kernel.cc +template +void ReshapeCsrKernel(const Context& dev_ctx, + const SparseCsrTensor& x, + const std::vector& new_shape, + SparseCsrTensor* out) { + /*将csr格式转化为coo格式后处理*/ +const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); +SparseCooTensor out_coo; +ReshapeCooKernel(dev_ctx, x_coo, new_shape, &out_coo); +CooToCsrKernel(dev_ctx, out_coo, out); +} + +} // namespace sparse +} // namespace phi + +PD_REGISTER_KERNEL(reshape_coo, + GPU, + ALL_LAYOUT, + phi::sparse::ReshapeCooKernel, + phi::dtype::float16, + float, + double, + int8_t, + uint8_t, + int16_t, + int, + int64_t, + bool) {} + +PD_REGISTER_KERNEL(reshape_csr, + GPU, + ALL_LAYOUT, + phi::sparse::ReshapeCsrKernel, + phi::dtype::float16, + float, + double, + int8_t, + uint8_t, + int16_t, + int, + int64_t, + bool) {} diff --git a/paddle/phi/kernels/sparse/gpu/sparse_utils_kernel.cu b/paddle/phi/kernels/sparse/gpu/sparse_utils_kernel.cu index c037f6b1b8360..c72a38cd8fd32 100644 --- a/paddle/phi/kernels/sparse/gpu/sparse_utils_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/sparse_utils_kernel.cu @@ -539,7 +539,8 @@ PD_REGISTER_KERNEL(csr_to_coo, int8_t, int16_t, int, - int64_t) {} + int64_t, + bool) {} PD_REGISTER_KERNEL(coo_to_csr, GPU, @@ -552,7 +553,8 @@ PD_REGISTER_KERNEL(coo_to_csr, int8_t, int16_t, int, - int64_t) {} + int64_t, + bool) {} PD_REGISTER_KERNEL(dense_to_csr, GPU, diff --git a/paddle/phi/kernels/sparse/unary_grad_kernel.h b/paddle/phi/kernels/sparse/unary_grad_kernel.h index eb2cf9ed697e9..a7060602f11cc 100644 --- a/paddle/phi/kernels/sparse/unary_grad_kernel.h +++ b/paddle/phi/kernels/sparse/unary_grad_kernel.h @@ -77,5 +77,19 @@ void CastCsrGradKernel(const Context& dev_ctx, DataType value_dtype, SparseCsrTensor* dx); +template +void ReshapeCooGradKernel(const Context& dev_ctx, + const SparseCooTensor& x, + const SparseCooTensor& dout, + // const std::vector& perm, + SparseCooTensor* dx); + +template +void ReshapeCsrGradKernel(const Context& dev_ctx, + const SparseCsrTensor& x, + const SparseCsrTensor& dout, + // const std::vector& perm, + SparseCsrTensor* dx); + } // namespace sparse } // namespace phi diff --git a/paddle/phi/kernels/sparse/unary_kernel.h b/paddle/phi/kernels/sparse/unary_kernel.h index fdb6b21a44427..003ae432d4d65 100644 --- a/paddle/phi/kernels/sparse/unary_kernel.h +++ b/paddle/phi/kernels/sparse/unary_kernel.h @@ -113,5 +113,50 @@ SparseCooTensor ReluCsr(const Context& dev_ctx, const SparseCooTensor& x) { return csr; } +template +void ReshapeCooKernel(const Context &dev_ctx, + const SparseCooTensor &x, + const std::vector & new_shape, + SparseCooTensor *out); + +template +void ReshapeCsrKernel(const Context &dev_ctx, + const SparseCsrTensor &x, + const std::vector & new_shape, + SparseCsrTensor *out); + +template +SparseCooTensor ReshapeCoo(const Context &dev_ctx, + const SparseCooTensor &x, + const std::vector& new_shape) { + //TODO: product(new_shape) must be equal with x.dims().numel() + // // PADDLE_ENFORCE_EQ(x.sparse_dim(), + // PADDLE_ENFORCE_EQ(x.dims().size(), + // new_shape.size(), + // phi::errors::InvalidArgument( + // "size of perm must be equal with the x.dims().size()")); + + SparseCooTensor coo; + ReshapeCooKernel(dev_ctx, x, new_shape, &coo); + return coo; +} + +template +SparseCsrTensor ReshapeCsr(const Context &dev_ctx, + const SparseCsrTensor &x, + const std::vector & new_shape) { + PADDLE_ENFORCE_LE( + 2, + new_shape.size(), + phi::errors::InvalidArgument("size of new_shape must be equal to 2 or 3")); + PADDLE_ENFORCE_GE( + 3, + new_shape.size(), + phi::errors::InvalidArgument("size of new_shape must be equal to 2 or 3")); + SparseCsrTensor csr; + ReshapeCsrKernel(dev_ctx, x, new_shape, &csr); + return csr; +} + } // namespace sparse } // namespace phi diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py new file mode 100644 index 0000000000000..84dd736c43fd1 --- /dev/null +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -0,0 +1,78 @@ +# 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. + +import paddle +import numpy as np +import unittest +from paddle.fluid.framework import _test_eager_guard + + +class TestReshape(unittest.TestCase): + # x: sparse, out: sparse + def check_result(self, x_shape, new_shape, format): + with _test_eager_guard(): + mask = paddle.randint(0, 2, x_shape).astype("float32") + origin_x = paddle.rand(x_shape, dtype='float32') * mask + dense_x = origin_x.detach() + dense_x.stop_gradient = False + # dense_out = paddle.transpose(dense_x, dims) + dense_out = paddle.reshape(dense_x, new_shape) + + if format == "coo": + sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) + else: + sp_x = origin_x.detach().to_sparse_csr() + sp_x.stop_gradient = False + # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) + sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) + + np.testing.assert_allclose(sp_out.to_dense().numpy(), + dense_out.numpy(), + rtol=1e-05) + dense_out.backward() + sp_out.backward() + np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), + # (dense_x.grad * mask).numpy(), + dense_x.grad.numpy(), + rtol=1e-05) + + def test_reshape_2d(self): + self.check_result([2, 5], [10,], 'coo') + self.check_result([10, 5], [2, 25], 'csr') + self.check_result([12, 5], [15, 4], 'coo') + self.check_result([9, 8], [18, 4], 'csr') + + def test_transpose_3d(self): + self.check_result([6, 2, 3], [6, 2, 3], 'coo') + self.check_result([6, 2, 3], [6, 2, 3], 'csr') + self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') + self.check_result([6, 2, 3], [6, 3, 2], 'csr') + # self.check_result([6, 2, 3], [1, 0, 2], 'coo') + # self.check_result([6, 2, 3], [1, 0, 2], 'csr') + # self.check_result([6, 2, 3], [2, 0, 1], 'coo') + # self.check_result([6, 2, 3], [2, 0, 1], 'csr') + # self.check_result([6, 2, 3], [2, 1, 0], 'coo') + # self.check_result([6, 2, 3], [2, 1, 0], 'csr') + # self.check_result([6, 2, 3], [1, 2, 0], 'coo') + # self.check_result([6, 2, 3], [1, 2, 0], 'csr') + + def test_transpose_nd(self): + self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') + # Randint now only supports access to dimension 0 to 9. + # self.check_result([i % 3 + 2 for i in range(9)], + # [(i + 2) % 9 for i in range(9)], 'coo') + + +if __name__ == "__main__": + unittest.main() diff --git a/python/paddle/incubate/sparse/__init__.py b/python/paddle/incubate/sparse/__init__.py index de89f46438a35..003ffcfdcd7ba 100644 --- a/python/paddle/incubate/sparse/__init__.py +++ b/python/paddle/incubate/sparse/__init__.py @@ -34,6 +34,7 @@ from .unary import deg2rad from .unary import rad2deg from .unary import expm1 +from .unary import reshape from .binary import mv from .binary import matmul @@ -79,4 +80,5 @@ 'divide', 'coalesce', 'is_same_shape', + 'reshape' ] diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index 621e31bc3e834..c173aaacac51c 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -608,3 +608,26 @@ def expm1(x, name=None): out = paddle.incubate.sparse.expm1(sparse_x) """ return _C_ops.sparse_expm1(x) + +@dygraph_only +def reshape(x, new_shape, name=None): + """ + Changes the dims order of ``x`` without changing its data, requiring x to be a SparseCooTensor or SparseCsrTensor. + .. math:: + out = transpose(x, dims) + Parameters: + x (Tensor): The input Sparse Tensor with data type float32, float64. + dims (list[int]): new dims. + name (str, optional): Name for the operation (optional, default is None). + For more information, please refer to :ref:`api_guide_Name`. + Returns: + A transposed Sparse Tensor with the same data type as ``x``. + Examples: + .. code-block:: python + import paddle + dense_x = paddle.to_tensor([[-2., 0.], [1., 2.]]) + sparse_x = dense_x.to_sparse_coo(1) + out = paddle.incubate.sparse.transpose(sparse_x, [1, 0]) + """ + return _C_ops.sparse_reshape(x, new_shape) + \ No newline at end of file From 5277531ec22a8d380c571d5b82ada49884ae33a7 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 21 Sep 2022 09:44:39 +0800 Subject: [PATCH 02/55] change the dtype in all test cases to int64 --- .../paddle/fluid/tests/unittests/test_sparse_reshape_op.py | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 84dd736c43fd1..008c93926c4c3 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -22,8 +22,11 @@ class TestReshape(unittest.TestCase): # x: sparse, out: sparse def check_result(self, x_shape, new_shape, format): with _test_eager_guard(): - mask = paddle.randint(0, 2, x_shape).astype("float32") - origin_x = paddle.rand(x_shape, dtype='float32') * mask + # mask = paddle.randint(0, 2, x_shape).astype("float32") + # origin_x = paddle.rand(x_shape, dtype='float32') * mask + mask = paddle.randint(0, 2, x_shape) + origin_x = paddle.randint(-100, 100, x_shape) * mask + dense_x = origin_x.detach() dense_x.stop_gradient = False # dense_out = paddle.transpose(dense_x, dims) From ab3e87188a015e63c0b8bfa03f94fc40dde6e2d7 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 21 Sep 2022 17:29:04 +0800 Subject: [PATCH 03/55] just one test case --- .../tests/unittests/test_sparse_reshape_op.py | 42 +++++++++---------- 1 file changed, 21 insertions(+), 21 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 008c93926c4c3..afc24ff61c85c 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -52,29 +52,29 @@ def check_result(self, x_shape, new_shape, format): def test_reshape_2d(self): self.check_result([2, 5], [10,], 'coo') - self.check_result([10, 5], [2, 25], 'csr') - self.check_result([12, 5], [15, 4], 'coo') - self.check_result([9, 8], [18, 4], 'csr') + # self.check_result([10, 5], [2, 25], 'csr') + # self.check_result([12, 5], [15, 4], 'coo') + # self.check_result([9, 8], [18, 4], 'csr') - def test_transpose_3d(self): - self.check_result([6, 2, 3], [6, 2, 3], 'coo') - self.check_result([6, 2, 3], [6, 2, 3], 'csr') - self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') - self.check_result([6, 2, 3], [6, 3, 2], 'csr') - # self.check_result([6, 2, 3], [1, 0, 2], 'coo') - # self.check_result([6, 2, 3], [1, 0, 2], 'csr') - # self.check_result([6, 2, 3], [2, 0, 1], 'coo') - # self.check_result([6, 2, 3], [2, 0, 1], 'csr') - # self.check_result([6, 2, 3], [2, 1, 0], 'coo') - # self.check_result([6, 2, 3], [2, 1, 0], 'csr') - # self.check_result([6, 2, 3], [1, 2, 0], 'coo') - # self.check_result([6, 2, 3], [1, 2, 0], 'csr') + # def test_transpose_3d(self): + # self.check_result([6, 2, 3], [6, 2, 3], 'coo') + # self.check_result([6, 2, 3], [6, 2, 3], 'csr') + # self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') + # self.check_result([6, 2, 3], [6, 3, 2], 'csr') + # # self.check_result([6, 2, 3], [1, 0, 2], 'coo') + # # self.check_result([6, 2, 3], [1, 0, 2], 'csr') + # # self.check_result([6, 2, 3], [2, 0, 1], 'coo') + # # self.check_result([6, 2, 3], [2, 0, 1], 'csr') + # # self.check_result([6, 2, 3], [2, 1, 0], 'coo') + # # self.check_result([6, 2, 3], [2, 1, 0], 'csr') + # # self.check_result([6, 2, 3], [1, 2, 0], 'coo') + # # self.check_result([6, 2, 3], [1, 2, 0], 'csr') - def test_transpose_nd(self): - self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') - # Randint now only supports access to dimension 0 to 9. - # self.check_result([i % 3 + 2 for i in range(9)], - # [(i + 2) % 9 for i in range(9)], 'coo') + # def test_transpose_nd(self): + # self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') + # # Randint now only supports access to dimension 0 to 9. + # # self.check_result([i % 3 + 2 for i in range(9)], + # # [(i + 2) % 9 for i in range(9)], 'coo') if __name__ == "__main__": From a8a4960f1fc57cfe22cbcd9a614ac455d5712c14 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Thu, 22 Sep 2022 11:31:06 +0800 Subject: [PATCH 04/55] modify comments --- .../phi/kernels/sparse/cpu/reshape_kernel.cc | 35 ++++--------------- python/paddle/incubate/sparse/unary.py | 16 ++++----- 2 files changed, 15 insertions(+), 36 deletions(-) diff --git a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc index 0a51afc04137c..c147cc444ee73 100644 --- a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc +++ b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc @@ -14,9 +14,7 @@ #include "paddle/phi/kernels/sparse/unary_kernel.h" -// #include "paddle/phi/core/ddim.cc" #include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" -// #include "paddle/phi/kernels/sparse/cpu/sparse_utils_kernel.cc" #include "paddle/phi/backends/cpu/cpu_context.h" #include "paddle/phi/core/kernel_registry.h" @@ -32,7 +30,6 @@ namespace sparse { template void ReshapeCooKernel(const Context& dev_ctx, const SparseCooTensor& x, - // const std::vector& sparse_part_permutation, const std::vector& new_shape, SparseCooTensor* out) { /* @@ -40,9 +37,8 @@ void ReshapeCooKernel(const Context& dev_ctx, */ // create "out" sparse tensor int64_t x_nnz = x.nnz(); - // DDim out_dims = x.dims().transpose(perm); DDim out_dims = phi::make_ddim(new_shape); - ////get sparse part dimensions of x and out + // get sparse part dimensions of x and out std::vector x_sparse_part_dims; std::vector out_sparse_part_dims; for (int i = 0; i < x.sparse_dim(); ++i) { @@ -51,7 +47,6 @@ void ReshapeCooKernel(const Context& dev_ctx, for (int i = 0; i < out_dims.size() - x.dense_dim(); ++i) { out_sparse_part_dims.push_back(out_dims[i]); } - // DenseTensor out_indices = EmptyLike(dev_ctx, x.indices()); DenseTensor out_indices = Empty(dev_ctx, {static_cast(out_sparse_part_dims.size()), x_nnz} ); @@ -62,24 +57,8 @@ void ReshapeCooKernel(const Context& dev_ctx, const DenseTensor& x_indices = x.indices(); const auto* x_indices_data = x_indices.data(); auto* out_indices_data = out_indices.data(); -// //i 表示 indices 的 行标 -// for (unsigned int i = 0; i < perm.size(); ++i) { -// // j 表示 indices 的 列标 -// for (int64_t j = 0; j < x_nnz; ++j) { -// // 修改 out indices 的索引为 (i, j)的元素值 -// /* Caution : 这是原来的计算逻辑，我认为是 错误的， -// 这里计算逻辑是： 原tensor的shape是 (10, 20, 30, 40, 50) -// 一个非零元素的索引为 (1, 2, 3, 4, 5) -// 进行transpose 后, tensor的shape 是 (30, 10, 50, 20, 40) -// 这里的计算逻辑就认为该非零元素的新索引就是 (3, 1, 5, 2, 4) -// 没错，这就是transpose的计算逻辑，transpose后元素在内存中的位置改变了 -// 你更改的逻辑其实是 reshape的计算逻辑，reshape后所有元素在内存中的位置均不变 -// */ -// out_indices_data[j + i * x_nnz] = x_indices_data[j + perm[i] * x_nnz]; -// } -// } - // 我的更改后的计算逻辑如下： + // 我的更改后的计算逻辑如下 const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); int64_t location = 0; @@ -103,11 +82,11 @@ void ReshapeCsrKernel(const Context& dev_ctx, const SparseCsrTensor& x, const std::vector& new_shape, SparseCsrTensor* out) { - /*将csr格式转化为coo格式后处理*/ -const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); -SparseCooTensor out_coo; -ReshapeCooKernel(dev_ctx, x_coo, new_shape, &out_coo); -CooToCsrKernel(dev_ctx, out_coo, out); + /*将csr格式转化为coo格式后处理*/ + const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); + SparseCooTensor out_coo; + ReshapeCooKernel(dev_ctx, x_coo, new_shape, &out_coo); + CooToCsrKernel(dev_ctx, out_coo, out); } } // namespace sparse diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index c173aaacac51c..e14250a8795de 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -612,22 +612,22 @@ def expm1(x, name=None): @dygraph_only def reshape(x, new_shape, name=None): """ - Changes the dims order of ``x`` without changing its data, requiring x to be a SparseCooTensor or SparseCsrTensor. + Changes the shape of ``x`` without changing its value, requiring x to be a SparseCooTensor or SparseCsrTensor. .. math:: - out = transpose(x, dims) + out = reshape(x, new_shape) Parameters: - x (Tensor): The input Sparse Tensor with data type float32, float64. - dims (list[int]): new dims. + x (Tensor): The input Sparse Tensor with data type float32, float64, int32, int64 and so on. + new_shape (list[int]): new shape. name (str, optional): Name for the operation (optional, default is None). For more information, please refer to :ref:`api_guide_Name`. Returns: - A transposed Sparse Tensor with the same data type as ``x``. + A reshaped Sparse Tensor with the same data type as ``x``. Examples: .. code-block:: python import paddle - dense_x = paddle.to_tensor([[-2., 0.], [1., 2.]]) - sparse_x = dense_x.to_sparse_coo(1) - out = paddle.incubate.sparse.transpose(sparse_x, [1, 0]) + dense_x = paddle.to_tensor([[-2., 0., 7.9], [1., 2., 0.]]) + sparse_x = dense_x.to_sparse_coo(2) + out = paddle.incubate.sparse.reshape(sparse_x, [6, ]) """ return _C_ops.sparse_reshape(x, new_shape) \ No newline at end of file From 234122f25563cdd29dd8a5215621bd4de2b36b68 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Thu, 22 Sep 2022 17:29:25 +0800 Subject: [PATCH 05/55] Update test_sparse_reshape_op.py --- .../fluid/tests/unittests/test_sparse_reshape_op.py | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index afc24ff61c85c..6073750162367 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -43,12 +43,12 @@ def check_result(self, x_shape, new_shape, format): np.testing.assert_allclose(sp_out.to_dense().numpy(), dense_out.numpy(), rtol=1e-05) - dense_out.backward() - sp_out.backward() - np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - # (dense_x.grad * mask).numpy(), - dense_x.grad.numpy(), - rtol=1e-05) + # dense_out.backward() + # sp_out.backward() + # np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), + # # (dense_x.grad * mask).numpy(), + # dense_x.grad.numpy(), + # rtol=1e-05) def test_reshape_2d(self): self.check_result([2, 5], [10,], 'coo') From 64f98b0fa340230e051ce73bc0de563fb77721b7 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Fri, 23 Sep 2022 11:06:50 +0800 Subject: [PATCH 06/55] chang the type of "shape" from vector to IntArray --- paddle/phi/api/yaml/sparse_backward.yaml | 2 +- paddle/phi/api/yaml/sparse_ops.yaml | 4 +- paddle/phi/infermeta/sparse/unary.cc | 8 -- paddle/phi/infermeta/sparse/unary.h | 4 - .../kernels/sparse/cpu/reshape_grad_kernel.cc | 25 ++---- .../phi/kernels/sparse/cpu/reshape_kernel.cc | 51 +++++------ .../kernels/sparse/gpu/reshape_grad_kernel.cu | 45 ++-------- .../phi/kernels/sparse/gpu/reshape_kernel.cu | 88 ++++++------------- paddle/phi/kernels/sparse/unary_grad_kernel.h | 4 +- paddle/phi/kernels/sparse/unary_kernel.h | 34 ++++--- python/paddle/incubate/sparse/unary.py | 10 ++- 11 files changed, 92 insertions(+), 183 deletions(-) diff --git a/paddle/phi/api/yaml/sparse_backward.yaml b/paddle/phi/api/yaml/sparse_backward.yaml index b0f4b22302738..eb50d0d9d8d93 100644 --- a/paddle/phi/api/yaml/sparse_backward.yaml +++ b/paddle/phi/api/yaml/sparse_backward.yaml @@ -407,7 +407,7 @@ data_type: query - backward_op : reshape_grad - forward : reshape(Tensor x, int64_t[] new_shape) -> Tensor(out) + forward : reshape(Tensor x, IntArray shape) -> Tensor(out) args : (Tensor x, Tensor out_grad) output : Tensor(x_grad) infer_meta : diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index 3258f58f9ce55..7bc299aa08681 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -459,10 +459,10 @@ backward: mv_grad - op : reshape - args : (Tensor x, int64_t[] new_shape) + args : (Tensor x, IntArray shape) output : Tensor(out) infer_meta : - func : sparse::ReshapeInferMeta + func : ReshapeInferMeta kernel : func : reshape_coo{sparse_coo -> sparse_coo}, reshape_csr{sparse_csr -> sparse_csr} diff --git a/paddle/phi/infermeta/sparse/unary.cc b/paddle/phi/infermeta/sparse/unary.cc index 6037c38fd96ed..45cb4f75e38c9 100644 --- a/paddle/phi/infermeta/sparse/unary.cc +++ b/paddle/phi/infermeta/sparse/unary.cc @@ -32,13 +32,5 @@ void ValuesInferMeta(const MetaTensor& x, MetaTensor* out) { out->set_layout(x.layout()); } -void ReshapeInferMeta(const MetaTensor& x, - const std::vector& new_shape, - MetaTensor* out) { - out->set_dtype(x.dtype()); - out->set_dims(phi::make_ddim(new_shape)); - out->set_layout(x.layout()); -} - } // namespace sparse } // namespace phi diff --git a/paddle/phi/infermeta/sparse/unary.h b/paddle/phi/infermeta/sparse/unary.h index 360852c559fe4..880e90b7ae697 100644 --- a/paddle/phi/infermeta/sparse/unary.h +++ b/paddle/phi/infermeta/sparse/unary.h @@ -24,9 +24,5 @@ void IndicesInferMeta(const MetaTensor& x, MetaTensor* out); void ValuesInferMeta(const MetaTensor& x, MetaTensor* out); -void ReshapeInferMeta(const MetaTensor& x, - const std::vector& new_shape, - MetaTensor* out); - } // namespace sparse } // namespace phi diff --git a/paddle/phi/kernels/sparse/cpu/reshape_grad_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_grad_kernel.cc index 57ae338af8fe9..fc843f81c31ee 100644 --- a/paddle/phi/kernels/sparse/cpu/reshape_grad_kernel.cc +++ b/paddle/phi/kernels/sparse/cpu/reshape_grad_kernel.cc @@ -25,33 +25,24 @@ namespace sparse { template void ReshapeCooGradKernel(const Context& dev_ctx, - const SparseCooTensor& x, - const SparseCooTensor& dout, - // const std::vector& perm, - SparseCooTensor* dx) { + const SparseCooTensor& x, + const SparseCooTensor& dout, + SparseCooTensor* dx) { EmptyLikeCooKernel(dev_ctx, x, dx); - std::vector x_shape(x.dims().size()); - for (int i=0; i(dev_ctx, dout, x_shape, dx); } template void ReshapeCsrGradKernel(const Context& dev_ctx, - const SparseCsrTensor& x, - const SparseCsrTensor& dout, - // const std::vector& perm, - SparseCsrTensor* dx) { + const SparseCsrTensor& x, + const SparseCsrTensor& dout, + SparseCsrTensor* dx) { EmptyLikeCsrKernel(dev_ctx, x, dx); - std::vector x_shape(x.dims().size()); - for (int i=0; i(dev_ctx, dout, x_shape, dx); } - } // namespace sparse } // namespace phi diff --git a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc index c147cc444ee73..48ee987c4ea69 100644 --- a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc +++ b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc @@ -15,6 +15,7 @@ #include "paddle/phi/kernels/sparse/unary_kernel.h" #include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" +#include "paddle/phi/core/ddim.h" #include "paddle/phi/backends/cpu/cpu_context.h" #include "paddle/phi/core/kernel_registry.h" @@ -29,15 +30,12 @@ namespace sparse { template void ReshapeCooKernel(const Context& dev_ctx, - const SparseCooTensor& x, - const std::vector& new_shape, - SparseCooTensor* out) { - /* - 目前只能针对 sparse part dims 部分进行reshape - */ - // create "out" sparse tensor + const SparseCooTensor& x, + const phi::IntArray& shape, + SparseCooTensor* out) { + // TODO: Currently, reshape is only applicable to sparse dims int64_t x_nnz = x.nnz(); - DDim out_dims = phi::make_ddim(new_shape); + DDim out_dims = phi::make_ddim(shape.GetData()); // get sparse part dimensions of x and out std::vector x_sparse_part_dims; std::vector out_sparse_part_dims; @@ -58,34 +56,31 @@ void ReshapeCooKernel(const Context& dev_ctx, const auto* x_indices_data = x_indices.data(); auto* out_indices_data = out_indices.data(); - // 我的更改后的计算逻辑如下 - const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); - const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); - int64_t location = 0; - - for (int64_t j = 0; j < x_nnz; ++j) { - location = 0; - for (int i = 0; i < x.sparse_dim(); ++i) { - location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; - } - for (size_t i = 0; i < out_sparse_part_dims.size(); ++i) { - out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; - location %= out_sparse_part_strides[i]; - } - } - + const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); + const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); + int64_t location = 0; + for (int64_t j = 0; j < x_nnz; ++j) { + location = 0; + for (int i = 0; i < x.sparse_dim(); ++i) { + location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; + } + for (size_t i = 0; i < out_sparse_part_dims.size(); ++i) { + out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; + location %= out_sparse_part_strides[i]; + } + } } template void ReshapeCsrKernel(const Context& dev_ctx, - const SparseCsrTensor& x, - const std::vector& new_shape, - SparseCsrTensor* out) { + const SparseCsrTensor& x, + const phi::IntArray& shape, + SparseCsrTensor* out) { /*将csr格式转化为coo格式后处理*/ const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); SparseCooTensor out_coo; - ReshapeCooKernel(dev_ctx, x_coo, new_shape, &out_coo); + ReshapeCooKernel(dev_ctx, x_coo, shape, &out_coo); CooToCsrKernel(dev_ctx, out_coo, out); } diff --git a/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu index fc24c6f39aca8..b67e79ef8efea 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu @@ -23,55 +23,26 @@ namespace phi { namespace sparse { -// std::vector get_gpu_grad_perm(std::vector perm) { -// std::vector grad_perm(perm.size()); -// for (unsigned int i = 0; i < perm.size(); ++i) { -// grad_perm[perm[i]] = i; -// } -// return grad_perm; -// } - - // just copy from paddle\phi\kernels\sparse\cpu\reshape_grad_kernel.cc template void ReshapeCooGradKernel(const Context& dev_ctx, - const SparseCooTensor& x, - const SparseCooTensor& dout, - // const std::vector& perm, - SparseCooTensor* dx) { + const SparseCooTensor& x, + const SparseCooTensor& dout, + SparseCooTensor* dx) { EmptyLikeCooKernel(dev_ctx, x, dx); - // std::vector grad_perm = get_gpu_grad_perm(perm); - // TransposeCooKernel(dev_ctx, dout, grad_perm, dx); - std::vector x_shape(x.dims().size()); - for (int i=0; i(dev_ctx, dout, x_shape, dx); } -// template -// void TransposeCsrGradKernel(const Context& dev_ctx, -// const SparseCsrTensor& x, -// const SparseCsrTensor& dout, -// const std::vector& perm, -// SparseCsrTensor* dx) { -// EmptyLikeCsrKernel(dev_ctx, x, dx); -// std::vector grad_perm = get_gpu_grad_perm(perm); -// TransposeCsrKernel(dev_ctx, dout, grad_perm, dx); -// } // just copy from paddle\phi\kernels\sparse\cpu\reshape_grad_kernel.cc template void ReshapeCsrGradKernel(const Context& dev_ctx, - const SparseCsrTensor& x, - const SparseCsrTensor& dout, - // const std::vector& perm, - SparseCsrTensor* dx) { + const SparseCsrTensor& x, + const SparseCsrTensor& dout, + SparseCsrTensor* dx) { EmptyLikeCsrKernel(dev_ctx, x, dx); - std::vector x_shape(x.dims().size()); - for (int i=0; i(dev_ctx, dout, x_shape, dx); } diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index 7cf8268be8ef1..bbb86788a0861 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -20,7 +20,6 @@ #include "paddle/phi/kernels/sparse/empty_kernel.h" #include "paddle/phi/kernels/sparse/impl/unary_kernel_impl.h" -// #include "paddle/phi/kernels/sparse/gpu/sparse_utils_kernel.cu" #include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" @@ -28,19 +27,12 @@ namespace phi { namespace sparse { __global__ void ReshapeCooCudaKernel(const int64_t *x_indices_data, - const int& num_x_sparse_part_dims, - // const int *perm, - const int& num_out_sparse_part_dims, - const int64_t& x_nnz, - const int64_t* x_sparse_part_strides, - const int64_t* out_sparse_part_strides, - int64_t *out_indices_data) { - - // for (std::size_t i = 0; i < n_dim; ++i) { - // CUDA_KERNEL_LOOP_TYPE(j, x_nnz, int64_t) { - // out_indices_data[j + i * x_nnz] = x_indices_data[j + perm[i] * x_nnz]; - // } - // } + const int& num_x_sparse_part_dims, + const int& num_out_sparse_part_dims, + const int64_t& x_nnz, + const int64_t* x_sparse_part_strides, + const int64_t* out_sparse_part_strides, + int64_t *out_indices_data) { CUDA_KERNEL_LOOP_TYPE(j, x_nnz, int64_t) { int64_t location = 0; @@ -55,18 +47,14 @@ __global__ void ReshapeCooCudaKernel(const int64_t *x_indices_data, } - template void ReshapeCooKernel(const Context &dev_ctx, - const SparseCooTensor &x, - const std::vector& new_shape, - SparseCooTensor *out) { - // create "out" sparse tensor + const SparseCooTensor &x, + const phi::IntArray& shape, + SparseCooTensor *out) { int64_t x_nnz = x.nnz(); - // DDim out_dims = x.dims().transpose(perm); - DDim out_dims = phi::make_ddim(new_shape); - - /////// get sparse part dimensions of x and out + DDim out_dims = phi::make_ddim(shape.GetData()); + // get sparse part dimensions of x and out std::vector x_sparse_part_dims; std::vector out_sparse_part_dims; for (int i = 0; i < x.sparse_dim(); ++i) { @@ -75,66 +63,44 @@ void ReshapeCooKernel(const Context &dev_ctx, for (int i = 0; i < out_dims.size() - x.dense_dim(); ++i) { out_sparse_part_dims.push_back(out_dims[i]); } - // DenseTensor out_indices = EmptyLike(dev_ctx, x.indices()); + DenseTensor out_indices = Empty(dev_ctx, {static_cast(out_sparse_part_dims.size()), x_nnz} ); - DenseTensor out_values(x.values()); - out->SetMember(out_indices, out_values, out_dims, x.coalesced()); - // compute values of indices + // compute values of out indices const DenseTensor& x_indices = x.indices(); const auto *x_indices_data = x_indices.data(); auto *out_indices_data = out_indices.data(); - - - - - -////////////// -// int *d_perm; -// #ifdef PADDLE_WITH_HIP -// hipMalloc(reinterpret_cast(&d_perm), sizeof(int) * perm.size()); -// hipMemcpy( -// d_perm, perm.data(), sizeof(int) * perm.size(), hipMemcpyHostToDevice); -// #else -// cudaMalloc(reinterpret_cast(&d_perm), sizeof(int) * perm.size()); -// cudaMemcpy( -// d_perm, perm.data(), sizeof(int) * perm.size(), cudaMemcpyHostToDevice); -// #endif -////////// - - - const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); - const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); + const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); + const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, x_nnz, 1); ReshapeCooCudaKernel<<>>( - // x_indices_data, d_perm, perm.size(), x_nnz, out_indices_data + config.thread_per_block.x, + 0, + dev_ctx.stream()>>>( x_indices_data, x_sparse_part_dims.size(), out_sparse_part_dims.size(), x_nnz, x_sparse_part_strides.Get(), out_sparse_part_strides.Get(), out_indices_data - ); + ); } // just copy from paddle\phi\kernels\sparse\cpu\reshape_kernel.cc template void ReshapeCsrKernel(const Context& dev_ctx, - const SparseCsrTensor& x, - const std::vector& new_shape, - SparseCsrTensor* out) { - /*将csr格式转化为coo格式后处理*/ -const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); -SparseCooTensor out_coo; -ReshapeCooKernel(dev_ctx, x_coo, new_shape, &out_coo); -CooToCsrKernel(dev_ctx, out_coo, out); + const SparseCsrTensor& x, + const phi::IntArray& shape, + SparseCsrTensor* out) { + /*transform csr format to coo format, and then use coo kernel*/ + const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); + SparseCooTensor out_coo; + ReshapeCooKernel(dev_ctx, x_coo, shape, &out_coo); + CooToCsrKernel(dev_ctx, out_coo, out); } } // namespace sparse diff --git a/paddle/phi/kernels/sparse/unary_grad_kernel.h b/paddle/phi/kernels/sparse/unary_grad_kernel.h index a7060602f11cc..2b0dd410975ff 100644 --- a/paddle/phi/kernels/sparse/unary_grad_kernel.h +++ b/paddle/phi/kernels/sparse/unary_grad_kernel.h @@ -16,6 +16,8 @@ #include "paddle/phi/core/sparse_coo_tensor.h" #include "paddle/phi/core/sparse_csr_tensor.h" +#include "paddle/phi/core/ddim.h" + namespace phi { namespace sparse { @@ -81,14 +83,12 @@ template void ReshapeCooGradKernel(const Context& dev_ctx, const SparseCooTensor& x, const SparseCooTensor& dout, - // const std::vector& perm, SparseCooTensor* dx); template void ReshapeCsrGradKernel(const Context& dev_ctx, const SparseCsrTensor& x, const SparseCsrTensor& dout, - // const std::vector& perm, SparseCsrTensor* dx); } // namespace sparse diff --git a/paddle/phi/kernels/sparse/unary_kernel.h b/paddle/phi/kernels/sparse/unary_kernel.h index 003ae432d4d65..c01994774b1f5 100644 --- a/paddle/phi/kernels/sparse/unary_kernel.h +++ b/paddle/phi/kernels/sparse/unary_kernel.h @@ -16,6 +16,9 @@ #include "paddle/phi/core/sparse_coo_tensor.h" #include "paddle/phi/core/sparse_csr_tensor.h" +#include "paddle/phi/common/int_array.h" +#include "paddle/phi/core/ddim.h" + namespace phi { namespace sparse { @@ -116,45 +119,38 @@ SparseCooTensor ReluCsr(const Context& dev_ctx, const SparseCooTensor& x) { template void ReshapeCooKernel(const Context &dev_ctx, const SparseCooTensor &x, - const std::vector & new_shape, + const phi::IntArray& shape, SparseCooTensor *out); template void ReshapeCsrKernel(const Context &dev_ctx, const SparseCsrTensor &x, - const std::vector & new_shape, + const phi::IntArray& shape, SparseCsrTensor *out); template SparseCooTensor ReshapeCoo(const Context &dev_ctx, - const SparseCooTensor &x, - const std::vector& new_shape) { - //TODO: product(new_shape) must be equal with x.dims().numel() - // // PADDLE_ENFORCE_EQ(x.sparse_dim(), - // PADDLE_ENFORCE_EQ(x.dims().size(), - // new_shape.size(), - // phi::errors::InvalidArgument( - // "size of perm must be equal with the x.dims().size()")); - + const SparseCooTensor &x, + const phi::IntArray& shape) { SparseCooTensor coo; - ReshapeCooKernel(dev_ctx, x, new_shape, &coo); + ReshapeCooKernel(dev_ctx, x, shape, &coo); return coo; } template SparseCsrTensor ReshapeCsr(const Context &dev_ctx, - const SparseCsrTensor &x, - const std::vector & new_shape) { + const SparseCsrTensor &x, + const phi::IntArray& shape) { PADDLE_ENFORCE_LE( 2, - new_shape.size(), - phi::errors::InvalidArgument("size of new_shape must be equal to 2 or 3")); + shape.size(), + phi::errors::InvalidArgument("size of shape must be equal to 2 or 3")); PADDLE_ENFORCE_GE( 3, - new_shape.size(), - phi::errors::InvalidArgument("size of new_shape must be equal to 2 or 3")); + shape.size(), + phi::errors::InvalidArgument("size of shape must be equal to 2 or 3")); SparseCsrTensor csr; - ReshapeCsrKernel(dev_ctx, x, new_shape, &csr); + ReshapeCsrKernel(dev_ctx, x, shape, &csr); return csr; } diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index e14250a8795de..8202aae3a413a 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -610,14 +610,16 @@ def expm1(x, name=None): return _C_ops.sparse_expm1(x) @dygraph_only -def reshape(x, new_shape, name=None): +def reshape(x, shape, name=None): """ Changes the shape of ``x`` without changing its value, requiring x to be a SparseCooTensor or SparseCsrTensor. + Currently this function can only reshape the sparse dims of ``x`` , but ``shape`` argument must + include all dims of the reshaped tensor. .. math:: - out = reshape(x, new_shape) + out = reshape(x, shape) Parameters: x (Tensor): The input Sparse Tensor with data type float32, float64, int32, int64 and so on. - new_shape (list[int]): new shape. + shape (list[int]): new shape. name (str, optional): Name for the operation (optional, default is None). For more information, please refer to :ref:`api_guide_Name`. Returns: @@ -629,5 +631,5 @@ def reshape(x, new_shape, name=None): sparse_x = dense_x.to_sparse_coo(2) out = paddle.incubate.sparse.reshape(sparse_x, [6, ]) """ - return _C_ops.sparse_reshape(x, new_shape) + return _C_ops.sparse_reshape(x, shape) \ No newline at end of file From 0f4660d188f2bf21b6185f33a1f0ad1e25900a78 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Fri, 23 Sep 2022 17:26:50 +0800 Subject: [PATCH 07/55] check whether sp_out.to_dense() is the cause of error --- .../fluid/tests/unittests/test_sparse_reshape_op.py | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 6073750162367..3aef9075845ac 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -40,9 +40,13 @@ def check_result(self, x_shape, new_shape, format): # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - np.testing.assert_allclose(sp_out.to_dense().numpy(), - dense_out.numpy(), - rtol=1e-05) + print(sp_out) + sp_out.to_dense() + print(sp_out.to_dense()) + # np.testing.assert_allclose(sp_out.to_dense().numpy(), + # dense_out.numpy(), + # rtol=1e-05) + # dense_out.backward() # sp_out.backward() # np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), From 4761494959b8463bb7db36e6ef276da8a3483ac8 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Fri, 23 Sep 2022 18:07:37 +0800 Subject: [PATCH 08/55] print sp_out --- .../tests/unittests/test_sparse_reshape_op.py | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 3aef9075845ac..fb6abef1f4945 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -40,13 +40,19 @@ def check_result(self, x_shape, new_shape, format): # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - print(sp_out) - sp_out.to_dense() - print(sp_out.to_dense()) + print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') + print("dense_out:", dense_out) + print("dense_out.numpy():", dense_out.numpy()) + print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') + print("sp_out:", sp_out) + print("sp_out.to_dense():", sp_out.to_dense()) + print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) + print(10*'=', "OccupyMars2025 the end", 10*'=') + # np.testing.assert_allclose(sp_out.to_dense().numpy(), # dense_out.numpy(), # rtol=1e-05) - + # dense_out.backward() # sp_out.backward() # np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), From aa72cc72855aba001ce161cd28550dae70457487 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Fri, 23 Sep 2022 21:49:15 +0800 Subject: [PATCH 09/55] Update reshape_kernel.cc --- paddle/phi/kernels/sparse/cpu/reshape_kernel.cc | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc index 48ee987c4ea69..cc6679317738b 100644 --- a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc +++ b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc @@ -77,7 +77,7 @@ void ReshapeCsrKernel(const Context& dev_ctx, const SparseCsrTensor& x, const phi::IntArray& shape, SparseCsrTensor* out) { - /*将csr格式转化为coo格式后处理*/ + /*transform csr format to coo format, and then use coo kernel*/ const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); SparseCooTensor out_coo; ReshapeCooKernel(dev_ctx, x_coo, shape, &out_coo); @@ -111,4 +111,5 @@ PD_REGISTER_KERNEL(reshape_csr, int16_t, int, int64_t, - bool) {} \ No newline at end of file + bool) {} + \ No newline at end of file From fb434fde66b1a4cafa25bd921004dfa1859923d8 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Fri, 23 Sep 2022 22:15:55 +0800 Subject: [PATCH 10/55] use numpy to generate the equal paddle tensor --- .../tests/unittests/test_sparse_reshape_op.py | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index fb6abef1f4945..b917e62668821 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -24,25 +24,30 @@ def check_result(self, x_shape, new_shape, format): with _test_eager_guard(): # mask = paddle.randint(0, 2, x_shape).astype("float32") # origin_x = paddle.rand(x_shape, dtype='float32') * mask - mask = paddle.randint(0, 2, x_shape) - origin_x = paddle.randint(-100, 100, x_shape) * mask + # mask = paddle.randint(0, 2, x_shape) + # origin_x = paddle.randint(-100, 100, x_shape) * mask - dense_x = origin_x.detach() + # dense_x = paddle.clone(origin_x.detach()) + mask = np.random.randint(0, 2, x_shape) + np_x = np.random.randint(-100, 100, x_shape) * mask + dense_x = paddle.to_tensor(np_x) dense_x.stop_gradient = False # dense_out = paddle.transpose(dense_x, dims) dense_out = paddle.reshape(dense_x, new_shape) if format == "coo": - sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) + # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) + sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) else: - sp_x = origin_x.detach().to_sparse_csr() + # sp_x = origin_x.detach().to_sparse_csr() + sp_x = paddle.to_tensor(np_x).to_sparse_csr() sp_x.stop_gradient = False # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') - print("dense_out:", dense_out) print("dense_out.numpy():", dense_out.numpy()) + print("dense_out:", dense_out) print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') print("sp_out:", sp_out) print("sp_out.to_dense():", sp_out.to_dense()) From 281a3d4c78509baf480e2cfe490127064666f0ca Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Fri, 23 Sep 2022 22:57:41 +0800 Subject: [PATCH 11/55] just check dense_tensor.numpy() --- .../tests/unittests/test_sparse_reshape_op.py | 39 +++++++++++-------- 1 file changed, 22 insertions(+), 17 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index b917e62668821..1f61d52366ddc 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -31,28 +31,33 @@ def check_result(self, x_shape, new_shape, format): mask = np.random.randint(0, 2, x_shape) np_x = np.random.randint(-100, 100, x_shape) * mask dense_x = paddle.to_tensor(np_x) + dense_x.numpy() + print(dense_x.numpy()) dense_x.stop_gradient = False + dense_x.numpy() # dense_out = paddle.transpose(dense_x, dims) dense_out = paddle.reshape(dense_x, new_shape) + dense_out.numpy() + print(dense_out.numpy()) - if format == "coo": - # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) - sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) - else: - # sp_x = origin_x.detach().to_sparse_csr() - sp_x = paddle.to_tensor(np_x).to_sparse_csr() - sp_x.stop_gradient = False - # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) - sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) + # if format == "coo": + # # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) + # sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) + # else: + # # sp_x = origin_x.detach().to_sparse_csr() + # sp_x = paddle.to_tensor(np_x).to_sparse_csr() + # sp_x.stop_gradient = False + # # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) + # sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') - print("dense_out.numpy():", dense_out.numpy()) - print("dense_out:", dense_out) - print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') - print("sp_out:", sp_out) - print("sp_out.to_dense():", sp_out.to_dense()) - print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) - print(10*'=', "OccupyMars2025 the end", 10*'=') + # print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') + # print("dense_out.numpy():", dense_out.numpy()) + # print("dense_out:", dense_out) + # print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') + # print("sp_out:", sp_out) + # print("sp_out.to_dense():", sp_out.to_dense()) + # print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) + # print(10*'=', "OccupyMars2025 the end", 10*'=') # np.testing.assert_allclose(sp_out.to_dense().numpy(), # dense_out.numpy(), From 64a3503cdc11ddf41d6f3884035afe54be3cbe7c Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Fri, 23 Sep 2022 23:11:00 +0800 Subject: [PATCH 12/55] check cpu and cuda versions --- .../tests/unittests/test_sparse_reshape_op.py | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 1f61d52366ddc..cf13b58431bdc 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -30,7 +30,20 @@ def check_result(self, x_shape, new_shape, format): # dense_x = paddle.clone(origin_x.detach()) mask = np.random.randint(0, 2, x_shape) np_x = np.random.randint(-100, 100, x_shape) * mask - dense_x = paddle.to_tensor(np_x) + ### cpu version + dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) + dense_x.numpy() + print(dense_x.numpy()) + dense_x.stop_gradient = False + dense_x.numpy() + # dense_out = paddle.transpose(dense_x, dims) + dense_out = paddle.reshape(dense_x, new_shape) + dense_out.numpy() + print(dense_out.numpy()) + + + ## cuda version + dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() print(dense_x.numpy()) dense_x.stop_gradient = False From 84f51dbf48997efe29eaa5dfb23722321d47e443 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sat, 24 Sep 2022 08:22:26 +0800 Subject: [PATCH 13/55] Update test_sparse_reshape_op.py --- .../tests/unittests/test_sparse_reshape_op.py | 87 ++++++++++++------- 1 file changed, 56 insertions(+), 31 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index cf13b58431bdc..880f90ff583dc 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -30,6 +30,8 @@ def check_result(self, x_shape, new_shape, format): # dense_x = paddle.clone(origin_x.detach()) mask = np.random.randint(0, 2, x_shape) np_x = np.random.randint(-100, 100, x_shape) * mask + + ### cpu version dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) dense_x.numpy() @@ -41,40 +43,63 @@ def check_result(self, x_shape, new_shape, format): dense_out.numpy() print(dense_out.numpy()) + if format == "coo": + # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) + sp_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()).to_sparse_coo(len(x_shape)) + else: + # sp_x = origin_x.detach().to_sparse_csr() + sp_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()).to_sparse_csr() + sp_x.stop_gradient = False + # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) + sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - ## cuda version - dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) - dense_x.numpy() - print(dense_x.numpy()) - dense_x.stop_gradient = False - dense_x.numpy() - # dense_out = paddle.transpose(dense_x, dims) - dense_out = paddle.reshape(dense_x, new_shape) - dense_out.numpy() - print(dense_out.numpy()) + print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') + print("dense_out.numpy():", dense_out.numpy()) + print("dense_out:", dense_out) + print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') + print("sp_out:", sp_out) + print("sp_out.to_dense():", sp_out.to_dense()) + print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) + print(10*'=', "OccupyMars2025 the end", 10*'=') + + np.testing.assert_allclose(sp_out.to_dense().numpy(), + dense_out.numpy(), + rtol=1e-05) - # if format == "coo": - # # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) - # sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) - # else: - # # sp_x = origin_x.detach().to_sparse_csr() - # sp_x = paddle.to_tensor(np_x).to_sparse_csr() - # sp_x.stop_gradient = False - # # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) - # sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) + if paddle.is_compiled_with_cuda(): + ## cuda version + dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) + dense_x.numpy() + print(dense_x.numpy()) + dense_x.stop_gradient = False + dense_x.numpy() + # dense_out = paddle.transpose(dense_x, dims) + dense_out = paddle.reshape(dense_x, new_shape) + dense_out.numpy() + print(dense_out.numpy()) - # print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') - # print("dense_out.numpy():", dense_out.numpy()) - # print("dense_out:", dense_out) - # print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') - # print("sp_out:", sp_out) - # print("sp_out.to_dense():", sp_out.to_dense()) - # print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) - # print(10*'=', "OccupyMars2025 the end", 10*'=') - - # np.testing.assert_allclose(sp_out.to_dense().numpy(), - # dense_out.numpy(), - # rtol=1e-05) + if format == "coo": + # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) + sp_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)).to_sparse_coo(len(x_shape)) + else: + # sp_x = origin_x.detach().to_sparse_csr() + sp_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)).to_sparse_csr() + sp_x.stop_gradient = False + # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) + sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) + + print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') + print("dense_out.numpy():", dense_out.numpy()) + print("dense_out:", dense_out) + print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') + print("sp_out:", sp_out) + print("sp_out.to_dense():", sp_out.to_dense()) + print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) + print(10*'=', "OccupyMars2025 the end", 10*'=') + + np.testing.assert_allclose(sp_out.to_dense().numpy(), + dense_out.numpy(), + rtol=1e-05) # dense_out.backward() # sp_out.backward() From 90bfea358acf470c8734d0b4fbfb226e4fe403fb Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sat, 24 Sep 2022 15:44:25 +0800 Subject: [PATCH 14/55] supply all test cases for cpu forward coo kernel --- .../tests/unittests/test_sparse_reshape_op.py | 46 ++++++++++--------- 1 file changed, 24 insertions(+), 22 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 880f90ff583dc..11d6f5db05e3b 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -21,7 +21,7 @@ class TestReshape(unittest.TestCase): # x: sparse, out: sparse def check_result(self, x_shape, new_shape, format): - with _test_eager_guard(): + # with _test_eager_guard(): # mask = paddle.randint(0, 2, x_shape).astype("float32") # origin_x = paddle.rand(x_shape, dtype='float32') * mask # mask = paddle.randint(0, 2, x_shape) @@ -53,20 +53,21 @@ def check_result(self, x_shape, new_shape, format): # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') - print("dense_out.numpy():", dense_out.numpy()) - print("dense_out:", dense_out) - print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') - print("sp_out:", sp_out) - print("sp_out.to_dense():", sp_out.to_dense()) - print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) - print(10*'=', "OccupyMars2025 the end", 10*'=') + # print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') + # print("dense_out.numpy():", dense_out.numpy()) + # print("dense_out:", dense_out) + # print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') + # print("sp_out:", sp_out) + # print("sp_out.to_dense():", sp_out.to_dense()) + # print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) + # print(10*'=', "OccupyMars2025 the end", 10*'=') np.testing.assert_allclose(sp_out.to_dense().numpy(), dense_out.numpy(), rtol=1e-05) - if paddle.is_compiled_with_cuda(): + # if paddle.is_compiled_with_cuda(): + if False: ## cuda version dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() @@ -104,32 +105,33 @@ def check_result(self, x_shape, new_shape, format): # dense_out.backward() # sp_out.backward() # np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - # # (dense_x.grad * mask).numpy(), - # dense_x.grad.numpy(), + # dense_x.grad.numpy() * mask, + # # dense_x.grad.numpy(), # rtol=1e-05) def test_reshape_2d(self): self.check_result([2, 5], [10,], 'coo') + self.check_result([12, 5], [15, 4], 'coo') # self.check_result([10, 5], [2, 25], 'csr') - # self.check_result([12, 5], [15, 4], 'coo') # self.check_result([9, 8], [18, 4], 'csr') - # def test_transpose_3d(self): - # self.check_result([6, 2, 3], [6, 2, 3], 'coo') + def test_transpose_3d(self): + self.check_result([6, 2, 3], [6, 2, 3], 'coo') + self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') + self.check_result([6, 2, 3], [1, 18, 2], 'coo') + self.check_result([6, 2, 3], [2, 9, 2], 'coo') + self.check_result([6, 2, 3], [2, 1, 18], 'coo') + self.check_result([6, 2, 3], [1, 2, 2, 3, 3], 'coo') + # self.check_result([6, 2, 3], [6, 2, 3], 'csr') - # self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') # self.check_result([6, 2, 3], [6, 3, 2], 'csr') - # # self.check_result([6, 2, 3], [1, 0, 2], 'coo') # # self.check_result([6, 2, 3], [1, 0, 2], 'csr') - # # self.check_result([6, 2, 3], [2, 0, 1], 'coo') # # self.check_result([6, 2, 3], [2, 0, 1], 'csr') - # # self.check_result([6, 2, 3], [2, 1, 0], 'coo') # # self.check_result([6, 2, 3], [2, 1, 0], 'csr') - # # self.check_result([6, 2, 3], [1, 2, 0], 'coo') # # self.check_result([6, 2, 3], [1, 2, 0], 'csr') - # def test_transpose_nd(self): - # self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') + def test_transpose_nd(self): + self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') # # Randint now only supports access to dimension 0 to 9. # # self.check_result([i % 3 + 2 for i in range(9)], # # [(i + 2) % 9 for i in range(9)], 'coo') From 21d35389cc75805c05843c8350779bbfe0a75461 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Thu, 29 Sep 2022 10:14:50 +0800 Subject: [PATCH 15/55] test forward coo cuda kernel --- paddle/phi/kernels/sparse/gpu/reshape_kernel.cu | 2 +- python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py | 3 ++- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index bbb86788a0861..2280292b822ff 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -48,7 +48,7 @@ __global__ void ReshapeCooCudaKernel(const int64_t *x_indices_data, template -void ReshapeCooKernel(const Context &dev_ctx, +void ReshapeCooKernel(const Context& dev_ctx, const SparseCooTensor &x, const phi::IntArray& shape, SparseCooTensor *out) { diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 11d6f5db05e3b..2c41bad045e0f 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -67,7 +67,8 @@ def check_result(self, x_shape, new_shape, format): rtol=1e-05) # if paddle.is_compiled_with_cuda(): - if False: + # if False: + if True: ## cuda version dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() From 687778aa0603e37c8a5f24f3d4d54ec0f8b47933 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Thu, 29 Sep 2022 10:25:29 +0800 Subject: [PATCH 16/55] change configuration of cuda kernel --- paddle/phi/kernels/sparse/gpu/reshape_kernel.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index 2280292b822ff..e13c50663e851 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -80,7 +80,8 @@ void ReshapeCooKernel(const Context& dev_ctx, auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, x_nnz, 1); ReshapeCooCudaKernel<<>>( x_indices_data, x_sparse_part_dims.size(), out_sparse_part_dims.size(), x_nnz, From 11ee8c3311e65981536c45850a2a3f51bd187cbd Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Thu, 29 Sep 2022 12:46:11 +0800 Subject: [PATCH 17/55] keep only one test case --- .../tests/unittests/test_sparse_reshape_op.py | 43 ++++++++++--------- 1 file changed, 22 insertions(+), 21 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 2c41bad045e0f..017570a0bb011 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -111,31 +111,32 @@ def check_result(self, x_shape, new_shape, format): # rtol=1e-05) def test_reshape_2d(self): - self.check_result([2, 5], [10,], 'coo') self.check_result([12, 5], [15, 4], 'coo') + # self.check_result([2, 5], [10,], 'coo') + # self.check_result([10, 5], [2, 25], 'csr') # self.check_result([9, 8], [18, 4], 'csr') - def test_transpose_3d(self): - self.check_result([6, 2, 3], [6, 2, 3], 'coo') - self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') - self.check_result([6, 2, 3], [1, 18, 2], 'coo') - self.check_result([6, 2, 3], [2, 9, 2], 'coo') - self.check_result([6, 2, 3], [2, 1, 18], 'coo') - self.check_result([6, 2, 3], [1, 2, 2, 3, 3], 'coo') - - # self.check_result([6, 2, 3], [6, 2, 3], 'csr') - # self.check_result([6, 2, 3], [6, 3, 2], 'csr') - # # self.check_result([6, 2, 3], [1, 0, 2], 'csr') - # # self.check_result([6, 2, 3], [2, 0, 1], 'csr') - # # self.check_result([6, 2, 3], [2, 1, 0], 'csr') - # # self.check_result([6, 2, 3], [1, 2, 0], 'csr') - - def test_transpose_nd(self): - self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') - # # Randint now only supports access to dimension 0 to 9. - # # self.check_result([i % 3 + 2 for i in range(9)], - # # [(i + 2) % 9 for i in range(9)], 'coo') + # def test_transpose_3d(self): + # self.check_result([6, 2, 3], [6, 2, 3], 'coo') + # self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') + # self.check_result([6, 2, 3], [1, 18, 2], 'coo') + # self.check_result([6, 2, 3], [2, 9, 2], 'coo') + # self.check_result([6, 2, 3], [2, 1, 18], 'coo') + # self.check_result([6, 2, 3], [1, 2, 2, 3, 3], 'coo') + + # # self.check_result([6, 2, 3], [6, 2, 3], 'csr') + # # self.check_result([6, 2, 3], [6, 3, 2], 'csr') + # # # self.check_result([6, 2, 3], [1, 0, 2], 'csr') + # # # self.check_result([6, 2, 3], [2, 0, 1], 'csr') + # # # self.check_result([6, 2, 3], [2, 1, 0], 'csr') + # # # self.check_result([6, 2, 3], [1, 2, 0], 'csr') + + # def test_transpose_nd(self): + # self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') + # # # Randint now only supports access to dimension 0 to 9. + # # # self.check_result([i % 3 + 2 for i in range(9)], + # # # [(i + 2) % 9 for i in range(9)], 'coo') if __name__ == "__main__": From de6d90351f61ac12849334606afee46dba2f1b1b Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Thu, 29 Sep 2022 15:27:22 +0800 Subject: [PATCH 18/55] test coo cpu kernel (forward and backward) --- .../tests/unittests/test_sparse_reshape_op.py | 23 +++++++++++-------- 1 file changed, 13 insertions(+), 10 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 017570a0bb011..ead52a2fcc01c 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -67,8 +67,8 @@ def check_result(self, x_shape, new_shape, format): rtol=1e-05) # if paddle.is_compiled_with_cuda(): - # if False: - if True: + if False: + # if True: ## cuda version dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() @@ -103,16 +103,19 @@ def check_result(self, x_shape, new_shape, format): dense_out.numpy(), rtol=1e-05) - # dense_out.backward() - # sp_out.backward() - # np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - # dense_x.grad.numpy() * mask, - # # dense_x.grad.numpy(), - # rtol=1e-05) + dense_out.backward() + sp_out.backward() + print("sp_x.grad.to_dense().numpy(): ", sp_x.grad.to_dense().numpy()) + print("dense_x.grad.numpy(): ", dense_x.grad.numpy()) + print("mask: ", mask) + np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), + dense_x.grad.numpy() * mask, + # dense_x.grad.numpy(), + rtol=1e-05) def test_reshape_2d(self): - self.check_result([12, 5], [15, 4], 'coo') - # self.check_result([2, 5], [10,], 'coo') + self.check_result([2, 5], [10,], 'coo') + # self.check_result([12, 5], [15, 4], 'coo') # self.check_result([10, 5], [2, 25], 'csr') # self.check_result([9, 8], [18, 4], 'csr') From 8b09f244d0de6c3c47c34b93460f4ec97a266780 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sat, 1 Oct 2022 16:43:40 +0800 Subject: [PATCH 19/55] row major or column major ??? --- paddle/phi/kernels/sparse/gpu/reshape_kernel.cu | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index e13c50663e851..0def787502a3c 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -37,10 +37,14 @@ __global__ void ReshapeCooCudaKernel(const int64_t *x_indices_data, CUDA_KERNEL_LOOP_TYPE(j, x_nnz, int64_t) { int64_t location = 0; for (int i = 0; i < num_x_sparse_part_dims; ++i) { - location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; + // location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; + // row major or column major ??? + location += x_indices_data[j * num_x_sparse_part_dims + i] * x_sparse_part_strides[i]; } for (int i = 0; i < num_out_sparse_part_dims; ++i) { - out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; + // out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; + // row major or column major ??? + out_indices_data[j * num_out_sparse_part_dims + i] = location / out_sparse_part_strides[i]; location %= out_sparse_part_strides[i]; } } From 4420f4fd8cb429190940e5799092299f73719f93 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sat, 1 Oct 2022 16:45:34 +0800 Subject: [PATCH 20/55] test cuda coo forward kernel --- .../tests/unittests/test_sparse_reshape_op.py | 22 +++++++++---------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index ead52a2fcc01c..fa46aeae9c74f 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -67,8 +67,8 @@ def check_result(self, x_shape, new_shape, format): rtol=1e-05) # if paddle.is_compiled_with_cuda(): - if False: - # if True: + # if False: + if True: ## cuda version dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() @@ -103,15 +103,15 @@ def check_result(self, x_shape, new_shape, format): dense_out.numpy(), rtol=1e-05) - dense_out.backward() - sp_out.backward() - print("sp_x.grad.to_dense().numpy(): ", sp_x.grad.to_dense().numpy()) - print("dense_x.grad.numpy(): ", dense_x.grad.numpy()) - print("mask: ", mask) - np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - dense_x.grad.numpy() * mask, - # dense_x.grad.numpy(), - rtol=1e-05) + # dense_out.backward() + # sp_out.backward() + # print("sp_x.grad.to_dense().numpy(): ", sp_x.grad.to_dense().numpy()) + # print("dense_x.grad.numpy(): ", dense_x.grad.numpy()) + # print("mask: ", mask) + # np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), + # dense_x.grad.numpy() * mask, + # # dense_x.grad.numpy(), + # rtol=1e-05) def test_reshape_2d(self): self.check_result([2, 5], [10,], 'coo') From d74bda932d285fc113a523ab15b9037bbac0bd55 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sat, 1 Oct 2022 17:50:15 +0800 Subject: [PATCH 21/55] complete declaration and registration --- paddle/phi/api/yaml/sparse_ops.yaml | 11 +++++++++++ python/paddle/incubate/sparse/__init__.py | 3 +++ 2 files changed, 14 insertions(+) diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index 1d7a4c0bafe53..c6ad1bfa583cb 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -469,3 +469,14 @@ transpose_csr{sparse_csr -> sparse_csr} layout : x backward : transpose_grad + +- op : reshape + args : (Tensor x, IntArray shape) + output : Tensor(out) + infer_meta : + func : ReshapeInferMeta + kernel : + func : reshape_coo{sparse_coo -> sparse_coo}, + reshape_csr{sparse_csr -> sparse_csr} + layout : x + backward : reshape_grad diff --git a/python/paddle/incubate/sparse/__init__.py b/python/paddle/incubate/sparse/__init__.py index 581310fbbd9b6..d4352ee2e9bde 100644 --- a/python/paddle/incubate/sparse/__init__.py +++ b/python/paddle/incubate/sparse/__init__.py @@ -12,6 +12,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from sympy import im from .creation import sparse_coo_tensor from .creation import sparse_csr_tensor @@ -35,6 +36,7 @@ from .unary import rad2deg from .unary import expm1 from .unary import transpose +from .unary import reshape from .binary import mv from .binary import matmul @@ -81,4 +83,5 @@ 'divide', 'coalesce', 'is_same_shape', + 'reshape' ] From 17ec7c31d4ef06df3501c20d9176b321b249a0de Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sat, 1 Oct 2022 18:13:39 +0800 Subject: [PATCH 22/55] Update __init__.py --- python/paddle/incubate/sparse/__init__.py | 1 - 1 file changed, 1 deletion(-) diff --git a/python/paddle/incubate/sparse/__init__.py b/python/paddle/incubate/sparse/__init__.py index d4352ee2e9bde..bbc110939c6a8 100644 --- a/python/paddle/incubate/sparse/__init__.py +++ b/python/paddle/incubate/sparse/__init__.py @@ -12,7 +12,6 @@ # See the License for the specific language governing permissions and # limitations under the License. -from sympy import im from .creation import sparse_coo_tensor from .creation import sparse_csr_tensor From ca08918b90201f9c627fc5052dadad0a8642fef7 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sat, 1 Oct 2022 20:17:11 +0800 Subject: [PATCH 23/55] rebuild --- python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index fa46aeae9c74f..d97904557724d 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -69,7 +69,7 @@ def check_result(self, x_shape, new_shape, format): # if paddle.is_compiled_with_cuda(): # if False: if True: - ## cuda version + ## cuda version dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() print(dense_x.numpy()) From cdd78a10a8fb2c6efbbcacbb0c30e7aa7bbb09c6 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sat, 1 Oct 2022 21:26:04 +0800 Subject: [PATCH 24/55] retrigger CI --- python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index d97904557724d..07b8ad19885af 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -15,7 +15,7 @@ import paddle import numpy as np import unittest -from paddle.fluid.framework import _test_eager_guard +# from paddle.fluid.framework import _test_eager_guard class TestReshape(unittest.TestCase): @@ -28,7 +28,7 @@ def check_result(self, x_shape, new_shape, format): # origin_x = paddle.randint(-100, 100, x_shape) * mask # dense_x = paddle.clone(origin_x.detach()) - mask = np.random.randint(0, 2, x_shape) + mask = np.random.randint( 0, 2, x_shape) np_x = np.random.randint(-100, 100, x_shape) * mask From 471c648f9254913628cc9eae5cf138679c08af89 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sun, 2 Oct 2022 03:25:11 +0800 Subject: [PATCH 25/55] add cudaMalloc and cudaMemcpy in ReshapeCooKernel and change back to row major order in a cuda dense tensor --- .../phi/kernels/sparse/gpu/reshape_kernel.cu | 64 +++++++++++++++---- 1 file changed, 51 insertions(+), 13 deletions(-) diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index 0def787502a3c..59ebf17105194 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -26,10 +26,10 @@ namespace phi { namespace sparse { -__global__ void ReshapeCooCudaKernel(const int64_t *x_indices_data, - const int& num_x_sparse_part_dims, - const int& num_out_sparse_part_dims, - const int64_t& x_nnz, +__global__ void ReshapeCooCudaKernel(const int64_t* x_indices_data, + const int num_x_sparse_part_dims, + const int num_out_sparse_part_dims, + const int64_t x_nnz, const int64_t* x_sparse_part_strides, const int64_t* out_sparse_part_strides, int64_t *out_indices_data) { @@ -37,14 +37,14 @@ __global__ void ReshapeCooCudaKernel(const int64_t *x_indices_data, CUDA_KERNEL_LOOP_TYPE(j, x_nnz, int64_t) { int64_t location = 0; for (int i = 0; i < num_x_sparse_part_dims; ++i) { - // location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; + location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; // row major or column major ??? - location += x_indices_data[j * num_x_sparse_part_dims + i] * x_sparse_part_strides[i]; + // location += x_indices_data[j * num_x_sparse_part_dims + i] * x_sparse_part_strides[i]; } for (int i = 0; i < num_out_sparse_part_dims; ++i) { - // out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; + out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; // row major or column major ??? - out_indices_data[j * num_out_sparse_part_dims + i] = location / out_sparse_part_strides[i]; + // out_indices_data[j * num_out_sparse_part_dims + i] = location / out_sparse_part_strides[i]; location %= out_sparse_part_strides[i]; } } @@ -57,6 +57,8 @@ void ReshapeCooKernel(const Context& dev_ctx, const phi::IntArray& shape, SparseCooTensor *out) { int64_t x_nnz = x.nnz(); + + // TODO: consider using "DDim DDim::reshape(std::vector& shape)" DDim out_dims = phi::make_ddim(shape.GetData()); // get sparse part dimensions of x and out std::vector x_sparse_part_dims; @@ -75,21 +77,57 @@ void ReshapeCooKernel(const Context& dev_ctx, out->SetMember(out_indices, out_values, out_dims, x.coalesced()); // compute values of out indices - const DenseTensor& x_indices = x.indices(); - const auto *x_indices_data = x_indices.data(); + // const DenseTensor& x_indices = x.indices(); + // const auto *x_indices_data = x_indices.data(); + const auto *x_indices_data = x.indices().data(); auto *out_indices_data = out_indices.data(); const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); + int64_t* destination_x_sparse_part_strides, destination_out_sparse_part_strides; + +#ifdef PADDLE_WITH_HIP + hipMalloc(reinterpret_cast(&destination_x_sparse_part_strides), + sizeof(int64_t) * x_sparse_part_strides.size()); + hipMemcpy( + destination_x_sparse_part_strides, + x_sparse_part_strides.Get(), + sizeof(int64_t) * x_sparse_part_strides.size(), + hipMemcpyHostToDevice); + hipMalloc(reinterpret_cast(&destination_out_sparse_part_strides), + sizeof(int64_t) * out_sparse_part_strides.size()); + hipMemcpy( + destination_out_sparse_part_strides, + out_sparse_part_strides.Get(), + sizeof(int64_t) * out_sparse_part_strides.size(), + hipMemcpyHostToDevice); +#else + cudaMalloc(reinterpret_cast(&destination_x_sparse_part_strides), + sizeof(int64_t) * x_sparse_part_strides.size()); + cudaMemcpy( + destination_x_sparse_part_strides, + x_sparse_part_strides.Get(), + sizeof(int64_t) * x_sparse_part_strides.size(), + cudaMemcpyHostToDevice); + cudaMalloc(reinterpret_cast(&destination_out_sparse_part_strides), + sizeof(int64_t) * out_sparse_part_strides.size()); + cudaMemcpy( + destination_out_sparse_part_strides, + out_sparse_part_strides.Get(), + sizeof(int64_t) * out_sparse_part_strides.size(), + cudaMemcpyHostToDevice); +#endif + auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, x_nnz, 1); ReshapeCooCudaKernel<<>>( x_indices_data, x_sparse_part_dims.size(), out_sparse_part_dims.size(), x_nnz, - x_sparse_part_strides.Get(), out_sparse_part_strides.Get(), + // x_sparse_part_strides.Get(), out_sparse_part_strides.Get(), + destination_x_sparse_part_strides, destination_out_sparse_part_strides, out_indices_data ); } From 5f332916a2d0ac7391f2f0ae48648ba92d42ab45 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sun, 2 Oct 2022 03:37:38 +0800 Subject: [PATCH 26/55] midify minor error --- paddle/phi/kernels/sparse/gpu/reshape_kernel.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index 59ebf17105194..815e14b410672 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -84,7 +84,7 @@ void ReshapeCooKernel(const Context& dev_ctx, const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); - int64_t* destination_x_sparse_part_strides, destination_out_sparse_part_strides; + int64_t* destination_x_sparse_part_strides, *destination_out_sparse_part_strides; #ifdef PADDLE_WITH_HIP hipMalloc(reinterpret_cast(&destination_x_sparse_part_strides), From f1a9d9b4cdff5683375521bf7091a8f7741b8500 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sun, 2 Oct 2022 08:13:58 +0800 Subject: [PATCH 27/55] test only cpu coo forward kernel --- python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 07b8ad19885af..968bd86cca260 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -67,8 +67,8 @@ def check_result(self, x_shape, new_shape, format): rtol=1e-05) # if paddle.is_compiled_with_cuda(): - # if False: - if True: + if False: + # if True: ## cuda version dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() From 6aa30618b2731e9ed8262c97187b57d6fc5a372e Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sun, 2 Oct 2022 16:29:52 +0800 Subject: [PATCH 28/55] add all test cases for coo forward kernel (both cpu and gpu) --- .../tests/unittests/test_sparse_reshape_op.py | 30 +++++++++---------- 1 file changed, 15 insertions(+), 15 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 968bd86cca260..56e48f4e8a6d6 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -67,8 +67,8 @@ def check_result(self, x_shape, new_shape, format): rtol=1e-05) # if paddle.is_compiled_with_cuda(): - if False: - # if True: + # if False: + if True: ## cuda version dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() @@ -115,18 +115,18 @@ def check_result(self, x_shape, new_shape, format): def test_reshape_2d(self): self.check_result([2, 5], [10,], 'coo') - # self.check_result([12, 5], [15, 4], 'coo') + self.check_result([12, 5], [15, 4], 'coo') # self.check_result([10, 5], [2, 25], 'csr') # self.check_result([9, 8], [18, 4], 'csr') - # def test_transpose_3d(self): - # self.check_result([6, 2, 3], [6, 2, 3], 'coo') - # self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') - # self.check_result([6, 2, 3], [1, 18, 2], 'coo') - # self.check_result([6, 2, 3], [2, 9, 2], 'coo') - # self.check_result([6, 2, 3], [2, 1, 18], 'coo') - # self.check_result([6, 2, 3], [1, 2, 2, 3, 3], 'coo') + def test_transpose_3d(self): + self.check_result([6, 2, 3], [6, 2, 3], 'coo') + self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') + self.check_result([6, 2, 3], [1, 18, 2], 'coo') + self.check_result([6, 2, 3], [2, 9, 2], 'coo') + self.check_result([6, 2, 3], [2, 1, 18], 'coo') + self.check_result([6, 2, 3], [1, 2, 2, 3, 3], 'coo') # # self.check_result([6, 2, 3], [6, 2, 3], 'csr') # # self.check_result([6, 2, 3], [6, 3, 2], 'csr') @@ -135,11 +135,11 @@ def test_reshape_2d(self): # # # self.check_result([6, 2, 3], [2, 1, 0], 'csr') # # # self.check_result([6, 2, 3], [1, 2, 0], 'csr') - # def test_transpose_nd(self): - # self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') - # # # Randint now only supports access to dimension 0 to 9. - # # # self.check_result([i % 3 + 2 for i in range(9)], - # # # [(i + 2) % 9 for i in range(9)], 'coo') + def test_transpose_nd(self): + self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') + # # Randint now only supports access to dimension 0 to 9. + # # self.check_result([i % 3 + 2 for i in range(9)], + # # [(i + 2) % 9 for i in range(9)], 'coo') if __name__ == "__main__": From 499b78a65edc56b3d6b9b9f1b266e4cc45ea8bb3 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sun, 2 Oct 2022 18:51:42 +0800 Subject: [PATCH 29/55] test all forward kernels (coo, csr; cpu, gpu) --- .../tests/unittests/test_sparse_reshape_op.py | 25 ++++++++----------- python/paddle/incubate/sparse/unary.py | 2 +- 2 files changed, 11 insertions(+), 16 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 56e48f4e8a6d6..62a269a5da839 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -34,14 +34,9 @@ def check_result(self, x_shape, new_shape, format): ### cpu version dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) - dense_x.numpy() - print(dense_x.numpy()) dense_x.stop_gradient = False - dense_x.numpy() # dense_out = paddle.transpose(dense_x, dims) dense_out = paddle.reshape(dense_x, new_shape) - dense_out.numpy() - print(dense_out.numpy()) if format == "coo": # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) @@ -67,8 +62,8 @@ def check_result(self, x_shape, new_shape, format): rtol=1e-05) # if paddle.is_compiled_with_cuda(): - # if False: - if True: + if False: + # if True: ## cuda version dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.numpy() @@ -117,8 +112,8 @@ def test_reshape_2d(self): self.check_result([2, 5], [10,], 'coo') self.check_result([12, 5], [15, 4], 'coo') - # self.check_result([10, 5], [2, 25], 'csr') - # self.check_result([9, 8], [18, 4], 'csr') + self.check_result([10, 5], [2, 25], 'csr') + self.check_result([9, 8], [18, 4], 'csr') def test_transpose_3d(self): self.check_result([6, 2, 3], [6, 2, 3], 'coo') @@ -128,12 +123,12 @@ def test_transpose_3d(self): self.check_result([6, 2, 3], [2, 1, 18], 'coo') self.check_result([6, 2, 3], [1, 2, 2, 3, 3], 'coo') - # # self.check_result([6, 2, 3], [6, 2, 3], 'csr') - # # self.check_result([6, 2, 3], [6, 3, 2], 'csr') - # # # self.check_result([6, 2, 3], [1, 0, 2], 'csr') - # # # self.check_result([6, 2, 3], [2, 0, 1], 'csr') - # # # self.check_result([6, 2, 3], [2, 1, 0], 'csr') - # # # self.check_result([6, 2, 3], [1, 2, 0], 'csr') + self.check_result([6, 2, 3], [6, 2, 3], 'csr') + self.check_result([6, 2, 3], [6, 3, 2], 'csr') + self.check_result([6, 2, 3], [2, 6, 3], 'csr') + self.check_result([6, 2, 3], [3, 6, 2], 'csr') + self.check_result([6, 2, 3], [4, 9, 1], 'csr') + self.check_result([6, 2, 3], [12, 1, 3], 'csr') def test_transpose_nd(self): self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index 09c6bc41be18a..81712d22b41d7 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -645,7 +645,7 @@ def reshape(x, shape, name=None): """ Changes the shape of ``x`` without changing its value, requiring x to be a SparseCooTensor or SparseCsrTensor. Currently this function can only reshape the sparse dims of ``x`` , but ``shape`` argument must - include all dims of the reshaped tensor. + include all dims of the reshaped tensor. Note: if x is a SparseCsrTensor, then len(shape) must be 2 or 3. .. math:: out = reshape(x, shape) Parameters: From 2e913357fec2f7a5e3852a2e5c1a48403c36cc0f Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Sun, 2 Oct 2022 20:18:18 +0800 Subject: [PATCH 30/55] add all test cases for all kinds of kernels --- .../tests/unittests/test_sparse_reshape_op.py | 125 +++++++----------- 1 file changed, 49 insertions(+), 76 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 62a269a5da839..c23be16b40506 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -21,92 +21,68 @@ class TestReshape(unittest.TestCase): # x: sparse, out: sparse def check_result(self, x_shape, new_shape, format): - # with _test_eager_guard(): - # mask = paddle.randint(0, 2, x_shape).astype("float32") - # origin_x = paddle.rand(x_shape, dtype='float32') * mask - # mask = paddle.randint(0, 2, x_shape) - # origin_x = paddle.randint(-100, 100, x_shape) * mask + mask = np.random.randint( 0, 2, x_shape) + np_x = np.random.randint(-100, 100, x_shape) * mask + + + ### cpu version + dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) + dense_x.stop_gradient = False + # dense_out = paddle.transpose(dense_x, dims) + dense_out = paddle.reshape(dense_x, new_shape) + + if format == "coo": + # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) + sp_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()).to_sparse_coo(len(x_shape)) + else: + # sp_x = origin_x.detach().to_sparse_csr() + sp_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()).to_sparse_csr() + sp_x.stop_gradient = False + # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) + sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) + + np.testing.assert_allclose(sp_out.to_dense().numpy(), + dense_out.numpy(), + rtol=1e-05) + + # check cpu backward computation + dense_out.backward() + sp_out.backward() + np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), + # dense_x.grad.numpy() * mask, + dense_x.grad.numpy() * np_x.astype('bool').astype('int'), + rtol=1e-05) - # dense_x = paddle.clone(origin_x.detach()) - mask = np.random.randint( 0, 2, x_shape) - np_x = np.random.randint(-100, 100, x_shape) * mask - - - ### cpu version - dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) + # if paddle.is_compiled_with_cuda(): + # if False: + if True: + ## cuda version + dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) dense_x.stop_gradient = False # dense_out = paddle.transpose(dense_x, dims) dense_out = paddle.reshape(dense_x, new_shape) if format == "coo": # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) - sp_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()).to_sparse_coo(len(x_shape)) + sp_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)).to_sparse_coo(len(x_shape)) else: # sp_x = origin_x.detach().to_sparse_csr() - sp_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()).to_sparse_csr() + sp_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)).to_sparse_csr() sp_x.stop_gradient = False # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - # print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') - # print("dense_out.numpy():", dense_out.numpy()) - # print("dense_out:", dense_out) - # print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') - # print("sp_out:", sp_out) - # print("sp_out.to_dense():", sp_out.to_dense()) - # print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) - # print(10*'=', "OccupyMars2025 the end", 10*'=') - np.testing.assert_allclose(sp_out.to_dense().numpy(), dense_out.numpy(), rtol=1e-05) - - # if paddle.is_compiled_with_cuda(): - if False: - # if True: - ## cuda version - dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) - dense_x.numpy() - print(dense_x.numpy()) - dense_x.stop_gradient = False - dense_x.numpy() - # dense_out = paddle.transpose(dense_x, dims) - dense_out = paddle.reshape(dense_x, new_shape) - dense_out.numpy() - print(dense_out.numpy()) - - if format == "coo": - # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) - sp_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)).to_sparse_coo(len(x_shape)) - else: - # sp_x = origin_x.detach().to_sparse_csr() - sp_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)).to_sparse_csr() - sp_x.stop_gradient = False - # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) - sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - - print(10*'=', "OccupyMars2025 the following is dense_out", 10*'=') - print("dense_out.numpy():", dense_out.numpy()) - print("dense_out:", dense_out) - print(10*'=', "OccupyMars2025 the following is sp_out", 10*'=') - print("sp_out:", sp_out) - print("sp_out.to_dense():", sp_out.to_dense()) - print("sp_out.to_dense().numpy():", sp_out.to_dense().numpy()) - print(10*'=', "OccupyMars2025 the end", 10*'=') - - np.testing.assert_allclose(sp_out.to_dense().numpy(), - dense_out.numpy(), - rtol=1e-05) - - # dense_out.backward() - # sp_out.backward() - # print("sp_x.grad.to_dense().numpy(): ", sp_x.grad.to_dense().numpy()) - # print("dense_x.grad.numpy(): ", dense_x.grad.numpy()) - # print("mask: ", mask) - # np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - # dense_x.grad.numpy() * mask, - # # dense_x.grad.numpy(), - # rtol=1e-05) + + # check gpu backward computation + dense_out.backward() + sp_out.backward() + np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), + # dense_x.grad.numpy() * mask, + dense_x.grad.numpy() * np_x.astype('bool').astype('int'), + rtol=1e-05) def test_reshape_2d(self): self.check_result([2, 5], [10,], 'coo') @@ -115,7 +91,7 @@ def test_reshape_2d(self): self.check_result([10, 5], [2, 25], 'csr') self.check_result([9, 8], [18, 4], 'csr') - def test_transpose_3d(self): + def test_reshape_3d(self): self.check_result([6, 2, 3], [6, 2, 3], 'coo') self.check_result([6, 2, 3], [2, 3, 3, 2], 'coo') self.check_result([6, 2, 3], [1, 18, 2], 'coo') @@ -130,12 +106,9 @@ def test_transpose_3d(self): self.check_result([6, 2, 3], [4, 9, 1], 'csr') self.check_result([6, 2, 3], [12, 1, 3], 'csr') - def test_transpose_nd(self): + def test_reshape_nd(self): self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') - # # Randint now only supports access to dimension 0 to 9. - # # self.check_result([i % 3 + 2 for i in range(9)], - # # [(i + 2) % 9 for i in range(9)], 'coo') - + self.check_result([3, 4, 4, 5, 7], [1, 12, 2, 5, 14], 'coo') if __name__ == "__main__": unittest.main() From cc21e67fb78b522d7eba220d5c1b8bdc551b8dfb Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 3 Oct 2022 08:43:40 +0800 Subject: [PATCH 31/55] just retrigger CI --- paddle/phi/kernels/sparse/gpu/reshape_kernel.cu | 1 - .../paddle/fluid/tests/unittests/test_sparse_reshape_op.py | 5 ++--- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index 815e14b410672..f04ec850bb867 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -57,7 +57,6 @@ void ReshapeCooKernel(const Context& dev_ctx, const phi::IntArray& shape, SparseCooTensor *out) { int64_t x_nnz = x.nnz(); - // TODO: consider using "DDim DDim::reshape(std::vector& shape)" DDim out_dims = phi::make_ddim(shape.GetData()); // get sparse part dimensions of x and out diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index c23be16b40506..bcea4aa3f10e1 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -23,9 +23,8 @@ class TestReshape(unittest.TestCase): def check_result(self, x_shape, new_shape, format): mask = np.random.randint( 0, 2, x_shape) np_x = np.random.randint(-100, 100, x_shape) * mask - - - ### cpu version + + # cpu version dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) dense_x.stop_gradient = False # dense_out = paddle.transpose(dense_x, dims) From b4e6d2f99571a7cdc4d86482383b680da7c2f396 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 13:55:51 +0800 Subject: [PATCH 32/55] Update sparse_ops.yaml --- paddle/phi/api/yaml/sparse_ops.yaml | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index c6ad1bfa583cb..a956644a075d8 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -470,6 +470,16 @@ layout : x backward : transpose_grad +- op : sync_batch_norm + args : (Tensor x, Tensor scale, Tensor bias, Tensor mean, Tensor variance, float momentum, float epsilon, str data_layout, bool is_test, bool use_global_stats, bool trainable_statistics, bool fuse_with_relu) + output : Tensor(out), Tensor(mean_out), Tensor(variance_out), Tensor(saved_mean), Tensor(saved_variance), Tensor(reserve_space) + infer_meta : + func : BatchNormInferMeta + kernel : + func : sync_batch_norm_coo{sparse_coo, dense, dense, dense, dense -> sparse_coo, dense, dense, dense, dense, dense} + data_type : x + backward : sync_batch_norm_grad + - op : reshape args : (Tensor x, IntArray shape) output : Tensor(out) From b5d6dbce747b26d115edf985b3298c05c8d1a1ea Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 13:58:32 +0800 Subject: [PATCH 33/55] Update sparse_ops.yaml --- paddle/phi/api/yaml/sparse_ops.yaml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index a956644a075d8..9d9d36983ae4d 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -479,7 +479,7 @@ func : sync_batch_norm_coo{sparse_coo, dense, dense, dense, dense -> sparse_coo, dense, dense, dense, dense, dense} data_type : x backward : sync_batch_norm_grad - + - op : reshape args : (Tensor x, IntArray shape) output : Tensor(out) @@ -489,4 +489,4 @@ func : reshape_coo{sparse_coo -> sparse_coo}, reshape_csr{sparse_csr -> sparse_csr} layout : x - backward : reshape_grad + backward : reshape_grad \ No newline at end of file From 46247a7aa7977146ec206f1fdd8ba3cbd9b70184 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 14:00:43 +0800 Subject: [PATCH 34/55] Update sparse_ops.yaml --- paddle/phi/api/yaml/sparse_ops.yaml | 13 +------------ 1 file changed, 1 insertion(+), 12 deletions(-) diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index 9d9d36983ae4d..9d52f319304b9 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -478,15 +478,4 @@ kernel : func : sync_batch_norm_coo{sparse_coo, dense, dense, dense, dense -> sparse_coo, dense, dense, dense, dense, dense} data_type : x - backward : sync_batch_norm_grad - -- op : reshape - args : (Tensor x, IntArray shape) - output : Tensor(out) - infer_meta : - func : ReshapeInferMeta - kernel : - func : reshape_coo{sparse_coo -> sparse_coo}, - reshape_csr{sparse_csr -> sparse_csr} - layout : x - backward : reshape_grad \ No newline at end of file + backward : sync_batch_norm_grad \ No newline at end of file From 0661bd2cfad5b76c112763b9ed6b1e3b137638ee Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 14:02:16 +0800 Subject: [PATCH 35/55] resolve conflicts --- paddle/phi/api/yaml/sparse_ops.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index 9d52f319304b9..b50d2e74904fb 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -478,4 +478,4 @@ kernel : func : sync_batch_norm_coo{sparse_coo, dense, dense, dense, dense -> sparse_coo, dense, dense, dense, dense, dense} data_type : x - backward : sync_batch_norm_grad \ No newline at end of file + backward : sync_batch_norm_grad From 3df911cff06c4db3fb1883a4fac9029d47c30f48 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 14:06:27 +0800 Subject: [PATCH 36/55] Update sparse_ops.yaml --- paddle/phi/api/yaml/sparse_ops.yaml | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index 43f8688fb81a0..6cdd0f6dbeb51 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -489,3 +489,15 @@ func : sync_batch_norm_coo{sparse_coo, dense, dense, dense, dense -> sparse_coo, dense, dense, dense, dense, dense} data_type : x backward : sync_batch_norm_grad + +- op : reshape + args : (Tensor x, IntArray shape) + output : Tensor(out) + infer_meta : + func : ReshapeInferMeta + kernel : + func : reshape_coo{sparse_coo -> sparse_coo}, + reshape_csr{sparse_csr -> sparse_csr} + layout : x + backward : reshape_grad + \ No newline at end of file From eaca3a21cf5a797f23512f0322f1285d00eb54d7 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 14:13:13 +0800 Subject: [PATCH 37/55] don't specify tensor place --- .../tests/unittests/test_sparse_reshape_op.py | 37 ++----------------- 1 file changed, 3 insertions(+), 34 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index bcea4aa3f10e1..cc9f03c519615 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -25,17 +25,17 @@ def check_result(self, x_shape, new_shape, format): np_x = np.random.randint(-100, 100, x_shape) * mask # cpu version - dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) + dense_x = paddle.to_tensor(np_x) dense_x.stop_gradient = False # dense_out = paddle.transpose(dense_x, dims) dense_out = paddle.reshape(dense_x, new_shape) if format == "coo": # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) - sp_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()).to_sparse_coo(len(x_shape)) + sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) else: # sp_x = origin_x.detach().to_sparse_csr() - sp_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()).to_sparse_csr() + sp_x = paddle.to_tensor(np_x).to_sparse_csr() sp_x.stop_gradient = False # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) @@ -52,37 +52,6 @@ def check_result(self, x_shape, new_shape, format): dense_x.grad.numpy() * np_x.astype('bool').astype('int'), rtol=1e-05) - # if paddle.is_compiled_with_cuda(): - # if False: - if True: - ## cuda version - dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) - dense_x.stop_gradient = False - # dense_out = paddle.transpose(dense_x, dims) - dense_out = paddle.reshape(dense_x, new_shape) - - if format == "coo": - # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) - sp_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)).to_sparse_coo(len(x_shape)) - else: - # sp_x = origin_x.detach().to_sparse_csr() - sp_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)).to_sparse_csr() - sp_x.stop_gradient = False - # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) - sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - - np.testing.assert_allclose(sp_out.to_dense().numpy(), - dense_out.numpy(), - rtol=1e-05) - - # check gpu backward computation - dense_out.backward() - sp_out.backward() - np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - # dense_x.grad.numpy() * mask, - dense_x.grad.numpy() * np_x.astype('bool').astype('int'), - rtol=1e-05) - def test_reshape_2d(self): self.check_result([2, 5], [10,], 'coo') self.check_result([12, 5], [15, 4], 'coo') From d80cf1a6ab3b54c4b74a88610f9140d1d628f4b4 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 15:24:04 +0800 Subject: [PATCH 38/55] new shape has -1 or 0 in it --- paddle/phi/kernels/sparse/cpu/reshape_kernel.cc | 5 ++++- paddle/phi/kernels/sparse/gpu/reshape_kernel.cu | 3 ++- .../tests/unittests/test_sparse_reshape_op.py | 15 +++++++++++++++ 3 files changed, 21 insertions(+), 2 deletions(-) diff --git a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc index cc6679317738b..d8096e819f71c 100644 --- a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc +++ b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc @@ -35,7 +35,10 @@ void ReshapeCooKernel(const Context& dev_ctx, SparseCooTensor* out) { // TODO: Currently, reshape is only applicable to sparse dims int64_t x_nnz = x.nnz(); - DDim out_dims = phi::make_ddim(shape.GetData()); + + // DDim out_dims = phi::make_ddim(shape.GetData()); + // Use DDim::reshape to handle -1 and 0 in the argument "shape" + DDim out_dims = x.dims().reshape(std::vector(shape.GetData().begin(), shape.GetData().end())); // get sparse part dimensions of x and out std::vector x_sparse_part_dims; std::vector out_sparse_part_dims; diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index f04ec850bb867..e7f3f5d4d42f7 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -58,7 +58,8 @@ void ReshapeCooKernel(const Context& dev_ctx, SparseCooTensor *out) { int64_t x_nnz = x.nnz(); // TODO: consider using "DDim DDim::reshape(std::vector& shape)" - DDim out_dims = phi::make_ddim(shape.GetData()); + // DDim out_dims = phi::make_ddim(shape.GetData()); + DDim out_dims = x.dims().reshape(std::vector(shape.GetData().begin(), shape.GetData().end())); // get sparse part dimensions of x and out std::vector x_sparse_part_dims; std::vector out_sparse_part_dims; diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index cc9f03c519615..339ed74eebe06 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -78,5 +78,20 @@ def test_reshape_nd(self): self.check_result([8, 3, 4, 4, 5, 3], [24, 8, 10, 3], 'coo') self.check_result([3, 4, 4, 5, 7], [1, 12, 2, 5, 14], 'coo') + def test_reshape_with_zero_or_minus_one_in_new_shape(self): + self.check_result([6, 2, 3], [-1, 0, 3], 'coo') + self.check_result([6, 2, 3], [2, 3, 0, -1], 'coo') + self.check_result([6, 2, 3], [1, -1, 2], 'coo') + self.check_result([6, 2, 3], [-1, 9, 2], 'coo') + self.check_result([6, 2, 3], [2, -1, 18], 'coo') + self.check_result([6, 2, 3], [1, 0, 2, -1, 3], 'coo') + + self.check_result([6, 2, 3], [0, 0, -1], 'csr') + self.check_result([6, 2, 3], [-1, 3, 2], 'csr') + self.check_result([6, 2, 3], [2, -1, 0], 'csr') + self.check_result([6, 2, 3], [-1, 6, 2], 'csr') + self.check_result([6, 2, 3], [-1, 9, 1], 'csr') + self.check_result([6, 2, 3], [-1, 1, 3], 'csr') + if __name__ == "__main__": unittest.main() From 1da6d604c7f35be4be439559452e272a2a9decd3 Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 15:49:24 +0800 Subject: [PATCH 39/55] Update unary_grad_kernel.h --- paddle/phi/kernels/sparse/unary_grad_kernel.h | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/paddle/phi/kernels/sparse/unary_grad_kernel.h b/paddle/phi/kernels/sparse/unary_grad_kernel.h index 933e1967e68c3..b446e1b99ed41 100644 --- a/paddle/phi/kernels/sparse/unary_grad_kernel.h +++ b/paddle/phi/kernels/sparse/unary_grad_kernel.h @@ -89,5 +89,17 @@ void TransposeCsrGradKernel(const Context& dev_ctx, const std::vector& perm, SparseCsrTensor* dx); +template +void ReshapeCooGradKernel(const Context& dev_ctx, + const SparseCooTensor& x, + const SparseCooTensor& dout, + SparseCooTensor* dx); + +template +void ReshapeCsrGradKernel(const Context& dev_ctx, + const SparseCsrTensor& x, + const SparseCsrTensor& dout, + SparseCsrTensor* dx); + } // namespace sparse } // namespace phi From bbf4b48efdc77d3c4c682321e14b5c5f9b36621a Mon Sep 17 00:00:00 2001 From: SmirnovKol <31559413+OccupyMars2025@users.noreply.github.com> Date: Mon, 10 Oct 2022 16:59:51 +0800 Subject: [PATCH 40/55] correct lvalue error --- paddle/phi/kernels/sparse/cpu/reshape_kernel.cc | 4 +++- paddle/phi/kernels/sparse/gpu/reshape_kernel.cu | 4 +++- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc index d8096e819f71c..19248869841fa 100644 --- a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc +++ b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc @@ -38,7 +38,9 @@ void ReshapeCooKernel(const Context& dev_ctx, // DDim out_dims = phi::make_ddim(shape.GetData()); // Use DDim::reshape to handle -1 and 0 in the argument "shape" - DDim out_dims = x.dims().reshape(std::vector(shape.GetData().begin(), shape.GetData().end())); + // DDim out_dims = x.dims().reshape(std::vector(shape.GetData().begin(), shape.GetData().end())); + std::vector new_shape(shape.GetData().begin(), shape.GetData().end()); + phi::DDim out_dims = x.dims().reshape(new_shape); // get sparse part dimensions of x and out std::vector x_sparse_part_dims; std::vector out_sparse_part_dims; diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index e7f3f5d4d42f7..eec00919847f3 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -59,7 +59,9 @@ void ReshapeCooKernel(const Context& dev_ctx, int64_t x_nnz = x.nnz(); // TODO: consider using "DDim DDim::reshape(std::vector& shape)" // DDim out_dims = phi::make_ddim(shape.GetData()); - DDim out_dims = x.dims().reshape(std::vector(shape.GetData().begin(), shape.GetData().end())); + // DDim out_dims = x.dims().reshape(std::vector(shape.GetData().begin(), shape.GetData().end())); + std::vector new_shape(shape.GetData().begin(), shape.GetData().end()); + phi::DDim out_dims = x.dims().reshape(new_shape); // get sparse part dimensions of x and out std::vector x_sparse_part_dims; std::vector out_sparse_part_dims; From 6db2391d6e23ebfb833bc5be5acb6723c2cab567 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <1945340016@qq.com> Date: Wed, 12 Oct 2022 12:13:20 +0800 Subject: [PATCH 41/55] code style --- .../phi/kernels/sparse/cpu/reshape_kernel.cc | 49 ++++--- .../phi/kernels/sparse/gpu/reshape_kernel.cu | 121 ++++++++---------- .../tests/unittests/test_sparse_reshape_op.py | 43 ++++--- 3 files changed, 102 insertions(+), 111 deletions(-) diff --git a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc index 19248869841fa..4f16515666810 100644 --- a/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc +++ b/paddle/phi/kernels/sparse/cpu/reshape_kernel.cc @@ -14,8 +14,8 @@ #include "paddle/phi/kernels/sparse/unary_kernel.h" -#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" #include "paddle/phi/core/ddim.h" +#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" #include "paddle/phi/backends/cpu/cpu_context.h" #include "paddle/phi/core/kernel_registry.h" @@ -33,12 +33,10 @@ void ReshapeCooKernel(const Context& dev_ctx, const SparseCooTensor& x, const phi::IntArray& shape, SparseCooTensor* out) { - // TODO: Currently, reshape is only applicable to sparse dims + // TODO(OccupyMars2025): Currently, reshape is only applicable to sparse dims int64_t x_nnz = x.nnz(); - - // DDim out_dims = phi::make_ddim(shape.GetData()); - // Use DDim::reshape to handle -1 and 0 in the argument "shape" - // DDim out_dims = x.dims().reshape(std::vector(shape.GetData().begin(), shape.GetData().end())); + + // Use DDim::reshape to handle -1 and 0 in the argument "shape" std::vector new_shape(shape.GetData().begin(), shape.GetData().end()); phi::DDim out_dims = x.dims().reshape(new_shape); // get sparse part dimensions of x and out @@ -50,9 +48,8 @@ void ReshapeCooKernel(const Context& dev_ctx, for (int i = 0; i < out_dims.size() - x.dense_dim(); ++i) { out_sparse_part_dims.push_back(out_dims[i]); } - DenseTensor out_indices = Empty(dev_ctx, - {static_cast(out_sparse_part_dims.size()), x_nnz} - ); + DenseTensor out_indices = Empty( + dev_ctx, {static_cast(out_sparse_part_dims.size()), x_nnz}); DenseTensor out_values(x.values()); out->SetMember(out_indices, out_values, out_dims, x.coalesced()); @@ -61,32 +58,33 @@ void ReshapeCooKernel(const Context& dev_ctx, const auto* x_indices_data = x_indices.data(); auto* out_indices_data = out_indices.data(); - const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); - const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); + const phi::DDim& x_sparse_part_strides = + phi::stride(phi::make_ddim(x_sparse_part_dims)); + const phi::DDim& out_sparse_part_strides = + phi::stride(phi::make_ddim(out_sparse_part_dims)); int64_t location = 0; for (int64_t j = 0; j < x_nnz; ++j) { - location = 0; - for (int i = 0; i < x.sparse_dim(); ++i) { - location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; - } - for (size_t i = 0; i < out_sparse_part_dims.size(); ++i) { - out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; - location %= out_sparse_part_strides[i]; - } + location = 0; + for (int i = 0; i < x.sparse_dim(); ++i) { + location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; + } + for (size_t i = 0; i < out_sparse_part_dims.size(); ++i) { + out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; + location %= out_sparse_part_strides[i]; + } } } - template void ReshapeCsrKernel(const Context& dev_ctx, const SparseCsrTensor& x, const phi::IntArray& shape, SparseCsrTensor* out) { - /*transform csr format to coo format, and then use coo kernel*/ - const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); - SparseCooTensor out_coo; - ReshapeCooKernel(dev_ctx, x_coo, shape, &out_coo); - CooToCsrKernel(dev_ctx, out_coo, out); + // transform csr format to coo format, and then use coo kernel + const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); + SparseCooTensor out_coo; + ReshapeCooKernel(dev_ctx, x_coo, shape, &out_coo); + CooToCsrKernel(dev_ctx, out_coo, out); } } // namespace sparse @@ -117,4 +115,3 @@ PD_REGISTER_KERNEL(reshape_csr, int, int64_t, bool) {} - \ No newline at end of file diff --git a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu index eec00919847f3..6e3a9842e8c30 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_kernel.cu @@ -22,7 +22,6 @@ #include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" - namespace phi { namespace sparse { @@ -32,34 +31,25 @@ __global__ void ReshapeCooCudaKernel(const int64_t* x_indices_data, const int64_t x_nnz, const int64_t* x_sparse_part_strides, const int64_t* out_sparse_part_strides, - int64_t *out_indices_data) { - - CUDA_KERNEL_LOOP_TYPE(j, x_nnz, int64_t) { - int64_t location = 0; - for (int i = 0; i < num_x_sparse_part_dims; ++i) { - location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; - // row major or column major ??? - // location += x_indices_data[j * num_x_sparse_part_dims + i] * x_sparse_part_strides[i]; - } - for (int i = 0; i < num_out_sparse_part_dims; ++i) { - out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; - // row major or column major ??? - // out_indices_data[j * num_out_sparse_part_dims + i] = location / out_sparse_part_strides[i]; - location %= out_sparse_part_strides[i]; - } + int64_t* out_indices_data) { + CUDA_KERNEL_LOOP_TYPE(j, x_nnz, int64_t) { + int64_t location = 0; + for (int i = 0; i < num_x_sparse_part_dims; ++i) { + location += x_indices_data[i * x_nnz + j] * x_sparse_part_strides[i]; } + for (int i = 0; i < num_out_sparse_part_dims; ++i) { + out_indices_data[i * x_nnz + j] = location / out_sparse_part_strides[i]; + location %= out_sparse_part_strides[i]; + } + } } - template void ReshapeCooKernel(const Context& dev_ctx, - const SparseCooTensor &x, + const SparseCooTensor& x, const phi::IntArray& shape, - SparseCooTensor *out) { + SparseCooTensor* out) { int64_t x_nnz = x.nnz(); - // TODO: consider using "DDim DDim::reshape(std::vector& shape)" - // DDim out_dims = phi::make_ddim(shape.GetData()); - // DDim out_dims = x.dims().reshape(std::vector(shape.GetData().begin(), shape.GetData().end())); std::vector new_shape(shape.GetData().begin(), shape.GetData().end()); phi::DDim out_dims = x.dims().reshape(new_shape); // get sparse part dimensions of x and out @@ -72,80 +62,75 @@ void ReshapeCooKernel(const Context& dev_ctx, out_sparse_part_dims.push_back(out_dims[i]); } - DenseTensor out_indices = Empty(dev_ctx, - {static_cast(out_sparse_part_dims.size()), x_nnz} - ); + DenseTensor out_indices = Empty( + dev_ctx, {static_cast(out_sparse_part_dims.size()), x_nnz}); DenseTensor out_values(x.values()); out->SetMember(out_indices, out_values, out_dims, x.coalesced()); // compute values of out indices - // const DenseTensor& x_indices = x.indices(); - // const auto *x_indices_data = x_indices.data(); - const auto *x_indices_data = x.indices().data(); - auto *out_indices_data = out_indices.data(); - const phi::DDim& x_sparse_part_strides = phi::stride(phi::make_ddim(x_sparse_part_dims)); - const phi::DDim& out_sparse_part_strides = phi::stride(phi::make_ddim(out_sparse_part_dims)); + const auto* x_indices_data = x.indices().data(); + auto* out_indices_data = out_indices.data(); + const phi::DDim& x_sparse_part_strides = + phi::stride(phi::make_ddim(x_sparse_part_dims)); + const phi::DDim& out_sparse_part_strides = + phi::stride(phi::make_ddim(out_sparse_part_dims)); - int64_t* destination_x_sparse_part_strides, *destination_out_sparse_part_strides; + int64_t *destination_x_sparse_part_strides, + *destination_out_sparse_part_strides; #ifdef PADDLE_WITH_HIP - hipMalloc(reinterpret_cast(&destination_x_sparse_part_strides), + hipMalloc(reinterpret_cast(&destination_x_sparse_part_strides), sizeof(int64_t) * x_sparse_part_strides.size()); - hipMemcpy( - destination_x_sparse_part_strides, - x_sparse_part_strides.Get(), - sizeof(int64_t) * x_sparse_part_strides.size(), - hipMemcpyHostToDevice); - hipMalloc(reinterpret_cast(&destination_out_sparse_part_strides), + hipMemcpy(destination_x_sparse_part_strides, + x_sparse_part_strides.Get(), + sizeof(int64_t) * x_sparse_part_strides.size(), + hipMemcpyHostToDevice); + hipMalloc(reinterpret_cast(&destination_out_sparse_part_strides), sizeof(int64_t) * out_sparse_part_strides.size()); - hipMemcpy( - destination_out_sparse_part_strides, - out_sparse_part_strides.Get(), - sizeof(int64_t) * out_sparse_part_strides.size(), - hipMemcpyHostToDevice); + hipMemcpy(destination_out_sparse_part_strides, + out_sparse_part_strides.Get(), + sizeof(int64_t) * out_sparse_part_strides.size(), + hipMemcpyHostToDevice); #else - cudaMalloc(reinterpret_cast(&destination_x_sparse_part_strides), - sizeof(int64_t) * x_sparse_part_strides.size()); - cudaMemcpy( - destination_x_sparse_part_strides, - x_sparse_part_strides.Get(), - sizeof(int64_t) * x_sparse_part_strides.size(), - cudaMemcpyHostToDevice); - cudaMalloc(reinterpret_cast(&destination_out_sparse_part_strides), - sizeof(int64_t) * out_sparse_part_strides.size()); - cudaMemcpy( - destination_out_sparse_part_strides, - out_sparse_part_strides.Get(), - sizeof(int64_t) * out_sparse_part_strides.size(), - cudaMemcpyHostToDevice); + cudaMalloc(reinterpret_cast(&destination_x_sparse_part_strides), + sizeof(int64_t) * x_sparse_part_strides.size()); + cudaMemcpy(destination_x_sparse_part_strides, + x_sparse_part_strides.Get(), + sizeof(int64_t) * x_sparse_part_strides.size(), + cudaMemcpyHostToDevice); + cudaMalloc(reinterpret_cast(&destination_out_sparse_part_strides), + sizeof(int64_t) * out_sparse_part_strides.size()); + cudaMemcpy(destination_out_sparse_part_strides, + out_sparse_part_strides.Get(), + sizeof(int64_t) * out_sparse_part_strides.size(), + cudaMemcpyHostToDevice); #endif auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, x_nnz, 1); ReshapeCooCudaKernel<<>>( - x_indices_data, - x_sparse_part_dims.size(), out_sparse_part_dims.size(), x_nnz, - // x_sparse_part_strides.Get(), out_sparse_part_strides.Get(), - destination_x_sparse_part_strides, destination_out_sparse_part_strides, - out_indices_data - ); + x_indices_data, + x_sparse_part_dims.size(), + out_sparse_part_dims.size(), + x_nnz, + destination_x_sparse_part_strides, + destination_out_sparse_part_strides, + out_indices_data); } - // just copy from paddle\phi\kernels\sparse\cpu\reshape_kernel.cc template void ReshapeCsrKernel(const Context& dev_ctx, const SparseCsrTensor& x, const phi::IntArray& shape, SparseCsrTensor* out) { - /*transform csr format to coo format, and then use coo kernel*/ + // transform csr format to coo format, and then use coo kernel const SparseCooTensor x_coo = CsrToCoo(dev_ctx, x); SparseCooTensor out_coo; ReshapeCooKernel(dev_ctx, x_coo, shape, &out_coo); - CooToCsrKernel(dev_ctx, out_coo, out); + CooToCsrKernel(dev_ctx, out_coo, out); } } // namespace sparse diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 339ed74eebe06..385c2338a7e63 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -15,45 +15,53 @@ import paddle import numpy as np import unittest -# from paddle.fluid.framework import _test_eager_guard class TestReshape(unittest.TestCase): - # x: sparse, out: sparse + """ + Test the API paddle.incubate.sparse.reshape on some sparse tensors. + x: sparse, out: sparse + """ + def check_result(self, x_shape, new_shape, format): - mask = np.random.randint( 0, 2, x_shape) - np_x = np.random.randint(-100, 100, x_shape) * mask - - # cpu version + """ + x_shape: original shape + new_shape: new shape + format: "coo" or "csr" + Transform a sparse tensor with shape "x_shape" to + a sparse tensor with shape "new_shape". + Compare the output of paddle.reshape and the output of + paddle.incubate.sparse.reshape. + """ + mask = np.random.randint(0, 2, x_shape) + np_x = np.random.randint(-100, 100, x_shape) * mask + dense_x = paddle.to_tensor(np_x) dense_x.stop_gradient = False - # dense_out = paddle.transpose(dense_x, dims) dense_out = paddle.reshape(dense_x, new_shape) if format == "coo": - # sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) else: - # sp_x = origin_x.detach().to_sparse_csr() sp_x = paddle.to_tensor(np_x).to_sparse_csr() sp_x.stop_gradient = False - # sp_out = paddle.incubate.sparse.transpose(sp_x, dims) sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) np.testing.assert_allclose(sp_out.to_dense().numpy(), - dense_out.numpy(), - rtol=1e-05) + dense_out.numpy(), + rtol=1e-05) - # check cpu backward computation dense_out.backward() sp_out.backward() np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - # dense_x.grad.numpy() * mask, - dense_x.grad.numpy() * np_x.astype('bool').astype('int'), - rtol=1e-05) + dense_x.grad.numpy() * + np_x.astype('bool').astype('int'), + rtol=1e-05) def test_reshape_2d(self): - self.check_result([2, 5], [10,], 'coo') + self.check_result([2, 5], [ + 10, + ], 'coo') self.check_result([12, 5], [15, 4], 'coo') self.check_result([10, 5], [2, 25], 'csr') @@ -93,5 +101,6 @@ def test_reshape_with_zero_or_minus_one_in_new_shape(self): self.check_result([6, 2, 3], [-1, 9, 1], 'csr') self.check_result([6, 2, 3], [-1, 1, 3], 'csr') + if __name__ == "__main__": unittest.main() From b01ebb98823ab81f9a69403fb0e55e72d4684f0a Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 12 Oct 2022 19:35:00 +0800 Subject: [PATCH 42/55] Update sparse_backward.yaml --- paddle/phi/api/yaml/sparse_backward.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/paddle/phi/api/yaml/sparse_backward.yaml b/paddle/phi/api/yaml/sparse_backward.yaml index 2bd366531618a..60a2137c72343 100644 --- a/paddle/phi/api/yaml/sparse_backward.yaml +++ b/paddle/phi/api/yaml/sparse_backward.yaml @@ -450,4 +450,4 @@ param : [x] kernel : func : reshape_coo_grad {sparse_coo, sparse_coo -> sparse_coo}, - reshape_csr_grad {sparse_csr, sparse_csr -> sparse_csr} \ No newline at end of file + reshape_csr_grad {sparse_csr, sparse_csr -> sparse_csr} From dd30c2e7bf9855de88eda346a72f8ecb985b9681 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 12 Oct 2022 19:36:42 +0800 Subject: [PATCH 43/55] Update sparse_ops.yaml --- paddle/phi/api/yaml/sparse_ops.yaml | 1 - 1 file changed, 1 deletion(-) diff --git a/paddle/phi/api/yaml/sparse_ops.yaml b/paddle/phi/api/yaml/sparse_ops.yaml index 6cdd0f6dbeb51..a5db3f959a80c 100644 --- a/paddle/phi/api/yaml/sparse_ops.yaml +++ b/paddle/phi/api/yaml/sparse_ops.yaml @@ -500,4 +500,3 @@ reshape_csr{sparse_csr -> sparse_csr} layout : x backward : reshape_grad - \ No newline at end of file From 64d7de22ebcdb4320a1f72452bc92316f771d821 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 12 Oct 2022 19:43:10 +0800 Subject: [PATCH 44/55] Update unary_kernel.h --- paddle/phi/kernels/sparse/unary_kernel.h | 25 ++++++++++++------------ 1 file changed, 12 insertions(+), 13 deletions(-) diff --git a/paddle/phi/kernels/sparse/unary_kernel.h b/paddle/phi/kernels/sparse/unary_kernel.h index fbe1a916d9823..a81e724d1fe48 100644 --- a/paddle/phi/kernels/sparse/unary_kernel.h +++ b/paddle/phi/kernels/sparse/unary_kernel.h @@ -14,11 +14,10 @@ #pragma once -#include "paddle/phi/core/sparse_coo_tensor.h" -#include "paddle/phi/core/sparse_csr_tensor.h" #include "paddle/phi/common/int_array.h" #include "paddle/phi/core/ddim.h" - +#include "paddle/phi/core/sparse_coo_tensor.h" +#include "paddle/phi/core/sparse_csr_tensor.h" namespace phi { namespace sparse { @@ -159,20 +158,20 @@ SparseCooTensor ReluCsr(const Context& dev_ctx, const SparseCooTensor& x) { } template -void ReshapeCooKernel(const Context &dev_ctx, - const SparseCooTensor &x, +void ReshapeCooKernel(const Context& dev_ctx, + const SparseCooTensor& x, const phi::IntArray& shape, - SparseCooTensor *out); + SparseCooTensor* out); template -void ReshapeCsrKernel(const Context &dev_ctx, - const SparseCsrTensor &x, +void ReshapeCsrKernel(const Context& dev_ctx, + const SparseCsrTensor& x, const phi::IntArray& shape, - SparseCsrTensor *out); + SparseCsrTensor* out); template -SparseCooTensor ReshapeCoo(const Context &dev_ctx, - const SparseCooTensor &x, +SparseCooTensor ReshapeCoo(const Context& dev_ctx, + const SparseCooTensor& x, const phi::IntArray& shape) { SparseCooTensor coo; ReshapeCooKernel(dev_ctx, x, shape, &coo); @@ -180,8 +179,8 @@ SparseCooTensor ReshapeCoo(const Context &dev_ctx, } template -SparseCsrTensor ReshapeCsr(const Context &dev_ctx, - const SparseCsrTensor &x, +SparseCsrTensor ReshapeCsr(const Context& dev_ctx, + const SparseCsrTensor& x, const phi::IntArray& shape) { PADDLE_ENFORCE_LE( 2, From 118c429bdd24481ab14444fe014d8a830172ace4 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 12 Oct 2022 20:27:42 +0800 Subject: [PATCH 45/55] Update unary.py --- python/paddle/incubate/sparse/unary.py | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index fa9b41f4b8a36..b1e2a3caa0b52 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -644,17 +644,17 @@ def expm1(x, name=None): def reshape(x, shape, name=None): """ Changes the shape of ``x`` without changing its value, requiring x to be a SparseCooTensor or SparseCsrTensor. - Currently this function can only reshape the sparse dims of ``x`` , but ``shape`` argument must - include all dims of the reshaped tensor. Note: if x is a SparseCsrTensor, then len(shape) must be 2 or 3. + Currently this function can only reshape the sparse dims of ``x`` , but ``shape`` argument must be specified + as the shape of the reshaped tensor. Note: if x is a SparseCsrTensor, then len(shape) must be 2 or 3. .. math:: out = reshape(x, shape) Parameters: - x (Tensor): The input Sparse Tensor with data type float32, float64, int32, int64 and so on. + x (Tensor): The input sparse tensor with data type float32, float64, int32, int64 and so on. shape (list[int]): new shape. name (str, optional): Name for the operation (optional, default is None). For more information, please refer to :ref:`api_guide_Name`. Returns: - A reshaped Sparse Tensor with the same data type as ``x``. + A reshaped sparse tensor with the same data type as ``x``. Examples: .. code-block:: python import paddle @@ -663,4 +663,4 @@ def reshape(x, shape, name=None): out = paddle.incubate.sparse.reshape(sparse_x, [6, ]) """ return _C_ops.sparse_reshape(x, shape) - \ No newline at end of file + From 4a510d271f18601b1372f26f7418ab666dc4ec37 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 12 Oct 2022 20:32:11 +0800 Subject: [PATCH 46/55] Update sparse_backward.yaml --- paddle/phi/api/yaml/sparse_backward.yaml | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/paddle/phi/api/yaml/sparse_backward.yaml b/paddle/phi/api/yaml/sparse_backward.yaml index 5410660784e50..1a294c2ffb3ef 100644 --- a/paddle/phi/api/yaml/sparse_backward.yaml +++ b/paddle/phi/api/yaml/sparse_backward.yaml @@ -272,6 +272,17 @@ func : relu_coo_grad {sparse_coo, sparse_coo -> sparse_coo}, relu_csr_grad {sparse_csr, sparse_csr -> sparse_csr} +- backward_op : reshape_grad + forward : reshape(Tensor x, IntArray shape) -> Tensor(out) + args : (Tensor x, Tensor out_grad) + output : Tensor(x_grad) + infer_meta : + func : UnchangedInferMeta + param : [x] + kernel : + func : reshape_coo_grad {sparse_coo, sparse_coo -> sparse_coo}, + reshape_csr_grad {sparse_csr, sparse_csr -> sparse_csr} + - backward_op : scale_grad forward : scale(Tensor x, float scale, float bias, bool bias_after_scale) -> Tensor(out) args : (Tensor out_grad, float scale) @@ -440,14 +451,3 @@ func : fused_attention_csr_grad{dense, dense, dense, sparse_csr, dense -> dense, dense, dense} layout : softmax data_type: query - -- backward_op : reshape_grad - forward : reshape(Tensor x, IntArray shape) -> Tensor(out) - args : (Tensor x, Tensor out_grad) - output : Tensor(x_grad) - infer_meta : - func : UnchangedInferMeta - param : [x] - kernel : - func : reshape_coo_grad {sparse_coo, sparse_coo -> sparse_coo}, - reshape_csr_grad {sparse_csr, sparse_csr -> sparse_csr} From f011729c2d457bf694d84ed8ed617f804b863084 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <31559413+OccupyMars2025@users.noreply.github.com> Date: Wed, 12 Oct 2022 20:35:03 +0800 Subject: [PATCH 47/55] Update unary.py From 659c034c29588ebb280bda2abbf034572b9392e5 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <1945340016@qq.com> Date: Thu, 13 Oct 2022 07:28:19 +0800 Subject: [PATCH 48/55] code style --- paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu | 1 - python/paddle/incubate/sparse/unary.py | 5 ++++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu b/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu index b67e79ef8efea..bfc81676eb804 100644 --- a/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu +++ b/paddle/phi/kernels/sparse/gpu/reshape_grad_kernel.cu @@ -34,7 +34,6 @@ void ReshapeCooGradKernel(const Context& dev_ctx, ReshapeCooKernel(dev_ctx, dout, x_shape, dx); } - // just copy from paddle\phi\kernels\sparse\cpu\reshape_grad_kernel.cc template void ReshapeCsrGradKernel(const Context& dev_ctx, diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index b1e2a3caa0b52..44b08e387b88a 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -640,6 +640,7 @@ def expm1(x, name=None): """ return _C_ops.sparse_expm1(x) + @dygraph_only def reshape(x, shape, name=None): """ @@ -648,13 +649,16 @@ def reshape(x, shape, name=None): as the shape of the reshaped tensor. Note: if x is a SparseCsrTensor, then len(shape) must be 2 or 3. .. math:: out = reshape(x, shape) + Parameters: x (Tensor): The input sparse tensor with data type float32, float64, int32, int64 and so on. shape (list[int]): new shape. name (str, optional): Name for the operation (optional, default is None). For more information, please refer to :ref:`api_guide_Name`. + Returns: A reshaped sparse tensor with the same data type as ``x``. + Examples: .. code-block:: python import paddle @@ -663,4 +667,3 @@ def reshape(x, shape, name=None): out = paddle.incubate.sparse.reshape(sparse_x, [6, ]) """ return _C_ops.sparse_reshape(x, shape) - From 213145b8b387d72bbb35c9acefdc49b783f5506b Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <1945340016@qq.com> Date: Thu, 13 Oct 2022 14:11:36 +0800 Subject: [PATCH 49/55] code style --- paddle/phi/api/yaml/sparse_backward.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/paddle/phi/api/yaml/sparse_backward.yaml b/paddle/phi/api/yaml/sparse_backward.yaml index 1a294c2ffb3ef..3e30379faef68 100644 --- a/paddle/phi/api/yaml/sparse_backward.yaml +++ b/paddle/phi/api/yaml/sparse_backward.yaml @@ -282,7 +282,7 @@ kernel : func : reshape_coo_grad {sparse_coo, sparse_coo -> sparse_coo}, reshape_csr_grad {sparse_csr, sparse_csr -> sparse_csr} - + - backward_op : scale_grad forward : scale(Tensor x, float scale, float bias, bool bias_after_scale) -> Tensor(out) args : (Tensor out_grad, float scale) From 0d77ea7acf632bcf1bdbbd7ff9f4391a5e04df67 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <1945340016@qq.com> Date: Thu, 13 Oct 2022 14:30:29 +0800 Subject: [PATCH 50/55] code style --- python/paddle/incubate/sparse/__init__.py | 37 +++-------------------- 1 file changed, 5 insertions(+), 32 deletions(-) diff --git a/python/paddle/incubate/sparse/__init__.py b/python/paddle/incubate/sparse/__init__.py index bbc110939c6a8..8b6866fa4da0b 100644 --- a/python/paddle/incubate/sparse/__init__.py +++ b/python/paddle/incubate/sparse/__init__.py @@ -51,36 +51,9 @@ from . import nn __all__ = [ - 'sparse_coo_tensor', - 'sparse_csr_tensor', - 'sin', - 'tan', - 'asin', - 'atan', - 'sinh', - 'tanh', - 'asinh', - 'atanh', - 'sqrt', - 'square', - 'log1p', - 'abs', - 'pow', - 'cast', - 'neg', - 'deg2rad', - 'rad2deg', - 'expm1', - 'mv', - 'matmul', - 'masked_matmul', - 'addmm', - 'add', - 'subtract', - 'transpose', - 'multiply', - 'divide', - 'coalesce', - 'is_same_shape', - 'reshape' + 'sparse_coo_tensor', 'sparse_csr_tensor', 'sin', 'tan', 'asin', 'atan', + 'sinh', 'tanh', 'asinh', 'atanh', 'sqrt', 'square', 'log1p', 'abs', 'pow', + 'cast', 'neg', 'deg2rad', 'rad2deg', 'expm1', 'mv', 'matmul', + 'masked_matmul', 'addmm', 'add', 'subtract', 'transpose', 'multiply', + 'divide', 'coalesce', 'is_same_shape', 'reshape' ] From 67cfa5023021e4f4c0d7276d2cda092b0a0de911 Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <31559413+OccupyMars2025@users.noreply.github.com> Date: Fri, 14 Oct 2022 11:45:38 +0800 Subject: [PATCH 51/55] Update unary.py --- python/paddle/incubate/sparse/unary.py | 1 + 1 file changed, 1 insertion(+) diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index 44b08e387b88a..d30b692c91296 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -648,6 +648,7 @@ def reshape(x, shape, name=None): Currently this function can only reshape the sparse dims of ``x`` , but ``shape`` argument must be specified as the shape of the reshaped tensor. Note: if x is a SparseCsrTensor, then len(shape) must be 2 or 3. .. math:: + out = reshape(x, shape) Parameters: From 32cd23c84f712330ab03161e4b73696d74cd033f Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <1945340016@qq.com> Date: Fri, 14 Oct 2022 14:01:09 +0800 Subject: [PATCH 52/55] specify tensor place explicitly --- .../tests/unittests/test_sparse_reshape_op.py | 36 +++++++++++-- python/paddle/incubate/sparse/unary.py | 50 +++++++++++++++---- 2 files changed, 73 insertions(+), 13 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index 385c2338a7e63..e9ef737f7743d 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -36,14 +36,18 @@ def check_result(self, x_shape, new_shape, format): mask = np.random.randint(0, 2, x_shape) np_x = np.random.randint(-100, 100, x_shape) * mask - dense_x = paddle.to_tensor(np_x) + # check cpu kernel + dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) dense_x.stop_gradient = False dense_out = paddle.reshape(dense_x, new_shape) if format == "coo": - sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) + sp_x = paddle.to_tensor(np_x, + place=paddle.CPUPlace()).to_sparse_coo( + len(x_shape)) else: - sp_x = paddle.to_tensor(np_x).to_sparse_csr() + sp_x = paddle.to_tensor(np_x, + place=paddle.CPUPlace()).to_sparse_csr() sp_x.stop_gradient = False sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) @@ -58,6 +62,32 @@ def check_result(self, x_shape, new_shape, format): np_x.astype('bool').astype('int'), rtol=1e-05) + # check gpu kernel + if paddle.device.is_compiled_with_cuda(): + dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) + dense_x.stop_gradient = False + dense_out = paddle.reshape(dense_x, new_shape) + + if format == "coo": + sp_x = paddle.to_tensor( + np_x, place=paddle.CUDAPlace(0)).to_sparse_coo(len(x_shape)) + else: + sp_x = paddle.to_tensor( + np_x, place=paddle.CUDAPlace(0)).to_sparse_csr() + sp_x.stop_gradient = False + sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) + + np.testing.assert_allclose(sp_out.to_dense().numpy(), + dense_out.numpy(), + rtol=1e-05) + + dense_out.backward() + sp_out.backward() + np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), + dense_x.grad.numpy() * + np_x.astype('bool').astype('int'), + rtol=1e-05) + def test_reshape_2d(self): self.check_result([2, 5], [ 10, diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index d30b692c91296..a3880dd973608 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -646,25 +646,55 @@ def reshape(x, shape, name=None): """ Changes the shape of ``x`` without changing its value, requiring x to be a SparseCooTensor or SparseCsrTensor. Currently this function can only reshape the sparse dims of ``x`` , but ``shape`` argument must be specified - as the shape of the reshaped tensor. Note: if x is a SparseCsrTensor, then len(shape) must be 2 or 3. - .. math:: + as the shape of the reshaped tensor. - out = reshape(x, shape) + Note that if x is a SparseCsrTensor, then len(shape) must be 2 or 3. - Parameters: - x (Tensor): The input sparse tensor with data type float32, float64, int32, int64 and so on. - shape (list[int]): new shape. + There are some tricks when specifying the target shape. + + - 1. -1 means the value of this dimension is inferred from the total element number of x and remaining dimensions. Thus one and only one dimension can be set -1. + + - 2. 0 means the actual dimension value is going to be copied from the corresponding dimension of x. The indices of 0 in the target shape can not exceed the rank of x. + + Here are some examples to explain it. + + - 1. Given a 3-D tensor x with a shape [2, 4, 6], and the target shape is [6, 8], the reshape operator will transform x into a 2-D tensor with shape [6, 8] and leaving x's data unchanged. + + - 2. Given a 3-D tensor x with a shape [2, 4, 6], and the target shape is [2, 3, -1, 2], the reshape operator will transform x into a 4-D tensor with shape [2, 3, 4, 2] and leaving x's data unchanged. In this case, one dimension of the target shape is set to -1, the value of this dimension is inferred from the total element number of x and remaining dimensions. + + - 3. Given a 3-D tensor x with a shape [2, 4, 6], and the target shape is [-1, 0, 3, 2], the reshape operator will transform x into a 4-D tensor with shape [2, 4, 3, 2] and leaving x's data unchanged. In this case, besides -1, 0 means the actual dimension value is going to be copied from the corresponding dimension of x. + + Args: + x (Tensor): The input sparse tensor with data type ``float32``, ``float64``, ``int32``, ``int64`` or ``bool``. + shape (list|tuple): Define the target shape. At most one dimension of the target shape can be -1. + The data type is ``int32``. name (str, optional): Name for the operation (optional, default is None). For more information, please refer to :ref:`api_guide_Name`. Returns: - A reshaped sparse tensor with the same data type as ``x``. + Tensor: A reshaped Tensor with the same data type as ``x``. Examples: .. code-block:: python + import paddle - dense_x = paddle.to_tensor([[-2., 0., 7.9], [1., 2., 0.]]) - sparse_x = dense_x.to_sparse_coo(2) - out = paddle.incubate.sparse.reshape(sparse_x, [6, ]) + import numpy as np + + x_shape = [6, 2, 3] + new_shape = [1, 0, 2, -1, 3] + format = "coo" + + mask = np.random.randint(0, 2, x_shape) + np_x = np.random.randint(-100, 100, x_shape) * mask + + if format == "coo": + sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) + else: + sp_x = paddle.to_tensor(np_x).to_sparse_csr() + sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) + + print(sp_out) + # the shape of sp_out is [1, 2, 2, 3, 3] + """ return _C_ops.sparse_reshape(x, shape) From a96baa6813bd66b35ede962a78544d299afece2c Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <1945340016@qq.com> Date: Fri, 14 Oct 2022 15:36:42 +0800 Subject: [PATCH 53/55] do not use numpy array --- .../tests/unittests/test_sparse_reshape_op.py | 43 +++---------------- 1 file changed, 6 insertions(+), 37 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index e9ef737f7743d..d153b1f37ab4d 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -33,21 +33,17 @@ def check_result(self, x_shape, new_shape, format): Compare the output of paddle.reshape and the output of paddle.incubate.sparse.reshape. """ - mask = np.random.randint(0, 2, x_shape) - np_x = np.random.randint(-100, 100, x_shape) * mask + mask = paddle.randint(0, 2, x_shape).astype("float32") + origin_x = (paddle.rand(x_shape, dtype='float32') + 1) * mask - # check cpu kernel - dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) + dense_x = origin_x.detach() dense_x.stop_gradient = False dense_out = paddle.reshape(dense_x, new_shape) if format == "coo": - sp_x = paddle.to_tensor(np_x, - place=paddle.CPUPlace()).to_sparse_coo( - len(x_shape)) + sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) else: - sp_x = paddle.to_tensor(np_x, - place=paddle.CPUPlace()).to_sparse_csr() + sp_x = origin_x.detach().to_sparse_csr() sp_x.stop_gradient = False sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) @@ -58,36 +54,9 @@ def check_result(self, x_shape, new_shape, format): dense_out.backward() sp_out.backward() np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - dense_x.grad.numpy() * - np_x.astype('bool').astype('int'), + (dense_x.grad * mask).numpy(), rtol=1e-05) - # check gpu kernel - if paddle.device.is_compiled_with_cuda(): - dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) - dense_x.stop_gradient = False - dense_out = paddle.reshape(dense_x, new_shape) - - if format == "coo": - sp_x = paddle.to_tensor( - np_x, place=paddle.CUDAPlace(0)).to_sparse_coo(len(x_shape)) - else: - sp_x = paddle.to_tensor( - np_x, place=paddle.CUDAPlace(0)).to_sparse_csr() - sp_x.stop_gradient = False - sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) - - np.testing.assert_allclose(sp_out.to_dense().numpy(), - dense_out.numpy(), - rtol=1e-05) - - dense_out.backward() - sp_out.backward() - np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - dense_x.grad.numpy() * - np_x.astype('bool').astype('int'), - rtol=1e-05) - def test_reshape_2d(self): self.check_result([2, 5], [ 10, From 8c7f144ac7cec22e967d8e99b73087ce08a32ebe Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <1945340016@qq.com> Date: Fri, 14 Oct 2022 18:51:19 +0800 Subject: [PATCH 54/55] use numpy array in unit test again --- .../tests/unittests/test_sparse_reshape_op.py | 43 ++++++++++++++++--- 1 file changed, 37 insertions(+), 6 deletions(-) diff --git a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py index d153b1f37ab4d..e9ef737f7743d 100644 --- a/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py +++ b/python/paddle/fluid/tests/unittests/test_sparse_reshape_op.py @@ -33,17 +33,21 @@ def check_result(self, x_shape, new_shape, format): Compare the output of paddle.reshape and the output of paddle.incubate.sparse.reshape. """ - mask = paddle.randint(0, 2, x_shape).astype("float32") - origin_x = (paddle.rand(x_shape, dtype='float32') + 1) * mask + mask = np.random.randint(0, 2, x_shape) + np_x = np.random.randint(-100, 100, x_shape) * mask - dense_x = origin_x.detach() + # check cpu kernel + dense_x = paddle.to_tensor(np_x, place=paddle.CPUPlace()) dense_x.stop_gradient = False dense_out = paddle.reshape(dense_x, new_shape) if format == "coo": - sp_x = origin_x.detach().to_sparse_coo(len(x_shape)) + sp_x = paddle.to_tensor(np_x, + place=paddle.CPUPlace()).to_sparse_coo( + len(x_shape)) else: - sp_x = origin_x.detach().to_sparse_csr() + sp_x = paddle.to_tensor(np_x, + place=paddle.CPUPlace()).to_sparse_csr() sp_x.stop_gradient = False sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) @@ -54,9 +58,36 @@ def check_result(self, x_shape, new_shape, format): dense_out.backward() sp_out.backward() np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), - (dense_x.grad * mask).numpy(), + dense_x.grad.numpy() * + np_x.astype('bool').astype('int'), rtol=1e-05) + # check gpu kernel + if paddle.device.is_compiled_with_cuda(): + dense_x = paddle.to_tensor(np_x, place=paddle.CUDAPlace(0)) + dense_x.stop_gradient = False + dense_out = paddle.reshape(dense_x, new_shape) + + if format == "coo": + sp_x = paddle.to_tensor( + np_x, place=paddle.CUDAPlace(0)).to_sparse_coo(len(x_shape)) + else: + sp_x = paddle.to_tensor( + np_x, place=paddle.CUDAPlace(0)).to_sparse_csr() + sp_x.stop_gradient = False + sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) + + np.testing.assert_allclose(sp_out.to_dense().numpy(), + dense_out.numpy(), + rtol=1e-05) + + dense_out.backward() + sp_out.backward() + np.testing.assert_allclose(sp_x.grad.to_dense().numpy(), + dense_x.grad.numpy() * + np_x.astype('bool').astype('int'), + rtol=1e-05) + def test_reshape_2d(self): self.check_result([2, 5], [ 10, From 518038c8a4309e46cdceaf94eff8136a2359edec Mon Sep 17 00:00:00 2001 From: OccupyMars2025 <1945340016@qq.com> Date: Mon, 17 Oct 2022 11:34:03 +0800 Subject: [PATCH 55/55] modify example code in docstring --- python/paddle/incubate/sparse/unary.py | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/python/paddle/incubate/sparse/unary.py b/python/paddle/incubate/sparse/unary.py index a3880dd973608..62ab27d0c04c6 100644 --- a/python/paddle/incubate/sparse/unary.py +++ b/python/paddle/incubate/sparse/unary.py @@ -678,19 +678,17 @@ def reshape(x, shape, name=None): .. code-block:: python import paddle - import numpy as np x_shape = [6, 2, 3] new_shape = [1, 0, 2, -1, 3] format = "coo" - mask = np.random.randint(0, 2, x_shape) - np_x = np.random.randint(-100, 100, x_shape) * mask + dense_x = paddle.randint(-100, 100, x_shape) * paddle.randint(0, 2, x_shape) if format == "coo": - sp_x = paddle.to_tensor(np_x).to_sparse_coo(len(x_shape)) + sp_x = dense_x.to_sparse_coo(len(x_shape)) else: - sp_x = paddle.to_tensor(np_x).to_sparse_csr() + sp_x = dense_x.to_sparse_csr() sp_out = paddle.incubate.sparse.reshape(sp_x, new_shape) print(sp_out)