From 7449c1d30a00d1ffdb00e2239a27c3d952a39c1e Mon Sep 17 00:00:00 2001
From: Scott Cutler <scutler@nvidia.com>
Date: Thu, 23 Apr 2026 21:24:11 -0700
Subject: [PATCH] various small cleanups

---
 ggml/src/ggml-cuda/allreduce.cu | 5 +++--
 ggml/src/ggml-cuda/ggml-cuda.cu | 3 +++
 2 files changed, 6 insertions(+), 2 deletions(-)

diff --git a/ggml/src/ggml-cuda/allreduce.cu b/ggml/src/ggml-cuda/allreduce.cu
index 270a3d0e7f..4a72d91899 100644
--- a/ggml/src/ggml-cuda/allreduce.cu
+++ b/ggml/src/ggml-cuda/allreduce.cu
@@ -350,12 +350,14 @@ ggml_cuda_ar_pipeline * ggml_cuda_ar_pipeline_init(
     for (int i = 0; i < n_devices; ++i) {
         ggml_cuda_set_device(p->devices[i]);
 
-        if (cudaStreamCreateWithFlags(&p->streams[i], cudaStreamNonBlocking) != cudaSuccess) {
+        cudaStream_t stream = nullptr;
+        if (cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking) != cudaSuccess) {
             GGML_LOG_ERROR("%s: cudaStreamCreateWithFlags failed for device %d\n",
                            __func__, p->devices[i]);
             ggml_cuda_ar_pipeline_free(p);
             return nullptr;
         }
+        p->streams[i] = stream;
 
         p->ev_pool[i] = new ggml_cuda_ar_event_slot[GGML_CUDA_AR_POOL_SIZE]();
         for (int s = 0; s < GGML_CUDA_AR_POOL_SIZE; ++s) {
@@ -490,7 +492,6 @@ bool ggml_cuda_ar_allreduce(
         ggml_cuda_ar_pipeline * p,
         ggml_backend_t        * backends,
         ggml_tensor           ** tensors) {
-    //printf("ggml_cuda_ar_allreduce\n");
     GGML_ASSERT(p != nullptr);
 
     const int n = p->n_devices;
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index d779c41adc..e545ca1cca 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -1241,6 +1241,9 @@ static void * ggml_backend_cuda_comm_init(ggml_backend_t * backends, size_t n_ba
 
         GGML_LOG_ERROR("%s: internal AllReduce pipeline init failed\n", __func__);
 #ifdef GGML_USE_NCCL
+        // Clear any sticky CUDA error left over from the failed pipeline init
+        // so NCCL's own error-check on entry doesn't observe it.
+        (void) cudaGetLastError();
         ret->preferred_provider = GGML_CUDA_ALLREDUCE_NCCL;
 #else
         delete ret;