diff --git a/tests/reviewed_contract_history.py b/tests/reviewed_contract_history.py index 5443a2793..315e54a99 100644 --- a/tests/reviewed_contract_history.py +++ b/tests/reviewed_contract_history.py @@ -112,76 +112,23 @@ def known_feedforward_geometry(source): # Scope-20 device boundaries are approved contract changes, not permission to # alter task math or guards. Full current/before bytes bound every inverse. -TRITON_DEVICE_BOUNDARY_FILES = {'tasks/triton2triton/vllm/triton_apply_grammar_bitmask/README.md': {'after_sha256': '802aec8ebf39be022fded4abf8b88815b251ff2374dd12eec071deb2e9f7df4b', - 'before_sha256': '38dfc5fb94d122c45228aa5803cb1953e7b1dff899a21c1cf94d4df8cb638218', - 'reverse_edits': [(56, - 70, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `_apply_grammar_bitmask_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, - 'tasks/triton2triton/vllm/triton_apply_grammar_bitmask/config.yaml': {'after_sha256': 'fb54d74399b7e25da2e50e4ad8a7f6a95d0cbc8f0ea76d61b7b6b10567d5028e', - 'before_sha256': '6b72732ff602cbe75aeb064f4979a3fc7f0c13fa0bdbe7acb8fd04d9d2a85a26', - 'reverse_edits': [(6, - 11, - [' - path: ' - 'source/triton_apply_grammar_bitmask.py\n', - ' scope: ' - 'symbols\n', - ' symbols:\n', - ' - ' - '_apply_grammar_bitmask_kernel\n', - ' ' - 'allow_new_helpers: ' - 'true\n'], - [' - ' - 'source/triton_apply_grammar_bitmask.py\n'])]}, - 'tasks/triton2triton/vllm/triton_awq_gemm/README.md': {'after_sha256': '0be43d97152f9a3dacb1808b2033f852277cae1d6c528488e76736d6ec15d5f7', - 'before_sha256': '9522b039a0b59837f313888ef7470015eef03c6aca0a5b8ddf2ad94b86d1cc7e', - 'reverse_edits': [(51, - 65, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `awq_gemm_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, +# Only contract files are pinned; README text is documentation and may change. +TRITON_DEVICE_BOUNDARY_FILES = {'tasks/triton2triton/vllm/triton_apply_grammar_bitmask/config.yaml': {'after_sha256': 'fb54d74399b7e25da2e50e4ad8a7f6a95d0cbc8f0ea76d61b7b6b10567d5028e', + 'before_sha256': '6b72732ff602cbe75aeb064f4979a3fc7f0c13fa0bdbe7acb8fd04d9d2a85a26', + 'reverse_edits': [(6, + 11, + [' - path: ' + 'source/triton_apply_grammar_bitmask.py\n', + ' scope: ' + 'symbols\n', + ' symbols:\n', + ' - ' + '_apply_grammar_bitmask_kernel\n', + ' ' + 'allow_new_helpers: ' + 'true\n'], + [' - ' + 'source/triton_apply_grammar_bitmask.py\n'])]}, 'tasks/triton2triton/vllm/triton_awq_gemm/config.yaml': {'after_sha256': '0de7eb1b20549c759c057d3cf3d34b664afd9337036cf3224dc0f4f58ad1819f', 'before_sha256': 'cdba9bac6e646921796eb59af68d9668af0447d58d0f8dc76a9c133331098fce', 'reverse_edits': [(6, @@ -195,86 +142,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_awq_gemm.py\n'])]}, - 'tasks/triton2triton/vllm/triton_bmm/README.md': {'after_sha256': 'c375293e91620c08c65ba24362db407cd12df59a11f7aa2db4d82c74048cb742', - 'before_sha256': '688c9a2650e4e2ec1b3cf9a2e1a72c79c65ed8fe870e5e1cea5de9d4b5dceaa2', - 'reverse_edits': [(41, - 90, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `bmm_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n', - '\n', - '\n', - '### Correctness-only domain controls\n', - '\n', - 'The protected `bmm_triton` wrapper defines the public domain. It accepts ' - 'two\n', - '3-D tensors with matching batch and inner dimensions and one shared dtype:\n', - 'float16, bfloat16 or float32, the only keys of its launch-configuration ' - 'table.\n', - 'It forwards all six operand strides to the kernel without copying, so\n', - 'transposed, offset, padded, element-strided and batch-broadcast (stride-0)\n', - 'views are valid inputs. It always allocates a new contiguous `(B, M, N)` ' - 'output\n', - 'of the input dtype on `a.device`; both operands must be on that GPU. ' - 'Zero-sized\n', - 'dimensions and the int64 path for tensors over 2**31 elements are not ' - 'exercised.\n', - '\n', - 'Eleven unscored `control-domain-*` rows (`case_index` 20000-20010, checks\n', - '`[correctness]`) come from the protected `_domain_controls.py` table; the\n', - 'adapter verifies that table and these rows against `workloads.json`. They ' - 'cover\n', - 'bfloat16 and float32 as well as float16, and B, M, N and K values that are ' - 'not\n', - 'multiples of the wrapper tiles (128/128/32 for 16-bit dtypes, 64/64/32 for\n', - 'float32), including `M = K = 1` and other small odd shapes. Operand ' - 'layouts\n', - 'include transposed A, B or both, a batch-inner `(M, B, K).transpose(0, 1)` ' - 'view,\n', - 'offset views of padded storage, element stride 2, batch stride 2 and a\n', - 'batch-broadcast B. Padding and stride gaps hold a sentinel value, so a ' - 'read\n', - 'outside a view changes the result.\n', - '\n', - 'Each control draws seeded N(0, 1) inputs on the CPU. It computes an FP64 ' - 'CPU\n', - '`torch.bmm` reference from pristine clones before calling the candidate. ' - 'The\n', - 'input views and their complete backing storage must be unchanged. The ' - 'output\n', - 'must have the exact shape, dtype and device, contiguous `(M*N, N, 1)` ' - 'strides\n', - 'and finite values, and match within `atol=rtol=1e-2` for every dtype. For\n', - 'float16 this is the original gate. One bfloat16 ulp is at most 0.78% of a\n', - 'value. Float32 keeps the bound that protected `_arena_checks.py` already ' - 'applies\n', - 'to every call, including these controls; layout, tail and dtype defects ' - 'give\n', - 'much larger errors at these magnitudes. The five scored float16 cases, ' - 'their\n', - 'seeds, tolerances, warmups, samples and timing are unchanged, and no ' - 'control\n', - 'is timed.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_bmm/config.yaml': {'after_sha256': '6e04236bbc34f35a60c1dcaaa47f218ac2e42d75a4cf85778240feaa8088462e', 'before_sha256': 'c7315e3dbb45c3e4e40f2d713e67b8071a14fd1c77d1f600bc047174b21d48c2', 'reverse_edits': [(6, @@ -285,40 +152,6 @@ def known_feedforward_geometry(source): ' - bmm_kernel\n', ' allow_new_helpers: true\n'], [' - source/triton_bmm.py\n'])]}, - 'tasks/triton2triton/vllm/triton_causal_conv1d_update/README.md': {'after_sha256': 'e01416edaebd6602a48872b9fbfdaea252a393896daed2890ed5953ba1f802f1', - 'before_sha256': 'c9bfde244528d690541c3e7027e283eaeafcc480f911cb6ed815cc88f1b96c42', - 'reverse_edits': [(54, - 63, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`_causal_conv1d_update_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch and ' - 'pointer preparation ' - 'remain protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and ' - 'benchmark protocol. ' - 'The structural\n', - 'guard enforces this ' - 'boundary; it does ' - 'not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_causal_conv1d_update/config.yaml': {'after_sha256': '7baf927cf33928b34680f96aed96049cb846b5af01d35a60f1bc691a577eddb5', 'before_sha256': '104ccd038ca4820dfa7cffc67f4f24875c8e3bfd8a6d2ebc882baf6ac1022b1d', 'reverse_edits': [(6, @@ -335,60 +168,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_causal_conv1d_update.py\n'])]}, - 'tasks/triton2triton/vllm/triton_chunked_prefill_paged_decode/README.md': {'after_sha256': 'f9680f7f601eba320655fd3874d3e340b1ac8ee45c057f8b8e6430f14a4fc071', - 'before_sha256': 'ee4c03b2a1cf9116985bb9f14bc12e978bf1c09cf841cd07be68e82fa3d33b0f', - 'reverse_edits': [(82, - 91, - ['\n', - '\n', - '### ' - 'Editable ' - 'device ' - 'implementation ' - 'boundary\n', - '\n', - 'The ' - 'editable ' - 'device ' - 'functions ' - 'are ' - '`cdiv_fn`, ' - '`kernel_paged_attention_2d`.\n', - 'Their ' - 'bodies and ' - 'additional ' - 'implementation ' - 'helpers may ' - 'be ' - 'optimized. ' - 'Module\n', - 'imports, ' - 'initialization, ' - 'host ' - 'dispatch ' - 'and pointer ' - 'preparation ' - 'remain ' - 'protected,\n', - 'along with ' - 'the ' - 'existing ' - 'numerical ' - 'checks and ' - 'benchmark ' - 'protocol. ' - 'The ' - 'structural\n', - 'guard ' - 'enforces ' - 'this ' - 'boundary; ' - 'it does not ' - 'isolate ' - 'arbitrary ' - 'Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_chunked_prefill_paged_decode/config.yaml': {'after_sha256': 'acba87fb1d4ca5620a8560a0751a8dd4650df4dc353b4b6f5798c66d049f9fa7', 'before_sha256': '9c2299ba86847d816573c626c13821ac00d26ed527a83001ef2491b04a5e8865', 'reverse_edits': [(6, @@ -409,40 +188,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_chunked_prefill_paged_decode.py\n'])]}, - 'tasks/triton2triton/vllm/triton_correct_attn_cp_out/README.md': {'after_sha256': 'dc3283ad5b4a43e8c46c26a7a20f6123854ed2f68cc8dcec88ab0fb97537757f', - 'before_sha256': 'b620b666d73be6b5346a22849ebf56639eb411b8153fff009b4c607832354232', - 'reverse_edits': [(54, - 63, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`_correct_attn_cp_out_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. Module\n', - 'imports, ' - 'initialization, host ' - 'dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and benchmark ' - 'protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does ' - 'not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_correct_attn_cp_out/config.yaml': {'after_sha256': '66ac82fce41b23f5d0aef546168b6523996140bf820c7b3812424b01873e4994', 'before_sha256': 'd734d0a7b8f9ddf17111dec1ee435e81fc68865079a8d524099b7486e9b1914f', 'reverse_edits': [(6, @@ -459,52 +204,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_correct_attn_cp_out.py\n'])]}, - 'tasks/triton2triton/vllm/triton_decode_attn_grouped_stage1/README.md': {'after_sha256': '0f2a551ce51340a3709600bc78d91b0acc2d2205c10756e4dfe45d5b6d380d7d', - 'before_sha256': 'ada3cf20f1eda4171fc77a6249f9aedd11890a2741fb8967134f0ca52dd9afc6', - 'reverse_edits': [(89, - 98, - ['\n', - '\n', - '### Editable ' - 'device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable ' - 'device ' - 'functions are ' - '`tanh`, ' - '`_fwd_grouped_kernel_stage1`.\n', - 'Their bodies ' - 'and ' - 'additional ' - 'implementation ' - 'helpers may ' - 'be optimized. ' - 'Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch ' - 'and pointer ' - 'preparation ' - 'remain ' - 'protected,\n', - 'along with ' - 'the existing ' - 'numerical ' - 'checks and ' - 'benchmark ' - 'protocol. The ' - 'structural\n', - 'guard ' - 'enforces this ' - 'boundary; it ' - 'does not ' - 'isolate ' - 'arbitrary ' - 'Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_decode_attn_grouped_stage1/config.yaml': {'after_sha256': 'd5abee78e0d739495acc39027a22d30462c42f37c1d86332403476af0c83fdc6', 'before_sha256': '24c98bbcaff663b18709845a25e03d1608cf50e6068e65caacace55e37e47dbd', 'reverse_edits': [(6, @@ -524,178 +223,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/_fwd_grouped_kernel_stage1.py\n'])]}, - 'tasks/triton2triton/vllm/triton_decode_attn_stage1/README.md': {'after_sha256': 'bf9984dd87ef182778e4063ee626cb1d740a41654593909bd2d8ba74b0428792', - 'before_sha256': '9e164b707e474647d6d64bd32612a5961b9cd4a405982c671cb320e6195f7578', - 'reverse_edits': [(28, - 29, - ['- Must handle the input domain stated under "Supported input domain" ' - 'below\n'], - ['- Must handle arbitrary sequence lengths, head dimensions, and page ' - 'sizes\n']), - (62, - 63, - ['The five original scored cases remain unchanged. The three ' - '`contract_case`\n'], - ['The five original scored cases remain unchanged. Additional ' - 'correctness-only\n']), - (88, - 187, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `tanh`, `_fwd_kernel_stage1`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n', - '\n', - '### Supported input domain\n', - '\n', - 'The protected wrapper launches one program per (request, query head, split), ' - 'uses a\n', - 'fixed `BLOCK_N = 64` token tile, and pads the K and V head dimensions to the ' - 'next\n', - 'power of two with masking. Within the limits below, a candidate must produce ' - 'the\n', - 'correct partial output and log-sum-exp for every active split:\n', - '\n', - '- Sequence lengths: at least 1 token per request, chosen independently per ' - 'request.\n', - ' Lengths need not be multiples of the page size or of `BLOCK_N`. Request ' - '`b` uses\n', - ' row `b` of `Req_to_tokens` and `B_Seqlen[b]`.\n', - "- Page size: any positive integer. Entry `p` of a request's `Req_to_tokens` " - 'row is\n', - " the physical page that holds that request's tokens `p * page_size` " - 'through\n', - ' `(p + 1) * page_size - 1`. The row needs at least `ceil(L / page_size)` ' - 'entries,\n', - ' and each referenced entry must be a valid page of `k_buffer`/`v_buffer`. ' - 'Entries\n', - " after a request's last logical page, unreferenced pages, and slots past " - 'a\n', - " request's length in its final page must not affect the result.\n", - '- Splits: any `num_kv_splits >= 1`. With `C = ceil(L / num_kv_splits)`, ' - 'split `s`\n', - ' covers tokens `[s * C, min(L, (s + 1) * C))`. A split with no tokens is ' - 'inactive:\n', - " the kernel does not write its `att_out` slot, and the caller's bytes there " - 'must\n', - ' remain unchanged.\n', - '- Head dimensions: the Q/K head dimension `Lk` and the V head dimension `Lv` ' - 'are\n', - ' independent integers from 16 through 256. They need not be powers of two ' - 'or\n', - ' equal. `att_out` has shape `[batch, num_heads, num_kv_splits, Lv + 1]`. ' - 'Head\n', - ' dimensions outside 16-256 are outside this contract.\n', - '- Heads: `num_heads` must be a multiple of `num_kv_heads`. Query head `h` ' - 'reads KV\n', - ' head `h // (num_heads // num_kv_heads)`.\n', - '- Data: `q`, `k_buffer` and `v_buffer` share one dtype, float16 or ' - 'bfloat16.\n', - ' `att_out` is float32, and `Req_to_tokens` and `B_Seqlen` are int32. A\n', - ' `logit_cap` of 0 disables capping; a positive value applies\n', - ' `logit_cap * tanh(logits / logit_cap)` before softmax.\n', - '- Layout: the last dimension of `q`, `k_buffer`, `v_buffer`, `att_out` and\n', - ' `Req_to_tokens` has stride 1, and `B_Seqlen` is contiguous. Each tensor ' - 'must\n', - ' hold fewer than 2^31 elements, because the kernel computes element offsets ' - 'in\n', - ' 32-bit integers.\n', - '\n', - '### Domain controls\n', - '\n', - 'Seven correctness-only `control-domain-*` rows (case indices 20000-20006) ' - 'sample\n', - 'this domain beyond the scored cases and earlier controls. The protected\n', - '`_domain_controls.py` holds their generator and oracle. `_arena_eval.py` ' - 'checks its\n', - "`CASES` table against the manifest's `domain_controls` list and the exact " - 'rows,\n', - 'and `scripts/task_runner.py` dispatches case indices of 20000 and above to ' - 'it.\n', - '\n', - '| Control | Lk / Lv | Page size | Splits | Lengths | Q / KV heads | Dtype, ' - 'cap |\n', - '|---|---|---|---|---|---|---|\n', - '| `page1_hd16_short_inactive` | 16 / 16 | 1 | 4 | 1, 2, 5, 70 | 4 / 4 | ' - 'fp16, 0 |\n', - '| `page3_hd96_group3` | 96 / 96 | 3 | 3 | 100, 7, 64 | 6 / 2 | fp16, 0 |\n', - '| `page48_hd32_mqa_one_split_capped` | 32 / 32 | 48 | 1 | 129, 47, 48 | 8 / ' - '1 | fp16, 1.0 |\n', - '| `page128_k192_v128` | 192 / 128 | 128 | 2 | 300, 129 | 4 / 2 | fp16, 0 |\n', - '| `page64_k80_v256` | 80 / 256 | 64 | 4 | 65, 191 | 2 / 1 | fp16, 0 |\n', - '| `long_page16_splits16` | 128 / 128 | 16 | 16 | 7999, 4099 | 8 / 2 | fp16, ' - '0 |\n', - '| `bf16_page5_hd112_splits7_capped` | 112 / 112 | 5 | 7 | 1, 33, 257 | 4 / 1 ' - '| bf16, 2.5 |\n', - '\n', - 'Together with the scored cases and earlier controls, the checked points are ' - 'head\n', - 'dimensions 16, 32, 64, 80, 96, 112, 128, 192 and 256 (with `Lk` both larger ' - 'and\n', - 'smaller than `Lv`), page sizes 1, 3, 5, 8, 16, 32, 48, 64 and 128, split ' - 'counts 1,\n', - '2, 3, 4, 5, 7, 8 and 16, and lengths from 1 to 7999 tokens.\n', - '\n', - 'Each domain control works as follows:\n', - '\n', - '- Pages: a seeded random permutation assigns physical pages. Table padding ' - 'entries\n', - ' point to unused decoy pages. Every K/V slot that no active position ' - 'references\n', - ' holds V = 32, so attending to such a slot changes the result.\n', - '- Oracle: before each call, an independent FP64 CPU oracle is computed ' - 'from\n', - ' pristine clones. The call also goes through the installed correctness ' - 'guard\n', - ' used by the other controls: the Triton dispatch audit and its FP32 ' - 'reference\n', - " (the original harness reference without a cap, the guard's capped " - 'reference\n', - ' otherwise).\n', - '- Outputs: every active output element is poisoned with NaN and every ' - 'inactive\n', - ' split slot is filled with 23.5. After the call, active elements must be ' - 'finite\n', - ' and match the oracle within atol = rtol = 0.01, and inactive slots must ' - 'be\n', - ' byte-identical.\n', - '- Inputs: any change to the bytes or metadata of `q`, `k_buffer`, ' - '`v_buffer`,\n', - ' `Req_to_tokens` or `B_Seqlen` is rejected, as is any change to the ' - 'shape,\n', - ' stride, dtype, device or storage of `att_out`.\n', - '- Second call: the wrapper runs again on the same tensors, with lengths ' - 'rotated\n', - ' across requests and all data and the page permutation redrawn from the ' - 'second\n', - ' declared seed.\n', - '\n', - 'On CPU, `validate-task` also checks that the FP64 oracle matches the ' - 'original\n', - 'harness reference for uncapped controls, and that it rejects zero output,\n', - 'uncapped logits and shifted page routing. These rows have no graph replay ' - 'or\n', - 'timing. The performance action still runs only `perf1`-`perf5`, and no ' - 'score\n', - 'weight, tolerance, warmup, sample count or timing boundary changes.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_decode_attn_stage1/config.yaml': {'after_sha256': 'c69bd6c3c573230c1358f6fb000a3ea82a56510fd921d5b5209e454fee63f684', 'before_sha256': '793e6e10b6227e67ce8d50ca18f88ca87deb888631af904c5a4998bfea6768ba', 'reverse_edits': [(6, @@ -713,33 +240,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/_fwd_kernel_stage1.py\n'])]}, - 'tasks/triton2triton/vllm/triton_decode_attn_stage2/README.md': {'after_sha256': 'b2ffabd5f50a3f17b8b427c4e2735f3f18de113e22f724ecf85c38ddb8265f1e', - 'before_sha256': '02b7883d55eaffb47865a326df31a2420408a8e395ad35cd4c43da20ea9ef6fc', - 'reverse_edits': [(79, - 93, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `_fwd_kernel_stage2`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_decode_attn_stage2/config.yaml': {'after_sha256': 'b3810ddb5da0e7632f526c54399db2023ac6ae54a25801ff3d13dfa7fd5047c9', 'before_sha256': '7e1e7018cd00382729b75e71be8f2c211c146e957381472ddcae0d61ae7e6866', 'reverse_edits': [(6, @@ -756,33 +256,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/_fwd_kernel_stage2.py\n'])]}, - 'tasks/triton2triton/vllm/triton_ep_gather/README.md': {'after_sha256': 'dc45435441bf951eb8eb501fc6145aef50ecaab21c42a40c896b89b6a7e696da', - 'before_sha256': '7fa777a1203ec1836570002443f93c14973a2362ee0f16652dc24dbf466b6cf7', - 'reverse_edits': [(44, - 58, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `_fwd_kernel_ep_gather`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_ep_gather/config.yaml': {'after_sha256': 'a0d64d529fd939f9367280f7b714b1f7710286a3042b7e5aec6ab55f98c72bab', 'before_sha256': '1c3a174ddf8be7fe280178935745ad80ae0eb4b97c1cc8fe62ceadce6738e771', 'reverse_edits': [(6, @@ -797,35 +270,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_ep_gather.py\n'])]}, - 'tasks/triton2triton/vllm/triton_ep_scatter_1/README.md': {'after_sha256': '1f3f587569cec719ec6df7df8a1a276b9815ddcb7ade33accf5aa38326ea2996', - 'before_sha256': 'b92081f80227ebd2be5086adea4e1da6794e7350ce3028eed4b92a5ff2c25693', - 'reverse_edits': [(44, - 53, - ['\n', - '\n', - '### Editable device ' - 'implementation boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`round_up_128`, ' - '`_fwd_kernel_ep_scatter_1`.\n', - 'Their bodies and additional ' - 'implementation helpers may ' - 'be optimized. Module\n', - 'imports, initialization, ' - 'host dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the existing ' - 'numerical checks and ' - 'benchmark protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does not ' - 'isolate arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_ep_scatter_1/config.yaml': {'after_sha256': 'b938b81659606fdf16fc26341d225ebebf74fa4ca06ad11e0e8560312772f805', 'before_sha256': '0b8c3d2c6db69fa55132fd26e1647845c6fe12669e5ccde2bfecdb6b2f2009f6', 'reverse_edits': [(6, @@ -841,33 +285,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_ep_scatter_1.py\n'])]}, - 'tasks/triton2triton/vllm/triton_expert_kernel/README.md': {'after_sha256': '830635586e7bf1a59134b851eb8034395175cd4113d3d52097e540c1275368da', - 'before_sha256': '14d610a41846b32393589aa59acd137fa83854c884c626e99af86802736ef7ea', - 'reverse_edits': [(47, - 61, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `expert_triton_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_expert_kernel/config.yaml': {'after_sha256': 'c65abb1bc437b26f4a62e68d6aa19b430f4bea874db35ecfc9f2e7aa84ac4ea7', 'before_sha256': '713b77fabbe6bd2775f79b76fc497f06d20093988ddc5438d52aa452c8efee4f', 'reverse_edits': [(6, @@ -882,39 +299,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_expert_kernel.py\n'])]}, - 'tasks/triton2triton/vllm/triton_fla_cumsum_scalar/README.md': {'after_sha256': '6d4aa1f9463abe19a5de42e29635ebf10def1284e5b6992e753965daa6e65055', - 'before_sha256': '768a3276cc87535c61847e5c788ac4aa294e3ed38e2944a93d090ed91f79d12e', - 'reverse_edits': [(45, - 54, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`chunk_local_cumsum_scalar_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation helpers ' - 'may be optimized. ' - 'Module\n', - 'imports, ' - 'initialization, host ' - 'dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and benchmark ' - 'protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does not ' - 'isolate arbitrary ' - 'Python execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_fla_cumsum_scalar/config.yaml': {'after_sha256': 'dd352194470d6bb09e72ba256bdd8ddda5b2ff2337f1a63049824b5d189bea52', 'before_sha256': '70d159272b991575cfebbe5dd0c9691153523882aa3d81295282e5c37b1e9db8', 'reverse_edits': [(6, @@ -930,39 +314,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_fla_cumsum_scalar.py\n'])]}, - 'tasks/triton2triton/vllm/triton_fla_cumsum_vector/README.md': {'after_sha256': 'c4f4a3d533d49e0d32abbdd3fd4f9c1ecc8641371da10c4fe4fdf4b0c351c5fd', - 'before_sha256': 'c0389d8bf0a67ef033edf35a528ff04d7aee3c2a26765fd13a6ea6af94de158e', - 'reverse_edits': [(45, - 54, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`chunk_local_cumsum_vector_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation helpers ' - 'may be optimized. ' - 'Module\n', - 'imports, ' - 'initialization, host ' - 'dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and benchmark ' - 'protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does not ' - 'isolate arbitrary ' - 'Python execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_fla_cumsum_vector/config.yaml': {'after_sha256': 'd1f55ec961bbe5bda4d98131e20fcf95f387ad4650ad4d191f95895fbfa4789b', 'before_sha256': '125e25621fe54e1b123e0cc3fdfcb12ebcf0e16cd262b187d197c2d36849ec2b', 'reverse_edits': [(6, @@ -978,35 +329,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_fla_cumsum_vector.py\n'])]}, - 'tasks/triton2triton/vllm/triton_fla_layernorm/README.md': {'after_sha256': 'd7d057a38a4390a58817653f635aa44a1c7acca0376d2aaa181c4b0fadea3132', - 'before_sha256': 'ee1f1ae727e1a49a156e4043618ad2ae7a3deac2f4cc4c6c663ad4e33b59f8f3', - 'reverse_edits': [(41, - 50, - ['\n', - '\n', - '### Editable device ' - 'implementation boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`layer_norm_fwd_kernel`.\n', - 'Their bodies and ' - 'additional implementation ' - 'helpers may be optimized. ' - 'Module\n', - 'imports, initialization, ' - 'host dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the existing ' - 'numerical checks and ' - 'benchmark protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does not ' - 'isolate arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_fla_layernorm/config.yaml': {'after_sha256': '0d27a2c8ca2b5d1edd587c771a1e93721fe0f3c199d87e9314e45262db373146', 'before_sha256': 'c86e2519f055c0271a9201c3818be96fb11bbd8163cd039e34b1836fdf348871', 'reverse_edits': [(6, @@ -1021,33 +343,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_fla_layernorm.py\n'])]}, - 'tasks/triton2triton/vllm/triton_flash_prefill_attention/README.md': {'after_sha256': '34cd567a7af6a77398f08b7fa1f288498d3a504ae9fe3f134f1e79e53d8b334d', - 'before_sha256': 'cca8c62515eb874e93a9e50f5371a5a50c5edf78f8aff2f14e37bf1b5490ca32', - 'reverse_edits': [(76, - 90, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `_fwd_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_flash_prefill_attention/config.yaml': {'after_sha256': 'dca4c1e0037ce58f17d85d6fd5159ef850130e2c74b79afc42d57d88803995b6', 'before_sha256': '2fbb4c96d9705a85dd42753c9a1638b67212d87639dbc53e36a072b9bee90fff', 'reverse_edits': [(6, @@ -1064,33 +359,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_flash_prefill_attention.py\n'])]}, - 'tasks/triton2triton/vllm/triton_fused_gdn_gating/README.md': {'after_sha256': '07ed66c5642f9bfe2943e38f8826a2d408897dab377713d1774c812ba49ef3bd', - 'before_sha256': '224db382310fe35792d0afcbb60d2c37e32888f7bdb1af02799693a61ed4d4b3', - 'reverse_edits': [(62, - 76, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `fused_gdn_gating_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_fused_gdn_gating/config.yaml': {'after_sha256': '00e08420c025b289d4059a9fb1cd65920dbd0122b017c3f8e9b490d5feb10245', 'before_sha256': '171a150edf22833e18a3e6956e9e403743bff25fb2762cba3d067a244f7641db', 'reverse_edits': [(6, @@ -1106,34 +374,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_fused_gdn_gating.py\n'])]}, - 'tasks/triton2triton/vllm/triton_fused_moe_lora/README.md': {'after_sha256': '77eafcbed9378290dcda73308b2ece126f0d4a67584e671f27ec2fd144931273', - 'before_sha256': '8fec6fffc0da84e702baaa4e258f746a7b796d60749c18a7b19d8515189d1ed2', - 'reverse_edits': [(84, - 98, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `_get_lora_id`, `_get_expert_id`, ' - '`_get_token_offs`, `fused_moe_lora_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_fused_moe_lora/config.yaml': {'after_sha256': '579e981824b47adb3e515709798ec2f419dd79f543f477e15001969911ddbc48', 'before_sha256': 'd55709edbab279c8bfa2e8d2b1de6d34df4ab10d1a717c1b3b23e8d87314349b', 'reverse_edits': [(6, @@ -1151,35 +391,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_fused_moe_lora.py\n'])]}, - 'tasks/triton2triton/vllm/triton_gumbel_sample/README.md': {'after_sha256': '0829c043a8b293dec4049a63a72f5e6f867001e6a09a05d20ae6fb7d7100a9cd', - 'before_sha256': 'd90372720a66a694251a764f8c1ec4d66acf68eee25cf1c325f75e84e3147799', - 'reverse_edits': [(46, - 55, - ['\n', - '\n', - '### Editable device ' - 'implementation boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`_gumbel_sample_kernel`.\n', - 'Their bodies and ' - 'additional implementation ' - 'helpers may be optimized. ' - 'Module\n', - 'imports, initialization, ' - 'host dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the existing ' - 'numerical checks and ' - 'benchmark protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does not ' - 'isolate arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_gumbel_sample/config.yaml': {'after_sha256': '3fac26ea5bb59a212e9d1b19202ee9dadb828e8115e2ae4f85e38f3bba5c6551', 'before_sha256': '719857b33e405fc18855ed590f16daca552bd9ce9ecbfa0d1f75417a5e664c94', 'reverse_edits': [(6, @@ -1194,49 +405,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_gumbel_sample.py\n'])]}, - 'tasks/triton2triton/vllm/triton_lightning_attn_kv_reduce/README.md': {'after_sha256': '33a07c73e464adb2c03e3fbc57050f6182f945da06ba9f56ed520dcec42dedae', - 'before_sha256': 'e8d3dd2ff72794c0797698f506dba6ff3b0112acf0cbcb7bfe98b9378b057839', - 'reverse_edits': [(53, - 62, - ['\n', - '\n', - '### Editable ' - 'device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable ' - 'device ' - 'functions are ' - '`_fwd_kv_reduce`.\n', - 'Their bodies ' - 'and additional ' - 'implementation ' - 'helpers may be ' - 'optimized. ' - 'Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch ' - 'and pointer ' - 'preparation ' - 'remain ' - 'protected,\n', - 'along with the ' - 'existing ' - 'numerical ' - 'checks and ' - 'benchmark ' - 'protocol. The ' - 'structural\n', - 'guard enforces ' - 'this boundary; ' - 'it does not ' - 'isolate ' - 'arbitrary ' - 'Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_lightning_attn_kv_reduce/config.yaml': {'after_sha256': '56cdeefe603318820ff265adc4276a0f7a52b508322a5c9cdcdd78c523ac1cf0', 'before_sha256': '170bdeb7d769ba61c2b3c053bab56dbb66f65d234ea109cf0e6c107fd25ec4ef', 'reverse_edits': [(6, @@ -1254,34 +422,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_lightning_attn_kv_reduce.py\n'])]}, - 'tasks/triton2triton/vllm/triton_lora_expand/README.md': {'after_sha256': '710f51d0aa63746a19d8f7316e0ffea67ccdb7ce4d9ed129aa990296cf2701de', - 'before_sha256': 'e1de456de9737753bf5059cdbe746b06bce0fd10672d628e1113792d4e258f3a', - 'reverse_edits': [(57, - 71, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `mm_k`, `do_expand_kernel`, ' - '`_lora_expand_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_lora_expand/config.yaml': {'after_sha256': '5599d6b8efcdc61ac58b8738c781bc8692527b25fa26ff62f6cb67d104d14ec1', 'before_sha256': '440ee2bc0389a8f18b1b9082fb00fe104a869619819309b0431e126178602131', 'reverse_edits': [(6, @@ -1298,34 +438,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_lora_expand.py\n'])]}, - 'tasks/triton2triton/vllm/triton_lora_shrink/README.md': {'after_sha256': 'a6a8373539b88204bf41f89598476d9c67c4f9a9d520b8887dafa9129776ccec', - 'before_sha256': 'a11e56f3f3f7fc809aa99f322b9d1afb4a6d00252f568fd24c37ade7d02f0dd1', - 'reverse_edits': [(79, - 93, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `mm_k`, `do_shrink_kernel`, ' - '`_lora_shrink_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_lora_shrink/config.yaml': {'after_sha256': 'bfdb6bb0a29cb205f22647f5176caf55c31ed3db818248f0868448d6c8a8fdf0', 'before_sha256': '313c1c2ccb404c3f9ca3045735eaa78d7224d5189ad221a89633e1830c29dbb5', 'reverse_edits': [(6, @@ -1342,34 +454,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_lora_shrink.py\n'])]}, - 'tasks/triton2triton/vllm/triton_matmul_persistent/README.md': {'after_sha256': '5ea2244d9b6a75857fb1c9ded4b724056ed4c02b21cbf1e9d4e64c165e179140', - 'before_sha256': '0e46ab7d08ae6b77aa256fa89b5fcbad1a8659af32cf5f344fd76911b638995e', - 'reverse_edits': [(63, - 77, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `_compute_pid`, ' - '`matmul_kernel_persistent`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_matmul_persistent/config.yaml': {'after_sha256': '173346b0733047ebc9be30b5b2af7c6f20dda5127fc9051890c5c270767821cf', 'before_sha256': '902b575005ab4bb8a563ea1ed7bf68e4f3239473fbf64eb8c282dec59aaa8d72', 'reverse_edits': [(6, @@ -1386,33 +470,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_matmul_persistent.py\n'])]}, - 'tasks/triton2triton/vllm/triton_mean/README.md': {'after_sha256': '339016b7be69cb6a98e98b3cbff6f5f15fa8432b823a0b1cde79ade62b9b815f', - 'before_sha256': '8bd9e87ca515a8e2caadbd1828d2ded335f531f22308ed16f5a22c9e9f99672e', - 'reverse_edits': [(60, - 74, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `mean_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_mean/config.yaml': {'after_sha256': 'ee104296178d1ccf897e4e50770c1348e5c63b9979b91a02ce1691335429299b', 'before_sha256': 'c3f57fb29709e2fb9c7f5c6dce970f8ac0d474ac143be383dc21b74e5dbb26a4', 'reverse_edits': [(6, @@ -1423,40 +480,6 @@ def known_feedforward_geometry(source): ' - mean_kernel\n', ' allow_new_helpers: true\n'], [' - source/triton_mean.py\n'])]}, - 'tasks/triton2triton/vllm/triton_merge_16x16_to_32x32/README.md': {'after_sha256': 'e5d4f924c5a65adec6d0373037b8fe8150c08f29584b619fedcbd2c4f981aabd', - 'before_sha256': '935b705e42500ef092e764087f293f71806a789001ed6b5e8e739f9cf38fae54', - 'reverse_edits': [(55, - 64, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`merge_16x16_to_32x32_inverse_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch and ' - 'pointer preparation ' - 'remain protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and ' - 'benchmark protocol. ' - 'The structural\n', - 'guard enforces this ' - 'boundary; it does ' - 'not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_merge_16x16_to_32x32/config.yaml': {'after_sha256': '216eb7997a713f039541f31b30e506a4acf96640dcedf247d3e0d49c80dfd878', 'before_sha256': 'e366a210f4501524d3ea499824bb46c82760c105cd1a895ae36f59ba9caba953', 'reverse_edits': [(6, @@ -1473,33 +496,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_merge_16x16_to_32x32.py\n'])]}, - 'tasks/triton2triton/vllm/triton_merge_attn_states/README.md': {'after_sha256': 'dc4ed0f79ef07c837098149f12b9db7b31f05428546c2a678782c0902f440787', - 'before_sha256': '6f11a6d78252446ffe1f4bebe1d235d29e8bfce8155601bbbe71535e6e13259f', - 'reverse_edits': [(44, - 58, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `merge_attn_states_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_merge_attn_states/config.yaml': {'after_sha256': 'b81487fe1dcc58b1a79a6f71c1dd8d78f86adf98b09824541fccd0d204a07982', 'before_sha256': 'e7d6291904f5a87e3b14a561e60a781b7074391c863c3568e74dca4709e43d9b', 'reverse_edits': [(6, @@ -1515,30 +511,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_merge_attn_states.py\n'])]}, - 'tasks/triton2triton/vllm/triton_min_p/README.md': {'after_sha256': '6e2a8bdb32719268b795cb177a25a01367faca40d45b5d6bda31a01bfacf13ec', - 'before_sha256': '704b6b17285d21fd3d609df5884b8ba9bfc0799aad869604829250a52c88aa5c', - 'reverse_edits': [(30, - 39, - ['\n', - '\n', - '### Editable device implementation ' - 'boundary\n', - '\n', - 'The editable device functions are ' - '`_min_p_kernel`.\n', - 'Their bodies and additional ' - 'implementation helpers may be ' - 'optimized. Module\n', - 'imports, initialization, host ' - 'dispatch and pointer preparation ' - 'remain protected,\n', - 'along with the existing numerical ' - 'checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it ' - 'does not isolate arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_min_p/config.yaml': {'after_sha256': '451ecdc34ee5ce1c1269bb617dfaa5835d8845a1f8c4d68843b5d1035018d163', 'before_sha256': 'cbfcd4e4f1e7e24da8ed20e8ee8a1264d3035115642ba2577cec70e4274545b3', 'reverse_edits': [(6, @@ -1550,30 +522,6 @@ def known_feedforward_geometry(source): ' - _min_p_kernel\n', ' allow_new_helpers: true\n'], [' - source/triton_min_p.py\n'])]}, - 'tasks/triton2triton/vllm/triton_mrope/README.md': {'after_sha256': 'bd741fe955f9323dbd4699deccd52d9cb3356d4b5097eb9034609f54fc605a68', - 'before_sha256': '6c6a3a34e733c8a6c9fedfa1acb8cd88068ef63fd9647b71f2898018c354f9f2', - 'reverse_edits': [(59, - 68, - ['\n', - '\n', - '### Editable device implementation ' - 'boundary\n', - '\n', - 'The editable device functions are ' - '`_triton_mrope_forward`.\n', - 'Their bodies and additional ' - 'implementation helpers may be ' - 'optimized. Module\n', - 'imports, initialization, host ' - 'dispatch and pointer preparation ' - 'remain protected,\n', - 'along with the existing numerical ' - 'checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it ' - 'does not isolate arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_mrope/config.yaml': {'after_sha256': '578f4fed58c2c6bb8cadf907064c765a627a03112024921cc5e4aa435f07c0fc', 'before_sha256': 'bd290946f28272245199bd7977af5f94c31fd8f29f9ede48deadac81d86403be', 'reverse_edits': [(6, @@ -1585,124 +533,6 @@ def known_feedforward_geometry(source): ' - _triton_mrope_forward\n', ' allow_new_helpers: true\n'], [' - source/triton_mrope.py\n'])]}, - 'tasks/triton2triton/vllm/triton_paged_prefix_prefill/README.md': {'after_sha256': '91f48bc2f53f36373d1b66a21a1ba325d8bd702c0e790a496d27638736d801f5', - 'before_sha256': '9b9f0668b7673c7acf3d85d6cb9724a850e6da0da039499e88a45960b62e1adb', - 'reverse_edits': [(33, - 34, - ['- Must handle the wrapper\'s full "Supported input domain" below, not only ' - 'the scored shapes\n'], - ['- Must handle arbitrary context lengths, query lengths, and head ' - 'dimensions\n']), - (92, - 160, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `_fwd_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n', - '\n', - '\n', - '### Supported input domain\n', - '\n', - 'The protected `context_attention_fwd` launch defines the inputs the kernel ' - 'must\n', - 'handle. The scored rows sample only part of this domain:\n', - '\n', - "- Data: float16 `q`, `k`, `v`, caches and output in the harness's " - 'contiguous\n', - ' layouts; int32 page table, start locations and sequence lengths. The 5D ' - 'K\n', - ' cache uses vector width x = 8 (16 bytes).\n', - '- Head dimension: a multiple of 8 from 16 through 128. The wrapper pads it ' - 'to\n', - ' the next power of two (`BLOCK_DMODEL_PADDED`) and the kernel masks the ' - 'padded\n', - ' lanes, so non-power-of-two sizes such as 24, 80, 96 and 112 are ' - 'included.\n', - ' Head dimensions below 16 or above 128 are outside the evaluated contract; ' - 'the\n', - ' fixed launch configuration has not been qualified for them.\n', - '- Heads: `num_heads` is a multiple of `num_kv_heads` (MHA, GQA with any ' - 'integer\n', - ' group size, or MQA).\n', - '- Physical page size: any positive size. Power-of-two pages launch\n', - ' `BLOCK_M=128, BLOCK_N=64`; other sizes launch `BLOCK_M=BLOCK_N=32`. ' - 'Context is\n', - ' read in 32-token tiles, so a page may be smaller or larger than a tile. ' - 'The\n', - ' page table may map logical pages to any physical pages and needs at ' - 'least\n', - ' `ceil(context_length / page_size)` entries per sequence.\n', - '- Lengths: per sequence, a context length >= 0 and a query length >= 0 ' - '(empty\n', - ' query segments and single-token queries included), with\n', - ' `b_seq_len = context + query` and `max_input_len` at least the largest ' - 'query\n', - ' length. Lengths need not be multiples of the context tile, the query tiles ' - 'or\n', - ' the page size.\n', - '- Scale and window: `sm_scale` omitted (`1/sqrt(head_dim)`) or any ' - 'explicit\n', - ' value; `sliding_window` omitted or <= 0 (disabled), or any positive ' - 'width,\n', - ' including widths longer than the sequence.\n', - '\n', - 'Seven further correctness-only rows sample this domain beyond the original\n', - 'cases and controls. Their declarations live in the protected\n', - '`_domain_controls.py`, which `_arena_checks.py` merges into the same ' - 'control\n', - 'table. They therefore use the same manifest binding and public-wrapper ' - 'checks\n', - 'described above: the FP32 oracle computed from pristine inputs, which ' - 'rebuilds\n', - 'context from the actual cache bytes through the page table; read-only input ' - 'and\n', - 'metadata checks; full-output 0.01 gates; and the unscored captured-replay\n', - 'diagnostic. Each row uses a non-identity page table, finite data in unused ' - 'pages\n', - 'and page tails, and an explicit nondefault scale. No row is scored or timed ' - 'in\n', - 'the performance action.\n', - '\n', - '| Control | head_dim | Page | Tiles M/N | Heads q/kv | Context lengths | ' - 'Query lengths | Scale | Window |\n', - '| --- | --- | --- | --- | --- | --- | --- | --- | --- |\n', - '| `domain_hd80_page8_two_m_tiles` | 80 | 8 | 128/64 | 8/2 | 0, 33, 130 | ' - '129, 1, 17 | 0.07 | none |\n', - '| `domain_hd96_page64_mqa_window` | 96 | 64 | 128/64 | 4/1 | 95, 64, 1 | 70, ' - '64, 2 | 0.25 | 40 |\n', - '| `domain_hd112_page12_gqa3_window` | 112 | 12 | 32/32 | 6/2 | 47, 0, 25 | ' - '33, 31, 1 | 0.06 | 3 |\n', - '| `domain_hd24_page40_wide_window` | 24 | 40 | 32/32 | 8/8 | 300, 81 | 65, ' - '96 | 0.3 | 4096 |\n', - '| `domain_hd16_page16_empty_query` | 16 | 16 | 128/64 | 2/1 | 5, 200, 31 | ' - '0, 3, 40 | 0.5 | none |\n', - '| `domain_hd128_page32_long_ragged` | 128 | 32 | 128/64 | 8/4 | 257, 0, 513 ' - '| 255, 7, 1 | 0.05 | none |\n', - '| `domain_hd32_page128_window` | 32 | 128 | 128/64 | 4/4 | 1, 129 | 128, 63 ' - '| 0.1 | 16 |\n', - '\n', - 'The default scale remains covered by the scored rows with head dimensions ' - '64\n', - 'and 128. These rows sample the domain; they do not enumerate it.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_paged_prefix_prefill/config.yaml': {'after_sha256': '7cc272bea549b0eca3fdf675b9a0654e6870167a0bfedc966fd518518d8b097a', 'before_sha256': '6140359155f3f66e00bb28378ae9675f04b48a3c51a1ab21ff546d5b963a8269', 'reverse_edits': [(6, @@ -1719,33 +549,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_paged_prefix_prefill.py\n'])]}, - 'tasks/triton2triton/vllm/triton_paged_prefix_prefill_alibi/README.md': {'after_sha256': '7df4e123e68e24b850a0e091f2614516829d4967fd65dc1a8ce4185d03e429ec', - 'before_sha256': '765fb8a7564c72024f9595a0aedc69bcc82bf96bb6cf59db0709003dc29699a2', - 'reverse_edits': [(112, - 126, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `_fwd_kernel_alibi`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_paged_prefix_prefill_alibi/config.yaml': {'after_sha256': '1459ef42022d411b6c358cd1ce4697b2c12c1c53d400c74b94f7e20d71f855cb', 'before_sha256': 'b6b06a0041f97d27c626a5c477ae7751cdaa29683130ebe9be687c91e888eb40', 'reverse_edits': [(6, @@ -1763,31 +566,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_paged_prefix_prefill_alibi.py\n'])]}, - 'tasks/triton2triton/vllm/triton_penalties/README.md': {'after_sha256': '2a8327217f233f1b239ad448d97c27c504ae9785cfe0109f0c747ee514380d51', - 'before_sha256': 'e5582672f76bfd4ef57a396c38b6859d99747eb8105fb0fdad5a98fcfa60c3f2', - 'reverse_edits': [(43, - 52, - ['\n', - '\n', - '### Editable device ' - 'implementation boundary\n', - '\n', - 'The editable device functions ' - 'are `_penalties_kernel`.\n', - 'Their bodies and additional ' - 'implementation helpers may be ' - 'optimized. Module\n', - 'imports, initialization, host ' - 'dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the existing ' - 'numerical checks and benchmark ' - 'protocol. The structural\n', - 'guard enforces this boundary; ' - 'it does not isolate arbitrary ' - 'Python execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_penalties/config.yaml': {'after_sha256': '4a21f1dd179fbe50b9c9e97a7478f3c4e6b53e4af0c57c2db7807fe20a38f4b5', 'before_sha256': 'c87835714a24223c95940ed03d26116831111a96fb7198b2a3c746ecb60a6ba1', 'reverse_edits': [(6, @@ -1801,66 +579,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_penalties.py\n'])]}, - 'tasks/triton2triton/vllm/triton_per_token_group_quant_fp8_colmajor/README.md': {'after_sha256': '7558ab9bf1409a7c61f0759f7e8d623a87b9260d070b137f458005023699320f', - 'before_sha256': 'f8f2975deae4f6a891bc876ab93cf5e2dfde09273a206377c105a333122b6382', - 'reverse_edits': [(38, - 47, - ['\n', - '\n', - '### ' - 'Editable ' - 'device ' - 'implementation ' - 'boundary\n', - '\n', - 'The ' - 'editable ' - 'device ' - 'functions ' - 'are ' - '`_per_token_group_quant_fp8_colmajor`.\n', - 'Their ' - 'bodies ' - 'and ' - 'additional ' - 'implementation ' - 'helpers ' - 'may ' - 'be ' - 'optimized. ' - 'Module\n', - 'imports, ' - 'initialization, ' - 'host ' - 'dispatch ' - 'and ' - 'pointer ' - 'preparation ' - 'remain ' - 'protected,\n', - 'along ' - 'with ' - 'the ' - 'existing ' - 'numerical ' - 'checks ' - 'and ' - 'benchmark ' - 'protocol. ' - 'The ' - 'structural\n', - 'guard ' - 'enforces ' - 'this ' - 'boundary; ' - 'it ' - 'does ' - 'not ' - 'isolate ' - 'arbitrary ' - 'Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_per_token_group_quant_fp8_colmajor/config.yaml': {'after_sha256': 'b95c4b7d2c45a36517fd45eef363e46665fb6de21f968cd41d6be2adf970d9c4', 'before_sha256': '15a3725e42c8657785196afd604fee95f3b9c5bb96df4483b6ba5f7271842746', 'reverse_edits': [(6, @@ -1881,51 +599,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_per_token_group_quant_fp8_colmajor.py\n'])]}, - 'tasks/triton2triton/vllm/triton_per_token_group_quant_int8/README.md': {'after_sha256': '8f8665c9d509a8aa46c73cc767124a37121e3523f58cd99e7153d37763198fb7', - 'before_sha256': '42e1099f5b25787703190f5fd026cdea54fecd105891f767928b481f86ead9be', - 'reverse_edits': [(42, - 51, - ['\n', - '\n', - '### Editable ' - 'device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable ' - 'device ' - 'functions are ' - '`_per_token_group_quant_int8`.\n', - 'Their bodies ' - 'and ' - 'additional ' - 'implementation ' - 'helpers may ' - 'be optimized. ' - 'Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch ' - 'and pointer ' - 'preparation ' - 'remain ' - 'protected,\n', - 'along with ' - 'the existing ' - 'numerical ' - 'checks and ' - 'benchmark ' - 'protocol. The ' - 'structural\n', - 'guard ' - 'enforces this ' - 'boundary; it ' - 'does not ' - 'isolate ' - 'arbitrary ' - 'Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_per_token_group_quant_int8/config.yaml': {'after_sha256': 'ff25dceeed235597a061a14e39b34a9e8c1a76a902f13617d676e865ddedd334', 'before_sha256': '5263a42ab1e4694b592059b8a745f1b2a3d48b403b70ac40805e4ac6772fdc25', 'reverse_edits': [(6, @@ -1943,40 +616,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_per_token_group_quant_int8.py\n'])]}, - 'tasks/triton2triton/vllm/triton_per_token_quant_int8/README.md': {'after_sha256': '5b1cdd96658e25336816f3dec8ce1b0699cef0d4931c0b7e68d16c35c3551cb0', - 'before_sha256': '52052c7e60fc57229a2b86e7736f01f768cf68b4556f8377595f7ef7680b473e', - 'reverse_edits': [(38, - 47, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`_per_token_quant_int8`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch and ' - 'pointer preparation ' - 'remain protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and ' - 'benchmark protocol. ' - 'The structural\n', - 'guard enforces this ' - 'boundary; it does ' - 'not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_per_token_quant_int8/config.yaml': {'after_sha256': '39ae15d077be4a29d58b60f1c308019484616112070674622150145badc60035', 'before_sha256': '441a8106d7efc9f0c0dc7f6c36c90a28efeb621c87451d0692e16a9e363ee831', 'reverse_edits': [(6, @@ -1993,38 +632,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_per_token_quant_int8.py\n'])]}, - 'tasks/triton2triton/vllm/triton_rejection_sample/README.md': {'after_sha256': '28eb4b42d001e80476767edbf85efc79768418095ed8181e4dde26576a8e8bb7', - 'before_sha256': '4cfc7d6f346342032ac3ab04e77da9d3fd2f2ed07e3d9cb85d9024526771a96f', - 'reverse_edits': [(39, - 48, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`_rejection_sample_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation helpers ' - 'may be optimized. ' - 'Module\n', - 'imports, ' - 'initialization, host ' - 'dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the existing ' - 'numerical checks and ' - 'benchmark protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does not ' - 'isolate arbitrary ' - 'Python execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_rejection_sample/config.yaml': {'after_sha256': '09ad9ce0ff1ba90775985212fbecba18233307a30f3bfc479e5b293a0ff72bdc', 'before_sha256': 'ed7bee1536ec258197e7e2123dae1b5efb1e2b4189f42589f09d226373300679', 'reverse_edits': [(6, @@ -2040,47 +647,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_rejection_sample.py\n'])]}, - 'tasks/triton2triton/vllm/triton_sample_recovered_tokens/README.md': {'after_sha256': 'a513a667069f51bfe2d50c930e3288e7beb72b19fd1119ba359531b5f8878629', - 'before_sha256': '00a8f44d6cc7f9d73e67baef45da8ea7d3788e082429a53cbdbaa43791455539', - 'reverse_edits': [(44, - 53, - ['\n', - '\n', - '### Editable ' - 'device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable ' - 'device functions ' - 'are ' - '`sample_recovered_tokens_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. ' - 'Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch ' - 'and pointer ' - 'preparation ' - 'remain ' - 'protected,\n', - 'along with the ' - 'existing ' - 'numerical checks ' - 'and benchmark ' - 'protocol. The ' - 'structural\n', - 'guard enforces ' - 'this boundary; ' - 'it does not ' - 'isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_sample_recovered_tokens/config.yaml': {'after_sha256': 'beabe1b62489b63d88942bf44dbe2b76d1249f56ed40abb890fc8a55c946c79f', 'before_sha256': '254909e7c398b0621191ef88f04695334251954f1541ae5eace0accc2b4fadcf', 'reverse_edits': [(6, @@ -2097,33 +663,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_sample_recovered_tokens.py\n'])]}, - 'tasks/triton2triton/vllm/triton_scaled_mm/README.md': {'after_sha256': '2f723d1eea6493bae0e6016f0cfa2a1198165e9801923f4f2c5b4e94d2b4bc3b', - 'before_sha256': '34668c3f22b9ef5977f6211907001f507699d04557808afd8fcf16d0e7e35c75', - 'reverse_edits': [(46, - 60, - ['\n', - '\n', - '### Editable device implementation boundary\n', - '\n', - 'The editable device functions are `scaled_mm_kernel`.\n', - 'Their bodies and additional implementation helpers may be optimized. ' - 'Module\n', - 'imports, initialization, host dispatch and pointer preparation remain ' - 'protected,\n', - 'along with the existing numerical checks and benchmark protocol. The ' - 'structural\n', - 'guard enforces this boundary; it does not isolate arbitrary Python ' - 'execution.\n', - '\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_scaled_mm/config.yaml': {'after_sha256': 'ac99a1b5b0cea20678b3fed86404b5fd0943e80cd6b9fed86562068156840f51', 'before_sha256': '0bd6449f5a1ca2e4f282e97a453f9ff7cf1f2f9137a621991943a66590a8e23f', 'reverse_edits': [(6, @@ -2137,45 +676,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_scaled_mm.py\n'])]}, - 'tasks/triton2triton/vllm/triton_silu_mul_fp8_quant_dg/README.md': {'after_sha256': '2813243834774a3de6dfcffd2924902615491efce579dfdf05d660314ce5ce69', - 'before_sha256': '4470271985b81190278a4304fc468d91d32a177b3e5a3f8d208f60717325fa48', - 'reverse_edits': [(59, - 68, - ['\n', - '\n', - '### Editable ' - 'device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable ' - 'device functions ' - 'are ' - '`_silu_mul_fp8_quant_deep_gemm`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. ' - 'Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch and ' - 'pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and ' - 'benchmark ' - 'protocol. The ' - 'structural\n', - 'guard enforces ' - 'this boundary; it ' - 'does not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_silu_mul_fp8_quant_dg/config.yaml': {'after_sha256': '09006b9821ed8345abebb2b2cc60b7de1398039e3033b437292504e760e0c722', 'before_sha256': '6708ba2b1970ed5c44736eadc153b95914ff030f4302de32b7b2e4971b66a9ed', 'reverse_edits': [(6, @@ -2192,39 +692,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_silu_mul_fp8_quant_dg.py\n'])]}, - 'tasks/triton2triton/vllm/triton_silu_mul_quant_fp8/README.md': {'after_sha256': '5d288e485ea72c598a18fd805b5ec9be986a02d533767dde9ddf6239116d2426', - 'before_sha256': 'f0c2b1acdb30f1c5f831e6fa26ff184220991352bd7e09381f8f31deaa09953b', - 'reverse_edits': [(43, - 52, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`_silu_mul_per_token_group_quant_fp8_colmajor`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. Module\n', - 'imports, ' - 'initialization, host ' - 'dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and benchmark ' - 'protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does not ' - 'isolate arbitrary ' - 'Python execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_silu_mul_quant_fp8/config.yaml': {'after_sha256': '1031f26d34c3e6a88429cc747fb79089c78fc27517f4a1a6caca34e0fb10a809', 'before_sha256': '53a3f188d15a7127ca07fa6b8ca21710077c77ba4a93c4fef9745594dbdb6fae', 'reverse_edits': [(6, @@ -2241,38 +708,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_silu_mul_quant_fp8.py\n'])]}, - 'tasks/triton2triton/vllm/triton_solve_tril_16x16/README.md': {'after_sha256': '215ae38f167f9dc1a87980996aa2cc89eb630f86d98a9e6e18565405714688a5', - 'before_sha256': 'f849b191e367faa0973525ef2b7b96944ed865520fd9eb48c35a6c12db4f9c80', - 'reverse_edits': [(47, - 56, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`solve_tril_16x16_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation helpers ' - 'may be optimized. ' - 'Module\n', - 'imports, ' - 'initialization, host ' - 'dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the existing ' - 'numerical checks and ' - 'benchmark protocol. The ' - 'structural\n', - 'guard enforces this ' - 'boundary; it does not ' - 'isolate arbitrary ' - 'Python execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_solve_tril_16x16/config.yaml': {'after_sha256': '5ea41f7f632b542dbc3b8fb117dfd7091147fa52d5329231b8bba07bcb75170e', 'before_sha256': '6569a133d70f6ca9587a903dcbe56f27dac0f53ada1f01da3c4f6cc7e99e755c', 'reverse_edits': [(6, @@ -2288,31 +723,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_solve_tril_16x16.py\n'])]}, - 'tasks/triton2triton/vllm/triton_topk_topp/README.md': {'after_sha256': '17a0d182ddc31aa3456bf55f70e69a6986845e295cadb1081072c0aae4a7562a', - 'before_sha256': '54030c630f2211cece8808956895ec5056937c455d116974d94b93296e4824be', - 'reverse_edits': [(41, - 50, - ['\n', - '\n', - '### Editable device ' - 'implementation boundary\n', - '\n', - 'The editable device functions ' - 'are `_topk_topp_kernel`.\n', - 'Their bodies and additional ' - 'implementation helpers may be ' - 'optimized. Module\n', - 'imports, initialization, host ' - 'dispatch and pointer ' - 'preparation remain ' - 'protected,\n', - 'along with the existing ' - 'numerical checks and benchmark ' - 'protocol. The structural\n', - 'guard enforces this boundary; ' - 'it does not isolate arbitrary ' - 'Python execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_topk_topp/config.yaml': {'after_sha256': 'e969d682b8029f4aa47f34eeaec705a32f79ae69857d7e5914f6c0bda5a45434', 'before_sha256': '9411f9918be12ff15202181b56a0c4c75e05fdd381411551ae460ab56872f39a', 'reverse_edits': [(6, @@ -2326,43 +736,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_topk_topp.py\n'])]}, - 'tasks/triton2triton/vllm/triton_unified_attention_2d/README.md': {'after_sha256': '03f52cab38e8b23ce442033ce9e08eedd84819c178e81e80a7e3a23cd88e9c4d', - 'before_sha256': '59e557a89fb25f195699afbae737407f7d2325351e4ef62c056a2dc104bbd335', - 'reverse_edits': [(75, - 84, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`cdiv_fn`, ' - '`apply_softcap`, ' - '`find_seq_idx`, ' - '`kernel_unified_attention_2d`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch and ' - 'pointer preparation ' - 'remain protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and ' - 'benchmark protocol. ' - 'The structural\n', - 'guard enforces this ' - 'boundary; it does ' - 'not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_unified_attention_2d/config.yaml': {'after_sha256': '3b645127e8b8b9f929e6d102a9f043099935413cb174267c9a4436bf92c4481b', 'before_sha256': '31b6b7d72a8c6703ad5eb6f956febd0b5f4f64c682f11b643aed2328dcc2a294', 'reverse_edits': [(6, @@ -2384,43 +757,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_unified_attention_2d.py\n'])]}, - 'tasks/triton2triton/vllm/triton_unified_attention_3d/README.md': {'after_sha256': '07be2b136d75106b315575bba812a31b20431a60cd943fdb6e486725bf44632c', - 'before_sha256': 'e56fa3e3752b2fa440a65b5b1ff5b222a798fc961ef9c5d874fbdf1daf8d50bc', - 'reverse_edits': [(75, - 84, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`cdiv_fn`, ' - '`apply_softcap`, ' - '`find_seq_idx`, ' - '`kernel_unified_attention_3d`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch and ' - 'pointer preparation ' - 'remain protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and ' - 'benchmark protocol. ' - 'The structural\n', - 'guard enforces this ' - 'boundary; it does ' - 'not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_unified_attention_3d/config.yaml': {'after_sha256': 'c30d7d6826e2eeaf0a1a80c0b553848367d05b123e87e31f43867ac85f5502df', 'before_sha256': 'b20d3a7328cf9dfdb3a6c404c86fd8d1431c23e55def172ba5e372bdf3016156', 'reverse_edits': [(6, @@ -2442,46 +778,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_unified_attention_3d.py\n'])]}, - 'tasks/triton2triton/vllm/triton_w8a8_block_int8_matmul/README.md': {'after_sha256': 'fb32276a001f986fee3ef7b961471657661a613564943e172555ee2d2d1b41bf', - 'before_sha256': '5109f1f44ae579b00a5086e5fefa48a195995015217dc749e29dc2c100627a6b', - 'reverse_edits': [(63, - 72, - ['\n', - '\n', - '### Editable ' - 'device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable ' - 'device functions ' - 'are ' - '`_w8a8_block_int8_matmul`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. ' - 'Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch and ' - 'pointer ' - 'preparation ' - 'remain ' - 'protected,\n', - 'along with the ' - 'existing ' - 'numerical checks ' - 'and benchmark ' - 'protocol. The ' - 'structural\n', - 'guard enforces ' - 'this boundary; it ' - 'does not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_w8a8_block_int8_matmul/config.yaml': {'after_sha256': '5f483f9292f468cd97adc7024c0453f7e40686dac93042210cd1fed56ea4e9e2', 'before_sha256': '769566d34e93447a4ebb1053710768c17a17c00c62cbb8aef895b99e195d6ef9', 'reverse_edits': [(6, @@ -2498,40 +794,6 @@ def known_feedforward_geometry(source): 'true\n'], [' - ' 'source/triton_w8a8_block_int8_matmul.py\n'])]}, - 'tasks/triton2triton/vllm/triton_wy_fast_recompute_wu/README.md': {'after_sha256': 'ea7cf3e5da9544062e0c09852a5f5da98da42edd97c5c91907f61e25e2288c17', - 'before_sha256': '5fe011ce830fc585fc7c00832b068a75739c2037079a17db140248d506d657f7', - 'reverse_edits': [(45, - 54, - ['\n', - '\n', - '### Editable device ' - 'implementation ' - 'boundary\n', - '\n', - 'The editable device ' - 'functions are ' - '`recompute_w_u_fwd_kernel`.\n', - 'Their bodies and ' - 'additional ' - 'implementation ' - 'helpers may be ' - 'optimized. Module\n', - 'imports, ' - 'initialization, ' - 'host dispatch and ' - 'pointer preparation ' - 'remain protected,\n', - 'along with the ' - 'existing numerical ' - 'checks and ' - 'benchmark protocol. ' - 'The structural\n', - 'guard enforces this ' - 'boundary; it does ' - 'not isolate ' - 'arbitrary Python ' - 'execution.\n'], - [])]}, 'tasks/triton2triton/vllm/triton_wy_fast_recompute_wu/config.yaml': {'after_sha256': '897ba21b4c27ade847aa27d9cd09973c52d907734c65fb6f3d9cfb267d8aca65', 'before_sha256': 'bdcd2bcf65cf8406df3f9d6f196874d203e1866c3e111f9890dad55fa4be96e6', 'reverse_edits': [(6, @@ -2549,494 +811,268 @@ def known_feedforward_geometry(source): [' - ' 'source/triton_wy_fast_recompute_wu.py\n'])]}} -ATTENTION_SAMPLING_COVERAGE_FILES = {'tasks/triton2triton/vllm/triton_correct_attn_cp_out/README.md': {'after_sha256': '0b874a9c778ced794e4e258f145d03ba70f9e2b1437e4eca836fd1ddc1378dd5', - 'before_sha256': 'dc3283ad5b4a43e8c46c26a7a20f6123854ed2f68cc8dcec88ab0fb97537757f', - 'reverse_edits': [(63, - 75, - ['\n', - '### Additional ' - 'unscored interface ' - 'coverage\n', - '\n', - 'Three explicit ' - 'correctness-only ' - 'rows exercise base-2 ' - 'exponent/logarithm ' - 'correction,\n', - 'nonzero current-rank ' - 'indices, and both ' - 'accepted ' - 'four-dimensional LSE ' - 'layouts\n', - '`[N,B,H,1]` / ' - '`[N,1,B,H]` with ' - '`[B,1,H,D]` output ' - 'inputs. Returned ' - 'corrected output\n', - 'and global LSE ' - 'retain normalized ' - 'shapes `[B,H,D]` / ' - '`[B,H]`. Each real ' - 'wrapper call\n', - 'uses the protected ' - 'pristine reference, ' - 'full-output checks, ' - 'immutable inputs ' - 'and\n', - 'Triton dispatch ' - 'audit. Manifest ' - 'parameters are ' - 'matched to the ' - 'actual generator.\n', - 'These are finite ' - 'representative ' - 'controls, not an ' - 'arbitrary-layout ' - 'guarantee.\n', - 'All five original ' - 'scored shapes, ' - 'seeds, values, ' - 'tolerances, ten ' - 'warmups, 100 ' - 'samples,\n', - 'reference and timing ' - 'code are unchanged; ' - 'the three new rows ' - 'do not affect ' - 'scoring.\n'], - [])]}, - 'tasks/triton2triton/vllm/triton_correct_attn_cp_out/workloads.json': {'after_sha256': 'a456cd982e5a9327d077d2b6c54fcb2d03488eddf5828d1cfe0adebfe08d0e52', - 'before_sha256': 'c3dbea0d0cf60a99c050e4795a8ab4d10411ec2f8ea01e5e59e3ea50155502a1', - 'reverse_edits': [(138, - 271, - [' },\n', - ' {\n', - ' ' - '"test_case_id": ' - '"base2_nonzero_rank",\n', - ' "dtype": ' - '"float16,float32",\n', - ' "params": ' - '{\n', - ' ' - '"contract_case": ' - '"base2_nonzero_rank",\n', - ' "B": ' - '2,\n', - ' "H": ' - '3,\n', - ' "D": ' - '64,\n', - ' "N": ' - '4,\n', - ' ' - '"lse_idx": 2,\n', - ' ' - '"is_base_e": ' - 'false,\n', - ' ' - '"out_layout": ' - '"BHD",\n', - ' ' - '"lse_layout": ' - '"NBH",\n', - ' "seed": ' - '1901,\n', - ' ' - '"input_shapes": ' - '{\n', - ' ' - '"out": [\n', - ' ' - '2,\n', - ' ' - '3,\n', - ' ' - '64\n', - ' ],\n', - ' ' - '"lses": [\n', - ' ' - '4,\n', - ' ' - '2,\n', - ' 3\n', - ' ]\n', - ' },\n', - ' ' - '"output_shapes": ' - '{\n', - ' ' - '"corrected": ' - '[\n', - ' ' - '2,\n', - ' ' - '3,\n', - ' ' - '64\n', - ' ],\n', - ' ' - '"final_lse": ' - '[\n', - ' ' - '2,\n', - ' 3\n', - ' ]\n', - ' }\n', - ' },\n', - ' "checks": ' - '[\n', - ' ' - '"correctness"\n', - ' ],\n', - ' "status": ' - '"PASS"\n', - ' },\n', - ' {\n', - ' ' - '"test_case_id": ' - '"basee_4d_trailing_lse",\n', - ' "dtype": ' - '"float16,float32",\n', - ' "params": ' - '{\n', - ' ' - '"contract_case": ' - '"basee_4d_trailing_lse",\n', - ' "B": ' - '2,\n', - ' "H": ' - '3,\n', - ' "D": ' - '64,\n', - ' "N": ' - '4,\n', - ' ' - '"lse_idx": 3,\n', - ' ' - '"is_base_e": ' - 'true,\n', - ' ' - '"out_layout": ' - '"B1HD",\n', - ' ' - '"lse_layout": ' - '"NBH1",\n', - ' "seed": ' - '1902,\n', - ' ' - '"input_shapes": ' - '{\n', - ' ' - '"out": [\n', - ' ' - '2,\n', - ' ' - '1,\n', - ' ' - '3,\n', - ' ' - '64\n', - ' ],\n', - ' ' - '"lses": [\n', - ' ' - '4,\n', - ' ' - '2,\n', - ' ' - '3,\n', - ' 1\n', - ' ]\n', - ' },\n', - ' ' - '"output_shapes": ' - '{\n', - ' ' - '"corrected": ' - '[\n', - ' ' - '2,\n', - ' ' - '3,\n', - ' ' - '64\n', - ' ],\n', - ' ' - '"final_lse": ' - '[\n', - ' ' - '2,\n', - ' 3\n', - ' ]\n', - ' }\n', - ' },\n', - ' "checks": ' - '[\n', - ' ' - '"correctness"\n', - ' ],\n', - ' "status": ' - '"PASS"\n', - ' },\n', - ' {\n', - ' ' - '"test_case_id": ' - '"base2_4d_middle_lse",\n', - ' "dtype": ' - '"float16,float32",\n', - ' "params": ' - '{\n', - ' ' - '"contract_case": ' - '"base2_4d_middle_lse",\n', - ' "B": ' - '2,\n', - ' "H": ' - '3,\n', - ' "D": ' - '64,\n', - ' "N": ' - '2,\n', - ' ' - '"lse_idx": 1,\n', - ' ' - '"is_base_e": ' - 'false,\n', - ' ' - '"out_layout": ' - '"B1HD",\n', - ' ' - '"lse_layout": ' - '"N1BH",\n', - ' "seed": ' - '1903,\n', - ' ' - '"input_shapes": ' - '{\n', - ' ' - '"out": [\n', - ' ' - '2,\n', - ' ' - '1,\n', - ' ' - '3,\n', - ' ' - '64\n', - ' ],\n', - ' ' - '"lses": [\n', - ' ' - '2,\n', - ' ' - '1,\n', - ' ' - '2,\n', - ' 3\n', - ' ]\n', - ' },\n', - ' ' - '"output_shapes": ' - '{\n', - ' ' - '"corrected": ' - '[\n', - ' ' - '2,\n', - ' ' - '3,\n', - ' ' - '64\n', - ' ],\n', - ' ' - '"final_lse": ' - '[\n', - ' ' - '2,\n', - ' 3\n', - ' ]\n', - ' }\n', - ' },\n', - ' "checks": ' - '[\n', - ' ' - '"correctness"\n', - ' ],\n', - ' "status": ' - '"PASS"\n'], - [])]}, - 'tasks/triton2triton/vllm/triton_decode_attn_grouped_stage1/README.md': {'after_sha256': '2cf4c9422ca383f9a7a89bb781ddc897f299d2c389bd3308b177845beb367a7e', - 'before_sha256': '0f2a551ce51340a3709600bc78d91b0acc2d2205c10756e4dfe45d5b6d380d7d', - 'reverse_edits': [(12, - 18, - ['head are processed together in tiles of BLOCK_H=16 heads:\n', - '- Uses a grid of (batch, num_kv_heads * ceil(kv_group_num / BLOCK_H), ' - 'num_kv_splits)\n', - '- Each program loads up to BLOCK_H Q vectors of one KV group and computes\n', - " attention against that shared KV head; a group's last tile is masked at " - 'the\n', - ' group boundary\n', - '- Reads K/V from a paged KV cache through the Req_to_tokens logical-page ' - 'table\n'], - ['head are processed together using BLOCK_H:\n', - '- Uses a grid of (batch, ceil(num_heads / BLOCK_H), num_kv_splits)\n', - '- Each program loads BLOCK_H Q vectors and computes attention against\n', - ' the shared KV head\n', - '- Reads K/V from a paged KV cache via Req_to_tokens mapping\n']), - (19, - 22, - ['- The kernel body keeps the upstream optional positional-encoding split\n', - ' (BLOCK_DPE, used upstream for DeepSeek MLA), but the protected wrapper ' - 'always\n', - ' launches with BLOCK_DPE=0, so that branch is unreachable through this ' - 'contract\n'], - ['- Supports optional positional encoding (BLOCK_DPE) for architectures like\n', - ' DeepSeek MLA\n']), - (34, - 36, - ['- Must be correct over the input domain listed under "Supported input ' - 'domain"\n', - ' below\n'], - ['- Must handle arbitrary sequence lengths, head dimensions, page sizes,\n', - ' and group sizes\n']), - (100, - 174, - ['\n', - 'The protected wrapper also fixes the launch configuration: the grid and the ' - 'tile\n', - 'sizes, `num_warps` and `num_stages` it passes at launch. Optimization notes ' - 'above\n', - 'that mention tuning those values refer to the original whole-file contract; ' - 'under\n', - 'this boundary, optimize the kernel bodies and implementation helpers ' - 'instead.\n', - '\n', - '### Group boundaries crossing a 16-head tile\n', - '\n', - 'Two additional correctness-only controls use 2 KV heads with 34 or 48 query ' - 'heads\n', - '(group sizes 17 and 24), ragged lengths 63/37 and nonidentity page routing. ' - 'The\n', - 'second also uses logit capping. They reuse the existing independent FP32 ' - 'reference,\n', - 'full active-output/unchanged inactive-slot checks and unscored graph ' - 'replay.\n', - 'The launch grid assigns `ceil(group_size/16)` blocks to each KV head ' - 'independently;\n', - "query-head offsets start at that KV group's base and the final block is " - 'masked at\n', - 'the group boundary. This prevents cross-group routing when a group needs a ' - 'partial\n', - 'second block. The original group sizes 1/4/8 have exactly the same launch ' - 'grid and\n', - 'head assignments as before. All original five scored cases and the three ' - 'earlier\n', - 'controls remain, along with seeds, inputs, tolerances, ten warmups and 100 ' - 'samples.\n', - 'No score weight or timing boundary changes. Kernel source changes require ' - 'fresh\n', - 'baseline/candidate measurements; historical GPU qualification does not ' - 'transfer.\n', - '\n', - '### Supported input domain\n', - '\n', - 'The protected wrapper fixes BLOCK_N=32, BLOCK_H=16, BLOCK_DPE=0, ' - '`num_warps=4`\n', - 'and `num_stages=2`, and derives every other launch value from its ' - 'arguments.\n', - 'The kernel must be correct over this domain:\n', - '\n', - '| Input | Contract in the protected wrapper and kernel | Checked values ' - '(scored cases and controls) |\n', - '| --- | --- | --- |\n', - "| Head widths | Lk is `k_buffer.shape[-1]` (and q's last dimension), Lv is " - '`v_buffer.shape[-1]`, and they may differ. `att_out` has Lv + 1 entries per ' - 'slot. Each width is padded to one next-power-of-two tile and masked; there ' - 'is no positional-encoding split. | Lk/Lv 16/16, 40/72, 64/64, 80/80, 96/64, ' - '128/128, 256/256 |\n', - '| Query and KV heads | `num_heads` must be a multiple of `num_kv_heads`. ' - '`kv_group_num = num_heads // num_kv_heads`, and query head h reads KV head ' - '`h // kv_group_num`. Each KV head gets `ceil(kv_group_num / 16)` head tiles. ' - 'With an indivisible head count the trailing query heads are never written. | ' - 'groups 1, 3, 4, 5, 7, 8, 12, 16, 17, 24, 33; KV heads 1, 2, 3, 4, 8 |\n', - '| Page size | Any positive integer. Position p of request b reads physical ' - 'page `Req_to_tokens[b, p // page_size]` at offset `p % page_size`. | 1, 2, ' - '3, 5, 12, 16, 32, 64 |\n', - '| KV splits | Any positive integer. With length L and S splits, split s ' - 'covers `[s*ceil(L/S), min((s+1)*ceil(L/S), L))`. Splits starting at or after ' - 'L are inactive and their slots are left untouched. | 1 through 8 |\n', - '| Sequence lengths | Per request, `1 <= b_seqlen[b] <= page_size * table ' - 'columns`. Lengths need not be multiples of the page size or of BLOCK_N. | 1 ' - 'to 512, including 1, 2, 3, 7, 21, 31, 33, 37, 45, 53, 63, 70, 77, 100 and ' - '130 |\n', - '| Data types and layout | The task uses float16 q, `k_buffer` and ' - '`v_buffer`, float32 `att_out`, and int32 `Req_to_tokens` and `B_Seqlen`. The ' - 'kernel addresses the last dimension of every tensor with unit stride and ' - 'receives the other strides from the wrapper. | contiguous generator layouts ' - '|\n', - '\n', - 'Outside the checked domain: widths whose padded tile is narrower than 16 ' - 'or\n', - "wider than 256 (for example MLA's 576/512 widths, which upstream splits " - 'with\n', - 'BLOCK_DPE), other data types, and zero-length requests (the kernel writes ' - 'none\n', - "of that request's slots).\n", - '\n', - '### Domain controls\n', - '\n', - 'Seven further correctness-only controls cover that domain beyond the cases\n', - 'above:\n', - '\n', - '| Control | Batch | Q/KV heads (group) | Lk/Lv | Page | Splits | Lengths | ' - 'Cap | Covers |\n', - '| --- | --- | --- | --- | --- | --- | --- | --- | --- |\n', - '| `domain_group1_page1` | 2 | 3/3 (1) | 64/64 | 1 | 3 | 45, 7 | 0 | MHA with ' - 'an odd head count, one-token pages, partial tile |\n', - '| `domain_group3_page5_dim80` | 2 | 6/2 (3) | 80/80 | 5 | 2 | 53, 21 | 1.5 | ' - 'non-power-of-two width, odd page size, capping |\n', - '| `domain_group5_page3_k96_v64` | 2 | 15/3 (5) | 96/64 | 3 | 5 | 70, 3 | 0 | ' - 'Lk > Lv, three KV heads, inactive splits |\n', - '| `domain_group7_page64_dim256` | 2 | 14/2 (7) | 256/256 | 64 | 1 | 100, 33 ' - '| 0 | widest tile, page larger than BLOCK_N, one split |\n', - '| `domain_group33_page12_k40_v72` | 2 | 66/2 (33) | 40/72 | 12 | 6 | 77, 1 | ' - '2.5 | three head tiles per group, Lk < Lv, length 1 |\n', - '| `domain_mqa_group12_page2_dim16` | 3 | 12/1 (12) | 16/16 | 2 | 4 | 31, 64, ' - '2 | 0 | MQA, narrowest tile, three requests |\n', - '| `domain_group16_page16_dim128` | 2 | 32/2 (16) | 128/128 | 16 | 7 | 130, ' - '40 | 1.5 | group exactly filling one tile, seven splits |\n', - '\n', - 'Each control draws q, K and V from a seeded CPU generator, so every device ' - 'sees\n', - 'the same bytes. Each request routes through a seeded permutation of ' - 'physical\n', - 'pages; the table has one spare column, and one physical page is never\n', - 'referenced. `att_out` starts at the finite sentinel 23.5. The protected\n', - '`_arena_domain_controls.py` declares the table, the generator and the ' - 'oracle,\n', - 'and the adapter checks the manifest rows against that table.\n', - '\n', - 'The oracle is an independent float64 evaluation of the pristine input ' - 'clones\n', - 'taken before the candidate call. Per request it gathers tokens through the ' - 'page\n', - 'table, maps query head h to KV head `h // kv_group_num`, and computes each\n', - "split's logsumexp and softmax-weighted values, applying `cap * tanh(x / " - 'cap)`\n', - 'before softmax when the cap is positive. The controls reuse every check of ' - 'the\n', - 'earlier controls: the launch audit; NaN poisoning of every active output\n', - 'element; read-only input, metadata and alias checks; byte-exact inactive ' - 'slots;\n', - 'finite outputs within atol = rtol = 0.01; and the unscored ' - '10-warmup/100-sample\n', - 'graph replay with perturbed inputs. Each control declares only ' - '`correctness`. The performance action, the five\n', - 'scored cases and their inputs, tolerances, warmups, samples and timing are\n', - 'unchanged.\n'], - [])]}, +ATTENTION_SAMPLING_COVERAGE_FILES = {'tasks/triton2triton/vllm/triton_correct_attn_cp_out/workloads.json': {'after_sha256': 'a456cd982e5a9327d077d2b6c54fcb2d03488eddf5828d1cfe0adebfe08d0e52', + 'before_sha256': 'c3dbea0d0cf60a99c050e4795a8ab4d10411ec2f8ea01e5e59e3ea50155502a1', + 'reverse_edits': [(138, + 271, + [' },\n', + ' {\n', + ' ' + '"test_case_id": ' + '"base2_nonzero_rank",\n', + ' "dtype": ' + '"float16,float32",\n', + ' "params": ' + '{\n', + ' ' + '"contract_case": ' + '"base2_nonzero_rank",\n', + ' "B": ' + '2,\n', + ' "H": ' + '3,\n', + ' "D": ' + '64,\n', + ' "N": ' + '4,\n', + ' ' + '"lse_idx": 2,\n', + ' ' + '"is_base_e": ' + 'false,\n', + ' ' + '"out_layout": ' + '"BHD",\n', + ' ' + '"lse_layout": ' + '"NBH",\n', + ' "seed": ' + '1901,\n', + ' ' + '"input_shapes": ' + '{\n', + ' ' + '"out": [\n', + ' ' + '2,\n', + ' ' + '3,\n', + ' ' + '64\n', + ' ],\n', + ' ' + '"lses": [\n', + ' ' + '4,\n', + ' ' + '2,\n', + ' 3\n', + ' ]\n', + ' },\n', + ' ' + '"output_shapes": ' + '{\n', + ' ' + '"corrected": ' + '[\n', + ' ' + '2,\n', + ' ' + '3,\n', + ' ' + '64\n', + ' ],\n', + ' ' + '"final_lse": ' + '[\n', + ' ' + '2,\n', + ' 3\n', + ' ]\n', + ' }\n', + ' },\n', + ' "checks": ' + '[\n', + ' ' + '"correctness"\n', + ' ],\n', + ' "status": ' + '"PASS"\n', + ' },\n', + ' {\n', + ' ' + '"test_case_id": ' + '"basee_4d_trailing_lse",\n', + ' "dtype": ' + '"float16,float32",\n', + ' "params": ' + '{\n', + ' ' + '"contract_case": ' + '"basee_4d_trailing_lse",\n', + ' "B": ' + '2,\n', + ' "H": ' + '3,\n', + ' "D": ' + '64,\n', + ' "N": ' + '4,\n', + ' ' + '"lse_idx": 3,\n', + ' ' + '"is_base_e": ' + 'true,\n', + ' ' + '"out_layout": ' + '"B1HD",\n', + ' ' + '"lse_layout": ' + '"NBH1",\n', + ' "seed": ' + '1902,\n', + ' ' + '"input_shapes": ' + '{\n', + ' ' + '"out": [\n', + ' ' + '2,\n', + ' ' + '1,\n', + ' ' + '3,\n', + ' ' + '64\n', + ' ],\n', + ' ' + '"lses": [\n', + ' ' + '4,\n', + ' ' + '2,\n', + ' ' + '3,\n', + ' 1\n', + ' ]\n', + ' },\n', + ' ' + '"output_shapes": ' + '{\n', + ' ' + '"corrected": ' + '[\n', + ' ' + '2,\n', + ' ' + '3,\n', + ' ' + '64\n', + ' ],\n', + ' ' + '"final_lse": ' + '[\n', + ' ' + '2,\n', + ' 3\n', + ' ]\n', + ' }\n', + ' },\n', + ' "checks": ' + '[\n', + ' ' + '"correctness"\n', + ' ],\n', + ' "status": ' + '"PASS"\n', + ' },\n', + ' {\n', + ' ' + '"test_case_id": ' + '"base2_4d_middle_lse",\n', + ' "dtype": ' + '"float16,float32",\n', + ' "params": ' + '{\n', + ' ' + '"contract_case": ' + '"base2_4d_middle_lse",\n', + ' "B": ' + '2,\n', + ' "H": ' + '3,\n', + ' "D": ' + '64,\n', + ' "N": ' + '2,\n', + ' ' + '"lse_idx": 1,\n', + ' ' + '"is_base_e": ' + 'false,\n', + ' ' + '"out_layout": ' + '"B1HD",\n', + ' ' + '"lse_layout": ' + '"N1BH",\n', + ' "seed": ' + '1903,\n', + ' ' + '"input_shapes": ' + '{\n', + ' ' + '"out": [\n', + ' ' + '2,\n', + ' ' + '1,\n', + ' ' + '3,\n', + ' ' + '64\n', + ' ],\n', + ' ' + '"lses": [\n', + ' ' + '2,\n', + ' ' + '1,\n', + ' ' + '2,\n', + ' 3\n', + ' ]\n', + ' },\n', + ' ' + '"output_shapes": ' + '{\n', + ' ' + '"corrected": ' + '[\n', + ' ' + '2,\n', + ' ' + '3,\n', + ' ' + '64\n', + ' ],\n', + ' ' + '"final_lse": ' + '[\n', + ' ' + '2,\n', + ' 3\n', + ' ]\n', + ' }\n', + ' },\n', + ' "checks": ' + '[\n', + ' ' + '"correctness"\n', + ' ],\n', + ' "status": ' + '"PASS"\n'], + [])]}, 'tasks/triton2triton/vllm/triton_decode_attn_grouped_stage1/source/_fwd_grouped_kernel_stage1.py': {'after_sha256': 'edd7803c83a25802b72386b774651218aedb6bdd6880d45f270ac4d49527ae12', 'before_sha256': '3f0d728ac011a7b9bd5eee1ae9d3eb69ef1ad91367918c4660c05a743cac7e4f', 'reverse_edits': [(53, @@ -3131,65 +1167,7 @@ def known_feedforward_geometry(source): [' ' 'triton.cdiv(head_num, ' 'min(BLOCK_H, ' - 'kv_group_num)),\n'])]}, - 'tasks/triton2triton/vllm/triton_rejection_sample/README.md': {'after_sha256': 'a44594d5c324187274f560d787fb2d30532ba2c396f10c3d0301a8ad203fe0b7', - 'before_sha256': '28eb4b42d001e80476767edbf85efc79768418095ed8181e4dde26576a8e8bb7', - 'reverse_edits': [(48, - 62, - ['\n', - '### Ragged ' - 'speculative-decoding ' - 'offsets\n', - '\n', - 'Two additional ' - 'correctness-only rows ' - 'use per-request ' - 'positive spans ' - '`[1,3,5,2]`\n', - 'and `[4,2,6,1]` in ' - 'cumulative logits ' - 'offsets. The output ' - 'capacity stays ' - 'uniform\n', - 'while each active ' - 'request may have a ' - 'different number of ' - 'draft/target logits.\n', - 'They cover bonus-only ' - 'requests, full ' - 'acceptance, and ' - 'early/late rejection. ' - 'In\n', - 'addition to the ' - 'retained protected ' - 'reference and read-only ' - 'checks, an ' - 'independent\n', - 'scalar oracle selects ' - 'the first mismatch and ' - 'checks every accepted ' - 'prefix and\n', - 'exact count. Undefined ' - 'unused output tails ' - 'remain unchecked. The ' - 'original five\n', - 'scored cases and ' - 'earlier controls, ' - 'kernel, seeds, integer ' - 'comparison, ten ' - 'warmups,\n', - '100 samples and timer ' - 'are unchanged. These ' - 'added cases do not ' - 'contribute a score.\n', - 'A zero speculative-step ' - 'request has one bonus ' - 'logit; no zero-logit ' - 'request is\n', - 'introduced by these ' - 'active-request ' - 'controls.\n'], - [])]}} + 'kv_group_num)),\n'])]}} def _exact_reviewed_file_inverse(text, record): @@ -3210,11 +1188,4 @@ def historical_attention_sampling_file(relative, text): def historical_triton_device_boundary_file(relative, text): record = TRITON_DEVICE_BOUNDARY_FILES.get(relative) - if record is None: - return text - # Only three later coverage-doc appendices are composed before the exact - # source20 boundary inverse. Both complete versions are individually pinned. - coverage = ATTENTION_SAMPLING_COVERAGE_FILES.get(relative) - if coverage is not None and hashlib.sha256(text.encode()).hexdigest() == coverage['after_sha256']: - text = _exact_reviewed_file_inverse(text, coverage) - return _exact_reviewed_file_inverse(text, record) + return _exact_reviewed_file_inverse(text, record) if record is not None else text diff --git a/tests/test_attention_sampling_coverage.py b/tests/test_attention_sampling_coverage.py index 93da3253b..221f30a3d 100644 --- a/tests/test_attention_sampling_coverage.py +++ b/tests/test_attention_sampling_coverage.py @@ -259,8 +259,15 @@ def check(*,case_index): assert result['status']=='FAIL' and 'manifest' in result['reason'].lower() +def test_reviewed_history_pins_contract_files_not_readme_text(): + from reviewed_contract_history import ATTENTION_SAMPLING_COVERAGE_FILES,TRITON_DEVICE_BOUNDARY_FILES + assert len(TRITON_DEVICE_BOUNDARY_FILES)==45 + assert all(relative.endswith('/config.yaml') for relative in TRITON_DEVICE_BOUNDARY_FILES) + assert not any(relative.endswith('/README.md') for relative in ATTENTION_SAMPLING_COVERAGE_FILES) + + @pytest.mark.parametrize('relative',__import__('reviewed_contract_history').TRITON_DEVICE_BOUNDARY_FILES) -def test_all_45_boundary_config_and_doc_inverses_require_exact_current_bytes(relative): +def test_all_45_boundary_config_inverses_require_exact_current_bytes(relative): import hashlib from reviewed_contract_history import TRITON_DEVICE_BOUNDARY_FILES,historical_triton_device_boundary_file source=(ROOT/relative).read_text()