“AI-generated code is allowed. You are 100% responsible for every line, however it was produced.” — llama.cpp AI Usage Policy

---

Every ten minutes, my server died.

That’s a setting. --sleep-idle-seconds 600 — ten minutes idle, and llama-server would hang.

No crash. No restart. Just dead.

Oh, it would try to spin up a new thread, just to give me hope! Then the thread would sit there staring at me. Any time I stepped away to think, or grab coffee, I’d come back to a corpse.

The fix was always the same. Stop the server. Start it again. Watch it come back to life like nothing happened.

I did that dance for two months.

I finally tracked down the mechanism. It only hits when you run with --flash-attn on and --sleep-idle-seconds set to anything. On wake, the server never recalculates its memory allocation before trying to spin up a new thread. So it hangs.

Somebody on GitHub had already done the hard diagnostic work. User christophkogler traced the whole thing to a stale process-static flag that survives a CUDA context reset:

“The sequence I traced is: On the initial load, ggml_cuda_flash_attn_ext_mma_f16_case() calls cudaFuncSetAttribute(...). A process-static shared_memory_limit_raised[device] flag records that the attribute was set. Entering the server sleeping state destroys the model and backend objects. During reload, ggml_backend_cuda_device_get_memory() calls cudaDeviceReset() when the CUDA backend active count is zero. The CUDA context and its function attributes are reset, but the process-static flag remains true. Flash-attention setup therefore skips cudaFuncSetAttribute() in the new context. cudaOccupancyMaxActiveBlocksPerMultiprocessor() evaluates the kernel with the default dynamic shared-memory limit and returns zero active blocks, leading to the assertion.”

A real bug. A real root cause. Reproducible 100% of the time by at least one other person in the thread.

Two weeks later, it got closed as “Not planned.”

I get it, a little. llama.cpp has 692 open issues and over 1,300 open pull requests right now. Not everything gets looked at, but this one had a traced call stack and a patch attached to it.

And it kept getting in my way for a reason bigger than my own annoyance. I’d built llama-server-manager, a wrapper whose whole premise is that you shouldn’t have to compile llama.cpp yourself. It’s supposed to detect your hardware, download the right binary out of the 28 archives llama.cpp ships per release, and update itself.

This bug meant I couldn’t update llama.cpp without hand-compiling a patched fork. That’s exactly the chore the wrapper exists to eliminate. It wasn’t just annoying, it was undermining the whole point of the thing I’d built.

So on Saturday, I sat down and made up my mind to fix it myself.

Five Hours With a Junior Developer

I had a working copy of Gemma 4 12B QAT. What more did I need?

I put OpenCode in Plan mode and started asking it questions. Some of llama.cpp’s files run past 5,000 lines. I’d read the code a few times already and knew I wasn’t going to absorb a codebase that size fast enough on my own.

So I leaned on the model to navigate it.

“How do I access the command line parameters at this point in the code, Gemma?” I knew the flags mattered, but I had no idea how to get them into the CUDA code where the actual hang happens.

Gemma found it. ggml/src/ggml-cuda/ggml-cuda.cu:4807 already had a params argument sitting in the function signature, unused for what I needed.

static ggml_backend_t ggml_backend_cuda_device_init_backend(ggml_backend_dev_t dev, const char * params) {

That’s where I started.

I found the duplicated logic christophkogler’s patch had worked around. The same shared-memory guard, copy-pasted across fattn-mma-f16.cuh and common.cuh. Instead of just stripping the guard and reapplying the fix unconditionally on every kernel launch, I consolidated it into one macro in common.cuh, with a force_reapply flag:

#define CUDA_SET_SHARED_MEMORY_LIMIT_FORCE(kernel, nbytes, force_reapply)                                  \
    do {                                                                                                   \
        static bool shared_memory_limit_raised[GGML_CUDA_MAX_DEVICES] = { false };                         \
        const int   id                                                = ggml_cuda_get_device();            \
        if (!shared_memory_limit_raised[id] || (force_reapply)) {                                          \
            CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes)); \
            shared_memory_limit_raised[id] = true;                                                         \
        }                                                                                                  \
    } while (0)

Called from fattn-mma-f16.cuh, gated on the actual failure condition instead of firing every time:

CUDA_SET_SHARED_MEMORY_LIMIT_FORCE(reinterpret_cast<fattn_kernel_ptr_t>(fattn_kernel), nbytes_shared_total,
    (ctx.flash_attn_type != 0 && ctx.sleep_idle_seconds != 0));

I explained what I wanted. Gemma wrote it. I tweaked it by hand to clean it up.

None of it came free. There were compile errors Gemma didn’t catch on the first pass. I’d read the error, point it at the offending line, let it try again.

At one point I dropped a TODO comment on a single line and pointed it there directly, so it would expand the parser to handle every way --flash-attn can actually show up on a command line: --flash-attn, or the short flag -fa.

static void ggml_backend_cuda_parse_params(const char * params, ggml_backend_cuda_context * ctx) {
    if (params == nullptr) {
        return;
    }
    // Look for --flash-attn on
    bool flash_attn = (strstr(params, "--flash-attn ") && !strstr(params, "--flash-attn off")) ||
                       (strstr(params, "-fa ") && !strstr(params, "-fa off"));
    if (flash_attn) {
        ctx->flash_attn_type = 1;
    }

    // Look for --sleep-idle-seconds <value>
    const char * sleep_idle = strstr(params, "--sleep-idle-seconds");
    if (sleep_idle != nullptr) {
        const char * value_ptr = strchr(sleep_idle, ' ');
        if (value_ptr != nullptr) {
            ctx->sleep_idle_seconds = atoi(value_ptr + 1);
        }
    }
}

That’s just how it goes with Gemma. It’s a good junior developer that occasionally needs its nose pointed at the exact line.

Five hours in, I had something that worked and, more importantly, looked sane.

Then came the part that actually humbled me. The paperwork.

The Two-Backend Problem

The Contributing doc says a CUDA change needs testing on at least two backends. Fine. CUDA and Vulkan both run on my RTX. That’s the two.

I found the CI tests, ran the CUDA suite. Slow, but it finished. Ran the Vulkan suite. Big nope. Turns out having an Nvidia card doesn’t mean you have Vulkan support. I’d never installed any of it.

It was 2 a.m. I went to bed.

Sunday

“Hey Gemini, what gives with this error?” “Oh hi Josh. Yeah you got some dependencies to install my dude.”

Gemini gave me a package name. That package didn’t exist. It gave me another. Then another. Eventually, after enough rounds of trial and error, I had this:

sudo apt install vulkan-tools shaderc glslc spirv-headers

And a build command to match:

rm -rf build/ && cmake -B build -DGGML_CUDA=ON -DGGML_VULKAN=ON -DCMAKE_BUILD_TYPE=Release && cmake --build build --config Release --parallel

It compiled. I kicked off the test suite, and it ran for almost two hours. I read Slashdot. I read the news. I drank some coffee. I had lunch.

The output was full of success messages the whole way through, so I had a decent read on where it was headed before it finished cleanly.

I fixed it. Time to prove it properly!

Pull the latest master, recompile, reproduce the hang on stock master, then switch to my branch and show the fix actually mattered.

The hang didn’t happen on master.

W. T. F.

I was certain I’d screwed something up, so I ran it again. Then again, for good measure.

My disbelief didn’t go away after two clean runs, so I uninstalled the Vulkan packages entirely and recompiled, on the theory that maybe the binary was quietly picking up Vulkan tooling even though I’d explicitly told it --device cuda0. Two more clean runs. Four total.

Then I tested the official prebuilt Vulkan binary, the one I never touched. Also clean.

I’d had this bug for months. I’d spent a weekend rewriting the fix myself. And somewhere between my Saturday afternoon with Gemma and testing it Sunday afternoon, it just stopped happening.

I went back through the files I’d edited and diffed them against what changed on master since I last pulled. Nothing obviously related. No major system updates on my end either.

christophkogler’s own writeup includes a line that’s been rattling around in my head since:

“This will probably not manifest with every model. It requires selection of a kernel whose dynamic shared-memory requirement exceeds the default limit; kernels that fit within the default can hide the stale-state problem.”

Maybe that’s it. Maybe something upstream shifted which kernel gets selected, and the bug is just hiding again, waiting. Maybe it’ll be back tomorrow. I don’t know yet, and I’m not going to pretend I do.

I posted my branch to the GitHub issue yesterday and tagged the two people who’d done the original diagnostic work, asking a specific question: are they running llama-server in router mode? I noticed christophkogler’s command runs a single model directly. No router:

./llama.cpp/build/bin/llama-server \
    -hf unsloth/gemma-4-12B-it-qat-GGUF:UD-Q4_K_XL \
    --spec-type draft-mtp \
    --spec-draft-n-max 4 \
    -ngl 999 \
    -fa on \
    --sleep-idle-seconds 300

I always run in router mode, because it’s the only way to swap models without restarting the server. That’s the entire reason llama-server-manager exists.

If router mode is the differentiator, that’s worth knowing. If it’s not, that’s worth knowing too. Either way, I haven’t heard back yet.

What I Actually Learned

I wasn’t sure a 12-billion-parameter model running on a 12GB consumer card could hold something the size of llama.cpp in its head. It did better than I expected.

It blew through its context window a few times and had to compact. Dump what it had, reload the relevant files, keep going. And it recovered every time instead of losing the thread. That’s a real data point, not a hopeful one.

llama.cpp’s AI Usage Policy says AI-generated code is allowed, and that you’re 100% responsible for every line, however it was produced. Here’s the honest accounting of this weekend: the AI wrote the code. I updated the code. I showed it exactly where it needed to expand for robustness, the flag parsing, the compile errors, the line where the fix actually needed to live.

I still own every bit of it.

Most people ask the same question about AI right now, whether they’re talking about code or anything else. Does it work? Is it a magic box?

I don’t have a clean answer. I fixed a real bug with real help from a model that occasionally needed its nose pointed at the right line. And then the bug I fixed might not have needed fixing anymore, for reasons I still don’t understand.

It isn’t resolved; that’s exactly why I wanted to write it down.