Fill LSTM CUDA operator opset gap: extend coverage from opset 14 to opset 22

[Copilot]

Description

Extends LSTM CUDA kernel registration from opset 14 to opset 22.

• lstm.cc: Cap existing opset 14 kernel to versioned 14–21, add new non-versioned kernel at opset 22
• cuda_execution_provider.cc: Update forward declarations and BuildKernelCreateInfo entries accordingly (versioned 14–21 + non-versioned 22) for all three types (float, double, MLFloat16)
• deep_cpu_lstm_op_test.cc: Add ONNXRuntime_TestLSTMForward_OpSet22_CUDA test targeting the new registration
• docs/OperatorKernels.md: Update CUDA LSTM entry from 14+ to [14, 21] and 22+

No spec-level behavior changes between opsets 14 and 22 for LSTM — this is purely a registration gap fill so the CUDA EP correctly claims nodes exported at newer opset versions.

Motivation and Context

LSTM CUDA kernel was registered only up to opset 14 while the ONNX spec defines LSTM through opset 22. Models exported at opset ≥15 would fall back to CPU. Follows the same pattern established by other opset gap PRs (ConvTranspose, MaxPool, Pad, etc.) referenced in #27729.

---

📍 Connect Copilot coding agent with Jira, Azure Boards or Linear to delegate work to Copilot in one click without leaving your project management tool.

[Tianlei Wu (tianleiwu)]

/azp run Windows GPU Doc Gen CI Pipeline

[azure-pipelines]

Azure Pipelines successfully started running 1 pipeline(s).

[Tianlei Wu (tianleiwu)]

Copilot, please update docs/OperatorKernels.md, and merge latest main branch to this branch.

[Copilot]

Copilot, please update docs/OperatorKernels.md, and merge latest main branch to this branch.

Updated docs/OperatorKernels.md to split the CUDA LSTM entry from 14+ into [14, 21] and 22+, and merged latest main. See commit 757c186.

[github-actions]

You can commit the suggested changes from lintrunner.

[Tianlei Wu (tianleiwu)]

Review Summary

Clean, low-risk opset gap-fill that extends CUDA LSTM kernel registration from opset 14 to cover opsets 14–21 (versioned) and 22+ (non-versioned). The pattern exactly matches the existing GRU and RNN registrations. Since the ONNX LSTM spec has no functional changes between opset 14 and 22 (the opset-22 doc string is literally kDoc_LSTM_ver14), the same kernel implementation handles all versions correctly.

Positives:

• Registration macros in lstm.cc follow the established three-macro structure from gru.cc and rnn.cc (REGISTER_KERNEL_VERSIONED_TYPED for 7–13, REGISTER_KERNEL_VERSIONED_TYPED_14 for 14–21, REGISTER_KERNEL_TYPED for 22+).
• Forward declarations and BuildKernelCreateInfo entries in cuda_execution_provider.cc are placed adjacent to GRU/RNN entries, maintaining the ordering convention.
• docs/OperatorKernels.md correctly updated.
• Test properly guards with DefaultCudaExecutionProvider() check and validates behavior equivalence against opset 14 expected values.

[Dmitri Smirnov (yuslepukhin)]

In cudnn_rnn_base.cc, multiple size computations use raw int64_t multiplication with no overflow guard:

// Line ~109: weight buffer size
int64_t w_size = num_directions_ * (number * hidden_size_ * (input_size + hidden_size_ + 2));

// Line ~254: output buffer size
int64_t output_size = seq_length * num_directions_ * batch_size * hidden_size_;

// Line ~240: reverse buffer
GetScratchBuffer(seq_length * batch_size * input_size, ...)

[Dmitri Smirnov (yuslepukhin)]

The code uses gsl::narrow_cast<int32_t>(seq_length) etc. throughout ComputeInternal. narrow_cast is an unchecked cast — it does NOT throw on truncation (unlike gsl::narrow). If seq_length, batch_size, or hidden_size exceed INT32_MAX, this silently truncates. Given these values come from user-controlled tensor shapes, gsl::narrow would be safer.

[Dmitri Smirnov (yuslepukhin)]

No explicit rank validation on input X.
f X has rank < 3, this is an out-of-bounds read on the shape vector. The ONNX schema should enforce rank-3, but an explicit ORT_RETURN_IF(X->Shape().NumDimensions() != 3, ...) guard would be defensive

int64_t seq_length = X->Shape()[0];
int64_t batch_size = X->Shape()[1];
int64_t input_size = X->Shape()[2];

[Dmitri Smirnov (yuslepukhin)]

SetZeroSequences takes zero_seq_index_cache by value.

This copies the entire vector on every call. Should be const std::vector<int32_t>&. This is a performance issue, not a correctness bug.

Coding standards require passing gsl::span in such cases.

[Dmitri Smirnov (yuslepukhin)]

Critical: The CUDA EP silently produces wrong results when input_forget=1 or when peephole weights P are provided. It should either implement the feature or return INVALID_ARGUMENT / decline the node so it falls back to CPU. This is the most serious functional correctness gap — the kernel claims the node but computes the wrong result.

[Dmitri Smirnov (yuslepukhin)]

CPU Implementation:

double type registered but throws ORT_NOT_IMPLEMENTED at runtime. CUDA supports it, we need parity.

[Tianlei Wu (tianleiwu)]

Addressing Review Feedback

All four review-level concerns from Dmitri Smirnov (@yuslepukhin) on cudnn_rnn_base.cc and the test have been addressed:

1. Overflow in size computations

Added #include "core/common/safeint.h" and wrapped overflow-prone multiplications with SafeInt<int64_t>:

• w_size computation in ReorganizeWeights
• output_size computation in ComputeInternal
• Scratch buffer size calculations for GetScratchBuffer<T> and ReverseBySequence

2. gsl::narrow_cast → gsl::narrow

Replaced all gsl::narrow_cast<int32_t> with checked gsl::narrow<int32_t> throughout ComputeInternal and SetZeroSequences. This covers seq_length, batch_size, input_size, and hidden_size_ casts.

3. Rank validation on input X

Added an explicit rank check immediately after the null check:

ORT_RETURN_IF(X->Shape().NumDimensions() != 3,
             "Input X must be 3-D [seq_length, batch_size, input_size], got rank ",
             X->Shape().NumDimensions());

4. SetZeroSequences takes vector by value

Changed the signature from const std::vector<int32_t> zero_seq_index_cache to gsl::span<const int32_t> zero_seq_index_cache in both the header and implementation. Call sites now construct a span over the meaningful portion of the cache.

5. Test improvements (addressing Tianlei Wu (@tianleiwu) feedback as well)

• Fixed a bug: input_size was 2 but W_data only had 8 elements for a [1, 8, 2] = 16 element tensor. Changed input_size to 1 so the shapes are consistent.
• Added Y output validation: The test now validates both Y (full sequence output) and Y_h (final hidden state). Expected values were computed via a reference LSTM implementation.

[Dmitri Smirnov (yuslepukhin)]

Both CPU and CUDA: OOB read before rank check on X
In lstm_base.cc:43-48:

int seq_length = narrow(X_shape[0]); // OOB if rank < 3
int batch_size = narrow(X_shape[1]);
int input_size = narrow(X_shape[2]);

Status status = ValidateInputs(X, ...); // rank check is HERE, too late

Same pattern in cudnn_rnn_base.cc:190-192:

int64_t seq_length = X->Shape()[0]; // No rank check anywhere
int64_t batch_size = X->Shape()[1];
int64_t input_size = X->Shape()[2];

[Dmitri Smirnov (yuslepukhin)]

Neither implementation validates W or R shapes:

W and R are not validated. No check that:

W.rank == 3, W[0] == num_directions, W[1] == 4hidden_size, W[2] == input_size
R.rank == 3, R[0] == num_directions, R[1] == 4hidden_size, R[2] == hidden_size.

In deep_cpu_lstm.cc:300-301, W_shape[1], W_shape[2] are accessed without rank validation — OOB if W is rank < 3.

[Dmitri Smirnov (yuslepukhin)]

CUDA: No ValidateInputs call exists at all. In ReorganizeWeights, W->Shape()[2] is accessed without checking rank.

[github-actions]

You can commit the suggested changes from lintrunner.

[Dmitri Smirnov (yuslepukhin)]

reorganized_w_data_size_in_bytes = w_size * sizeof(T);

Multiplication is not SafeInt-guarded (it's int64_t * size_t — standard arithmetic). Should be SafeInt<size_t>(w_size) * sizeof(T).

---

Refers to: onnxruntime/core/providers/cuda/rnn/cudnn_rnn_base.cc:112 in a08329d. [](commit_id = a08329d, deletion_comment = False)

[Dmitri Smirnov (yuslepukhin)]

  CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(y_data, y_reorganized_data.get(), output_size * sizeof(T),

Same issue — output_size is SafeInt-computed but the final * sizeof(T) is raw multiplication.

---

Refers to: onnxruntime/core/providers/cuda/rnn/cudnn_rnn_base.cc:396 in a08329d. [](commit_id = a08329d, deletion_comment = False)

[Dmitri Smirnov (yuslepukhin)]

ORT_RETURN_IF_ERROR(tmp_rnn_desc.Set(W->Shape()[2],  // input_size

CacheCudnnRnnWeights accesses W->Shape()[2] without rank check — cudnn_rnn_base.cc:153.

The W/R rank check was added in ReorganizeWeights, but CacheCudnnRnnWeights also accesses W->Shape()[2] directly (line 153: tmp_rnn_desc.Set(W->Shape()[2], ...)) without going through ReorganizeWeights first.

If W is a constant input with rank < 3, this is an OOB read in the constructor path.

---

Refers to: onnxruntime/core/providers/cuda/rnn/cudnn_rnn_base.cc:148 in a08329d. [](commit_id = a08329d, deletion_comment = False)

[Tianlei Wu (tianleiwu)]

Addressing all review feedback from Dmitri Smirnov (@yuslepukhin) — here's the status of each item on the current HEAD (9aeafcf):

Feedback	Status
SafeInt overflow guards for size computations	✅ Fixed — w_size, output_size, scratch buffer sizes, and count in SetWeightBias all use SafeInt
gsl::narrow_cast → checked gsl::narrow	✅ Fixed — both ComputeInternal int32 casts and CudnnRNN::Set int casts
Rank-3 validation on input X	✅ Fixed — ORT_RETURN_IF guard before shape access
SetZeroSequences pass-by-value	✅ Fixed — changed to gsl::span<const int32_t>
input_forget=1 / peephole P silently ignored	✅ Fixed — input_forget validated in LSTM ctor (gated on attribute presence for opset compat); peephole P rejected at runtime in ComputeInternal
W/R shape not validated / OOB in ReorganizeWeights	✅ Fixed — rank-3 checks added at top of ReorganizeWeights
CUDA no ValidateInputs / W->Shape()[2] without rank check	✅ Fixed — covered by the ReorganizeWeights rank checks above
Loop for (int i = 0; i < batch_size) type mismatch	✅ Fixed — changed to int64_t i with gsl::narrow<int32_t> for the assignment
cudaStreamSynchronize return unchecked	✅ Fixed — wrapped in CUDA_RETURN_IF_ERROR
CUDA_CALL_THROW in Status-returning SetWeightBias	✅ Fixed — changed to CUDA_RETURN_IF_ERROR
int offsets in SetCudnnRnnWeightBias can wrap	✅ Fixed — changed to size_t
CPU double ORT_NOT_IMPLEMENTED parity	⏭️ Out of scope — pre-existing CPU issue, not introduced by this PR
CPU lstm_base.cc OOB before rank check	⏭️ Out of scope — pre-existing CPU issue in a different file

[Dmitri Smirnov (yuslepukhin)]

LGTM