Valid kernel JSON is not a repaired kernel

The interface worked. The source did not change.

The three model generations were actual GPU runs; the byte, JSON and provenance checks were performed locally on retained records.^[2]

We made three development calls to Qwen2.5-Coder-7B-Instruct on one deliberately modified RMSNorm kernel. The first response failed the JSON interface. The second ended inside a source string. The third produced a valid submission after JSON repacking and copied the input CUDA source exactly.^[1]

A carrier packages the CUDA source and launch descriptor in JSON. Accepting that wrapper is useful evidence about the interface. It is not a repaired kernel. Here the launch descriptor was fixed, so the unchanged source also preserved the two authored defects. We did not compile or execute any generated CUDA in these three attempts.^[1]

Keep every attempt visible

Attempt 1 used the initial prompt and a configured 60-second soft generation limit. The model returned 602 tokens, ending with an end-of-sequence token (EOS). Its response was wrapped in Markdown fences, so strict JSON parsing rejected it. It also used the wrong outer fields. Natural completion did not make it a valid submission.^[1]

For attempt 2, we named the exact protocol, source and launch fields, showed the complete input carrier, and requested a full translation unit with its includes, extern-C declaration and notices. The prompt grew from 1,670 to 1,953 tokens. The model returned 765 tokens without EOS and stopped inside the source string. The recorded reason is soft_time_or_other_stop: the exact stopping criterion was not captured.^[1]

Attempt 3 kept that prompt and raised the soft generation limit from 60 to 180 seconds. It returned 1,190 tokens, ending with EOS. The unchanged extractor accepted the object and repacked its JSON serialization. It did not strip fences, repair keys or edit CUDA. The resulting carrier and decoded source were byte-identical to the input.^[1]

What stayed fixed and what changed

All three calls used the same pinned Qwen model and four weight shards, BF16 weights, eager attention, greedy single-sequence decoding and a 4,096-new-token limit. The observed software was Transformers 4.57.1, Tokenizers 0.22.1 and Torch 2.9.0+cu130 on an NVIDIA L4. Later calls reused the initial runtime image; this does not establish the same physical GPU.^[1][6]

The source fixture, mathematical contract, fixed launch and strict extractor stayed unchanged. The schema instructions changed after attempt 1. The soft time limit and generation-timer instrumentation changed after attempt 2. These are three adaptive development attempts, not three repetitions of a frozen test or evidence that one prompt change caused the result.^[1]

The fixture came from a FlashInfer RMSNorm adaptation. We deliberately removed a block barrier and replaced the normalization divisor H with blockDim.x. The prompt described those authored changes; it did not contain authentic compiler or runtime feedback. This exposed development example is excluded from training and held-out confirmation claims.^[2][5]

Keep three checks separate

First, validate the carrier: is there exactly one JSON object with the required fields, protocol and allowed launch? Preserve the raw response and record any permitted serialization change. A lossless JSON repack should never silently become a source edit.^[1]

Second, compare the proposed change. For this fixed-launch source-repair task, compare the decoded source with the input. A matching hash makes the no-change result explicit. In tasks that allow launch changes, inspect those fields too; a source hash alone would miss that kind of repair.

Third, judge the program independently. Compilation, numerical correctness, memory safety and synchronization need their own evidence. A changed source file does not pass those checks merely because it is different. This calibration stopped before that stage.

What this result can support

The observed result is narrow: two format rejections, then one valid carrier with no source change. It does not estimate the model's repair rate, demonstrate training improvement or establish kernel correctness. No throughput, provider comparison or cost-per-result claim follows from it.^[1]

The downloadable aggregate records all three attempts, exact model and software pins, prompt and decoder changes, and source identities. The provenance record credits FlashInfer and Qwen. The aggregate summarizes the evidence; the separate archive retains the selected source, prompts and responses. These are checks of retained records and source identities, not independent re-execution.^[1][2]

Download the saved-byte package and run python3 -I -B verify_saved_bytes.py using only the Python standard library. The release check returned two JSON rejections and one valid carrier with unchanged source. It checks the exact selected file hashes and saved JSON/source relationships. It does not rerun Qwen, rerun tokenizer decoding, compile CUDA or grade correctness.^[3][4]

For a kernel-data evaluation, make format validity, the proposed change and independently judged correctness separate fields. That gives a lab a clear account of what each record actually demonstrates.