KNOWN_ISSUES.md

# Known limitations — tritllm-kernel

Items previously raised in code review have been addressed:

- The implicit one-warp-per-block launch contract in the educational kernels
  is now an early-return guard: kernels return without writing if launched
  with `blockDim.x != 32` or `in_features % 64 != 0`.
- The dead `trit_pipeline` / `k_v29_pipeline` path was removed.
- The C API now validates pointers, dimensions, and the
  `cols / GROUP_SIZE == num_groups` invariant, and reports the result via
  `trit_gemv_get_last_error()`. CUDA launch errors are captured into the same
  channel.
- `get_gpu_name(buf, buflen)` now refuses null pointers and `buflen <= 0`.

This document lists what remains.

## Design tradeoff (not a bug)

### Lane-0 scale-and-add in `trit_gemv_uniform` / `trit_gemv_variable`
**Where:** [`trit_gemv.cu:223-232, 279-286`](trit_gemv.cu#L223)

After the warp reduction in the educational kernels, only lane 0 multiplies
the group sum by the scale and accumulates into `row_acc`. The other 31 lanes
are idle for the scale/add path. This is correct, just slow — the published
paper benchmarks are produced by the deferred-reduction kernel
`k_d3_hardened` in `trit_gemv_standalone.cu`, which does not have this
limitation.

The `trit_gemv_uniform` / `trit_gemv_variable` kernels in `trit_gemv.cu` are
kept as a smaller, single-file reference implementation that is easier to read
and reason about. If you need maximum throughput, use the C API in
`trit_gemv_standalone.cu`.

## Future cleanup

The C API in `trit_gemv_standalone.cu` exposes several historical kernel
variants (`v9`, `v27`, `v28`, `v29`, plus `k_d3_hardened` via
`trit_gemv_d3_int8_dp4a`). They all work, but the public API is wider than
needed. A future release will trim to one canonical entry point per depth
(`trit_gemv_d1`, `trit_gemv_d2`, `trit_gemv_d3`, `trit_gemv_d4`).