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,