COLIDE

The problem

Intrusion detection on IoT traffic wants two things that pull against each other: a model small enough to run on a constrained GPU at batch size one, and an explanation a human operator can act on.

The second requirement is usually solved by calling a cloud API, which is exactly what you cannot do on an air-gapped sensor. The first is usually solved by handing the model to TensorRT or torch.compile and trusting the framework — which turns out to work poorly for this shape of model. Production inference stacks are tuned for large graphs and large batches. A sub-1M-parameter model with a bidirectional LSTM in the middle of it is close to their worst case: dynamic recurrent control flow, tiny tensors, and per-launch overhead that dominates the actual arithmetic.

COLIDE asks what you get back by writing the forward path yourself, and it tries to answer that question in a way that survives review.

What was built

• A CNN-BiLSTM classifier (~530K parameters) trained on BoT-IoT and ToN-IoT, with knowledge distillation from a classical tree-ensemble teacher and focal loss for the minority attack classes.
• Four fused CUDA C++ kernels replacing the PyTorch operators for the projection/conv block, the second conv block, the BiLSTM, and the dense head. Each kernel is standalone — no build system, it compiles with a single nvcc invocation and benchmarks itself.
• A statistical benchmark harness that runs each configuration as independent subprocess trials and records the distribution, not a point estimate — because a single run of this workload is not reproducible enough to publish.
• An asynchronous alert-dispatch prototype that hands detections to a local quantized LLM (TinyLlama 1.1B, Q4, ~0.77 GB) through a bounded queue, so generation never blocks detection.
• A claim verifier. scripts/verify_claims.py walks every bold number in the repository’s documentation and re-derives it from the JSON artifact it came from; check_stale_claims.py fails the build if a withdrawn figure reappears anywhere. That tooling exists because a number was withdrawn once already, and I wanted the repo to catch the next one instead of me.

What is measured

Detection, on BoT-IoT, under a sealed multi-seed protocol: macro-F1 0.9780 ± 0.0033 across five seeds. That is competitive with the classical baselines rather than better than all of them — protocol LightGBM still leads on pure F1 at 0.9818. The argument is accuracy per unit of latency, not accuracy supremacy, and the write-up says so.

Per-block, on an RTX 3050, matched operator against matched operator: the custom kernels win clearly on Blocks 1, 2 and 4 (3.24×–6.55×). Block 3 — the BiLSTM — is the hard one, and its optimisation history is the interesting part: naive → precomputed input projection → transposed recurrent weights for coalesced access → CUDA Graphs → FP16 half2 gate packing, cumulatively 7.55×–9.50× over the naive kernel, and 1.30×–1.47× over cuDNN across five independent sessions on that laptop.

The full-model framework latencies were measured too, as absolute multi-session ranges on the same box: ORT CPU 487–699 µs, torch.compile 1,519–1,777 µs, eager PyTorch 2,050–2,247 µs, TensorRT FP16 2,427–2,966 µs, ORT GPU 3,862–4,652 µs. Two things in that list are worth sitting with. TensorRT FP16 is slower than eager PyTorch here, on this model size, on this hardware. And ORT on the CPU beats several GPU paths, because the model is small enough that launch overhead dominates.

Cross-hardware runs on V100S and A100 nodes at a university computing centre give a separate full-model batch-1 matrix, where TensorRT native is the fastest GPU path (528 µs V100S / 588 µs A100).

Alert dispatch costs 16.60 µs at p99. Generation itself takes roughly 8.5 seconds per alert and happens entirely off the detection path — that separation is the whole point of the design, and conflating the two numbers would be the easiest way to make this project look better than it is.

What is not claimed

This section is load-bearing, not modesty.

• There is no full-pipeline speedup number. The custom CUDA path covers fused Blocks 1–4 only. The full model also has attention, LayerNorm, residual connections, temporal pooling and the classifier, none of which are in the kernel chain. So the partial pipeline sum (594–675 µs) is reported as an absolute latency with its scope attached, and is never divided by a full-model framework latency. “Custom CUDA is N× faster than TensorRT” would be an easy sentence to write and it would be false.
• A negative result stands unretracted. On the cluster GPUs, matching PyTorch Block 3 measured faster than the CUDA Block 3 (~363 vs ~513 µs on V100S; ~385–391 vs ~667–671 µs on A100). Those are historical wall-clock numbers from before a kernel fix, and they have not been re-run on the server since. Rather than quietly dropping the comparison, the active claim set treats server Block 3 as unresolved until a fresh run exists.
• The recurrent architecture is not defended as necessary. An MLP ablation reaches competitive macro-F1 on this feature set. The BiLSTM is kept because it is a genuinely hard case for compilers and kernels — which is what the project is actually studying — not because the data demanded it.
• The LLM explanations are not validated. The dispatch mechanism is measured; free-form output quality is weak under the explainability checks, and the work is not titled as validated LLM-based explainability.
• An earlier ToN-IoT result was withdrawn. A 26-feature “clean” run reported much stronger numbers, and was invalid: the loader kept a target-derived column in the feature matrix, fit encoders before the split, and applied SMOTE to integer-encoded categoricals. The corrected leakage-safe rerun is weaker and honest — CNN 0.8075 against a same-split random forest at 0.9626 — and the old figures are tombstoned in the repository rather than deleted.

What measuring this taught me

Two findings changed how I benchmark anything now.

The first: a kernel that failed numerical validation about six times in thirty was blamed, in my own notes, on floating-point accumulation order. That explanation was wrong, and it was checkably wrong — accumulation order is deterministic, and the same seeded input through the same binary was giving different output every run. compute-sanitizer --tool racecheck found a real shared-memory race between a timestep’s hidden-state write and the next timestep’s read of it, across an intervening __syncthreads(). Double-buffering the hidden state fixed it: zero hazards under racecheck, 100/100 runs passing, and 20/20 passing at a tolerance three orders of magnitude tighter than the one it had been scraping by on.

The second: re-running the same benchmark suite twice in one sitting moved framework latencies by 14–17%, and moved individual kernel configurations by up to 27% — while each session’s own coefficient of variation looked perfectly tight. Within-session variance understates the real uncertainty on a WSL2 development box. Every latency figure above is therefore a range across independent sessions, and single-run tables are treated as unpublishable.

Stack

CUDA C++ · PyTorch · knowledge distillation · TensorRT / ONNX Runtime (baselines) · TinyLlama 1.1B Q4 · BoT-IoT / ToN-IoT · compute-sanitizer · a claim-verification harness