# Two CUDA PRs into llama.cpp, and what they taught me about a 100K-star codebase

My first pull request to llama.cpp was 95 lines of CUDA. Georgi Gerganov, the guy who started the project, merged it the next day.

The second one was even smaller. Nine new lines of code, and the change that mattered most was deleting a condition. Honestly, that one taught me more.

Neither of these is a big PR. But llama.cpp has over 100,000 stars and pretty much everyone running LLMs on their own machine depends on it, so even a small change forces you to understand how the thing works underneath. This post is what I picked up along the way: how ggml decides which device runs an op, what my two changes actually did, how testing and review went, and the boring little things that decide whether a PR gets merged or sits there.

## Finding something to work on

llama.cpp sits on top of **ggml**, a tensor library with backends for nearly everything: CPU, CUDA, Metal, Vulkan, SYCL, WebGPU and a few more. Every backend supports some set of ops. None of them support all of them.

Two things made it easy to find useful work:

1. **`docs/ops.md`**: a big table of every op against every backend, marked ✅ (supported), 🟡 (partially supported) or ❌ (not supported).
2. **Issue [#14909](https://github.com/ggml-org/llama.cpp/issues/14909), "Implement missing ops from backends"**, opened by Aman Gupta (am17an). It's an umbrella issue with a simple deal: pick an ❌, implement it, test it with `test-backend-ops`, send a PR.

For a first contribution you can't really ask for more. The task is clear, there's a test harness, and there's almost always a similar op already implemented that you can learn from.

## How ggml decides where an op runs

Before writing any kernel, I spent time on the part of the CUDA backend that routes everything. In `ggml/src/ggml-cuda/ggml-cuda.cu`, two functions do most of the work:

- **`ggml_backend_cuda_device_supports_op`** answers one question: can CUDA run this op, with these types and shapes?
- **`ggml_cuda_compute_forward`** is a big `switch` over op types that calls the actual CUDA implementation.

When the scheduler splits the graph across devices, it asks each backend whether it supports each op. If CUDA says no, **that op just runs on the CPU instead**, with copies to and from the GPU around it. Nothing crashes. Nothing warns you. The output is still correct, only slower.

Keep that in mind, because it's basically the whole story of my second PR.

## PR 1: POOL_1D on CUDA ([#27573](https://github.com/ggml-org/llama.cpp/pull/27573))

`GGML_OP_POOL_1D` is 1D pooling. You slide a window of size `k` with stride `s` across each row (with `p` elements of padding on each side) and output either the average or the max of each window. CUDA already had `POOL_2D`, but not the 1D version, so it was sitting at ❌ in the table.

### The kernel

I copied the structure of `pool2d.cu`, which is how most CUDA ops in ggml are organised: a `.cuh` header, a `.cu` file with the kernel and a launcher, and two small hooks in `ggml-cuda.cu` (one in `supports_op`, one in `compute_forward`).

Each thread computes one output element. With `nr` rows and `OW` outputs per row, that's `nr * OW` threads, in blocks of 256:

```cpp
static __global__ void pool1d_nchw_kernel(
        const int iw, const int ow,
        const int kw, const int sw, const int pw,
        const int parallel_elements,
        const float * src, float * dst, const enum ggml_op_pool op) {
    const int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx >= parallel_elements) {
        return;
    }

    const int nc = idx / ow;        // which row
    const int cur_ow = idx % ow;    // which output in that row

    const float * i_ptr = src + nc * iw;
    float * o_ptr = dst + nc * ow;

    const int start = cur_ow * sw - pw;
    const int b = max(0, start);
    const int e = min(iw, start + kw);

    float res;
    switch (op) {
        case GGML_OP_POOL_AVG: res = 0.0f;     break;
        case GGML_OP_POOL_MAX: res = -FLT_MAX; break;
        default: return;
    }

    int count = 0;
    for (int i = b; i < e; i++) {
        float cur = __ldg(i_ptr + i);   // simplified: the real code falls back to a plain load below sm_35
        switch (op) {
            case GGML_OP_POOL_AVG: res += cur;          break;
            case GGML_OP_POOL_MAX: res = max(res, cur); break;
            default: break;
        }
        count++;
    }

    if (op == GGML_OP_POOL_AVG) {
        res = (count > 0) ? (res / count) : 0.0f;
    }

    o_ptr[cur_ow] = res;
}
```

It's a simple kernel, and that's fine. Pooling is memory-bound: each output reads at most `k` values and does almost no math, so shared-memory tiling wouldn't buy anything here. The one small thing is `__ldg`, which sends the reads through the read-only data cache.

### The part that actually mattered: matching the CPU

The tricky question wasn't speed. It was **what "average" means at the edges.**

With padding, windows near the start and end of a row hang off the edge. There are two sensible ways to average those:

- divide by `k`, treating the padding as zeros, or
- divide by the number of real elements that fall inside the window.

Different frameworks pick different answers. In ggml there's only one right answer: **whatever the CPU backend does**, because the CPU implementation is the reference every other backend gets tested against. So before writing any CUDA, I read `ggml_compute_forward_pool_1d` in `ggml-cpu/ops.cpp`:

```cpp
for (int ki = 0; ki < k; ++ki) {
    const int j = base + ki;
    if (j < 0 || j >= (int) IW) {
        continue;               // padded positions are skipped...
    }
    ...
    ++count;                    // ...and not counted
}
...
case GGML_OP_POOL_AVG: res = (count > 0) ? (res / count) : 0.0f; break;
```

The CPU skips padded positions and divides by how many real elements it saw. That's why my kernel clamps the window to `[max(0, start), min(iw, start + kw))` and divides by `count` instead of `kw`. If you get this wrong, max pooling still passes everything, and average pooling fails only on the windows that touch the padding. That's exactly the kind of bug that slips through when you only test the easy shapes.

### Testing with `test-backend-ops`

ggml has a really good tool for this. `test-backend-ops` builds the graph for an op, runs it on the CPU and on the backend you're testing, and compares the outputs.

The POOL_1D cases come from a nested loop in `tests/test-backend-ops.cpp`:

- 2 pooling modes (avg, max)
- × 3 kernel sizes (1, 2, 3)
- × 3 strides (1, 2, 3)
- × 4 paddings (0, 1, 2, 3)
- × 3 input shapes (`[10,3,2,1]`, `[11,1,3,2]`, `[128,2,1,3]`)

That's **216 cases**, and they're well chosen: an odd width of 11, padding larger than the kernel, strides that skip elements entirely. I work on a Mac with no NVIDIA GPU, so I ran everything on Kaggle's free 2× T4 instances:

```
./build/bin/test-backend-ops -o POOL_1D
```

216 out of 216 passed.

### The review

This is the bit people usually skip when they write about open source, and it's where I learned the most.

- **A bot stopped me first.** Right after I opened the PR, llama.cpp's bot pointed out that I had more than one PR open, and new contributors only get one at a time. So I turned it into a draft. Big projects guard their reviewers' time, and after seeing how many PRs come in, I get why.
- **The only change I was asked for was a newline.** Aman Gupta (am17an) reviewed it, approved it, and pointed me to a failing CI job: my two new files didn't end with a newline, which fails the repo's editorconfig check. One tiny commit fixed it.
- **Not every red check is yours.** Two WebGPU checks failed on code I hadn't touched. I said that in the PR, so the reviewer didn't have to go figure it out.
- am17an moved it out of draft and added the "merge ready" label, and **Georgi Gerganov merged it on August 23**, less than a day after I opened it.

The kernel took the most thinking. Getting it merged mostly came down to following the process and making the reviewer's life easy.

## PR 2: the silent CPU fallback ([#28897](https://github.com/ggml-org/llama.cpp/pull/28897))

My second PR started from a 🟡 in that same ops table. CUDA's `DUP` (copy a tensor, possibly into a different layout) was only marked as partially supported, and I wanted to know which cases were missing. The ops table is generated from `docs/ops/CUDA.csv`, which lists every single test case and whether CUDA supports it, so I searched it for `DUP`. Every row was `yes` except the ones with `type=i32` and `type=i16`. That told me exactly where to look: whatever was saying no had to be checking the type.

Sure enough, in `ggml_backend_cuda_device_supports_op`:

```cpp
case GGML_OP_DUP:
    {
        ggml_type src0_type = op->src[0]->type;
        return src0_type != GGML_TYPE_I32 && src0_type != GGML_TYPE_I16;
    } break;
```

So CUDA refused `DUP` for 32-bit and 16-bit integer tensors, and those quietly ran on the CPU. Then I grepped `cpy.cu` for `int32_t` expecting to write a new kernel, and found something I didn't expect: **an i32 → i32 copy path already existed.** It just never got used for `DUP`, because the gate said no. Only i16 → i16 was actually missing.

The fix had two parts. First, add the i16 branch to `ggml_cuda_cpy`, mirroring the i32 one right above it (including the transposed-copy variant):

```cpp
} else if (src0->type == GGML_TYPE_I16 && src1->type == GGML_TYPE_I16) {
    if (can_be_transposed) {
        ggml_cpy_scalar_cuda<int16_t, int16_t, true>
            (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
             ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
    } else {
        ggml_cpy_scalar_cuda<int16_t, int16_t>
            (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
             ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
    }
}
```

Second, take the gate out:

```cpp
case GGML_OP_DUP:
    return true;
```

Eight lines added in `cpy.cu`, and a four-line condition replaced with one.

### A small change still needs real testing

A tiny diff doesn't mean tiny risk. Copies are used all over the graph, so I didn't stop at the op's own tests:

- `test-backend-ops -o DUP`: **10/10** cases passed (f32, f16, i32 and i16, contiguous and permuted).
- The **full `test-backend-ops` suite: 16,097/16,097 passed, no regressions**, on 2× Tesla T4 (compute capability 7.5).

### The docs I chose not to regenerate

That CSV is generated from test runs, so the normal thing is to regenerate it along with `docs/ops.md`.

I didn't, and I explained why in the PR. A T4 is compute capability 7.5. Regenerating on it would have marked ops that need newer GPUs as unsupported, **quietly downgrading rows that had nothing to do with my change**. So I edited just the four `DUP` rows for i32 and i16 from `no` to `yes`, and changed the `DUP` cell in `ops.md` from 🟡 to ✅.

It's a small call, but it's the kind reviewers notice: don't touch what you can't check.

### The review

Georgi Gerganov approved it on September 14. Aman Gupta approved it too and **merged it on September 15**.

## What I took away from these two

1. **A silent fallback is a real bug, even when nothing breaks.** If `supports_op` says no, your op runs on the CPU and nobody finds out unless they go looking. If something on the GPU is slower than it should be, check the gate before you start profiling the kernel.
2. **The reference implementation is the spec.** For POOL_1D, "average" meant exactly what the CPU does at the padded edges. Read the reference before you write the kernel.
3. **If you can enumerate the edge cases, do it.** 216 cases sounds like a lot until you realise the bugs that matter live in maybe 20 of them.
4. **Run the whole suite, not just your op.** 16,097 tests for a 9-line change sounds like overkill. It isn't, when the thing you touched is used everywhere.
5. **Make the reviewer's job easy.** Give the exact test command, the GPU, the result, and anything you deliberately didn't do. Respect the process (drafts, CI rules). People maintaining a project this size review a lot of PRs, and the clear ones get merged.
6. **Small is fine.** Neither PR is impressive on its own. Both are merged code in something a huge number of people rely on, and each one taught me a part of the codebase I didn't understand before.

## What's next

I'm now working on FP8 in vLLM. Right now, block-FP8 models on Ada GPUs (L4, L40S, RTX 4090) fall back to a slower weight-only path, because vLLM's CUTLASS blockwise kernels need hardware that only exists on Hopper and newer. The issue is [vllm#58241](https://github.com/vllm-project/vllm/issues/58241), and I'll write it up once I have real numbers.

If you want to make your first llama.cpp contribution, start with [issue #14909](https://github.com/ggml-org/llama.cpp/issues/14909) and the ops table. Pick an ❌, find the closest op that's already implemented, read the CPU reference, and run `test-backend-ops`.

**Links**
- PR #27573, CUDA POOL_1D: https://github.com/ggml-org/llama.cpp/pull/27573
- PR #28897, CUDA DUP for i16/i32: https://github.com/ggml-org/llama.cpp/pull/28897
- My GitHub: https://github.com/amankarki151