From 014ad9f61d682c278ea417fdd90325cdd29eb755 Mon Sep 17 00:00:00 2001
From: Scott Cutler <scutler@nvidia.com>
Date: Thu, 23 Apr 2026 14:47:49 -0700
Subject: [PATCH] remove dead warmup code

---
 ggml/src/ggml-cuda/allreduce.cu | 65 ---------------------------------
 1 file changed, 65 deletions(-)

diff --git a/ggml/src/ggml-cuda/allreduce.cu b/ggml/src/ggml-cuda/allreduce.cu
index 61612510a4..a9dad4c13a 100644
--- a/ggml/src/ggml-cuda/allreduce.cu
+++ b/ggml/src/ggml-cuda/allreduce.cu
@@ -420,71 +420,6 @@ ggml_cuda_ar_pipeline * ggml_cuda_ar_pipeline_init(
     }
 #endif
 
-#if 0
-    // Warmup: run the kernel N times to pay first-use driver / PCIe /
-    // page-mapping costs during model load and encourage the GPU clock
-    // governor to boost before inference begins.
-    if (n_devices == 2) {
-        printf("ggml_cuda_ar_pipeline_init warmup\n");
-        
-        constexpr int    WARMUP_ITERS = 64;
-        constexpr size_t WARMUP_COUNT = 8192;  // 32 KB of fp32
-        constexpr size_t WARMUP_BYTES = WARMUP_COUNT * sizeof(float);
-
-        float * dev_buf[2] = {};
-        bool warmup_ok = true;
-        for (int i = 0; i < 2; ++i) {
-            ggml_cuda_set_device(p->devices[i]);
-            if (cudaMalloc(reinterpret_cast<void **>(&dev_buf[i]), WARMUP_BYTES) != cudaSuccess) {
-                GGML_LOG_WARN("%s: warmup alloc failed for device %d, skipping\n",
-                              __func__, p->devices[i]);
-                warmup_ok = false;
-                break;
-            }
-        }
-
-        if (warmup_ok) {
-            // Warmup uses float kernel.
-            for (int iter = 0; iter < WARMUP_ITERS; ++iter) {
-                for (int r = 0; r < 2; ++r) {
-                    *ggml_cuda_ar_arrival_ptr(p, /*slot=*/0, r) = 0;
-                }
-                for (int r = 0; r < 2; ++r) {
-                    ggml_cuda_set_device(p->devices[r]);
-                    ggml_cuda_ar_kernel<float><<<dim3(1), dim3(256), 0, p->streams[r]>>>(
-                        dev_buf[r], dev_buf[r],
-                        reinterpret_cast<float *>(p->host_buf[r]),
-                        reinterpret_cast<const float *>(p->host_buf[1 - r]),
-                        static_cast<int>(WARMUP_COUNT),
-                        ggml_cuda_ar_arrival_ptr(p, /*slot=*/0, r),
-                        ggml_cuda_ar_arrival_ptr(p, /*slot=*/0, 1 - r)
-#if GGML_CUDA_AR_WATCHDOG
-                       ,p->debug_ring[r],
-                        0,  // max_spin = 0 (no limit during warmup)
-                        r,
-                        0   // slot = 0
-#endif
-                        );
-                }
-            }
-            for (int i = 0; i < 2; ++i) {
-                ggml_cuda_set_device(p->devices[i]);
-                cudaStreamSynchronize(p->streams[i]);
-            }
-            GGML_LOG_DEBUG("%s: warmup complete (%d iters x %zu KB)\n",
-                           __func__, WARMUP_ITERS, WARMUP_BYTES >> 10);
-        }
-
-        for (int i = 0; i < 2; ++i) {
-            if (dev_buf[i]) {
-                ggml_cuda_set_device(p->devices[i]);
-                cudaFree(dev_buf[i]);
-            }
-        }
-        
-        printf("ggml_cuda_ar_pipeline_init warmup finished\n");
-    }
-#endif
     GGML_LOG_INFO("%s: initialized AllReduce pipeline: %d GPUs, "
                   "%zu KB staging per GPU\n",
                   __func__, n_devices, max_bytes >> 10);