[Copilot]

Description

CUDA EP's Slice always allocated an output buffer and launched the per-element _SliceKernel, even when the output is a contiguous subregion of the input (e.g. tensor[i:j, :]). This adds a fast path that copies such slices with a single cudaMemcpyAsync (DeviceToDevice) from input_ptr + offset, skipping the per-element index decomposition.

• TryComputeContiguousSliceOffset (core/providers/cuda/tensor/slice.cc): detects the contiguous case and returns the start element offset. Conditions: all steps == 1; scanning from the rightmost axis, axes are fully included up to the first trimmed (pivot) axis; every axis left of the pivot selects exactly one element. Offset is computed from row-major strides.
• Slice::CallSliceImp: takes the memcpy path when the slice is contiguous, otherwise falls back to the existing kernel. Placed in CallSliceImp rather than the shared ComputeInternal, so training SliceGrad (which overrides CallSliceImp with scatter semantics) is unaffected.
• Tests (slice_op.test.cc): added Slice3D_ContiguousLeadingAxis and Slice3D_ContiguousSingleLeadingIndex, exercising the path across CPU/CUDA EPs.

Covered patterns: leading-axis slicing, batch splitting on axis 0, and single-index leading selection. Strided, middle-axis, and inner-axis slices remain on the kernel path.

Motivation and Context

Models with repeated leading-axis slicing (e.g. per-layer embedding extraction) incur one kernel launch per Slice. Replacing the trivially-contiguous cases with a single memcpy removes the per-element index computation and launch overhead, matching how Reshape/Squeeze/Flatten avoid copies. Full zero-copy aliasing would require allocation-planner support for offset aliases; this is the lower-risk memcpy alternative noted in the issue.

[Copilot]

Pull request overview

This PR optimizes the CUDA Execution Provider’s Slice implementation by adding a fast path for step-1 slices that map to a single contiguous subregion of the input, enabling a single cudaMemcpyAsync(DeviceToDevice) instead of launching the per-element slice kernel.

Changes:

• Added contiguous-slice detection (TryComputeContiguousSliceOffset) and a memcpy-based fast path in Slice::CallSliceImp for eligible cases.
• Added new Slice tests intended to exercise contiguous leading-axis slicing behavior across providers.

Reviewed changes

Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.

File	Description
onnxruntime/core/providers/cuda/tensor/slice.cc	Adds contiguous-slice detection and a cudaMemcpyAsync fast path in the CUDA Slice kernel implementation.
onnxruntime/test/providers/cpu/tensor/slice_op.test.cc	Adds new Slice test cases targeting contiguous leading-axis slice patterns.

Suppressed comments (1)

onnxruntime/core/providers/cuda/tensor/slice.cc:291

• In the memcpy fast path, the byte offset and copy size are derived from int64_t output_size/offset_in_elements without guarding against size_t overflow. If either multiplication overflows, cudaMemcpyAsync will copy the wrong region/size. Since this is an optimization, add conservative bounds checks and fall back to the kernel path when the copy cannot be represented safely.

      const char* input_data = static_cast<const char*>(input_tensor->DataRaw()) +
                               offset_in_elements * static_cast<int64_t>(element_size);
      CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(output_tensor->MutableDataRaw(), input_data,
                                           static_cast<size_t>(output_size) * element_size,
                                           cudaMemcpyDeviceToDevice, Stream(ctx)));

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[Tianlei Wu (tianleiwu)]

Reviewed the contiguous-region detection, checked offset/byte arithmetic, CUDA stream behavior, SliceGrad separation, and test coverage. The detector matched exhaustive row-major enumeration for 11,110 non-empty step-1 slices across ranks 1–4, and all required checks pass on this head. No actionable findings.