Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

refine the slice Op #24967

Merged
merged 5 commits into from
Jun 10, 2020
Merged

refine the slice Op #24967

Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
4f48d3f
refine the slice Op test=develo
wangchaochaohu Jun 5, 2020
6909b00
refine the code test=develop
wangchaochaohu Jun 8, 2020
a3431c9
refine the compiler error test=develop
wangchaochaohu Jun 8, 2020
8938483
refine the code test=develop
wangchaochaohu Jun 9, 2020
d7e3077
refine the code test=develop
wangchaochaohu Jun 9, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
29 changes: 17 additions & 12 deletions paddle/fluid/operators/slice_op.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,12 +13,12 @@ See the License for the specific language governing permissions and
limitations under the License. */

#include <thrust/device_vector.h>
#include "paddle/fluid/framework/tensor_util.h"
#include "paddle/fluid/operators/math/math_function.h"
#include "paddle/fluid/operators/slice_op.h"
#include "paddle/fluid/platform/cuda_device_function.h"
#include "paddle/fluid/platform/cuda_primitives.h"
#include "paddle/fluid/platform/float16.h"

namespace paddle {
namespace operators {

Expand Down Expand Up @@ -94,17 +94,22 @@ class SliceGradKernel<paddle::platform::CUDADeviceContext,
dim3 blocks((numel - 1) / PADDLE_CUDA_NUM_THREADS + 1);
dim3 threads(PADDLE_CUDA_NUM_THREADS);
auto stream = ctx.cuda_device_context().stream();

auto out_shape = framework::vectorize<int64_t>(out_dims);
thrust::device_vector<int64_t> out_dims_vec(out_shape.begin(),
out_shape.end());
auto in_shape = framework::vectorize<int64_t>(in_dims);
thrust::device_vector<int64_t> in_dims_vec(in_shape.begin(),
in_shape.end());
thrust::device_vector<int64_t> offsets_vec(offsets.begin(), offsets.end());
const int64_t* out_dims_ptr = thrust::raw_pointer_cast(out_dims_vec.data());
const int64_t* in_dims_ptr = thrust::raw_pointer_cast(in_dims_vec.data());
const int64_t* offsets_ptr = thrust::raw_pointer_cast(offsets_vec.data());
const std::vector<int64_t> out_shape =
framework::vectorize<int64_t>(out_dims);
const std::vector<int64_t> in_shape =
framework::vectorize<int64_t>(in_dims);

framework::Tensor out_dims_tensor;
framework::Tensor in_dims_tensor;
framework::Tensor offsets_tensor;
framework::TensorFromVector(out_shape, ctx.device_context(),
&out_dims_tensor);
framework::TensorFromVector(in_shape, ctx.device_context(),
&in_dims_tensor);
framework::TensorFromVector(offsets, ctx.device_context(), &offsets_tensor);
const int64_t* out_dims_ptr = out_dims_tensor.data<int64_t>();
const int64_t* in_dims_ptr = in_dims_tensor.data<int64_t>();
const int64_t* offsets_ptr = offsets_tensor.data<int64_t>();

switch (rank) {
case 1:
Expand Down