From 0b1e3f1e1012480587d42ad474613bbd4142ee0e Mon Sep 17 00:00:00 2001 From: Thiago Crepaldi Date: Wed, 1 Apr 2020 14:28:17 -0700 Subject: [PATCH] Revert _SliceKernel cuda implementation --- onnxruntime/core/providers/cuda/tensor/slice_impl.cu | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/onnxruntime/core/providers/cuda/tensor/slice_impl.cu b/onnxruntime/core/providers/cuda/tensor/slice_impl.cu index f98bc19837..0c5b3eb514 100644 --- a/onnxruntime/core/providers/cuda/tensor/slice_impl.cu +++ b/onnxruntime/core/providers/cuda/tensor/slice_impl.cu @@ -29,12 +29,15 @@ __global__ void _SliceKernel(const int32_t dimension_count, break; } - output_strides[dim].divmod(value, div, mod); - input_index += (starts[dim] + div * steps[dim]) * input_strides[dim]; + output_strides.data_[dim].divmod(value, div, mod); + input_index += (starts.data_[dim] + div * steps.data_[dim]) * input_strides.data_[dim]; value = mod; } - input_index += starts[dim] + mod * steps[dim]; - output_data[id] = input_data[input_index]; + input_index += starts.data_[dim] + mod * steps.data_[dim]; + if (is_grad) + output_data[input_index] = input_data[id]; + else + output_data[id] = input_data[input_index]; } Status SliceImpl(const size_t element_size,