[CUDA] Add cuDNN-free ArgMax/ArgMin/ReduceSum and fix LogSoftmax on plugin EP

.github/workflows/linux_cuda_no_cudnn.yml

@@ -93,15 +93,37 @@ jobs:
           ldd /build/Release/Release/libonnxruntime_providers_cuda.so | tee /tmp/ldd.txt
           ! grep -i cudnn /tmp/ldd.txt'
 
-      - name: Run no-cuDNN CUDA EP smoke test
+      - name: Run no-cuDNN CUDA EP op smoke test
         run: |
-          docker run --rm --gpus all \
+          # The build image runs as a non-root user by default, but this step
+          # deletes the system cuDNN libraries and rebuilds the dynamic linker
+          # cache (both require root). Run this disposable test container as root.
+          docker run --rm --gpus all --user root \
             -v "${{ runner.temp }}/Release:/build/Release" \
             "${{ steps.build_docker_image_step.outputs.full-image-name }}" \
             bash -lc 'set -e
           PATH=/opt/python/cp312-cp312/bin:$PATH
           LD_LIBRARY_PATH=/usr/local/cuda-13.0/lib64:${LD_LIBRARY_PATH:-}
           export PATH LD_LIBRARY_PATH
+
+          # The build image contains cuDNN. Remove its runtime libraries from this disposable
+          # test container so the smoke test proves that the provider works when cuDNN is
+          # physically unavailable, rather than merely relying on enable_cudnn=0.
+          for root in /usr /opt /lib /lib64; do
+            if [ -e "$root" ]; then
+              find "$root" -name "libcudnn*.so*" -print -delete 2>/dev/null || true
+            fi
+          done
+          ldconfig
+          if ldconfig -p | grep -qi cudnn; then
+            echo "cuDNN remains available in the dynamic linker cache" >&2
+            exit 1
+          fi
+          if find /usr /opt /lib /lib64 -name "libcudnn*.so*" -print -quit 2>/dev/null | grep -q .; then
+            echo "cuDNN runtime libraries remain in the no-cuDNN test container" >&2
+            exit 1
+          fi
+
           WHEEL_PATH=$(find /build/Release/Release/dist -type f -name "onnxruntime_gpu-*.whl" | head -n 1)
           if [ -z "$WHEEL_PATH" ]; then
             echo "No built onnxruntime GPU wheel found under /build/Release/Release/dist" >&2
@@ -111,21 +133,114 @@ jobs:
           python -m pip install --no-cache-dir --force-reinstall --no-deps numpy onnx "$WHEEL_PATH"
           python - <<"PY"
           import numpy as np
-          import onnx
           import onnxruntime as ort
           from onnx import TensorProto, helper
 
-          x = helper.make_tensor_value_info("x", TensorProto.FLOAT, [2, 3])
-          y = helper.make_tensor_value_info("y", TensorProto.FLOAT, [2, 3])
-          node = helper.make_node("Add", ["x", "x"], ["y"])
-          graph = helper.make_graph([node], "cuda_no_cudnn_smoke", [x], [y])
-          model = helper.make_model(graph, opset_imports=[helper.make_opsetid("", 21)])
-          model.ir_version = 10
-
           providers = [("CUDAExecutionProvider", {"enable_cudnn": "0"})]
-          sess = ort.InferenceSession(model.SerializeToString(), providers=providers)
-          data = np.arange(6, dtype=np.float32).reshape(2, 3)
-          result = sess.run(None, {"x": data})[0]
-          np.testing.assert_allclose(result, data + data)
-          print("CUDA no-cuDNN smoke test passed")
+
+
+          def run_op(name, model, feeds, expected, rtol=1e-4, atol=1e-4):
+              model.ir_version = 10
+              options = ort.SessionOptions()
+              options.add_session_config_entry("session.disable_cpu_ep_fallback", "1")
+              sess = ort.InferenceSession(model.SerializeToString(), sess_options=options, providers=providers)
+              got = sess.run(None, feeds)[0]
+              np.testing.assert_allclose(got, expected, rtol=rtol, atol=atol)
+              print("[no-cuDNN] " + name + " passed")
+
+
+          def make_model(nodes, inputs, outputs, opset, initializers=None):
+              graph = helper.make_graph(nodes, "no_cudnn_smoke", inputs, outputs, initializer=initializers or [])
+              return helper.make_model(graph, opset_imports=[helper.make_opsetid("", opset)])
+
+
+          def f32_in(shape):
+              return helper.make_tensor_value_info("x", TensorProto.FLOAT, shape)
+
+
+          data = np.random.rand(2, 3).astype(np.float32)
+
+          # Elementwise baseline (does not depend on cuDNN).
+          run_op(
+              "Add",
+              make_model(
+                  [helper.make_node("Add", ["x", "x"], ["y"])],
+                  [f32_in([2, 3])],
+                  [helper.make_tensor_value_info("y", TensorProto.FLOAT, [2, 3])],
+                  21,
+              ),
+              {"x": data},
+              data + data,
+          )
+
+          # LogSoftmax over the last axis (softmax kernel, no cuDNN).
+          shifted = data - data.max(axis=1, keepdims=True)
+          log_softmax = shifted - np.log(np.exp(shifted).sum(axis=1, keepdims=True))
+          run_op(
+              "LogSoftmax",
+              make_model(
+                  [helper.make_node("LogSoftmax", ["x"], ["y"], axis=1)],
+                  [f32_in([2, 3])],
+                  [helper.make_tensor_value_info("y", TensorProto.FLOAT, [2, 3])],
+                  13,
+              ),
+              {"x": data},
+              log_softmax,
+          )
+
+          # ReduceMean over the last axis (matrix-reduction path, no cuDNN).
+          run_op(
+              "ReduceMean",
+              make_model(
+                  [helper.make_node("ReduceMean", ["x"], ["y"], axes=[1], keepdims=1)],
+                  [f32_in([2, 3])],
+                  [helper.make_tensor_value_info("y", TensorProto.FLOAT, [2, 1])],
+                  13,
+              ),
+              {"x": data},
+              data.mean(axis=1, keepdims=True),
+          )
+
+          # ReduceSum over a middle axis, which requires the general no-cuDNN path.
+          reduce_sum_data = np.random.rand(2, 3, 4).astype(np.float32)
+          run_op(
+              "ReduceSum",
+              make_model(
+                  [helper.make_node("ReduceSum", ["x", "axes"], ["y"], keepdims=1)],
+                  [f32_in([2, 3, 4])],
+                  [helper.make_tensor_value_info("y", TensorProto.FLOAT, [2, 1, 4])],
+                  13,
+                  initializers=[helper.make_tensor("axes", TensorProto.INT64, [1], [1])],
+              ),
+              {"x": reduce_sum_data},
+              reduce_sum_data.sum(axis=1, keepdims=True),
+          )
+
+          # ArgMax over the last axis (custom cuDNN-free kernel).
+          run_op(
+              "ArgMax",
+              make_model(
+                  [helper.make_node("ArgMax", ["x"], ["y"], axis=1, keepdims=1)],
+                  [f32_in([2, 3])],
+                  [helper.make_tensor_value_info("y", TensorProto.INT64, [2, 1])],
+                  13,
+              ),
+              {"x": data},
+              np.argmax(data, axis=1).reshape(2, 1).astype(np.int64),
+          )
+
+          # ArgMin over the last axis (custom cuDNN-free kernel).
+          run_op(
+              "ArgMin",
+              make_model(
+                  [helper.make_node("ArgMin", ["x"], ["y"], axis=1, keepdims=1)],
+                  [f32_in([2, 3])],
+                  [helper.make_tensor_value_info("y", TensorProto.INT64, [2, 1])],
+                  13,
+              ),
+              {"x": data},
+              np.argmin(data, axis=1).reshape(2, 1).astype(np.int64),
+          )
+
+          print("CUDA no-cuDNN op smoke tests passed")
           PY'

docs/cuda_plugin_ep/QUICK_START.md

@@ -27,7 +27,7 @@ For local Linux CUDA 13 validation, use the no-cuDNN helper script. It keeps `CU
 bash .env/cuda_130_plugin_no_cudnn.sh --build --test_plugin
 ```
 
-The test mode sets `ORT_TEST_CUDA_PLUGIN_EP=1` and `ORT_TEST_CUDA_PLUGIN_NO_CUDNN=1`, which passes `enable_cudnn=0` to plugin sessions and skips plugin tests for operators that still require cuDNN, such as Conv, ConvTranspose, BatchNormalization, InstanceNormalization, LRN, ArgMax, reductions, Einsum, and cuDNN-backed pooling paths.
+The test mode sets `ORT_TEST_CUDA_PLUGIN_EP=1` and `ORT_TEST_CUDA_PLUGIN_NO_CUDNN=1`, which passes `enable_cudnn=0` to plugin sessions and skips plugin tests for operators that still require cuDNN, such as Conv, ConvTranspose, BatchNormalization, InstanceNormalization, LRN, Einsum, and cuDNN-backed pooling paths.
 
 ## Minimum ONNX Runtime Version
 

onnxruntime/core/providers/cuda/cuda_kernel.h

@@ -133,6 +133,10 @@ class CudaKernel : public OpKernel {
     return RequireCudnnHandle(GetCudnnHandle(static_cast<CudaStream*>(ctx->GetComputeStream())));
   }
 
+  inline cudnnHandle_t TryGetCudnnHandle(OpKernelContext* ctx) const {
+    return GetCudnnHandle(static_cast<CudaStream*>(ctx->GetComputeStream()));
+  }
+
   static inline cudnnHandle_t GetCudnnHandle(onnxruntime::CudaStream* stream) {
     return stream ? stream->cudnn_handle_ : nullptr;
   }

onnxruntime/core/providers/cuda/math/softmax.h

@@ -47,7 +47,7 @@ class Softmax final : public CudaKernel {
       }
     }
 
-    log_softmax_ = info.GetKernelDef().OpName() == "LogSoftmax";
+    log_softmax_ = node.OpType() == "LogSoftmax";
   }
 
   Status ComputeInternal(OpKernelContext* context) const override;

onnxruntime/core/providers/cuda/plugin/cuda_kernel_adapter.h

@@ -1064,8 +1064,30 @@ class CudaKernel : public OpKernel {
           std::string("cuDNN is unavailable or disabled for CUDA Plugin Execution Provider: ") +
               onnxruntime::cuda::CudnnLibrary::Get().Error()));
     }
-    if (handle != nullptr && stream != nullptr) {
-      CUDNN_CALL_THROW(cudnnSetStream(handle, stream));
+    // Bind the shared handle to the current compute stream. cudaStream_t 0/nullptr is the default
+    // stream, which is still a valid stream to bind, so do this unconditionally to avoid leaving
+    // the handle bound to a stale stream from a previous call.
+    CUDNN_CALL_THROW(cudnnSetStream(handle, stream));
+    return handle;
+  }
+
+  cudnnHandle_t TryGetCudnnHandle(OpKernelContext* ctx) const {
+    auto stream = Stream(ctx);
+    auto handle = GetCudnnHandle(stream);
+    if (handle != nullptr) {
+      return handle;
+    }
+
+    handle = DefaultCudnnHandle();
+    if (handle != nullptr) {
+      // Bind the shared handle to the current compute stream. cudaStream_t 0/nullptr is the default
+      // stream, which is still a valid stream to bind, so do this unconditionally to avoid leaving
+      // the handle bound to a stale stream from a previous call.
+      // Keep this accessor non-throwing: if the stream cannot be bound, treat it as "no cuDNN handle"
+      // so callers can fall back to a cuDNN-free path instead of failing.
+      if (!CUDNN_CALL(cudnnSetStream(handle, stream)).IsOK()) {
+        return nullptr;
+      }
     }
     return handle;
   }