upstream-archive/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-08-07 02:20:48 -05:00
Code Issues Packages Projects Releases Wiki Activity

remove dead warmup code

Browse Source
This commit is contained in:
Scott Cutler
2026-04-23 14:47:49 -07:00
parent ebc31bccfa
commit 014ad9f61d
1 changed files with 0 additions and 65 deletions
Download Patch File Download Diff File Expand all files Collapse all files

65
ggml/src/ggml-cuda/allreduce.cu
View File

@@ -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);