diff --git a/README.md b/README.md index 842a09b2f27ebbbf3eb694add76b01e9a8d9c72f..2e856277df863ed4ef7061aa090954923b2c3288 100644 --- a/README.md +++ b/README.md @@ -28,14 +28,14 @@ All-BFP8: every text matrix (attention, MLP, LM head) is stored as TTNN `bfloat8 `tensors/tensor_cache_bf16/*.tensorbin` holds the exact TTNN host tensors that the runtime uploads to device DRAM. The directory name comes from the runtime's tensor-cache writer. The files are not a regenerable cache: they are the checkpoint. Every file is hashed in the manifest, and the loader refuses to start if any file is missing or changed. The drafter (`gemma-4-12B-it-assistant/`) stores its tensors the same way. Its recorded dtype is `bfp8-lm4`: BFP8 weights and a BFP4 drafter LM head. -The loader requires two proofs, both bound to the manifests: +The loader requires two proofs, both bound to the manifests and recorded on the current runtime (p3d, see [Runtime](#runtime)): - [`gemma-4-12B-it/equivalence.json`](gemma-4-12B-it/equivalence.json): the full logits of the native reload are bit-exact to quantizing the original weights at load time, over 4,093 teacher-forced tokens (`exact_logits_equal: true`). - [`gemma-4-12B-it-assistant/spec_equivalence.json`](gemma-4-12B-it-assistant/spec_equivalence.json): greedy speculative decoding with draft length 5 produced token streams identical to non-speculative greedy decoding on the same proven runtime (`streams_identical: true`, 3,130 tokens compared). ## Quality versus the original BF16 model -Teacher-forced comparison against the original model run in BF16 on CPU ([`evidence/quality/quant_table.json`](evidence/quality/quant_table.json)). dPPL is the perplexity change relative to BF16, top-1 is argmax agreement with BF16, and KL is the mean KL divergence on the short-prompt set: +This table is the precision-selection study. It is a teacher-forced comparison against the original model run in BF16 on CPU, measured on the exact (unfused) TT build ([`evidence/quality/quant_table.json`](evidence/quality/quant_table.json)). For the shipped runtime's own gate, see [Runtime](#runtime). dPPL is the perplexity change relative to BF16, top-1 is argmax agreement with BF16, and KL is the mean KL divergence on the short-prompt set: | Plan | Weight read | chat dPPL / top-1 | code dPPL / top-1 | 16K book dPPL / top-1 | KL | GSM8K-50 | |---|---:|---:|---:|---:|---:|---:| @@ -46,24 +46,42 @@ Teacher-forced comparison against the original model run in BF16 on CPU ([`evide **Why BFP8 and not BFP4:** with the MLP in BFP4, perplexity rises about 5% and top-1 agreement falls to about 87%. All-BFP4 is worse still. BFP8 stays within about 1–2% of BF16 perplexity. -The shipped runtime adds batched speculative verification and a fused GELU×up kernel. It passed a decode-path quality gate against the exact (unfused) build ([`evidence/quality/runtime-gate-dtf-score.json`](evidence/quality/runtime-gate-dtf-score.json)): argmax agreement with the exact build was 98.8% on chat, 98.8% on code and 98.4% on long-book text. The NLL change versus the exact build was +0.0017 ± 0.0012 on chat, −0.0004 ± 0.0030 on code and −0.0061 ± 0.0096 on long-book. The two proofs above were recorded with this shipped runtime configuration. +## Runtime + +The current runtime is **p3d**: image `lottolabs/gemma4-12b-tt-p150:p3d`, ID `sha256:331e79b24fa2575c436158be994c6bc0e184d776d15305e82538d2e4d646f544`. Its Gemma 4 tree is at git commit `061e48d`. On top of the earlier runtime it adds batched speculative verification, a fused GELU×up kernel, a decode KV write through a rows kernel, and dual-NoC weight reads for the decode matmuls (the TTNN in1 reader overlay, `dn` branch `efc465f`). + +These fast paths are numerics switches, and the proofs bind them (`runtime_env`: `GEMMA4_VERIFY_SDPA=batched`, `GEMMA4_FUSE_GELU_MUL=1`, `GEMMA4_DECODE_KV_ROWS=1`, `GEMMA4_DECODE_MM_DN=1`). They change floating-point rounding relative to the previous runtime p3c, so greedy outputs are not token-identical to p3c's. Within p3d, speculative output is identical to non-speculative output, and HTTP output is identical to the direct runtime (see below). The checkpoint tensors are byte-identical to the p3c release. Only `equivalence.json` and `spec_equivalence.json` were re-recorded on p3d. + +**Quality gate for p3d.** Decode-path teacher forcing was compared against the original model in BF16 on CPU ([`evidence/quality/p3d-gate/`](evidence/quality/p3d-gate/)): + +| Domain | Positions | dPPL vs BF16 | top-1 vs BF16 | argmax agreement with exact build | +|---|---:|---:|---:|---:| +| chat | 5,888 | +1.00% | 98.2% | 98.5% | +| code | 3,328 | +0.61% | 98.7% | 98.9% | +| long book | 2,048 | +1.15% | 96.5% | 97.9% | + +GSM8K-100 scored 97/100 on p3d. The exact (unfused) build and the BF16 CPU reference also scored 97/100. + +**Previous runtime:** p3c, image `lottolabs/gemma4-12b-tt-p150:p3c`, ID `sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73`. It was published at repo revision [`3d9001d5`](https://huggingface.co/Lottolabs/gemma-4-12B-it-TT-BFP8-P150/tree/3d9001d5f52a5f3350b90e87703440a44255bd95), and its evidence is under [`evidence/previous-p3c/`](evidence/previous-p3c/). On that runtime, HTTP decode ran at 50.6 / 49.3 / 46.9 / 45.1 / 29.7 / 21.8 tok/s at the prompt lengths in the table below. To serve p3c, use the `launch.py` from that revision; each launcher pins its own runtime image. ## Performance -These numbers were measured over HTTP on one P150 with greedy decoding, 512-token outputs, chat prompts, one user and speculative decoding on ([`evidence/http-p3c/`](evidence/http-p3c/)): +These numbers were measured on runtime p3d over HTTP on one P150 with greedy decoding, 512-token outputs, chat prompts, one user and speculative decoding on ([`evidence/http-p3d/`](evidence/http-p3d/)): | Prompt tokens | 128 | 2K | 8K | 32K | 131K | 262K (261,632) | |---|---:|---:|---:|---:|---:|---:| -| Decode tok/s | 50.6 | 49.3 | 46.9 | 45.1 | 29.7 | 21.8 | -| TTFT (s) | 0.09 | 0.64 | 3.2 | 17.0 | 124 | 395 | +| Decode tok/s | 63.1 | 52.4 | 53.9 | 43.4 | 30.9 | 24.6 | +| TTFT (s) | 0.09 | 0.65 | 3.2 | 17.0 | 124 | 395 | -With the LocalMaxxing official prompt (a local run of the official protocol, not submitted), the server reached **49.9 output tok/s** with a **95 ms** TTFT (median of 5, 256 output tokens; [`evidence/localmaxxing/speed-test.json`](evidence/localmaxxing/speed-test.json)). +Decode speed with speculative decoding depends on how many drafted tokens the content lets the model accept. That is why it does not fall monotonically with prompt length. -## Serving correctness +**LocalMaxxing, 2K prompt (p3d).** The server reached **53.9 output tok/s** (median of 5), 3,063 tok/s prefill and a 651 ms TTFT, with 1,996 prompt and 256 output tokens. The run was submitted as `cmuosddj10kd1lq010lf2ahkj` and approved as an unverified run ([`evidence/localmaxxing/p3d-2k/speed-test.json`](evidence/localmaxxing/p3d-2k/speed-test.json); the prompt is included). -- **HTTP matches the direct runtime token for token.** All 28 of 28 checks passed ([`evidence/http-p3c/exact.json`](evidence/http-p3c/exact.json)). They cover chat and completion outputs on 10 prompts, request isolation, sampled requests followed by greedy ones, cancellation mid-prefill and mid-decode, and rejection of image input. The 512-token outputs at 32K and 261,632 prompt tokens were also identical to the direct runtime ([`evidence/http-p3c/tput_vs_direct.json`](evidence/http-p3c/tput_vs_direct.json)). -- **Long context works to the native limit.** Passkeys at the beginning, middle and end of the prompt were retrieved at 32,768, 131,072 and 262,016 prompt tokens. A request totalling 262,145 tokens was rejected with HTTP 400 before generation started ([`evidence/http-p3c/long.json`](evidence/http-p3c/long.json)). -- **The public package was verified end to end.** It was downloaded fresh and served on a P150 ([`evidence/public-download-verification.json`](evidence/public-download-verification.json), [`evidence/public-serving-smoke.json`](evidence/public-serving-smoke.json)). +## Serving correctness (p3d) + +- **HTTP matches the direct runtime token for token.** All 28 of 28 checks passed ([`evidence/http-p3d/exact.json`](evidence/http-p3d/exact.json)). They cover chat and completion outputs on 10 prompts, request isolation, sampled requests followed by greedy ones, cancellation mid-prefill and mid-decode, and rejection of image input. The 512-token outputs at 32K and 261,632 prompt tokens were also identical to the direct runtime ([`evidence/http-p3d/tput_vs_direct.json`](evidence/http-p3d/tput_vs_direct.json)). +- **Long context works to the native limit.** Passkeys at the beginning, middle and end of the prompt were retrieved at 32,768, 131,072 and 262,016 prompt tokens. A request totalling 262,145 tokens was rejected with HTTP 400 before generation started ([`evidence/http-p3d/long.json`](evidence/http-p3d/long.json)). +- **The public package was verified end to end.** It was downloaded fresh, anonymously, and served on a P150 ([`evidence/public-download-verification.json`](evidence/public-download-verification.json), [`evidence/public-serving-smoke.json`](evidence/public-serving-smoke.json)). ## Download and serve @@ -131,7 +149,13 @@ curl http://127.0.0.1:8000/v1/chat/completions -H 'Content-Type: application/jso - [`release-manifest.json`](release-manifest.json): the complete repository inventory with size and sha256 for every file. - [`SHA256SUMS`](SHA256SUMS): file checksums. - [`reproduction.json`](reproduction.json): source, runtime, serving and evidence contract. -- [`runtime/`](runtime/): source of the runtime image layers, recorded in [`runtime/source.json`](runtime/source.json). It contains the Gemma 4 TT model tree (`runtime/gemma4/`, git commit `ce67960382fffed8550b38717cedf574e821da06`, byte-identical to the tree in the image), the vLLM/TT-plugin serving overlay, the TTNN matmul overlay, the Dockerfile chain and the build script. This is provenance: serving uses the checksum-pinned `runtime-image.tar.gz`. +- [`runtime/`](runtime/): source of the runtime image layers, recorded in [`runtime/source.json`](runtime/source.json). It contains: + - the Gemma 4 TT model tree (`runtime/gemma4/`, git commit `061e48ddf6a0eab5dac1f958731fcb70768231d1`, byte-identical to the tree in the image); + - the TTNN overlay (`runtime/ttnn-overlay/`: 1D matmul factory and the dual-NoC in1 reader, `dn` commit `efc465f4b1a44025ae4a33e4dbc7d371b3f9266c`, byte-identical to the image); + - the vLLM/TT-plugin serving overlay; + - the Dockerfile chain (p2b → p2c → ttnn-dn1 → fast2 → p3d) and the build script. + + This is provenance: serving uses the checksum-pinned `runtime-image.tar.gz`. - [`native_checkpoint.py`](native_checkpoint.py) and [`provenance/build_native_checkpoint.py`](provenance/build_native_checkpoint.py): the manifest, proof and verification tool, and the checkpoint builder. ## License and attribution diff --git a/SHA256SUMS b/SHA256SUMS index 86c7693a6c3a7d6e2baa665bc28a3eae198dbc79..6ce17c989739d2051b786b8c033f9ff07053b8ac 100644 --- a/SHA256SUMS +++ b/SHA256SUMS @@ -1,39 +1,66 @@ cfc7749b96f63bd31c3c42b5c471bf756814053e847c10f3eb003417bc523d30 LICENSE -98ed9bdcd1983193f71e4d7e327eebe72a3be34c6d19e8f378aa5b7e5e370a54 README.md -dcd509a51cd400fd48f042a8147c99df5381132386b53347b7b37a5f48b16094 evidence/http-p3c/exact.json -67124b62e2fd13138863abda1b821e6e66e066541e572e5ec321a4b275b0e959 evidence/http-p3c/long.json -dfc6acee93d6a341f92e453592fd2ee5b52329a229be413fa4407f7bb681c7c0 evidence/http-p3c/tput-128/requests-128-c1.json -316fb623119a50661ae7e768c3bf8ad8733683ceb56953cdbdf2a2d963a5df9f evidence/http-p3c/tput-128/summary.json -44bdcbb47481b62ab615f72a2bc29c2a4d2524a4ec318f8fd19159ef4ac00c5f evidence/http-p3c/tput-128/warmup-128.json -fdb5caa830280f0ba36b1045a905cfeb02bee98e54d72bfa8a0f54f712ea4c20 evidence/http-p3c/tput-131072/requests-131072-c1.json -ea4fa34dd27a8011fe1cdad7cf26f5acd3af67c0d7cab5c1dd8a3f7c680005c4 evidence/http-p3c/tput-131072/summary.json -9aa37602c1c2ddfbe18be1c210ed50c953a20eda00f1a94b14fb3e2c249dedbd evidence/http-p3c/tput-131072/warmup-131072.json -d73bc532e74a494cb908c00f98f3745e465472e6fa66290b3de69ea0d673849a evidence/http-p3c/tput-2048/requests-2048-c1.json -b0b79e75d5f23f31261ff91863a7737dc9e0849476eac663e7577b587daa7cf8 evidence/http-p3c/tput-2048/summary.json -abc0819a5df996ee123cba52233fc2d81c16df49fffc0c2e1bd8e934eb28333c evidence/http-p3c/tput-2048/warmup-2048.json -e211831e0530c7b1cd1aa356a95beaec3a2215a7fe77d9908e858b9f3b453c05 evidence/http-p3c/tput-261632/requests-261632-c1.json -9ebdbacd8886c3faf8e2753fb010cd2a87d00d3da3a085051b9663265b6bcff5 evidence/http-p3c/tput-261632/summary.json -94a93fe15d162528cb287df6f4ff7f95244e05ea8cfaae89076b4b518eaf39d4 evidence/http-p3c/tput-261632/warmup-261632.json -37e1fc84216f03c5db76d35995054e367adc55c4fa6f2892dabb61ed6be1c12b evidence/http-p3c/tput-32768/requests-32768-c1.json -6862ec9e00ddad920c0324a442ee5d28737a2e839de8e89c693c0cd809df1774 evidence/http-p3c/tput-32768/summary.json -e4f3cc40e54e56410341d29bd26d9102f06dbd00e35901c4a05b2a05f085c3bd evidence/http-p3c/tput-32768/warmup-32768.json -e67e8fbe22e47880715b26cee62f0d999a6395388d36774117c880f6ff6e4842 evidence/http-p3c/tput-8192/requests-8192-c1.json -af278ee14869619894a8ff9a4c12c69f155d35c26f3b7868190d9559ec001981 evidence/http-p3c/tput-8192/summary.json -957010a8c724f2a1f82ce7ffc5dcbdabfeb2e1420c5a9274fda0faa8c85c05bb evidence/http-p3c/tput-8192/warmup-8192.json -1c7778066e729cf8e2f302418c63f1640eba7d88a6b943b9681f3819c1579adb evidence/http-p3c/tput_vs_direct.json -f03041863057e793e816ded6f048899034e04f11a3a4dd0fd102a04d40e7b5de evidence/localmaxxing/speed-test.json -92c92f98fb43b89ff04bd970974f23a3647264a722da2e674b033cf0d8523a32 evidence/public-download-verification.json -791d5b2dff0a130024200ab016e671c1d74495229b3461f83464bb35124478a9 evidence/public-serving-smoke.json -07b8133a548b886bc30794b9f984f0a5dee9b9301f7c90ff671806cbb4ce766c evidence/public-serving-smoke/requests-2048-c1.json -a386aef5e382c304cc3c5315b8021166dfadec3d256b6af7887acf35383262e2 evidence/public-serving-smoke/summary.json -3cb6e6002631824144551034aeed233081c44182771b51d8c3ade48555ff7426 evidence/public-serving-smoke/warmup-2048.json +fb17f137237deb11b801ab11ca09a446b02e6a35c9a278f294daee9a82faedfb README.md +2d2a52c8c9dab817a10c4bec4c9d483ee1214b3a22cc883fb754549dac8dbf0e evidence/http-p3d/exact.json +4ae89585c639211abecfb7347a476c18d43f809385b0eb2a8077d3e994e80c8e evidence/http-p3d/long.json +534ef3c7105f7f2d8fc7e0bad4836b6d5507139c0a18a0b97377d4a8051ba035 evidence/http-p3d/tput-128/requests-128-c1.json +4d213e69b02e4b71062c358b3cda4a8db94a615b43c16f3dfd619c053788dac6 evidence/http-p3d/tput-128/summary.json +7a94e4f69a9e5ac6de2d96cd7eb9647fda3c50ba406cdaa2c68bba7aba5c74ad evidence/http-p3d/tput-128/warmup-128.json +2bbf64804564dfaf71c0f5b78db64db026eaf16ab5d515125e295f145c137548 evidence/http-p3d/tput-131072/requests-131072-c1.json +3b5fc67ff3e95a6ba8fc59f721fb79ba6df82993af27b8f6b1ef280cdd134ddf evidence/http-p3d/tput-131072/summary.json +3b0071bad4dda913079bd8d853d419220e5c6cb58b34618f3059d3e4c305d0a4 evidence/http-p3d/tput-131072/warmup-131072.json +2f6176f102977e953ddd33a8d938fe0cbf1661c1466a287e406c2f1ac7215174 evidence/http-p3d/tput-2048/requests-2048-c1.json +4206f0885321e0cd121e2efbbb91daf2c87bb3343081e81709fb2f80b9d7fcb0 evidence/http-p3d/tput-2048/summary.json +cc417c28c2a3952197e27e516b10c358665bd9792ef1cebb990de8037287292f evidence/http-p3d/tput-2048/warmup-2048.json +1cfdf00f237dccc9e9ddfed7d0f64f224f574b048d5b395edb0f81ced7a0dfdb evidence/http-p3d/tput-261632/requests-261632-c1.json +bad5d6e144e98944b706ad12b62017d4355f4148b3ff2d6ee72543f3cefe72c8 evidence/http-p3d/tput-261632/summary.json +c6a74eafe2f81e64b185d4327d79ca49bcc52f52a9572d353d55cf054656c196 evidence/http-p3d/tput-261632/warmup-261632.json +91e11398058527c1d04f3911de6218b2f3f183dfadd3977280e738dd59907aa2 evidence/http-p3d/tput-32768/requests-32768-c1.json +89effe9fecb4560388df68a68cba91117c9861f68be7991c3f89773c1e515117 evidence/http-p3d/tput-32768/summary.json +7de3bc8246a36dfcd42b9802f4312874fa5a4919f93c01552326727bcb14b3e3 evidence/http-p3d/tput-32768/warmup-32768.json +7aa5cf9a79ccd791f80a4ddb1350d3da0ecc7a14901a4cee7ec81df8e2fc1fa3 evidence/http-p3d/tput-8192/requests-8192-c1.json +2657f8971531b9f9bb940856f83740aebf1fdd7f1c6f313356ad13c102b54ac3 evidence/http-p3d/tput-8192/summary.json +5df2100d05536b268b8a811e814d66628320eca6016f7968d46f9747eeba74a1 evidence/http-p3d/tput-8192/warmup-8192.json +1c7778066e729cf8e2f302418c63f1640eba7d88a6b943b9681f3819c1579adb evidence/http-p3d/tput_vs_direct.json +cd754ee70eca7fe4db4c243e6ca985e46e46dd89157735b8ebadd92dafcff4a1 evidence/localmaxxing/p3d-2k/prompt-2k.txt +e341a88dd41612504017df8c950b3670d8dc01b057887f697a8b761dd8a32106 evidence/localmaxxing/p3d-2k/speed-test.json +dcd509a51cd400fd48f042a8147c99df5381132386b53347b7b37a5f48b16094 evidence/previous-p3c/http/exact.json +67124b62e2fd13138863abda1b821e6e66e066541e572e5ec321a4b275b0e959 evidence/previous-p3c/http/long.json +dfc6acee93d6a341f92e453592fd2ee5b52329a229be413fa4407f7bb681c7c0 evidence/previous-p3c/http/tput-128/requests-128-c1.json +316fb623119a50661ae7e768c3bf8ad8733683ceb56953cdbdf2a2d963a5df9f evidence/previous-p3c/http/tput-128/summary.json +44bdcbb47481b62ab615f72a2bc29c2a4d2524a4ec318f8fd19159ef4ac00c5f evidence/previous-p3c/http/tput-128/warmup-128.json +fdb5caa830280f0ba36b1045a905cfeb02bee98e54d72bfa8a0f54f712ea4c20 evidence/previous-p3c/http/tput-131072/requests-131072-c1.json +ea4fa34dd27a8011fe1cdad7cf26f5acd3af67c0d7cab5c1dd8a3f7c680005c4 evidence/previous-p3c/http/tput-131072/summary.json +9aa37602c1c2ddfbe18be1c210ed50c953a20eda00f1a94b14fb3e2c249dedbd evidence/previous-p3c/http/tput-131072/warmup-131072.json +d73bc532e74a494cb908c00f98f3745e465472e6fa66290b3de69ea0d673849a evidence/previous-p3c/http/tput-2048/requests-2048-c1.json +b0b79e75d5f23f31261ff91863a7737dc9e0849476eac663e7577b587daa7cf8 evidence/previous-p3c/http/tput-2048/summary.json +abc0819a5df996ee123cba52233fc2d81c16df49fffc0c2e1bd8e934eb28333c evidence/previous-p3c/http/tput-2048/warmup-2048.json +e211831e0530c7b1cd1aa356a95beaec3a2215a7fe77d9908e858b9f3b453c05 evidence/previous-p3c/http/tput-261632/requests-261632-c1.json +9ebdbacd8886c3faf8e2753fb010cd2a87d00d3da3a085051b9663265b6bcff5 evidence/previous-p3c/http/tput-261632/summary.json +94a93fe15d162528cb287df6f4ff7f95244e05ea8cfaae89076b4b518eaf39d4 evidence/previous-p3c/http/tput-261632/warmup-261632.json +37e1fc84216f03c5db76d35995054e367adc55c4fa6f2892dabb61ed6be1c12b evidence/previous-p3c/http/tput-32768/requests-32768-c1.json +6862ec9e00ddad920c0324a442ee5d28737a2e839de8e89c693c0cd809df1774 evidence/previous-p3c/http/tput-32768/summary.json +e4f3cc40e54e56410341d29bd26d9102f06dbd00e35901c4a05b2a05f085c3bd evidence/previous-p3c/http/tput-32768/warmup-32768.json +e67e8fbe22e47880715b26cee62f0d999a6395388d36774117c880f6ff6e4842 evidence/previous-p3c/http/tput-8192/requests-8192-c1.json +af278ee14869619894a8ff9a4c12c69f155d35c26f3b7868190d9559ec001981 evidence/previous-p3c/http/tput-8192/summary.json +957010a8c724f2a1f82ce7ffc5dcbdabfeb2e1420c5a9274fda0faa8c85c05bb evidence/previous-p3c/http/tput-8192/warmup-8192.json +1c7778066e729cf8e2f302418c63f1640eba7d88a6b943b9681f3819c1579adb evidence/previous-p3c/http/tput_vs_direct.json +f03041863057e793e816ded6f048899034e04f11a3a4dd0fd102a04d40e7b5de evidence/previous-p3c/localmaxxing-speed-test.json +92c92f98fb43b89ff04bd970974f23a3647264a722da2e674b033cf0d8523a32 evidence/previous-p3c/public-download-verification.json +791d5b2dff0a130024200ab016e671c1d74495229b3461f83464bb35124478a9 evidence/previous-p3c/public-serving-smoke.json +07b8133a548b886bc30794b9f984f0a5dee9b9301f7c90ff671806cbb4ce766c evidence/previous-p3c/public-serving-smoke/requests-2048-c1.json +a386aef5e382c304cc3c5315b8021166dfadec3d256b6af7887acf35383262e2 evidence/previous-p3c/public-serving-smoke/summary.json +3cb6e6002631824144551034aeed233081c44182771b51d8c3ade48555ff7426 evidence/previous-p3c/public-serving-smoke/warmup-2048.json +b118bf712e7d70058213f378920c0b129a2366b554f23d489c441ecbb2e5f2ba evidence/previous-p3c/runtime-gate-dtf-score.json +156ecb1746f24b3c6c7d6143e74034caf533b286270baab8f9bb7ad117fc8750 evidence/quality/p3d-gate/book-2048-score.json +51c580756dd62962bcad40d1cf53100319f8d9615902f2a5db24f78af42f510a evidence/quality/p3d-gate/dtf-score.json +9bb6af5e16abd148f7b7d528e045587bec96b94249a89be227352840257fb703 evidence/quality/p3d-gate/gsm8k-100-exact-build.json +82c012ee2d4dd6f5986f7a7bbbec425589405c98e487b93747457b0180e220ab evidence/quality/p3d-gate/gsm8k-100.json 32904d8e6be2daea2b797cc71b8d587e06605dc23421d96e1912c8593ed4b991 evidence/quality/quant_table.json -b118bf712e7d70058213f378920c0b129a2366b554f23d489c441ecbb2e5f2ba evidence/quality/runtime-gate-dtf-score.json b6f19209588fcefe41f65b193fad6148446253c470d36e29441ecc5158a54e6d gemma-4-12B-it-assistant/config.json 2db4559e6ec51d81a7cd51b003743ca689fed5220d41b6a403b8d6947e1c74b7 gemma-4-12B-it-assistant/drafter_manifest.json 02b56bd11e1cd1e363e701a85a2fd7fbaa2992ec3358c1cd7cc44ead7208f505 gemma-4-12B-it-assistant/generation_config.json 3279c173daddd7186e79d652ad94022415736d3a1370625696c898429b06d6df gemma-4-12B-it-assistant/model.safetensors -032cc468d6bf33bca7267be0e70d1d6932e7dcc4dfb97acb5a75e736f2ee8d0d gemma-4-12B-it-assistant/spec_equivalence.json +1c09aae0f94d18ae7075800ccf55a7866769bec1fe41475ec781c054ad4ee489 gemma-4-12B-it-assistant/spec_equivalence.json 70eb60efff28b3606bbc91b35a3311e7b4ccae3b70e0f3787c9c0b32cc4ebf3c gemma-4-12B-it-assistant/tensors/assistant_tensor_cache_bfp8/final_norm/weight_dtype_BFLOAT16_layout_ROW_MAJOR.tensorbin 1179fa4a06ab8e4bae470dce5c4848d5b3b5eca1a3c87d8a02001bfb1aacf02f gemma-4-12B-it-assistant/tensors/assistant_tensor_cache_bfp8/layer_0/layer_0/input_layernorm/weight_dtype_BFLOAT16_layout_ROW_MAJOR.tensorbin bc68d0c960e2782993ff0d354800e3f4d67dd80d1b31f0775dc5574114402787 gemma-4-12B-it-assistant/tensors/assistant_tensor_cache_bfp8/layer_0/layer_0/mlp/down_proj.weight_bfp8_dtype_BFLOAT8_B_layout_TILE.tensorbin @@ -84,7 +111,7 @@ a926c8853869f484c379a9d5b3397488c04905b15d90082ac679ecf7e6466479 gemma-4-12B-it ba02924ee433b4f9689d504ddadc0e3247d23972021668dc03bd9e27c17670dc gemma-4-12B-it-assistant/tensors/assistant_tensor_cache_bfp8/pre_projection_weight_dtype_BFLOAT8_B_layout_TILE.tensorbin ae53464bf3be25802b3a5b37def7fd89667067d7577049b3b2d74c4d8de4c6d4 gemma-4-12B-it/chat_template.jinja 478c46e8d2c52d5c2d85bf67e3b3e8c90e7c9d91086cee27e3c267907e936bd9 gemma-4-12B-it/config.json -1bd70995c42ce9d484c717d4169a000df36ce5181798c6f54b26138128daf296 gemma-4-12B-it/equivalence.json +824b1705373431eb0a79dd3774eaf1890a1849cc1de06375b1506a2f2073abd2 gemma-4-12B-it/equivalence.json a8349d9bd64cc5841297fcb5002f0fdc4749c473c8f1b10ea337f9ce4ee7014e gemma-4-12B-it/generation_config.json 007d9f3da8d54363a75e8af84cf0c9b762eafd6edf23679de48b08c62d584d55 gemma-4-12B-it/native_manifest.json 612f5da6a29a50dea769db21923d2bd7b828a94f2f25cc41f28a33abac9cb990 gemma-4-12B-it/precision_plan.json @@ -622,17 +649,18 @@ e5656d7f88172fb4add6d5969a1c15d5d50f0eb3c5e0223a448c90816ec43e8e gemma-4-12B-it dfeaed45d94523cd991122388b8ca2e0cbeeb5fe34d3d3b37c047c26bfc715e3 gemma-4-12B-it/tensors/tensor_cache_bf16/lm_head.weight_bfp8_dtype_BFLOAT8_B_layout_TILE.tensorbin cc8d3a0ce36466ccc1278bf987df5f71db1719b9ca6b4118264f45cb627bfe0f gemma-4-12B-it/tokenizer.json a62f4e85a47c0c136edaaa3a4f591fd6783717299a9def47e5ad03a49f6a5eb9 gemma-4-12B-it/tokenizer_config.json -dd9ce6c7f2776f93acd7091dcd7c60eb0460a7d18e8c655e44e11846fc5d5df6 launch.py +22186ea9c0c9fb94823f666c740555459728d5750278315ac2ab364ccbb4cbfb launch.py cc507fe91c3279813aa4bdf697e7a2a811d5f67ea25afe90206830ec4c3db9cd native_checkpoint.py 3aa234a04e768e734a8ede9cdd9da1716366b06ea47b22a5ca466a49ad4fb56f provenance/build_native_checkpoint.py -1d0b9b438090dbbc63e1e40d5ef8e5a2cd0c5d8f91356d1d21c08db7bb4dea08 release-manifest.json -30578daa4ad6a0e529ac99470c5131c7738ed2a60e70fbaedb32ac0de5736d0a reproduction.json -782cab35025b19b55fe2212a2619cd97dfd5f6da6c04a00e68b9c69a8f92d2dd runtime-image.tar.gz -48fdc1e0b7cb088fcd12c56de79300f90cb6a5a79a49e029f3ec6deec565b66a runtime-release.json -d9e47832fb3354959883d3069a9f1c7e2cc8be0a6e2b9d37f22c36c154749fdd runtime/Dockerfile.fast1 +8d60ffcd6654cce863901b4f0a5f0d5fc25b46d0fa65c58911a1832486460f04 release-manifest.json +2f916687347fd9b47ddb7dc18585a80a7285aacf4bb727db4fe8b52218d203fa reproduction.json +59074c4431d616cd56e0a001684f320825cb3f5ade7d16d6e3e58667d127e6ce runtime-image.tar.gz +7d4c57f97b7ab4f2f353b4e4dffcd56899b24e73a459d5fc14578aaf48ae59e9 runtime-release.json +b99769e8cd55f06640377a142de2199d0a2cda7830081cab5a1974c15915d14d runtime/Dockerfile.fast2 46eb8d5cca3549b243d6014e9aa6eca8e1d7ac940c32d423f480eca69b96290c runtime/Dockerfile.p2b 91fdd9a2cae0bffb6715422bfda322e47e9b05d40d99e3a6cc962aba89132b9e runtime/Dockerfile.p2c -abaee1bb49548bf7b8abebe6519e68d9d7bc3fa9e86e3ef76f1ea9a86e64337c runtime/Dockerfile.p3c +70c93a67b96223c1f0bcfdee039ba1ffb2e17ed1725678d9bad98fb1dd19ee18 runtime/Dockerfile.p3d +005744b93e3b95864bd6152c3084226c448b48aa124311855fa4896dbfaa6fe6 runtime/Dockerfile.ttnn-dn1 93ce271364dd8bac7ccd0e01e1d20fa3f950ce279ca076b6a272ee3ea039073a runtime/build.sh 8a8a3b18e19e275cd5039ea176d0931a6a5f3797985472c8d03366613b99ecc2 runtime/gemma4/README.md 501d5b563ca412fa21db82bf39117b5480ecf99bcaf11996b9c3d5f1093bb9a8 runtime/gemma4/__init__.py @@ -679,15 +707,15 @@ adbf13ab702a97058e31b3822d2bede9aac7f76dd109c86286061e2522f4e392 runtime/gemma4 0e3247c34a359a73bd4c8ecc7365e028372084554eeeba41682d1fd818c296af runtime/gemma4/tt/assistant/model.py df24b2bf4311ffaee781692ff8f465dea0f415e2323cf87fba4ce22227edb8ae runtime/gemma4/tt/attention/__init__.py 01b0f3e14c31aae5ed01104aee7f4cb03913007c0d0018f9d35d341d3fe049ef runtime/gemma4/tt/attention/config.py -0ec303ff75ded2a03f773c22aba27c719e2f055c71b9560fe18661d0272fd3db runtime/gemma4/tt/attention/decode.py +df1f33dce05739d867ff0f4f11facabfa5625818118d8cf4d5fc8e97df36c4ac runtime/gemma4/tt/attention/decode.py 3a1e9f520cdde61dd49550dbfca1072d009ee34878bafe5923c655821c09bff7 runtime/gemma4/tt/attention/kv_cache.py 9074a2d1de4b3a65d11271370f1fce39a78e5523b841752bcaa87dccb4b56470 runtime/gemma4/tt/attention/kv_cache_hybrid.py -f5cda2ade4031a66e631a732199f3b4b9a00b45ef2e252aee0d3138406bbcc43 runtime/gemma4/tt/attention/operations.py +8e87380345b46db028f882df1df2ea05b59a4e0278ed8516a1f15b9663384a65 runtime/gemma4/tt/attention/operations.py 32c6e1c9423b9c22f14558c33d76d9164cb8a27c939cb5d1ccb6da039b7c831f runtime/gemma4/tt/attention/prefill.py 3ec94d60c6d1085bed0bf83fbc450aa9599a476341d3c6d77466b1b2e17c00d6 runtime/gemma4/tt/attention/weights.py 3580f2e25950381465e55dd2886d2832c33d8d6fc618b5a2b94088d0b1a91105 runtime/gemma4/tt/ccl.py e4369b067c3c2e7930fb538a34ab04b237940048e771f2f8f0560776f03a578f runtime/gemma4/tt/common.py -93ebedf2eb3b0c78ebbddfc00c630b13c726d9b13bb07fb53a073a207731bae4 runtime/gemma4/tt/decode_mm.py +85e5a625134c36925d559c4b7557b0139e824a77c4e724bf3b647015b77f7fac runtime/gemma4/tt/decode_mm.py 567aef533a30a692e4c0b28695c15d13e2cfc024ee621e4147c836b63220d5c3 runtime/gemma4/tt/experts/__init__.py 80bd7d18ffd0e0b6bd969d05a8e9972ff8473dc63631f1cb6de0bd619a50f33b runtime/gemma4/tt/experts/config.py 112b9df4b88e8395053fdbf69315e3a43e9f3f4877e699ca5a48462ccd8e804b runtime/gemma4/tt/experts/decode.py @@ -727,9 +755,11 @@ fb0ab8fe2667f410bbe1afa8f09d3a3bbd2d159908d39ceccc4dd07068ccbd05 runtime/servin 0925cd2054f881f59625c2cc31ddb8a8f7dc29a661acf720f303a2cb50366705 runtime/serving-overlay/home/container_app_user/vllm/vllm/config/speculative.py d0ab1698eab1d59203c777555e4457a6e11b3d141c898ffe2befcf8b972dbf65 runtime/serving-overlay/home/container_app_user/vllm/vllm/v1/engine/detokenizer.py e7e004902362d832cfea5f4e22eedcc530494c3e42ea1e720159712ae6d8e1d3 runtime/serving-overlay/home/container_app_user/vllm/vllm/v1/engine/output_processor.py -03bfbc377872223a5aa72bee2148b941831969065e9241aa5dedd0d8ddb36c61 runtime/source.json +dc0a865c919aecaee44bc113cced7f3be2d64172505f0a5c5fc166dcd8fb57f6 runtime/source.json 288c44040187729651e90de21e4d270750eda3d05be109a699404327738e8e4a runtime/tools/compare_gen.py cc507fe91c3279813aa4bdf697e7a2a811d5f67ea25afe90206830ec4c3db9cd runtime/tools/native_checkpoint.py -36485b68752f51c3642e3dd777f5da0cdb9de61ab819d5680f95402e1a24ede3 runtime/tools/tt_eval.py -1de00fc9d9d9ca537d7a90af22dd61b4723c54b6ecd236f7053124963789a62b runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/factory/matmul_multicore_reuse_mcast_1d_program_factory.cpp +505baaad9c88ccfb2e4f264f557477b30ae066673837b6c7fe56f0dc197f0934 runtime/tools/tt_eval.py +fa967085d4bccc78e448454911c38468415c1403e3138e11dbbf9d176ffd3c67 runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/factory/matmul_multicore_reuse_mcast_1d_program_factory.cpp +a78857488315cac5257d5e8a93f03800bbf40acb528ba0edd0ad67660a2cbad1 runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in0_sender_padding.cpp +3620262683632718607446e2a8eba75ea74cd714eee00886d3b2a00ad4f08dc4 runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in1_sender_writer_padding.cpp a9ae81f0f28d1f5c7c093817485cd72d28df4072714a1a30672c43e0166a88d9 serve_native.py diff --git a/evidence/http-p3d/exact.json b/evidence/http-p3d/exact.json new file mode 100644 index 0000000000000000000000000000000000000000..24183813f16a272927c6f1aca3cc326a6da298a3 --- /dev/null +++ b/evidence/http-p3d/exact.json @@ -0,0 +1,185 @@ +{ + "checks": [ + { + "name": "health", + "ok": true + }, + { + "name": "chat_exact_story", + "ok": true, + "n": 250, + "ref_n": 250, + "finish": "stop", + "events": 110, + "first_divergence": null + }, + { + "name": "chat_exact_explain", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 107, + "first_divergence": null + }, + { + "name": "chat_exact_code", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 86, + "first_divergence": null + }, + { + "name": "chat_exact_math", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 67, + "first_divergence": null + }, + { + "name": "chat_exact_logic", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 94, + "first_divergence": null + }, + { + "name": "chat_exact_code-lru", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 86, + "first_divergence": null + }, + { + "name": "chat_exact_code-c", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 74, + "first_divergence": null + }, + { + "name": "chat_exact_code-sql", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 91, + "first_divergence": null + }, + { + "name": "chat_exact_long-summary", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 115, + "first_divergence": null + }, + { + "name": "chat_exact_long-code", + "ok": true, + "n": 320, + "ref_n": 320, + "finish": "length", + "events": 106, + "first_divergence": null + }, + { + "name": "completion_exact_story", + "ok": true, + "n": 250 + }, + { + "name": "completion_exact_explain", + "ok": true, + "n": 320 + }, + { + "name": "completion_exact_code", + "ok": true, + "n": 320 + }, + { + "name": "completion_exact_math", + "ok": true, + "n": 320 + }, + { + "name": "completion_exact_logic", + "ok": true, + "n": 320 + }, + { + "name": "completion_exact_code-lru", + "ok": true, + "n": 320 + }, + { + "name": "completion_exact_code-c", + "ok": true, + "n": 320 + }, + { + "name": "completion_exact_code-sql", + "ok": true, + "n": 320 + }, + { + "name": "completion_exact_long-summary", + "ok": true, + "n": 320 + }, + { + "name": "completion_exact_long-code", + "ok": true, + "n": 320 + }, + { + "name": "isolation_ABA", + "ok": true, + "n": 128 + }, + { + "name": "sampled_request", + "ok": true, + "n": 64 + }, + { + "name": "greedy_after_sampled", + "ok": true + }, + { + "name": "cancel_mid_decode", + "ok": true, + "cancelled_after": 42, + "next_request_s": 2.9789515789889265 + }, + { + "name": "cancel_mid_prefill", + "ok": true, + "cancelled_after_s": 8.169508557009976, + "next_request_s": 11.751794009993318 + }, + { + "name": "image_rejected", + "ok": true, + "status": 400, + "body": "{\"error\":{\"message\":\"/model is not a multimodal model\",\"type\":\"BadRequestError\",\"param\":null,\"code\":400}}" + }, + { + "name": "health_end", + "ok": true + } + ], + "exact_seconds": 100.73274602601305 +} \ No newline at end of file diff --git a/evidence/http-p3d/long.json b/evidence/http-p3d/long.json new file mode 100644 index 0000000000000000000000000000000000000000..a9baddaafbac842403086ce5b13f90fb155635d9 --- /dev/null +++ b/evidence/http-p3d/long.json @@ -0,0 +1,231 @@ +[ + { + "name": "configuration", + "url": "http://127.0.0.1:8000", + "model": "google/gemma-4-12B-it", + "contexts": [ + 32768, + 131072, + 262016 + ], + "tokens": 128, + "native_max_len": 262144, + "request_timeout_seconds": 3600, + "cases": [ + { + "input_tokens": 32768, + "output_tokens": 128 + }, + { + "input_tokens": 131072, + "output_tokens": 128 + }, + { + "input_tokens": 262016, + "output_tokens": 128 + } + ], + "begin_key_token_region": [ + 0, + 43 + ], + "end_key_suffix_tokens": 65, + "timing_basis": "Client first/last non-empty SSE text arrival; MTP may emit multiple tokens per event" + }, + { + "name": "short_before", + "input_tokens": 25, + "requested_output_tokens": 32, + "requested_total_tokens": 57, + "ttft_seconds": 0.09304460400016978, + "decode_wall_seconds": 0.04858210598467849, + "seconds": 0.18983703898265958, + "decode_tokens_per_second": 144.08597276963692, + "finish_reason": "stop", + "sse_done": true, + "response": { + "text": "The capital of France is Paris.", + "usage": { + "prompt_tokens": 25, + "total_tokens": 33, + "completion_tokens": 8 + } + }, + "expected_keys": { + "known_answer": "Paris" + }, + "retrieved_keys": { + "known_answer": true + }, + "accuracy": true + }, + { + "name": "layout_32768", + "input_tokens": 32768, + "begin_key_offset": 35, + "middle_key_offset": 16021, + "end_key_offset": 32707 + }, + { + "name": "stream_32768_128", + "input_tokens": 32768, + "requested_output_tokens": 128, + "requested_total_tokens": 32896, + "ttft_seconds": 16.924654099013424, + "decode_wall_seconds": 0.43388441897695884, + "seconds": 17.35861225798726, + "decode_tokens_per_second": 78.36188282623154, + "finish_reason": "stop", + "sse_done": true, + "response": { + "text": "CEDAR-7419-QUARTZ\nHARBOR-5096-VIOLET\nORBIT-2863-MARBLE", + "usage": { + "prompt_tokens": 32768, + "total_tokens": 32803, + "completion_tokens": 35 + } + }, + "expected_keys": { + "begin": "CEDAR-7419-QUARTZ", + "middle": "HARBOR-5096-VIOLET", + "end": "ORBIT-2863-MARBLE" + }, + "retrieved_keys": { + "begin": true, + "middle": true, + "end": true + }, + "accuracy": true + }, + { + "name": "layout_131072", + "input_tokens": 131072, + "begin_key_offset": 35, + "middle_key_offset": 65832, + "end_key_offset": 131011 + }, + { + "name": "stream_131072_128", + "input_tokens": 131072, + "requested_output_tokens": 128, + "requested_total_tokens": 131200, + "ttft_seconds": 123.8322498589987, + "decode_wall_seconds": 0.5941902149934322, + "seconds": 124.42653353299829, + "decode_tokens_per_second": 57.22073360022567, + "finish_reason": "stop", + "sse_done": true, + "response": { + "text": "CEDAR-7419-QUARTZ\nHARBOR-5096-VIOLET\nORBIT-2863-MARBLE", + "usage": { + "prompt_tokens": 131072, + "total_tokens": 131107, + "completion_tokens": 35 + } + }, + "expected_keys": { + "begin": "CEDAR-7419-QUARTZ", + "middle": "HARBOR-5096-VIOLET", + "end": "ORBIT-2863-MARBLE" + }, + "retrieved_keys": { + "begin": true, + "middle": true, + "end": true + }, + "accuracy": true + }, + { + "name": "layout_262016", + "input_tokens": 262016, + "begin_key_offset": 35, + "middle_key_offset": 130632, + "end_key_offset": 261955 + }, + { + "name": "stream_262016_128", + "input_tokens": 262016, + "requested_output_tokens": 128, + "requested_total_tokens": 262144, + "ttft_seconds": 394.9077806400019, + "decode_wall_seconds": 0.9461079349857755, + "seconds": 395.9787403039809, + "decode_tokens_per_second": 35.9367031421327, + "finish_reason": "stop", + "sse_done": true, + "response": { + "text": "CEDAR-7419-QUARTZ\nHARBOR-5096-VIOLET\nORBIT-2863-MARBLE", + "usage": { + "prompt_tokens": 262016, + "total_tokens": 262051, + "completion_tokens": 35 + } + }, + "expected_keys": { + "begin": "CEDAR-7419-QUARTZ", + "middle": "HARBOR-5096-VIOLET", + "end": "ORBIT-2863-MARBLE" + }, + "retrieved_keys": { + "begin": true, + "middle": true, + "end": true + }, + "accuracy": true + }, + { + "name": "native_overflow_rejection", + "input_tokens": 262016, + "requested_output_tokens": 129, + "requested_total_tokens": 262145, + "native_max_len": 262144, + "status": 400, + "seconds": 0.014746526983799413, + "rejected_before_generation": true, + "response": { + "error": { + "message": "You passed 262016 input tokens and requested 129 output tokens. However, the model's context length is only 262144 tokens, resulting in a maximum input length of 262015 tokens. Please reduce the length of the input prompt. (parameter=input_tokens, value=262016)", + "type": "BadRequestError", + "param": "input_tokens", + "code": 400 + } + } + }, + { + "name": "mtp_acceptance", + "accepted_tokens_before": 10881.0, + "accepted_tokens_after": 10962.0, + "accepted_tokens_delta": 81.0 + }, + { + "name": "short_after", + "input_tokens": 25, + "requested_output_tokens": 32, + "requested_total_tokens": 57, + "ttft_seconds": 0.09321309201186523, + "decode_wall_seconds": 0.04820338898571208, + "seconds": 0.18969050902524032, + "decode_tokens_per_second": 145.21800535798144, + "finish_reason": "stop", + "sse_done": true, + "response": { + "text": "The capital of France is Paris.", + "usage": { + "prompt_tokens": 25, + "total_tokens": 33, + "completion_tokens": 8 + } + }, + "expected_keys": { + "known_answer": "Paris" + }, + "retrieved_keys": { + "known_answer": true + }, + "accuracy": true + }, + { + "name": "state_isolation", + "exact_text_match": true + } +] diff --git a/evidence/http-p3d/tput-128/requests-128-c1.json b/evidence/http-p3d/tput-128/requests-128-c1.json new file mode 100644 index 0000000000000000000000000000000000000000..7047bbe96af104a8e71795d1a519df1f1b810bb7 --- /dev/null +++ b/evidence/http-p3d/tput-128/requests-128-c1.json @@ -0,0 +1,34 @@ +[ + { + "index": 0, + "start_seconds": 1.8349994206801057e-05, + "end_seconds": 8.190938635001658, + "ttft_seconds": 0.09066972200525925, + "latency_seconds": 8.19092028500745, + "decode_tokens_per_second": 63.08490319983218, + "text_sha256": "8b03541774e1f6be7b4c4473ed2e299b30efe7d684e4e280632204176904b0e4", + "usage": { + "prompt_tokens": 128, + "total_tokens": 640, + "completion_tokens": 512 + }, + "error": null, + "text": "# Understanding Database Transactions: The Foundation of Data Integrity\n\nIn the realm of database management systems (DBMS), a **transaction** is the fundamental unit of work. It is a sequence of one or more operations (such as inserting, updating, deleting, or querying data) that are executed as a single logical unit. The primary purpose of a transaction is to ensure that even in the event of system failures, power outages, or concurrent access by multiple users, the database remains in a consistent and reliable state.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the lifecycle of a transaction, concurrency control, and practical real-world examples.\n\n---\n\n### 1. The ACID Properties: The Gold Standard\nFor a database to be considered reliable, every transaction must adhere to the ACID properties. These four principles ensure that data remains accurate and protected.\n\n#### A - Atomicity (\"All or Nothing\")\nAtomicity dictates that a transaction must be treated as an indivisible unit. Either every single operation within the transaction succeeds, or none of them do. If any part of the transaction fails (due to a crash, a constraint violation, or a power failure), the entire transaction is \"rolled back,\" and the database is returned to its state prior to the start of the transaction.\n\n* **Example:** Imagine a bank transfer where \\$100 is moved from Account A to Account B. This involves two steps: (1) Subtracting \\$100 from A, and (2) Adding \\$100 to B. If the system crashes after step 1 but before step 2, the money would vanish into thin air. Atomicity ensures that if step 2 fails, step 1 is undone.\n\n#### C - Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, foreign keys, and triggers. A transaction cannot leave the database in a \"broken\" state where data violates the business logic.\n\n* **Example:** If a database has a rule that an account balance cannot drop below zero, and a transaction attempts to withdraw \\$500 from an account containing only \\$200, the database will reject the transaction to maintain consistency.\n\n#### I - Isolation\nIn modern systems, thousands of transactions happen simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. The result of running multiple transactions at the same time should be the same as if they were executed sequentially (one after the" + }, + { + "index": 1, + "start_seconds": 8.190989373979392, + "end_seconds": 16.381804777978687, + "ttft_seconds": 0.09145842600264587, + "latency_seconds": 8.190815403999295, + "decode_tokens_per_second": 63.091868561597515, + "text_sha256": "8b03541774e1f6be7b4c4473ed2e299b30efe7d684e4e280632204176904b0e4", + "usage": { + "prompt_tokens": 128, + "total_tokens": 640, + "completion_tokens": 512 + }, + "error": null, + "text": "# Understanding Database Transactions: The Foundation of Data Integrity\n\nIn the realm of database management systems (DBMS), a **transaction** is the fundamental unit of work. It is a sequence of one or more operations (such as inserting, updating, deleting, or querying data) that are executed as a single logical unit. The primary purpose of a transaction is to ensure that even in the event of system failures, power outages, or concurrent access by multiple users, the database remains in a consistent and reliable state.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the lifecycle of a transaction, concurrency control, and practical real-world examples.\n\n---\n\n### 1. The ACID Properties: The Gold Standard\nFor a database to be considered reliable, every transaction must adhere to the ACID properties. These four principles ensure that data remains accurate and protected.\n\n#### A - Atomicity (\"All or Nothing\")\nAtomicity dictates that a transaction must be treated as an indivisible unit. Either every single operation within the transaction succeeds, or none of them do. If any part of the transaction fails (due to a crash, a constraint violation, or a power failure), the entire transaction is \"rolled back,\" and the database is returned to its state prior to the start of the transaction.\n\n* **Example:** Imagine a bank transfer where \\$100 is moved from Account A to Account B. This involves two steps: (1) Subtracting \\$100 from A, and (2) Adding \\$100 to B. If the system crashes after step 1 but before step 2, the money would vanish into thin air. Atomicity ensures that if step 2 fails, step 1 is undone.\n\n#### C - Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, foreign keys, and triggers. A transaction cannot leave the database in a \"broken\" state where data violates the business logic.\n\n* **Example:** If a database has a rule that an account balance cannot drop below zero, and a transaction attempts to withdraw \\$500 from an account containing only \\$200, the database will reject the transaction to maintain consistency.\n\n#### I - Isolation\nIn modern systems, thousands of transactions happen simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. The result of running multiple transactions at the same time should be the same as if they were executed sequentially (one after the" + } +] \ No newline at end of file diff --git a/evidence/http-p3d/tput-128/summary.json b/evidence/http-p3d/tput-128/summary.json new file mode 100644 index 0000000000000000000000000000000000000000..9b0c58ff08fb81043e0371d75bd418bb5ec4678a --- /dev/null +++ b/evidence/http-p3d/tput-128/summary.json @@ -0,0 +1,24 @@ +{ + "model": "google/gemma-4-12B-it", + "output_tokens": 512, + "method": "Closed loop; synchronized initial clients; max(min_requests,2*concurrency) requests; nearest-rank percentiles; aggregate includes prefill and queue drain", + "cases": [ + { + "input_tokens": 128, + "concurrency": 1, + "requests": 2, + "successes": 2, + "errors": 0, + "wall_seconds": 16.38178642798448, + "matching_single_request_outputs": 2, + "aggregate_output_tokens_per_second": 62.508445248116146, + "requests_per_second": 0.12208680712522685, + "ttft_seconds_p50": 0.09066972200525925, + "ttft_seconds_p95": 0.09145842600264587, + "latency_seconds_p50": 8.190815403999295, + "latency_seconds_p95": 8.19092028500745, + "decode_tokens_per_second_p50": 63.08490319983218, + "decode_tokens_per_second_p95": 63.091868561597515 + } + ] +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-128/warmup-128.json b/evidence/http-p3d/tput-128/warmup-128.json new file mode 100644 index 0000000000000000000000000000000000000000..28542025f2d7d7cde1c7e8999adb19bd80b7d71e --- /dev/null +++ b/evidence/http-p3d/tput-128/warmup-128.json @@ -0,0 +1,16 @@ +{ + "index": -1, + "start_seconds": 3.100140020251274e-07, + "end_seconds": 8.237141543999314, + "ttft_seconds": 0.09195410198299214, + "latency_seconds": 8.237141233985312, + "decode_tokens_per_second": 62.73699920467591, + "text_sha256": "8b03541774e1f6be7b4c4473ed2e299b30efe7d684e4e280632204176904b0e4", + "usage": { + "prompt_tokens": 128, + "total_tokens": 640, + "completion_tokens": 512 + }, + "error": null, + "text": "# Understanding Database Transactions: The Foundation of Data Integrity\n\nIn the realm of database management systems (DBMS), a **transaction** is the fundamental unit of work. It is a sequence of one or more operations (such as inserting, updating, deleting, or querying data) that are executed as a single logical unit. The primary purpose of a transaction is to ensure that even in the event of system failures, power outages, or concurrent access by multiple users, the database remains in a consistent and reliable state.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the lifecycle of a transaction, concurrency control, and practical real-world examples.\n\n---\n\n### 1. The ACID Properties: The Gold Standard\nFor a database to be considered reliable, every transaction must adhere to the ACID properties. These four principles ensure that data remains accurate and protected.\n\n#### A - Atomicity (\"All or Nothing\")\nAtomicity dictates that a transaction must be treated as an indivisible unit. Either every single operation within the transaction succeeds, or none of them do. If any part of the transaction fails (due to a crash, a constraint violation, or a power failure), the entire transaction is \"rolled back,\" and the database is returned to its state prior to the start of the transaction.\n\n* **Example:** Imagine a bank transfer where \\$100 is moved from Account A to Account B. This involves two steps: (1) Subtracting \\$100 from A, and (2) Adding \\$100 to B. If the system crashes after step 1 but before step 2, the money would vanish into thin air. Atomicity ensures that if step 2 fails, step 1 is undone.\n\n#### C - Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, foreign keys, and triggers. A transaction cannot leave the database in a \"broken\" state where data violates the business logic.\n\n* **Example:** If a database has a rule that an account balance cannot drop below zero, and a transaction attempts to withdraw \\$500 from an account containing only \\$200, the database will reject the transaction to maintain consistency.\n\n#### I - Isolation\nIn modern systems, thousands of transactions happen simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. The result of running multiple transactions at the same time should be the same as if they were executed sequentially (one after the" +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-131072/requests-131072-c1.json b/evidence/http-p3d/tput-131072/requests-131072-c1.json new file mode 100644 index 0000000000000000000000000000000000000000..6ce1ec63368875d4982d47b6703eb79dbe1c3a88 --- /dev/null +++ b/evidence/http-p3d/tput-131072/requests-131072-c1.json @@ -0,0 +1,34 @@ +[ + { + "index": 0, + "start_seconds": 3.928999649360776e-05, + "end_seconds": 140.3550663930073, + "ttft_seconds": 123.82656395100639, + "latency_seconds": 140.3550271030108, + "decode_tokens_per_second": 30.91648280188757, + "text_sha256": "372e489ce2a70556e81cd2e3a391d7f7ae934471533aa26b7720c790c8d94c3e", + "usage": { + "prompt_tokens": 131072, + "total_tokens": 131584, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is the fundamental unit of work. It is a sequence of one or more operations performed on a database that must be treated as a single, indivisible logical unit. The core philosophy behind transactions is simple: either the entire sequence of operations succeeds, or none of them do. This prevents the database from ending up in a \"half-baked\" or corrupted state where some data is updated while related data remains unchanged.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the **transaction lifecycle**, and real-world **concurrency scenarios**.\n\n---\n\n### 1. The Pillars of Transactions: ACID Properties\n\nThe reliability of any relational database management system (RDBMS) like MySQL, PostgreSQL, or Oracle rests on the ACID acronym. These four properties ensure that even in the event of a power failure, system crash, or simultaneous access by thousands of users, the data remains accurate.\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity guarantees that a transaction is treated as a single \"atom.\" If a transaction consists of five different SQL statements (e.g., deducting money from one account, adding it to another, and logging the history), and the system crashes after the third statement, the database must \"roll back\" the first two. The database should look as if nothing ever happened.\n* **Analogy:** Think of an ATM withdrawal. If the machine gives you the cash but fails to update your balance, the transaction failed. Atomicity ensures that if the balance isn't updated, the cash isn't dispensed.\n\n#### B. Consistency (\"Follow the Rules\")\nConsistency ensures that a transaction brings the database from one valid state to another. Every transaction must follow the predefined rules of the database, such as constraints (e.g., \"Age must be over 18\"), triggers, and cascades. If a transaction would result in a violation of these rules, the database will reject it.\n* **Example:** If a database rule states that a bank account balance cannot be negative, and a transaction tries to withdraw more money than is available, the transaction is aborted to maintain consistency.\n\n#### C. Isolation (\"Don't Peek\")\nIsolation ensures that concurrent transactions (multiple people using the database at the same time) do not interfere with each other. Even though many transactions might be running simultaneously, the result should be the same as if they were executed one" + }, + { + "index": 1, + "start_seconds": 140.35510921300738, + "end_seconds": 280.67249108300894, + "ttft_seconds": 123.78938921200461, + "latency_seconds": 140.31738187000155, + "decode_tokens_per_second": 30.91739205383985, + "text_sha256": "372e489ce2a70556e81cd2e3a391d7f7ae934471533aa26b7720c790c8d94c3e", + "usage": { + "prompt_tokens": 131072, + "total_tokens": 131584, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is the fundamental unit of work. It is a sequence of one or more operations performed on a database that must be treated as a single, indivisible logical unit. The core philosophy behind transactions is simple: either the entire sequence of operations succeeds, or none of them do. This prevents the database from ending up in a \"half-baked\" or corrupted state where some data is updated while related data remains unchanged.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the **transaction lifecycle**, and real-world **concurrency scenarios**.\n\n---\n\n### 1. The Pillars of Transactions: ACID Properties\n\nThe reliability of any relational database management system (RDBMS) like MySQL, PostgreSQL, or Oracle rests on the ACID acronym. These four properties ensure that even in the event of a power failure, system crash, or simultaneous access by thousands of users, the data remains accurate.\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity guarantees that a transaction is treated as a single \"atom.\" If a transaction consists of five different SQL statements (e.g., deducting money from one account, adding it to another, and logging the history), and the system crashes after the third statement, the database must \"roll back\" the first two. The database should look as if nothing ever happened.\n* **Analogy:** Think of an ATM withdrawal. If the machine gives you the cash but fails to update your balance, the transaction failed. Atomicity ensures that if the balance isn't updated, the cash isn't dispensed.\n\n#### B. Consistency (\"Follow the Rules\")\nConsistency ensures that a transaction brings the database from one valid state to another. Every transaction must follow the predefined rules of the database, such as constraints (e.g., \"Age must be over 18\"), triggers, and cascades. If a transaction would result in a violation of these rules, the database will reject it.\n* **Example:** If a database rule states that a bank account balance cannot be negative, and a transaction tries to withdraw more money than is available, the transaction is aborted to maintain consistency.\n\n#### C. Isolation (\"Don't Peek\")\nIsolation ensures that concurrent transactions (multiple people using the database at the same time) do not interfere with each other. Even though many transactions might be running simultaneously, the result should be the same as if they were executed one" + } +] \ No newline at end of file diff --git a/evidence/http-p3d/tput-131072/summary.json b/evidence/http-p3d/tput-131072/summary.json new file mode 100644 index 0000000000000000000000000000000000000000..292bda4573b8168c8dd2473e40f7469659109a20 --- /dev/null +++ b/evidence/http-p3d/tput-131072/summary.json @@ -0,0 +1,24 @@ +{ + "model": "google/gemma-4-12B-it", + "output_tokens": 512, + "method": "Closed loop; synchronized initial clients; max(min_requests,2*concurrency) requests; nearest-rank percentiles; aggregate includes prefill and queue drain", + "cases": [ + { + "input_tokens": 131072, + "concurrency": 1, + "requests": 2, + "successes": 2, + "errors": 0, + "wall_seconds": 280.67245179301244, + "matching_single_request_outputs": 2, + "aggregate_output_tokens_per_second": 3.6483808562557805, + "requests_per_second": 0.007125743859874571, + "ttft_seconds_p50": 123.78938921200461, + "ttft_seconds_p95": 123.82656395100639, + "latency_seconds_p50": 140.31738187000155, + "latency_seconds_p95": 140.3550271030108, + "decode_tokens_per_second_p50": 30.91648280188757, + "decode_tokens_per_second_p95": 30.91739205383985 + } + ] +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-131072/warmup-131072.json b/evidence/http-p3d/tput-131072/warmup-131072.json new file mode 100644 index 0000000000000000000000000000000000000000..5834d02f9c9fa7066ebd6db384045e13e88446de --- /dev/null +++ b/evidence/http-p3d/tput-131072/warmup-131072.json @@ -0,0 +1,16 @@ +{ + "index": -1, + "start_seconds": 1.00000761449337e-06, + "end_seconds": 143.09652435599128, + "ttft_seconds": 126.40863994599204, + "latency_seconds": 143.09652335598366, + "decode_tokens_per_second": 30.621160536435426, + "text_sha256": "372e489ce2a70556e81cd2e3a391d7f7ae934471533aa26b7720c790c8d94c3e", + "usage": { + "prompt_tokens": 131072, + "total_tokens": 131584, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is the fundamental unit of work. It is a sequence of one or more operations performed on a database that must be treated as a single, indivisible logical unit. The core philosophy behind transactions is simple: either the entire sequence of operations succeeds, or none of them do. This prevents the database from ending up in a \"half-baked\" or corrupted state where some data is updated while related data remains unchanged.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the **transaction lifecycle**, and real-world **concurrency scenarios**.\n\n---\n\n### 1. The Pillars of Transactions: ACID Properties\n\nThe reliability of any relational database management system (RDBMS) like MySQL, PostgreSQL, or Oracle rests on the ACID acronym. These four properties ensure that even in the event of a power failure, system crash, or simultaneous access by thousands of users, the data remains accurate.\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity guarantees that a transaction is treated as a single \"atom.\" If a transaction consists of five different SQL statements (e.g., deducting money from one account, adding it to another, and logging the history), and the system crashes after the third statement, the database must \"roll back\" the first two. The database should look as if nothing ever happened.\n* **Analogy:** Think of an ATM withdrawal. If the machine gives you the cash but fails to update your balance, the transaction failed. Atomicity ensures that if the balance isn't updated, the cash isn't dispensed.\n\n#### B. Consistency (\"Follow the Rules\")\nConsistency ensures that a transaction brings the database from one valid state to another. Every transaction must follow the predefined rules of the database, such as constraints (e.g., \"Age must be over 18\"), triggers, and cascades. If a transaction would result in a violation of these rules, the database will reject it.\n* **Example:** If a database rule states that a bank account balance cannot be negative, and a transaction tries to withdraw more money than is available, the transaction is aborted to maintain consistency.\n\n#### C. Isolation (\"Don't Peek\")\nIsolation ensures that concurrent transactions (multiple people using the database at the same time) do not interfere with each other. Even though many transactions might be running simultaneously, the result should be the same as if they were executed one" +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-2048/requests-2048-c1.json b/evidence/http-p3d/tput-2048/requests-2048-c1.json new file mode 100644 index 0000000000000000000000000000000000000000..6ce44410471f318dcb78253d00151291aebbe8d2 --- /dev/null +++ b/evidence/http-p3d/tput-2048/requests-2048-c1.json @@ -0,0 +1,34 @@ +[ + { + "index": 0, + "start_seconds": 1.8780003301799297e-05, + "end_seconds": 10.399938565999037, + "ttft_seconds": 0.6453376279969234, + "latency_seconds": 10.399919785995735, + "decode_tokens_per_second": 52.386027641385155, + "text_sha256": "02f0d03a55e50a19e8fd9e605f9812f55ef641b9e2ebacc042e79c1c4837431c", + "usage": { + "prompt_tokens": 2048, + "total_tokens": 2560, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is a fundamental concept that ensures data integrity and reliability. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and the changes are permanently saved, or if any single part fails, the entire transaction is undone, leaving the database in its original state.\n\nTo understand why this is critical, we must look at the **ACID properties**, the mechanisms that govern transactions, and real-world applications.\n\n---\n\n### 1. The ACID Properties\nThe reliability of a database system is measured by its adherence to the ACID model. These four properties ensure that even in the event of system crashes, power failures, or concurrent access by thousands of users, the data remains accurate.\n\n#### A. Atomicity\nAtomicity treats a transaction as an indivisible unit. Think of it like a physical atom\u2014it cannot be split. If a transaction involves five different SQL queries (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), atomicity guarantees that if the fourth query fails, the first three are \"rolled back\" (undone).\n* **Example:** If you are transferring money from Bank Account A to Bank Account B, two things must happen: money must be subtracted from A, and money must be added to B. If the system crashes after subtracting from A but before adding to B, the money would vanish. Atomicity prevents this by ensuring that if the addition to B fails, the subtraction from A is reversed.\n\n#### B. Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, cascades, and triggers. A transaction must not violate the \"rules\" of the database.\n* **Example:** Imagine a database rule that says \"Account balances cannot be negative.\" If a transaction attempts to withdraw \\$100 from an account that only has \\$50, the database will detect a violation of the consistency rule and abort the transaction.\n\n#### C. Isolation\nIn modern systems, thousands of transactions happen simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. Even though they are running at the same time, the result should be the same as if they were executed one after another (sequentially).\n* **" + }, + { + "index": 1, + "start_seconds": 10.399984395015053, + "end_seconds": 20.804596664005658, + "ttft_seconds": 0.6470787139842287, + "latency_seconds": 10.404612268990604, + "decode_tokens_per_second": 52.37015888910936, + "text_sha256": "02f0d03a55e50a19e8fd9e605f9812f55ef641b9e2ebacc042e79c1c4837431c", + "usage": { + "prompt_tokens": 2048, + "total_tokens": 2560, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is a fundamental concept that ensures data integrity and reliability. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and the changes are permanently saved, or if any single part fails, the entire transaction is undone, leaving the database in its original state.\n\nTo understand why this is critical, we must look at the **ACID properties**, the mechanisms that govern transactions, and real-world applications.\n\n---\n\n### 1. The ACID Properties\nThe reliability of a database system is measured by its adherence to the ACID model. These four properties ensure that even in the event of system crashes, power failures, or concurrent access by thousands of users, the data remains accurate.\n\n#### A. Atomicity\nAtomicity treats a transaction as an indivisible unit. Think of it like a physical atom\u2014it cannot be split. If a transaction involves five different SQL queries (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), atomicity guarantees that if the fourth query fails, the first three are \"rolled back\" (undone).\n* **Example:** If you are transferring money from Bank Account A to Bank Account B, two things must happen: money must be subtracted from A, and money must be added to B. If the system crashes after subtracting from A but before adding to B, the money would vanish. Atomicity prevents this by ensuring that if the addition to B fails, the subtraction from A is reversed.\n\n#### B. Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, cascades, and triggers. A transaction must not violate the \"rules\" of the database.\n* **Example:** Imagine a database rule that says \"Account balances cannot be negative.\" If a transaction attempts to withdraw \\$100 from an account that only has \\$50, the database will detect a violation of the consistency rule and abort the transaction.\n\n#### C. Isolation\nIn modern systems, thousands of transactions happen simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. Even though they are running at the same time, the result should be the same as if they were executed one after another (sequentially).\n* **" + } +] \ No newline at end of file diff --git a/evidence/http-p3d/tput-2048/summary.json b/evidence/http-p3d/tput-2048/summary.json new file mode 100644 index 0000000000000000000000000000000000000000..41830f19407e797ed3f25a36fa602872ab7f90a3 --- /dev/null +++ b/evidence/http-p3d/tput-2048/summary.json @@ -0,0 +1,24 @@ +{ + "model": "google/gemma-4-12B-it", + "output_tokens": 512, + "method": "Closed loop; synchronized initial clients; max(min_requests,2*concurrency) requests; nearest-rank percentiles; aggregate includes prefill and queue drain", + "cases": [ + { + "input_tokens": 2048, + "concurrency": 1, + "requests": 2, + "successes": 2, + "errors": 0, + "wall_seconds": 20.804577884002356, + "matching_single_request_outputs": 2, + "aggregate_output_tokens_per_second": 49.21993638656822, + "requests_per_second": 0.09613268825501606, + "ttft_seconds_p50": 0.6453376279969234, + "ttft_seconds_p95": 0.6470787139842287, + "latency_seconds_p50": 10.399919785995735, + "latency_seconds_p95": 10.404612268990604, + "decode_tokens_per_second_p50": 52.37015888910936, + "decode_tokens_per_second_p95": 52.386027641385155 + } + ] +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-2048/warmup-2048.json b/evidence/http-p3d/tput-2048/warmup-2048.json new file mode 100644 index 0000000000000000000000000000000000000000..64e4a73235ba276ba71be943c0dd2ed850fbd73e --- /dev/null +++ b/evidence/http-p3d/tput-2048/warmup-2048.json @@ -0,0 +1,16 @@ +{ + "index": -1, + "start_seconds": 4.00003045797348e-07, + "end_seconds": 10.159259158011992, + "ttft_seconds": 0.6454166269977577, + "latency_seconds": 10.159258758008946, + "decode_tokens_per_second": 53.711591469489335, + "text_sha256": "02f0d03a55e50a19e8fd9e605f9812f55ef641b9e2ebacc042e79c1c4837431c", + "usage": { + "prompt_tokens": 2048, + "total_tokens": 2560, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is a fundamental concept that ensures data integrity and reliability. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and the changes are permanently saved, or if any single part fails, the entire transaction is undone, leaving the database in its original state.\n\nTo understand why this is critical, we must look at the **ACID properties**, the mechanisms that govern transactions, and real-world applications.\n\n---\n\n### 1. The ACID Properties\nThe reliability of a database system is measured by its adherence to the ACID model. These four properties ensure that even in the event of system crashes, power failures, or concurrent access by thousands of users, the data remains accurate.\n\n#### A. Atomicity\nAtomicity treats a transaction as an indivisible unit. Think of it like a physical atom\u2014it cannot be split. If a transaction involves five different SQL queries (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), atomicity guarantees that if the fourth query fails, the first three are \"rolled back\" (undone).\n* **Example:** If you are transferring money from Bank Account A to Bank Account B, two things must happen: money must be subtracted from A, and money must be added to B. If the system crashes after subtracting from A but before adding to B, the money would vanish. Atomicity prevents this by ensuring that if the addition to B fails, the subtraction from A is reversed.\n\n#### B. Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, cascades, and triggers. A transaction must not violate the \"rules\" of the database.\n* **Example:** Imagine a database rule that says \"Account balances cannot be negative.\" If a transaction attempts to withdraw \\$100 from an account that only has \\$50, the database will detect a violation of the consistency rule and abort the transaction.\n\n#### C. Isolation\nIn modern systems, thousands of transactions happen simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. Even though they are running at the same time, the result should be the same as if they were executed one after another (sequentially).\n* **" +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-261632/requests-261632-c1.json b/evidence/http-p3d/tput-261632/requests-261632-c1.json new file mode 100644 index 0000000000000000000000000000000000000000..b2f7fad01574e85a5f43c1f8f8f21ec2d1ba7412 --- /dev/null +++ b/evidence/http-p3d/tput-261632/requests-261632-c1.json @@ -0,0 +1,34 @@ +[ + { + "index": 0, + "start_seconds": 1.9529979908838868e-05, + "end_seconds": 415.5554907670012, + "ttft_seconds": 394.7760664890229, + "latency_seconds": 415.55547123702127, + "decode_tokens_per_second": 24.59174523194995, + "text_sha256": "886da530a6b8ed5f967b0ceebe827d949fabefbce66a9a71c5076273db38c59d", + "usage": { + "prompt_tokens": 261632, + "total_tokens": 262144, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is the fundamental unit of work. It is not merely a single action, like inserting one row into a table; rather, it is a logical sequence of operations that must be treated as a single, indivisible unit. To ensure that data remains accurate and reliable\u2014especially in multi-user environments where thousands of people might be accessing the same information simultaneously\u2014databases rely on the concept of transactions.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the **transaction lifecycle**, and **concurrency control mechanisms**.\n\n---\n\n### 1. The Core Philosophy: The ACID Properties\nThe gold standard for any database transaction is the **ACID** acronym. If a system cannot guarantee these four properties, it cannot be considered a reliable relational database management system (RDBMS).\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as a single \"atom.\" Either every single operation within the transaction succeeds, or none of them do. If a transaction involves five steps and the fourth step fails (due to a power outage, a constraint violation, or a network error), the database must \"roll back\" the first three steps so that the data remains in its original state.\n\n* **Example:** Imagine you are transferring \\$100 from your Savings Account to your Checking Account.\n 1. Subtract \\$100 from Savings.\n 2. Add \\$100 to Checking.\n If the system crashes after step 1 but before step 2, your money would simply vanish. Atomicity prevents this by ensuring that if step 2 fails, step 1 is undone.\n\n#### B. Consistency (Rules of the Game)\nConsistency ensures that a transaction brings the database from one valid state to another. Every database has predefined rules (constraints), such as \"account balances cannot be negative\" or \"every order must have a valid Customer ID.\" A transaction must never violate these rules. If a transaction attempts to perform an action that would break a rule, the database will reject the entire transaction.\n\n* **Example:** If you try to transfer \\$500 from an account that only has \\$200, and the database has a \"minimum balance\" constraint, the transaction will fail because it would result in an inconsistent state (a negative balance).\n\n#### C. Isolation (The \"Privacy\" of Transactions)\n" + }, + { + "index": 1, + "start_seconds": 415.55553735699505, + "end_seconds": 831.1073085209937, + "ttft_seconds": 394.76245176198427, + "latency_seconds": 415.55177116399864, + "decode_tokens_per_second": 24.580007387981023, + "text_sha256": "886da530a6b8ed5f967b0ceebe827d949fabefbce66a9a71c5076273db38c59d", + "usage": { + "prompt_tokens": 261632, + "total_tokens": 262144, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is the fundamental unit of work. It is not merely a single action, like inserting one row into a table; rather, it is a logical sequence of operations that must be treated as a single, indivisible unit. To ensure that data remains accurate and reliable\u2014especially in multi-user environments where thousands of people might be accessing the same information simultaneously\u2014databases rely on the concept of transactions.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the **transaction lifecycle**, and **concurrency control mechanisms**.\n\n---\n\n### 1. The Core Philosophy: The ACID Properties\nThe gold standard for any database transaction is the **ACID** acronym. If a system cannot guarantee these four properties, it cannot be considered a reliable relational database management system (RDBMS).\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as a single \"atom.\" Either every single operation within the transaction succeeds, or none of them do. If a transaction involves five steps and the fourth step fails (due to a power outage, a constraint violation, or a network error), the database must \"roll back\" the first three steps so that the data remains in its original state.\n\n* **Example:** Imagine you are transferring \\$100 from your Savings Account to your Checking Account.\n 1. Subtract \\$100 from Savings.\n 2. Add \\$100 to Checking.\n If the system crashes after step 1 but before step 2, your money would simply vanish. Atomicity prevents this by ensuring that if step 2 fails, step 1 is undone.\n\n#### B. Consistency (Rules of the Game)\nConsistency ensures that a transaction brings the database from one valid state to another. Every database has predefined rules (constraints), such as \"account balances cannot be negative\" or \"every order must have a valid Customer ID.\" A transaction must never violate these rules. If a transaction attempts to perform an action that would break a rule, the database will reject the entire transaction.\n\n* **Example:** If you try to transfer \\$500 from an account that only has \\$200, and the database has a \"minimum balance\" constraint, the transaction will fail because it would result in an inconsistent state (a negative balance).\n\n#### C. Isolation (The \"Privacy\" of Transactions)\n" + } +] \ No newline at end of file diff --git a/evidence/http-p3d/tput-261632/summary.json b/evidence/http-p3d/tput-261632/summary.json new file mode 100644 index 0000000000000000000000000000000000000000..2ef5f3d08e3a37c4bab03ea52c729265af8d25eb --- /dev/null +++ b/evidence/http-p3d/tput-261632/summary.json @@ -0,0 +1,24 @@ +{ + "model": "google/gemma-4-12B-it", + "output_tokens": 512, + "method": "Closed loop; synchronized initial clients; max(min_requests,2*concurrency) requests; nearest-rank percentiles; aggregate includes prefill and queue drain", + "cases": [ + { + "input_tokens": 261632, + "concurrency": 1, + "requests": 2, + "successes": 2, + "errors": 0, + "wall_seconds": 831.1072889910138, + "matching_single_request_outputs": 2, + "aggregate_output_tokens_per_second": 1.2320912276478324, + "requests_per_second": 0.0024064281789996727, + "ttft_seconds_p50": 394.76245176198427, + "ttft_seconds_p95": 394.7760664890229, + "latency_seconds_p50": 415.55177116399864, + "latency_seconds_p95": 415.55547123702127, + "decode_tokens_per_second_p50": 24.580007387981023, + "decode_tokens_per_second_p95": 24.59174523194995 + } + ] +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-261632/warmup-261632.json b/evidence/http-p3d/tput-261632/warmup-261632.json new file mode 100644 index 0000000000000000000000000000000000000000..af611c8100ce1049a9d2a32e383379af18215860 --- /dev/null +++ b/evidence/http-p3d/tput-261632/warmup-261632.json @@ -0,0 +1,16 @@ +{ + "index": -1, + "start_seconds": 1.200009137392044e-06, + "end_seconds": 417.3713296529895, + "ttft_seconds": 396.3232168069808, + "latency_seconds": 417.3713284529804, + "decode_tokens_per_second": 24.27779241951352, + "text_sha256": "886da530a6b8ed5f967b0ceebe827d949fabefbce66a9a71c5076273db38c59d", + "usage": { + "prompt_tokens": 261632, + "total_tokens": 262144, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is the fundamental unit of work. It is not merely a single action, like inserting one row into a table; rather, it is a logical sequence of operations that must be treated as a single, indivisible unit. To ensure that data remains accurate and reliable\u2014especially in multi-user environments where thousands of people might be accessing the same information simultaneously\u2014databases rely on the concept of transactions.\n\nTo understand transactions deeply, we must explore the **ACID properties**, the **transaction lifecycle**, and **concurrency control mechanisms**.\n\n---\n\n### 1. The Core Philosophy: The ACID Properties\nThe gold standard for any database transaction is the **ACID** acronym. If a system cannot guarantee these four properties, it cannot be considered a reliable relational database management system (RDBMS).\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as a single \"atom.\" Either every single operation within the transaction succeeds, or none of them do. If a transaction involves five steps and the fourth step fails (due to a power outage, a constraint violation, or a network error), the database must \"roll back\" the first three steps so that the data remains in its original state.\n\n* **Example:** Imagine you are transferring \\$100 from your Savings Account to your Checking Account.\n 1. Subtract \\$100 from Savings.\n 2. Add \\$100 to Checking.\n If the system crashes after step 1 but before step 2, your money would simply vanish. Atomicity prevents this by ensuring that if step 2 fails, step 1 is undone.\n\n#### B. Consistency (Rules of the Game)\nConsistency ensures that a transaction brings the database from one valid state to another. Every database has predefined rules (constraints), such as \"account balances cannot be negative\" or \"every order must have a valid Customer ID.\" A transaction must never violate these rules. If a transaction attempts to perform an action that would break a rule, the database will reject the entire transaction.\n\n* **Example:** If you try to transfer \\$500 from an account that only has \\$200, and the database has a \"minimum balance\" constraint, the transaction will fail because it would result in an inconsistent state (a negative balance).\n\n#### C. Isolation (The \"Privacy\" of Transactions)\n" +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-32768/requests-32768-c1.json b/evidence/http-p3d/tput-32768/requests-32768-c1.json new file mode 100644 index 0000000000000000000000000000000000000000..128767f2980925ab611493f00f41b62d1b180294 --- /dev/null +++ b/evidence/http-p3d/tput-32768/requests-32768-c1.json @@ -0,0 +1,34 @@ +[ + { + "index": 0, + "start_seconds": 2.0428997231647372e-05, + "end_seconds": 28.759651961998316, + "ttft_seconds": 16.977761214016937, + "latency_seconds": 28.759631533001084, + "decode_tokens_per_second": 43.37198960875458, + "text_sha256": "1506d9360ed3ea7ef3017e332c703f80b9171cdd9b987634b75b4789218ad98d", + "usage": { + "prompt_tokens": 32768, + "total_tokens": 33280, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of software engineering and data management, a **database transaction** is a fundamental concept that ensures data integrity. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and is permanently saved to the database, or none of them are applied.\n\nTo understand why this is critical, we must look at the **ACID properties**, the mechanisms that enforce them, and real-world scenarios where transactions prevent catastrophic data loss.\n\n---\n\n### 1. The ACID Properties\nThe reliability of a database management system (DBMS) is measured by its adherence to the ACID model. These four properties are the \"gold standard\" for transaction management.\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as an indivisible unit. If a transaction consists of five different SQL statements (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), and the fourth statement fails due to a power outage or a constraint violation, the database must \"roll back\" the first three statements. The database should return to the exact state it was in before the transaction started.\n\n#### B. Consistency (Valid State to Valid State)\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including-defined constraints (like unique keys),s, and foreign key relationships. For example, if a rule states that an account balance cannot drop below zero, any transaction that would result in a negative balance must be rejected by the system.\n\n#### C. Isolation (Independence of Concurrent Transactions)\nIn modern applications, thousands of users may access a database simultaneously. Isolation ensures that the concurrent execution of transactions results in a system state that is the same as if the transactions were executed sequentially (one after the other). Without isolation, \"dirty reads\" or \"lost updates\" could occur\u2014where one user sees half-finished data from another user's ongoing transaction.\n\n#### D. Durability (Permanence)\nOnce a transaction has been committed (successfully completed), it must remain committed even in the event of a system failure, such as a crash or a power loss. This is usually achieved by recording the transaction in non-volatile memory (like a hard drive or SSD) via transaction logs before confirming success to the user." + }, + { + "index": 1, + "start_seconds": 28.759694711014163, + "end_seconds": 57.533709167008055, + "ttft_seconds": 17.023219919996336, + "latency_seconds": 28.774014455993893, + "decode_tokens_per_second": 43.48662823635663, + "text_sha256": "1506d9360ed3ea7ef3017e332c703f80b9171cdd9b987634b75b4789218ad98d", + "usage": { + "prompt_tokens": 32768, + "total_tokens": 33280, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of software engineering and data management, a **database transaction** is a fundamental concept that ensures data integrity. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and is permanently saved to the database, or none of them are applied.\n\nTo understand why this is critical, we must look at the **ACID properties**, the mechanisms that enforce them, and real-world scenarios where transactions prevent catastrophic data loss.\n\n---\n\n### 1. The ACID Properties\nThe reliability of a database management system (DBMS) is measured by its adherence to the ACID model. These four properties are the \"gold standard\" for transaction management.\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as an indivisible unit. If a transaction consists of five different SQL statements (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), and the fourth statement fails due to a power outage or a constraint violation, the database must \"roll back\" the first three statements. The database should return to the exact state it was in before the transaction started.\n\n#### B. Consistency (Valid State to Valid State)\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including-defined constraints (like unique keys),s, and foreign key relationships. For example, if a rule states that an account balance cannot drop below zero, any transaction that would result in a negative balance must be rejected by the system.\n\n#### C. Isolation (Independence of Concurrent Transactions)\nIn modern applications, thousands of users may access a database simultaneously. Isolation ensures that the concurrent execution of transactions results in a system state that is the same as if the transactions were executed sequentially (one after the other). Without isolation, \"dirty reads\" or \"lost updates\" could occur\u2014where one user sees half-finished data from another user's ongoing transaction.\n\n#### D. Durability (Permanence)\nOnce a transaction has been committed (successfully completed), it must remain committed even in the event of a system failure, such as a crash or a power loss. This is usually achieved by recording the transaction in non-volatile memory (like a hard drive or SSD) via transaction logs before confirming success to the user." + } +] \ No newline at end of file diff --git a/evidence/http-p3d/tput-32768/summary.json b/evidence/http-p3d/tput-32768/summary.json new file mode 100644 index 0000000000000000000000000000000000000000..80b204c537f2cc02f52072f6d0361a1b133df011 --- /dev/null +++ b/evidence/http-p3d/tput-32768/summary.json @@ -0,0 +1,24 @@ +{ + "model": "google/gemma-4-12B-it", + "output_tokens": 512, + "method": "Closed loop; synchronized initial clients; max(min_requests,2*concurrency) requests; nearest-rank percentiles; aggregate includes prefill and queue drain", + "cases": [ + { + "input_tokens": 32768, + "concurrency": 1, + "requests": 2, + "successes": 2, + "errors": 0, + "wall_seconds": 57.533688738010824, + "matching_single_request_outputs": 2, + "aggregate_output_tokens_per_second": 17.79826780554526, + "requests_per_second": 0.03476224180770559, + "ttft_seconds_p50": 16.977761214016937, + "ttft_seconds_p95": 17.023219919996336, + "latency_seconds_p50": 28.759631533001084, + "latency_seconds_p95": 28.774014455993893, + "decode_tokens_per_second_p50": 43.37198960875458, + "decode_tokens_per_second_p95": 43.48662823635663 + } + ] +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-32768/warmup-32768.json b/evidence/http-p3d/tput-32768/warmup-32768.json new file mode 100644 index 0000000000000000000000000000000000000000..6cccf76cbd62493c03554f9532531c19899e3253 --- /dev/null +++ b/evidence/http-p3d/tput-32768/warmup-32768.json @@ -0,0 +1,16 @@ +{ + "index": -1, + "start_seconds": 8.200004231184721e-07, + "end_seconds": 28.692602641007397, + "ttft_seconds": 16.91264510701876, + "latency_seconds": 28.692601821006974, + "decode_tokens_per_second": 43.37904968705863, + "text_sha256": "1506d9360ed3ea7ef3017e332c703f80b9171cdd9b987634b75b4789218ad98d", + "usage": { + "prompt_tokens": 32768, + "total_tokens": 33280, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of software engineering and data management, a **database transaction** is a fundamental concept that ensures data integrity. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and is permanently saved to the database, or none of them are applied.\n\nTo understand why this is critical, we must look at the **ACID properties**, the mechanisms that enforce them, and real-world scenarios where transactions prevent catastrophic data loss.\n\n---\n\n### 1. The ACID Properties\nThe reliability of a database management system (DBMS) is measured by its adherence to the ACID model. These four properties are the \"gold standard\" for transaction management.\n\n#### A. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as an indivisible unit. If a transaction consists of five different SQL statements (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), and the fourth statement fails due to a power outage or a constraint violation, the database must \"roll back\" the first three statements. The database should return to the exact state it was in before the transaction started.\n\n#### B. Consistency (Valid State to Valid State)\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including-defined constraints (like unique keys),s, and foreign key relationships. For example, if a rule states that an account balance cannot drop below zero, any transaction that would result in a negative balance must be rejected by the system.\n\n#### C. Isolation (Independence of Concurrent Transactions)\nIn modern applications, thousands of users may access a database simultaneously. Isolation ensures that the concurrent execution of transactions results in a system state that is the same as if the transactions were executed sequentially (one after the other). Without isolation, \"dirty reads\" or \"lost updates\" could occur\u2014where one user sees half-finished data from another user's ongoing transaction.\n\n#### D. Durability (Permanence)\nOnce a transaction has been committed (successfully completed), it must remain committed even in the event of a system failure, such as a crash or a power loss. This is usually achieved by recording the transaction in non-volatile memory (like a hard drive or SSD) via transaction logs before confirming success to the user." +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-8192/requests-8192-c1.json b/evidence/http-p3d/tput-8192/requests-8192-c1.json new file mode 100644 index 0000000000000000000000000000000000000000..cf4193d694b8d5516e9cd7e5d415b38e4b33faa4 --- /dev/null +++ b/evidence/http-p3d/tput-8192/requests-8192-c1.json @@ -0,0 +1,34 @@ +[ + { + "index": 0, + "start_seconds": 2.2149994038045406e-05, + "end_seconds": 12.67971995798871, + "ttft_seconds": 3.2031257979979273, + "latency_seconds": 12.679697807994671, + "decode_tokens_per_second": 53.92284405708368, + "text_sha256": "ed6c97ae9c6aa59779ca029eff53e380a46050f750027f929e9644ef5075218c", + "usage": { + "prompt_tokens": 8192, + "total_tokens": 8704, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is a fundamental concept that ensures data integrity and reliability. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and is permanently saved to the database, or, if any single part fails, the entire transaction is undone, leaving the database in its original state.\n\nTo understand why this is critical, we must look at the **ACID properties**, which serve as the gold standard for ensuring that database transactions are processed reliably.\n\n---\n\n### The ACID Properties\n\nFor a database management system (DBMS) to be considered reliable, it must adhere to the following four principles:\n\n#### 1. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as an indivisible unit. If a transaction consists of five different SQL queries (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), and the system crashes during the third query, the first two queries must be \"rolled back\" (undone). The database should not be left in a \"half-finished\" state.\n\n#### 2. Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, cascades, and triggers. For example, if a database rule states that an account balance cannot drop below zero, any transaction that would result in a negative balance must be rejected by the system.\n\n#### 3. Isolation\nIn modern applications, thousands of users may access a database simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. Even though multiple transactions are happening at the same time, the result should be the same as if they were executed one after another (sequentially). This prevents \"dirty reads\" (reading uncommitted data) or \"lost updates\" (where two people update the same row at the exact same millisecond and one update is overwritten).\n\n#### 4. Durability\nDurability guarantees that once a transaction has been committed (successfully completed), it will remain committed even in the event of a system failure, power outage, or crash. This is usually achieved by recording the transaction in non-volatile memory (like a hard drive or SSD) via a transaction log before confirming success to the user.\n\n---\n\n### Real-World Example: The" + }, + { + "index": 1, + "start_seconds": 12.679757648002123, + "end_seconds": 25.362715021008626, + "ttft_seconds": 3.206180963985389, + "latency_seconds": 12.682957373006502, + "decode_tokens_per_second": 53.92168694176561, + "text_sha256": "ed6c97ae9c6aa59779ca029eff53e380a46050f750027f929e9644ef5075218c", + "usage": { + "prompt_tokens": 8192, + "total_tokens": 8704, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is a fundamental concept that ensures data integrity and reliability. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and is permanently saved to the database, or, if any single part fails, the entire transaction is undone, leaving the database in its original state.\n\nTo understand why this is critical, we must look at the **ACID properties**, which serve as the gold standard for ensuring that database transactions are processed reliably.\n\n---\n\n### The ACID Properties\n\nFor a database management system (DBMS) to be considered reliable, it must adhere to the following four principles:\n\n#### 1. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as an indivisible unit. If a transaction consists of five different SQL queries (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), and the system crashes during the third query, the first two queries must be \"rolled back\" (undone). The database should not be left in a \"half-finished\" state.\n\n#### 2. Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, cascades, and triggers. For example, if a database rule states that an account balance cannot drop below zero, any transaction that would result in a negative balance must be rejected by the system.\n\n#### 3. Isolation\nIn modern applications, thousands of users may access a database simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. Even though multiple transactions are happening at the same time, the result should be the same as if they were executed one after another (sequentially). This prevents \"dirty reads\" (reading uncommitted data) or \"lost updates\" (where two people update the same row at the exact same millisecond and one update is overwritten).\n\n#### 4. Durability\nDurability guarantees that once a transaction has been committed (successfully completed), it will remain committed even in the event of a system failure, power outage, or crash. This is usually achieved by recording the transaction in non-volatile memory (like a hard drive or SSD) via a transaction log before confirming success to the user.\n\n---\n\n### Real-World Example: The" + } +] \ No newline at end of file diff --git a/evidence/http-p3d/tput-8192/summary.json b/evidence/http-p3d/tput-8192/summary.json new file mode 100644 index 0000000000000000000000000000000000000000..e2803f48ecc96c516b25464b790b9cc25c666454 --- /dev/null +++ b/evidence/http-p3d/tput-8192/summary.json @@ -0,0 +1,24 @@ +{ + "model": "google/gemma-4-12B-it", + "output_tokens": 512, + "method": "Closed loop; synchronized initial clients; max(min_requests,2*concurrency) requests; nearest-rank percentiles; aggregate includes prefill and queue drain", + "cases": [ + { + "input_tokens": 8192, + "concurrency": 1, + "requests": 2, + "successes": 2, + "errors": 0, + "wall_seconds": 25.362692871014588, + "matching_single_request_outputs": 2, + "aggregate_output_tokens_per_second": 40.374261724008996, + "requests_per_second": 0.07885597992970507, + "ttft_seconds_p50": 3.2031257979979273, + "ttft_seconds_p95": 3.206180963985389, + "latency_seconds_p50": 12.679697807994671, + "latency_seconds_p95": 12.682957373006502, + "decode_tokens_per_second_p50": 53.92168694176561, + "decode_tokens_per_second_p95": 53.92284405708368 + } + ] +} \ No newline at end of file diff --git a/evidence/http-p3d/tput-8192/warmup-8192.json b/evidence/http-p3d/tput-8192/warmup-8192.json new file mode 100644 index 0000000000000000000000000000000000000000..22a5a9904ca44434b406b365dab1a513f34041c2 --- /dev/null +++ b/evidence/http-p3d/tput-8192/warmup-8192.json @@ -0,0 +1,16 @@ +{ + "index": -1, + "start_seconds": 3.00002284348011e-07, + "end_seconds": 12.817421763000311, + "ttft_seconds": 3.203757772978861, + "latency_seconds": 12.817421462998027, + "decode_tokens_per_second": 53.15396023562198, + "text_sha256": "ed6c97ae9c6aa59779ca029eff53e380a46050f750027f929e9644ef5075218c", + "usage": { + "prompt_tokens": 8192, + "total_tokens": 8704, + "completion_tokens": 512 + }, + "error": null, + "text": "### Understanding Database Transactions: A Comprehensive Guide\n\nIn the world of data management, a **database transaction** is a fundamental concept that ensures data integrity and reliability. At its simplest, a transaction is a sequence of one or more operations performed as a single logical unit of work. The core philosophy of a transaction is \"all or nothing\": either every operation within the transaction succeeds and is permanently saved to the database, or, if any single part fails, the entire transaction is undone, leaving the database in its original state.\n\nTo understand why this is critical, we must look at the **ACID properties**, which serve as the gold standard for ensuring that database transactions are processed reliably.\n\n---\n\n### The ACID Properties\n\nFor a database management system (DBMS) to be considered reliable, it must adhere to the following four principles:\n\n#### 1. Atomicity (\"All or Nothing\")\nAtomicity ensures that a transaction is treated as an indivisible unit. If a transaction consists of five different SQL queries (e.g., updating a balance, creating a log entry, updating an inventory count, etc.), and the system crashes during the third query, the first two queries must be \"rolled back\" (undone). The database should not be left in a \"half-finished\" state.\n\n#### 2. Consistency\nConsistency ensures that a transaction brings the database from one valid state to another, maintaining all predefined rules, including constraints, cascades, and triggers. For example, if a database rule states that an account balance cannot drop below zero, any transaction that would result in a negative balance must be rejected by the system.\n\n#### 3. Isolation\nIn modern applications, thousands of users may access a database simultaneously. Isolation ensures that concurrent transactions do not interfere with each other. Even though multiple transactions are happening at the same time, the result should be the same as if they were executed one after another (sequentially). This prevents \"dirty reads\" (reading uncommitted data) or \"lost updates\" (where two people update the same row at the exact same millisecond and one update is overwritten).\n\n#### 4. Durability\nDurability guarantees that once a transaction has been committed (successfully completed), it will remain committed even in the event of a system failure, power outage, or crash. This is usually achieved by recording the transaction in non-volatile memory (like a hard drive or SSD) via a transaction log before confirming success to the user.\n\n---\n\n### Real-World Example: The" +} \ No newline at end of file diff --git a/evidence/http-p3c/tput_vs_direct.json b/evidence/http-p3d/tput_vs_direct.json similarity index 100% rename from evidence/http-p3c/tput_vs_direct.json rename to evidence/http-p3d/tput_vs_direct.json diff --git a/evidence/localmaxxing/p3d-2k/prompt-2k.txt b/evidence/localmaxxing/p3d-2k/prompt-2k.txt new file mode 100644 index 0000000000000000000000000000000000000000..290d0de1edc985f37bfb46e8a0cc53c73269a07b --- /dev/null +++ b/evidence/localmaxxing/p3d-2k/prompt-2k.txt @@ -0,0 +1,216 @@ +Summarize the following passage from Newton’s Opticks, then explain its main experiment in plain language. + + added +about twelve Years after to complete the Theory; except the third Book, +and the last Proposition of the Second, which were since put together +out of scatter'd Papers. To avoid being engaged in Disputes about these +Matters, I have hitherto delayed the printing, and should still have +delayed it, had not the Importunity of Friends prevailed upon me. If any +other Papers writ on this Subject are got out of my Hands they are +imperfect, and were perhaps written before I had tried all the +Experiments here set down, and fully satisfied my self about the Laws of +Refractions and Composition of Colours. I have here publish'd what I +think proper to come abroad, wishing that it may not be translated into +another Language without my Consent._ + +_The Crowns of Colours, which sometimes appear about the Sun and Moon, I +have endeavoured to give an Account of; but for want of sufficient +Observations leave that Matter to be farther examined. The Subject of +the Third Book I have also left imperfect, not having tried all the +Experiments which I intended when I was about these Matters, nor +repeated some of those which I did try, until I had satisfied my self +about all their Circumstances. To communicate what I have tried, and +leave the rest to others for farther Enquiry, is all my Design in +publishing these Papers._ + +_In a Letter written to Mr._ Leibnitz _in the year 1679, and published +by Dr._ Wallis, _I mention'd a Method by which I had found some general +Theorems about squaring Curvilinear Figures, or comparing them with the +Conic Sections, or other the simplest Figures with which they may be +compared. And some Years ago I lent out a Manuscript containing such +Theorems, and having since met with some Things copied out of it, I have +on this Occasion made it publick, prefixing to it an_ Introduction, _and +subjoining a_ Scholium _concerning that Method. And I have joined with +it another small Tract concerning the Curvilinear Figures of the Second +Kind, which was also written many Years ago, and made known to some +Friends, who have solicited the making it publick._ + + _I. N._ + +April 1, 1704. + + +Advertisement II + +_In this Second Edition of these Opticks I have omitted the Mathematical +Tracts publish'd at the End of the former Edition, as not belonging to +the Subject. And at the End of the Third Book I have added some +Questions. And to shew that I do not take Gravity for an essential +Property of Bodies, I have added one Question concerning its Cause, +chusing to propose it by way of a Question, because I am not yet +satisfied about it for want of Experiments._ + + _I. N._ + +July 16, 1717. + + +Advertisement to this Fourth Edition + +_This new Edition of Sir_ Isaac Newton's Opticks _is carefully printed +from the Third Edition, as it was corrected by the Author's own Hand, +and left before his Death with the Bookseller. Since Sir_ Isaac's +Lectiones Opticæ, _which he publickly read in the University of_ +Cambridge _in the Years 1669, 1670, and 1671, are lately printed, it has +been thought proper to make at the bottom of the Pages several Citations +from thence, where may be found the Demonstrations, which the Author +omitted in these_ Opticks. + + * * * * * + +Transcriber's Note: There are several greek letters used in the +descriptions of the illustrations. They are signified by [Greek: +letter]. Square roots are noted by the letters sqrt before the equation. + + * * * * * + +THE FIRST BOOK OF OPTICKS + + + + +_PART I._ + + +My Design in this Book is not to explain the Properties of Light by +Hypotheses, but to propose and prove them by Reason and Experiments: In +order to which I shall premise the following Definitions and Axioms. + + + + +_DEFINITIONS_ + + +DEFIN. I. + +_By the Rays of Light I understand its least Parts, and those as well +Successive in the same Lines, as Contemporary in several Lines._ For it +is manifest that Light consists of Parts, both Successive and +Contemporary; because in the same place you may stop that which comes +one moment, and let pass that which comes presently after; and in the +same time you may stop it in any one place, and let it pass in any +other. For that part of Light which is stopp'd cannot be the same with +that which is let pass. The least Light or part of Light, which may be +stopp'd alone without the rest of the Light, or propagated alone, or do +or suffer any thing alone, which the rest of the Light doth not or +suffers not, I call a Ray of Light. + + +DEFIN. II. + +_Refrangibility of the Rays of Light, is their Disposition to be +refracted or turned out of their Way in passing out of one transparent +Body or Medium into another. And a greater or less Refrangibility of +Rays, is their Disposition to be turned more or less out of their Way in +like Incidences on the same Medium._ Mathematicians usually consider the +Rays of Light to be Lines reaching from the luminous Body to the Body +illuminated, and the refraction of those Rays to be the bending or +breaking of those lines in their passing out of one Medium into another. +And thus may Rays and Refractions be considered, if Light be propagated +in an instant. But by an Argument taken from the Æquations of the times +of the Eclipses of _Jupiter's Satellites_, it seems that Light is +propagated in time, spending in its passage from the Sun to us about +seven Minutes of time: And therefore I have chosen to define Rays and +Refractions in such general terms as may agree to Light in both cases. + + +DEFIN. III. + +_Reflexibility of Rays, is their Disposition to be reflected or turned +back into the same Medium from any other Medium upon whose Surface they +fall. And Rays are more or less reflexible, which are turned back more +or less easily._ As if Light pass out of a Glass into Air, and by being +inclined more and more to the common Surface of the Glass and Air, +begins at length to be totally reflected by that Surface; those sorts of +Rays which at like Incidences are reflected most copiously, or by +inclining the Rays begin soonest to be totally reflected, are most +reflexible. + + +DEFIN. IV. + +_The Angle of Incidence is that Angle, which the Line described by the +incident Ray contains with the Perpendicular to the reflecting or +refracting Surface at the Point of Incidence._ + + +DEFIN. V. + +_The Angle of Reflexion or Refraction, is the Angle which the line +described by the reflected or refracted Ray containeth with the +Perpendicular to the reflecting or refracting Surface at the Point of +Incidence._ + + +DEFIN. VI. + +_The Sines of Incidence, Reflexion, and Refraction, are the Sines of the +Angles of Incidence, Reflexion, and Refraction._ + + +DEFIN. VII + +_The Light whose Rays are all alike Refrangible, I call Simple, +Homogeneal and Similar; and that whose Rays are some more Refrangible +than others, I call Compound, Heterogeneal and Dissimilar._ The former +Light I call Homogeneal, not because I would affirm it so in all +respects, but because the Rays which agree in Refrangibility, agree at +least in all those their other Properties which I consider in the +following Discourse. + + +DEFIN. VIII. + +_The Colours of Homogeneal Lights, I call Primary, Homogeneal and +Simple; and those of Heterogeneal Lights, Heterogeneal and Compound._ +For these are always compounded of the colours of Homogeneal Lights; as +will appear in the following Discourse. + + + + +_AXIOMS._ + + +AX. I. + +_The Angles of Reflexion and Refraction, lie in one and the same Plane +with the Angle of Incidence._ + + +AX. II. + +_The Angle of Reflexion is equal to the Angle of Incidence._ + + +AX. III. + +_If the refracted Ray be returned directly back to the Point of +Incidence, it shall be refracted into the Line before described by the +incident Ray._ + + +AX. IV. + +_Refraction out of the rarer Medium into the denser, is made towards the +Perpendicular; that is, so that the Angle of Refraction be less than the +Angle of Incidence._ + + +AX. V. + +_The Sine of Incidence is either accurately or very nearly in a given +Ratio to the Sine of Refraction._ + +Whence \ No newline at end of file diff --git a/evidence/localmaxxing/p3d-2k/speed-test.json b/evidence/localmaxxing/p3d-2k/speed-test.json new file mode 100644 index 0000000000000000000000000000000000000000..7b84f24352dc9e0772f7b38df5da7c9ba644b94f --- /dev/null +++ b/evidence/localmaxxing/p3d-2k/speed-test.json @@ -0,0 +1,212 @@ +{ + "agentFeedback": { + "benchmarkStatus": "completed", + "canApiValidate": true, + "canSubmit": true, + "engine": "vllm", + "message": "Speed-test payload is ready for API validation.", + "mode": "remote", + "nextCommand": "lmx speed-test dry-run /var/tmp/gemma4-12b/runs/lmx-2k-p3d/speed-test.json", + "outputPath": "/var/tmp/gemma4-12b/runs/lmx-2k-p3d/speed-test.json", + "outputPathAbsolute": "/var/tmp/gemma4-12b/runs/lmx-2k-p3d/speed-test.json", + "requiresMetrics": false, + "runPersisted": true, + "savedRunPath": "/var/tmp/gemma4-12b/runs/lmx-2k-p3d/runs/Lottolabs-gemma-4-12B-it-TT-BFP8-P150/20261001T001840Z.json", + "savedRunPathAbsolute": "/var/tmp/gemma4-12b/runs/lmx-2k-p3d/runs/Lottolabs-gemma-4-12B-it-TT-BFP8-P150/20261001T001840Z.json", + "savedRunPathRelative": "/var/tmp/gemma4-12b/runs/lmx-2k-p3d/runs/Lottolabs-gemma-4-12B-it-TT-BFP8-P150/20261001T001840Z.json", + "status": "ready_for_api_validation", + "submissionId": "cmuosddj10kd1lq010lf2ahkj", + "submissionStatus": "approved", + "submitCommand": "lmx speed-test submit /var/tmp/gemma4-12b/runs/lmx-2k-p3d/speed-test.json" + }, + "backend": "tt-metal", + "batchSize": 1, + "benchmarkMode": "remote", + "contextLength": 262144, + "detectedEngines": [ + { + "binaries": { + "llama-bench": "/home/lotto/llama.cpp/build/bin/llama-bench", + "llama-cli": "/home/lotto/llama.cpp/build/bin/llama-cli", + "llama-server": "/home/lotto/llama.cpp/build/bin/llama-server" + }, + "installed": true, + "name": "llama.cpp" + } + ], + "engineFlags": { + "baseUrl": "http://127.0.0.1:8002", + "concurrency": 1, + "iterations": 5, + "maxTokens": 256, + "mode": "remote", + "prefixCacheBust": "leading_nonce_per_request", + "promptFile": "/var/tmp/gemma4-12b/runs/prompt-2k.txt", + "promptSource": "file", + "servedModel": "google/gemma-4-12B-it", + "servedModelSource": "explicit", + "specDecoding": true, + "specDraftModel": "google/gemma-4-12B-it-assistant", + "specMethod": "assistant", + "specNumTokens": 5, + "stream": true, + "timeoutSeconds": 600, + "warmup": 2 + }, + "engineName": "vllm", + "hardware": { + "cpu": "AMD Ryzen 9 9950X", + "gpuCount": 1, + "gpuName": "Tenstorrent P150", + "hwClass": "DISCRETE_GPU", + "os": "Ubuntu 26.04 LTS", + "ramGb": 96, + "vramGb": 32 + }, + "hardwareSource": "file", + "hfId": "Lottolabs/gemma-4-12B-it-TT-BFP8-P150", + "id": "cmuosddj10kd1lq010lf2ahkj", + "metricSource": "remote_endpoint", + "modelResolution": { + "candidates": [ + { + "benchmarkCount": 38, + "displayName": "gemma-4-12B-it", + "family": "Gemma", + "hfId": "google/gemma-4-12B-it", + "params": 12 + }, + { + "benchmarkCount": 9, + "displayName": "gemma-4-12B-it-qat-q4_0-gguf", + "family": "Gemma", + "hfId": "google/gemma-4-12B-it-qat-q4_0-gguf", + "params": 12 + }, + { + "benchmarkCount": 0, + "displayName": "gemma-4-12B-it-qat-q4_0-unquantized", + "family": "Gemma", + "hfId": "google/gemma-4-12B-it-qat-q4_0-unquantized", + "params": 12 + } + ], + "hfId": "Lottolabs/gemma-4-12B-it-TT-BFP8-P150", + "searchCommand": "lmx model search google/gemma-4-12B-it", + "searchQuery": "google/gemma-4-12B-it", + "searchQuerySource": "served_model", + "servedModel": "google/gemma-4-12B-it", + "servedModelSource": "explicit", + "status": "alias" + }, + "modelRevision": "main", + "notes": "Single Tenstorrent P150. Lottolabs/gemma-4-12B-it-TT-BFP8-P150, runtime p3d (image sha256:331e79b24fa2575c436158be994c6bc0e184d776d15305e82538d2e4d646f544, dual-NoC weight reader): all-BFP8 TT-native weights with exact equivalence proof, gemma-4-12B-it-assistant drafter K=5 with exact greedy verification, native 262144 context, greedy. ~2K-token natural-text prompt (Newton, Opticks).", + "outputText": "### Summary of the Passage\n\nIn this introductory section of *Opticks*, Isaac Newton establishes the foundational framework for his study of light. He begins by clarifying his intent: he aims to explain the properties of light through **reason and experimentation** rather than through abstract hypotheses. \n\nTo do this, he provides a series of formal **Definitions** and **Axioms**:\n* **Definitions:** He defines \"Rays of Light\" as the smallest parts of light, \"Refrangibility\" (how light bends when moving between mediums), \"Reflexibility\" (how light bounces off surfaces), and the specific angles involved in these movements. Crucially, he distinguishes between **Homogeneal Light** (light where all rays bend the same way) and **Heterogeneal Light** (light made of different types of rays).\n* **Axioms:** He lists the fundamental rules of optics, such as the fact that the angle of reflection equals the angle of incidence, and that light bends toward the perpendicular when moving from a \"rarer\" (thinner) medium to a \"denser\" one.\n\n---\n\n### The Main Experiment (Explained in Plain Language)\n\nWhile the text provided is primarily a list of definitions and rules,", + "outputTokens": 256, + "prefillTokens": 0, + "prompt": "Summarize the following passage from Newton’s Opticks, then explain its main experiment in plain language.\n\n added\nabout twelve Years after to complete the Theory; except the third Book,\nand the last Proposition of the Second, which were since put together\nout of scatter'd Papers. To avoid being engaged in Disputes about these\nMatters, I have hitherto delayed the printing, and should still have\ndelayed it, had not the Importunity of Friends prevailed upon me. If any\nother Papers writ on this Subject are got out of my Hands they are\nimperfect, and were perhaps written before I had tried all the\nExperiments here set down, and fully satisfied my self about the Laws of\nRefractions and Composition of Colours. I have here publish'd what I\nthink proper to come abroad, wishing that it may not be translated into\nanother Language without my Consent._\n\n_The Crowns of Colours, which sometimes appear about the Sun and Moon, I\nhave endeavoured to give an Account of; but for want of sufficient\nObservations leave that Matter to be farther examined. The Subject of\nthe Third Book I have also left imperfect, not having tried all the\nExperiments which I intended when I was about these Matters, nor\nrepeated some of those which I did try, until I had satisfied my self\nabout all their Circumstances. To communicate what I have tried, and\nleave the rest to others for farther Enquiry, is all my Design in\npublishing these Papers._\n\n_In a Letter written to Mr._ Leibnitz _in the year 1679, and published\nby Dr._ Wallis, _I mention'd a Method by which I had found some general\nTheorems about squaring Curvilinear Figures, or comparing them with the\nConic Sections, or other the simplest Figures with which they may be\ncompared. And some Years ago I lent out a Manuscript containing such\nTheorems, and having since met with some Things copied out of it, I have\non this Occasion made it publick, prefixing to it an_ Introduction, _and\nsubjoining a_ Scholium _concerning that Method. And I have joined with\nit another small Tract concerning the Curvilinear Figures of the Second\nKind, which was also written many Years ago, and made known to some\nFriends, who have solicited the making it publick._\n\n _I. N._\n\nApril 1, 1704.\n\n\nAdvertisement II\n\n_In this Second Edition of these Opticks I have omitted the Mathematical\nTracts publish'd at the End of the former Edition, as not belonging to\nthe Subject. And at the End of the Third Book I have added some\nQuestions. And to shew that I do not take Gravity for an essential\nProperty of Bodies, I have added one Question concerning its Cause,\nchusing to propose it by way of a Question, because I am not yet\nsatisfied about it for want of Experiments._\n\n _I. N._\n\nJuly 16, 1717.\n\n\nAdvertisement to this Fourth Edition\n\n_This new Edition of Sir_ Isaac Newton's Opticks _is carefully printed\nfrom the Third Edition, as it was corrected by the Author's own Hand,\nand left before his Death with the Bookseller. Since Sir_ Isaac's\nLectiones Opticæ, _which he publickly read in the University of_\nCambridge _in the Years 1669, 1670, and 1671, are lately printed, it has\nbeen thought proper to make at the bottom of the Pages several Citations\nfrom thence, where may be found the Demonstrations, which the Author\nomitted in these_ Opticks.\n\n * * * * *\n\nTranscriber's Note: There are several greek letters used in the\ndescriptions of the illustrations. They are signified by [Greek:\nletter]. Square roots are noted by the letters sqrt before the equation.\n\n * * * * *\n\nTHE FIRST BOOK OF OPTICKS\n\n\n\n\n_PART I._\n\n\nMy Design in this Book is not to explain the Properties of Light by\nHypotheses, but to propose and prove them by Reason and Experiments: In\norder to which I shall premise the following Definitions and Axioms.\n\n\n\n\n_DEFINITIONS_\n\n\nDEFIN. I.\n\n_By the Rays of Light I understand its least Parts, and those as well\nSuccessive in the same Lines, as Contemporary in several Lines._ For it\nis manifest that Light consists of Parts, both Successive and\nContemporary; because in the same place you may stop that which comes\none moment, and let pass that which comes presently after; and in the\nsame time you may stop it in any one place, and let it pass in any\nother. For that part of Light which is stopp'd cannot be the same with\nthat which is let pass. The least Light or part of Light, which may be\nstopp'd alone without the rest of the Light, or propagated alone, or do\nor suffer any thing alone, which the rest of the Light doth not or\nsuffers not, I call a Ray of Light.\n\n\nDEFIN. II.\n\n_Refrangibility of the Rays of Light, is their Disposition to be\nrefracted or turned out of their Way in passing out of one transparent\nBody or Medium into another. And a greater or less Refrangibility of\nRays, is their Disposition to be turned more or less out of their Way in\nlike Incidences on the same Medium._ Mathematicians usually consider the\nRays of Light to be Lines reaching from the luminous Body to the Body\nilluminated, and the refraction of those Rays to be the bending or\nbreaking of those lines in their passing out of one Medium into another.\nAnd thus may Rays and Refractions be considered, if Light be propagated\nin an instant. But by an Argument taken from the Æquations of the times\nof the Eclipses of _Jupiter's Satellites_, it seems that Light is\npropagated in time, spending in its passage from the Sun to us about\nseven Minutes of time: And therefore I have chosen to define Rays and\nRefractions in such general terms as may agree to Light in both cases.\n\n\nDEFIN. III.\n\n_Reflexibility of Rays, is their Disposition to be reflected or turned\nback into the same Medium from any other Medium upon whose Surface they\nfall. And Rays are more or less reflexible, which are turned back more\nor less easily._ As if Light pass out of a Glass into Air, and by being\ninclined more and more to the common Surface of the Glass and Air,\nbegins at length to be totally reflected by that Surface; those sorts of\nRays which at like Incidences are reflected most copiously, or by\ninclining the Rays begin soonest to be totally reflected, are most\nreflexible.\n\n\nDEFIN. IV.\n\n_The Angle of Incidence is that Angle, which the Line described by the\nincident Ray contains with the Perpendicular to the reflecting or\nrefracting Surface at the Point of Incidence._\n\n\nDEFIN. V.\n\n_The Angle of Reflexion or Refraction, is the Angle which the line\ndescribed by the reflected or refracted Ray containeth with the\nPerpendicular to the reflecting or refracting Surface at the Point of\nIncidence._\n\n\nDEFIN. VI.\n\n_The Sines of Incidence, Reflexion, and Refraction, are the Sines of the\nAngles of Incidence, Reflexion, and Refraction._\n\n\nDEFIN. VII\n\n_The Light whose Rays are all alike Refrangible, I call Simple,\nHomogeneal and Similar; and that whose Rays are some more Refrangible\nthan others, I call Compound, Heterogeneal and Dissimilar._ The former\nLight I call Homogeneal, not because I would affirm it so in all\nrespects, but because the Rays which agree in Refrangibility, agree at\nleast in all those their other Properties which I consider in the\nfollowing Discourse.\n\n\nDEFIN. VIII.\n\n_The Colours of Homogeneal Lights, I call Primary, Homogeneal and\nSimple; and those of Heterogeneal Lights, Heterogeneal and Compound._\nFor these are always compounded of the colours of Homogeneal Lights; as\nwill appear in the following Discourse.\n\n\n\n\n_AXIOMS._\n\n\nAX. I.\n\n_The Angles of Reflexion and Refraction, lie in one and the same Plane\nwith the Angle of Incidence._\n\n\nAX. II.\n\n_The Angle of Reflexion is equal to the Angle of Incidence._\n\n\nAX. III.\n\n_If the refracted Ray be returned directly back to the Point of\nIncidence, it shall be refracted into the Line before described by the\nincident Ray._\n\n\nAX. IV.\n\n_Refraction out of the rarer Medium into the denser, is made towards the\nPerpendicular; that is, so that the Angle of Refraction be less than the\nAngle of Incidence._\n\n\nAX. V.\n\n_The Sine of Incidence is either accurately or very nearly in a given\nRatio to the Sine of Refraction._\n\nWhence", + "promptTokens": 1996, + "provenance": { + "benchmarkMode": "remote", + "cli": "localmaxxing-go", + "createdAt": "2026-10-01T00:18:40Z", + "metricSource": "remote_endpoint", + "timingSource": "client_observed_http", + "ttftSource": "stream_first_token" + }, + "quantization": "BFP8", + "quantizationResolution": { + "cli": "BFP8", + "status": "matched", + "trusted": "BFP8", + "trustedSource": "cli" + }, + "sampleStats": { + "tokSOut": { + "count": 5, + "max": 59.8, + "mean": 54.2, + "min": 48.5, + "p50": 53.9, + "stddev": 4.6 + }, + "tokSTotal": { + "count": 5, + "max": 458.4, + "mean": 420.18, + "min": 381.7, + "p50": 417.9, + "stddev": 31.11 + }, + "ttftMs": { + "count": 5, + "max": 652.07, + "mean": 651.48, + "min": 650.74, + "p50": 651.49, + "stddev": 0.5 + } + }, + "samples": [ + { + "iteration": 1, + "outputTokens": 256, + "promptTokens": 1993, + "request": 1, + "tokSOut": 53.9, + "tokSPrefill": 3062.7, + "ttftMs": 650.74 + }, + { + "iteration": 2, + "outputTokens": 256, + "promptTokens": 1998, + "request": 1, + "tokSOut": 48.5, + "tokSPrefill": 3065.3, + "ttftMs": 651.8 + }, + { + "iteration": 3, + "outputTokens": 256, + "promptTokens": 1994, + "request": 1, + "tokSOut": 57.6, + "tokSPrefill": 3060.7, + "ttftMs": 651.49 + }, + { + "iteration": 4, + "outputTokens": 256, + "promptTokens": 1997, + "request": 1, + "tokSOut": 51.2, + "tokSPrefill": 3062.6, + "ttftMs": 652.07 + }, + { + "iteration": 5, + "outputTokens": 256, + "promptTokens": 1996, + "request": 1, + "tokSOut": 59.8, + "tokSPrefill": 3064.6, + "ttftMs": 651.32 + } + ], + "submissionId": "cmuosddj10kd1lq010lf2ahkj", + "submissionStatus": "approved", + "submittedAt": "2026-10-01T00:18:41.725Z", + "timingSource": "client_observed_http", + "tokSOut": 53.9, + "tokSOutSource": "inter_token", + "tokSPrefill": 3062.7, + "tokSPrefillSource": "estimated_from_ttft", + "tokSTotal": 417.9, + "tokenSources": { + "output": "endpoint_usage", + "prompt": "endpoint_usage" + }, + "ttftMs": 651.49, + "ttftSource": "stream_first_token" +} diff --git a/evidence/http-p3c/exact.json b/evidence/previous-p3c/http/exact.json similarity index 100% rename from evidence/http-p3c/exact.json rename to evidence/previous-p3c/http/exact.json diff --git a/evidence/http-p3c/long.json b/evidence/previous-p3c/http/long.json similarity index 100% rename from evidence/http-p3c/long.json rename to evidence/previous-p3c/http/long.json diff --git a/evidence/http-p3c/tput-128/requests-128-c1.json b/evidence/previous-p3c/http/tput-128/requests-128-c1.json similarity index 100% rename from evidence/http-p3c/tput-128/requests-128-c1.json rename to evidence/previous-p3c/http/tput-128/requests-128-c1.json diff --git a/evidence/http-p3c/tput-128/summary.json b/evidence/previous-p3c/http/tput-128/summary.json similarity index 100% rename from evidence/http-p3c/tput-128/summary.json rename to evidence/previous-p3c/http/tput-128/summary.json diff --git a/evidence/http-p3c/tput-128/warmup-128.json b/evidence/previous-p3c/http/tput-128/warmup-128.json similarity index 100% rename from evidence/http-p3c/tput-128/warmup-128.json rename to evidence/previous-p3c/http/tput-128/warmup-128.json diff --git a/evidence/http-p3c/tput-131072/requests-131072-c1.json b/evidence/previous-p3c/http/tput-131072/requests-131072-c1.json similarity index 100% rename from evidence/http-p3c/tput-131072/requests-131072-c1.json rename to evidence/previous-p3c/http/tput-131072/requests-131072-c1.json diff --git a/evidence/http-p3c/tput-131072/summary.json b/evidence/previous-p3c/http/tput-131072/summary.json similarity index 100% rename from evidence/http-p3c/tput-131072/summary.json rename to evidence/previous-p3c/http/tput-131072/summary.json diff --git a/evidence/http-p3c/tput-131072/warmup-131072.json b/evidence/previous-p3c/http/tput-131072/warmup-131072.json similarity index 100% rename from evidence/http-p3c/tput-131072/warmup-131072.json rename to evidence/previous-p3c/http/tput-131072/warmup-131072.json diff --git a/evidence/http-p3c/tput-2048/requests-2048-c1.json b/evidence/previous-p3c/http/tput-2048/requests-2048-c1.json similarity index 100% rename from evidence/http-p3c/tput-2048/requests-2048-c1.json rename to evidence/previous-p3c/http/tput-2048/requests-2048-c1.json diff --git a/evidence/http-p3c/tput-2048/summary.json b/evidence/previous-p3c/http/tput-2048/summary.json similarity index 100% rename from evidence/http-p3c/tput-2048/summary.json rename to evidence/previous-p3c/http/tput-2048/summary.json diff --git a/evidence/http-p3c/tput-2048/warmup-2048.json b/evidence/previous-p3c/http/tput-2048/warmup-2048.json similarity index 100% rename from evidence/http-p3c/tput-2048/warmup-2048.json rename to evidence/previous-p3c/http/tput-2048/warmup-2048.json diff --git a/evidence/http-p3c/tput-261632/requests-261632-c1.json b/evidence/previous-p3c/http/tput-261632/requests-261632-c1.json similarity index 100% rename from evidence/http-p3c/tput-261632/requests-261632-c1.json rename to evidence/previous-p3c/http/tput-261632/requests-261632-c1.json diff --git a/evidence/http-p3c/tput-261632/summary.json b/evidence/previous-p3c/http/tput-261632/summary.json similarity index 100% rename from evidence/http-p3c/tput-261632/summary.json rename to evidence/previous-p3c/http/tput-261632/summary.json diff --git a/evidence/http-p3c/tput-261632/warmup-261632.json b/evidence/previous-p3c/http/tput-261632/warmup-261632.json similarity index 100% rename from evidence/http-p3c/tput-261632/warmup-261632.json rename to evidence/previous-p3c/http/tput-261632/warmup-261632.json diff --git a/evidence/http-p3c/tput-32768/requests-32768-c1.json b/evidence/previous-p3c/http/tput-32768/requests-32768-c1.json similarity index 100% rename from evidence/http-p3c/tput-32768/requests-32768-c1.json rename to evidence/previous-p3c/http/tput-32768/requests-32768-c1.json diff --git a/evidence/http-p3c/tput-32768/summary.json b/evidence/previous-p3c/http/tput-32768/summary.json similarity index 100% rename from evidence/http-p3c/tput-32768/summary.json rename to evidence/previous-p3c/http/tput-32768/summary.json diff --git a/evidence/http-p3c/tput-32768/warmup-32768.json b/evidence/previous-p3c/http/tput-32768/warmup-32768.json similarity index 100% rename from evidence/http-p3c/tput-32768/warmup-32768.json rename to evidence/previous-p3c/http/tput-32768/warmup-32768.json diff --git a/evidence/http-p3c/tput-8192/requests-8192-c1.json b/evidence/previous-p3c/http/tput-8192/requests-8192-c1.json similarity index 100% rename from evidence/http-p3c/tput-8192/requests-8192-c1.json rename to evidence/previous-p3c/http/tput-8192/requests-8192-c1.json diff --git a/evidence/http-p3c/tput-8192/summary.json b/evidence/previous-p3c/http/tput-8192/summary.json similarity index 100% rename from evidence/http-p3c/tput-8192/summary.json rename to evidence/previous-p3c/http/tput-8192/summary.json diff --git a/evidence/http-p3c/tput-8192/warmup-8192.json b/evidence/previous-p3c/http/tput-8192/warmup-8192.json similarity index 100% rename from evidence/http-p3c/tput-8192/warmup-8192.json rename to evidence/previous-p3c/http/tput-8192/warmup-8192.json diff --git a/evidence/previous-p3c/http/tput_vs_direct.json b/evidence/previous-p3c/http/tput_vs_direct.json new file mode 100644 index 0000000000000000000000000000000000000000..bba72e8c1da033d8d335021674fec33b24f75daa --- /dev/null +++ b/evidence/previous-p3c/http/tput_vs_direct.json @@ -0,0 +1,22 @@ +[ + { + "ctx": 32768, + "direct_tokens": 512, + "http_requests": 3, + "identical": [ + true, + true, + true + ] + }, + { + "ctx": 261632, + "direct_tokens": 512, + "http_requests": 3, + "identical": [ + true, + true, + true + ] + } +] \ No newline at end of file diff --git a/evidence/localmaxxing/speed-test.json b/evidence/previous-p3c/localmaxxing-speed-test.json similarity index 100% rename from evidence/localmaxxing/speed-test.json rename to evidence/previous-p3c/localmaxxing-speed-test.json diff --git a/evidence/public-download-verification.json b/evidence/previous-p3c/public-download-verification.json similarity index 100% rename from evidence/public-download-verification.json rename to evidence/previous-p3c/public-download-verification.json diff --git a/evidence/public-serving-smoke.json b/evidence/previous-p3c/public-serving-smoke.json similarity index 100% rename from evidence/public-serving-smoke.json rename to evidence/previous-p3c/public-serving-smoke.json diff --git a/evidence/public-serving-smoke/requests-2048-c1.json b/evidence/previous-p3c/public-serving-smoke/requests-2048-c1.json similarity index 100% rename from evidence/public-serving-smoke/requests-2048-c1.json rename to evidence/previous-p3c/public-serving-smoke/requests-2048-c1.json diff --git a/evidence/public-serving-smoke/summary.json b/evidence/previous-p3c/public-serving-smoke/summary.json similarity index 100% rename from evidence/public-serving-smoke/summary.json rename to evidence/previous-p3c/public-serving-smoke/summary.json diff --git a/evidence/public-serving-smoke/warmup-2048.json b/evidence/previous-p3c/public-serving-smoke/warmup-2048.json similarity index 100% rename from evidence/public-serving-smoke/warmup-2048.json rename to evidence/previous-p3c/public-serving-smoke/warmup-2048.json diff --git a/evidence/quality/runtime-gate-dtf-score.json b/evidence/previous-p3c/runtime-gate-dtf-score.json similarity index 100% rename from evidence/quality/runtime-gate-dtf-score.json rename to evidence/previous-p3c/runtime-gate-dtf-score.json diff --git a/evidence/quality/p3d-gate/book-2048-score.json b/evidence/quality/p3d-gate/book-2048-score.json new file mode 100644 index 0000000000000000000000000000000000000000..6f2d84e1b17e08e3895643e8ea1d396409d4cdd5 --- /dev/null +++ b/evidence/quality/p3d-gate/book-2048-score.json @@ -0,0 +1,15 @@ +{ + "long-book": { + "n": 2048, + "dNLL": 0.011399740353226662, + "se": 0.0044630522780970535, + "dPPL_pct": 1.1464953422546387, + "top1": 0.96484375, + "kl_mean": 0.007551407441496849, + "kl_p99": 0.08058559149503708, + "vs_base_dNLL": -0.003821163671091199, + "vs_base_se": 0.004111117531952485, + "vs_base_top1_pts": -0.29296875, + "argmax_agree_base": 0.978515625 + } +} \ No newline at end of file diff --git a/evidence/quality/p3d-gate/dtf-score.json b/evidence/quality/p3d-gate/dtf-score.json new file mode 100644 index 0000000000000000000000000000000000000000..f80d0aca1c2f2b09bf15f908d45063eed03a0eb6 --- /dev/null +++ b/evidence/quality/p3d-gate/dtf-score.json @@ -0,0 +1,41 @@ +{ + "chat": { + "n": 5888, + "dNLL": 0.009954653680324554, + "se": 0.0020533486807181715, + "dPPL_pct": 1.0004401206970215, + "top1": 0.9816576086956522, + "kl_mean": 0.0037782012950628996, + "kl_p99": 0.047721605747938156, + "vs_base_dNLL": 0.0008893578196875751, + "vs_base_se": 0.0011506362538229252, + "vs_base_top1_pts": 0.23777173913044347, + "argmax_agree_base": 0.9852241847826086 + }, + "code": { + "n": 3328, + "dNLL": 0.006125741638243198, + "se": 0.003328559730022777, + "dPPL_pct": 0.614464282989502, + "top1": 0.9870793269230769, + "kl_mean": 0.005452338140457869, + "kl_p99": 0.08425664901733398, + "vs_base_dNLL": 0.0008188523352146149, + "vs_base_se": 0.003045922652773373, + "vs_base_top1_pts": 0.15024038461537437, + "argmax_agree_base": 0.9885817307692307 + }, + "long-book": { + "n": 256, + "dNLL": 0.011955169960856438, + "se": 0.015475481748580933, + "dPPL_pct": 1.2026786804199219, + "top1": 0.95703125, + "kl_mean": 0.00987553782761097, + "kl_p99": 0.10478656738996506, + "vs_base_dNLL": 0.002162549179047346, + "vs_base_se": 0.008185286074876785, + "vs_base_top1_pts": -1.171875, + "argmax_agree_base": 0.98046875 + } +} \ No newline at end of file diff --git a/evidence/quality/p3d-gate/gsm8k-100-exact-build.json b/evidence/quality/p3d-gate/gsm8k-100-exact-build.json new file mode 100644 index 0000000000000000000000000000000000000000..dab1343ffbc45e1361ac711dcb407440fb2d8fbd --- /dev/null +++ b/evidence/quality/p3d-gate/gsm8k-100-exact-build.json @@ -0,0 +1,17 @@ +{ + "n": 100, + "candidate": "/var/tmp/gemma4-12b/runs/speed-gsm-base/generations.jsonl", + "reference": "/var/tmp/gemma4-12b/quality/gsm8k/ref-bf16-100.jsonl", + "candidate_accuracy": 0.97, + "reference_accuracy": 0.97, + "answer_agreement": 0.95, + "text_exact_match": 0.25, + "disagree_idx": [ + 5, + 62, + 85, + 87, + 89 + ], + "missing_in_candidate": [] +} diff --git a/evidence/quality/p3d-gate/gsm8k-100.json b/evidence/quality/p3d-gate/gsm8k-100.json new file mode 100644 index 0000000000000000000000000000000000000000..298cb8c50137c54877c69771f3eee99da232cb01 --- /dev/null +++ b/evidence/quality/p3d-gate/gsm8k-100.json @@ -0,0 +1,17 @@ +{ + "n": 100, + "candidate": "/var/tmp/gemma4-12b/runs/speed-g6/gsm/generations.jsonl", + "reference": "/var/tmp/gemma4-12b/quality/gsm8k/ref-bf16-100.jsonl", + "candidate_accuracy": 0.97, + "reference_accuracy": 0.97, + "answer_agreement": 0.95, + "text_exact_match": 0.22, + "disagree_idx": [ + 5, + 12, + 62, + 87, + 89 + ], + "missing_in_candidate": [] +} diff --git a/gemma-4-12B-it-assistant/spec_equivalence.json b/gemma-4-12B-it-assistant/spec_equivalence.json index 706189cba0bc317f2baa3189725a74fa2e933196..88703921d83a0f02ae20db00da11d8a3c5f946a3 100644 --- a/gemma-4-12B-it-assistant/spec_equivalence.json +++ b/gemma-4-12B-it-assistant/spec_equivalence.json @@ -3,16 +3,18 @@ "drafter_manifest_sha256": "2db4559e6ec51d81a7cd51b003743ca689fed5220d41b6a403b8d6947e1c74b7", "max_new": 320, "prompts_sha256": "12691a871d1878fcd5fd8fdaceacdf412a6300bc38ba3d61413a3cbe0ff76b3c", - "reference_metadata_sha256": "b1190856d991324991e0206cdbe0b2a3e5c9ab83351bf9f17cce6361aa6524fa", + "reference_metadata_sha256": "38667e8fbca1a26852a7e9361a1f56ec0f76bcae8b03e524e98d8867b91ce95f", "runtime_env": { + "GEMMA4_DECODE_KV_ROWS": "1", + "GEMMA4_DECODE_MM_DN": "1", "GEMMA4_FUSE_GELU_MUL": "1", "GEMMA4_VERIFY_SDPA": "batched" }, "runtime_identity": { "binaries": { - "build/lib/_ttnncpp.so": "72186b132bf7be85bbb6e1b54bf6de125fbd7a557b12711c84efc6d3e2e08f8d", + "build/lib/_ttnncpp.so": "90680abdecf91fff303ea439313eef14e2358ba4b47ca68971171f63dbb5de33", "build/lib/libtt_metal.so": "0497308232f180a0973406e2afd290b9eafa9e310a16b09bf69acc4434eb0e78", - "ttnn/ttnn/_ttnn.so": "a6448405abd714d6ba310dbbc409c01842f42017b8c1d6bb1c30503fb194c7b2" + "ttnn/ttnn/_ttnn.so": "0d4a1051d2d284d14be8d72d0d818436cb0c736f17bdf39a8e82026dfab43300" }, "gemma4_tree": { "__init__.py": "501d5b563ca412fa21db82bf39117b5480ecf99bcaf11996b9c3d5f1093bb9a8", @@ -26,15 +28,15 @@ "tt/assistant/model.py": "0e3247c34a359a73bd4c8ecc7365e028372084554eeeba41682d1fd818c296af", "tt/attention/__init__.py": "df24b2bf4311ffaee781692ff8f465dea0f415e2323cf87fba4ce22227edb8ae", "tt/attention/config.py": "01b0f3e14c31aae5ed01104aee7f4cb03913007c0d0018f9d35d341d3fe049ef", - "tt/attention/decode.py": "0ec303ff75ded2a03f773c22aba27c719e2f055c71b9560fe18661d0272fd3db", + "tt/attention/decode.py": "df1f33dce05739d867ff0f4f11facabfa5625818118d8cf4d5fc8e97df36c4ac", "tt/attention/kv_cache.py": "3a1e9f520cdde61dd49550dbfca1072d009ee34878bafe5923c655821c09bff7", "tt/attention/kv_cache_hybrid.py": "9074a2d1de4b3a65d11271370f1fce39a78e5523b841752bcaa87dccb4b56470", - "tt/attention/operations.py": "f5cda2ade4031a66e631a732199f3b4b9a00b45ef2e252aee0d3138406bbcc43", + "tt/attention/operations.py": "8e87380345b46db028f882df1df2ea05b59a4e0278ed8516a1f15b9663384a65", "tt/attention/prefill.py": "32c6e1c9423b9c22f14558c33d76d9164cb8a27c939cb5d1ccb6da039b7c831f", "tt/attention/weights.py": "3ec94d60c6d1085bed0bf83fbc450aa9599a476341d3c6d77466b1b2e17c00d6", "tt/ccl.py": "3580f2e25950381465e55dd2886d2832c33d8d6fc618b5a2b94088d0b1a91105", "tt/common.py": "e4369b067c3c2e7930fb538a34ab04b237940048e771f2f8f0560776f03a578f", - "tt/decode_mm.py": "93ebedf2eb3b0c78ebbddfc00c630b13c726d9b13bb07fb53a073a207731bae4", + "tt/decode_mm.py": "85e5a625134c36925d559c4b7557b0139e824a77c4e724bf3b647015b77f7fac", "tt/experts/__init__.py": "567aef533a30a692e4c0b28695c15d13e2cfc024ee621e4147c836b63220d5c3", "tt/experts/config.py": "80bd7d18ffd0e0b6bd969d05a8e9972ff8473dc63631f1cb6de0bd619a50f33b", "tt/experts/decode.py": "112b9df4b88e8395053fdbf69315e3a43e9f3f4877e699ca5a48462ccd8e804b", @@ -66,21 +68,21 @@ } }, "scope": "greedy token streams of speculative decoding equal non-speculative greedy decoding on the recorded prompts, same proven runtime", - "speculative_metadata_sha256": "8d73468115afbf703356394074e01a74a3cfcacac68650214bb52d19d0a3ed0d", + "speculative_metadata_sha256": "f36c0c488a67aca60982cc80dba2195a2614a890ec008c8b5c99ec4dbd6465a9", "streams_identical": true, "streams_sha256": { - "code": "e81a12880f67e38b9409ceb85846ca731a0cdb9efe4046d458bae25c1be09417", + "code": "271e958e56f958fee3cca5a57f6ac818589d3b55597f226a9a307b905b83734d", "code-c": "4755648cb1b633a42596b416aeae1d2e4d2519b29188d31d7be9f85a73eb05cc", - "code-lru": "2ddfbc830a509221a668dc6f430dd4c7f65d424f0a47ef0297fe9825ce64bcdd", - "code-sql": "36484b712be95b98a7fc1b0e06fd2b4a3d4f09017f3a6f791102529d803c37c8", - "explain": "c4c3e77a59cab64b75d4604f3dadcd75fa06cbcdfb77c4e6e3209a5f10691332", - "logic": "f28d9e2ee05d50b2225c186bfe2fdeb2c5464f5e4809ebe6fc4c70e7a72b62fa", - "long-code": "557d5b465ba6ba3b9260b33dd5b7d943c70599b8de06c669be70ec70178b50dd", - "long-summary": "f9a898b16dc9eed88e73e18d07c9abb019f937bea73287f1d101dea86272ff6e", + "code-lru": "151bcf66339d8e529eb18d84cd31d5d53cce7ef0adf86c8fd73730fcd01c38aa", + "code-sql": "3e0b762512eb0dcf5e3c67176e1ab27d379b0b652cd053fbdc8725f3a9e567e0", + "explain": "37af975a9113066fc1e70b41db996ea5e1a15d3277c06372bd88c373fc69199b", + "logic": "a343b9a1074418f3070fde6c6a7b00c297f2d53e15ac9bbd236e32b32fdeb8aa", + "long-code": "641bf0520a5d8ee9401188eba9b35095e5cd781b727f6de7196ea1368b603bc8", + "long-summary": "d81cecb93b441c4a652b4032cb15ffd43dc351716e0d500851c75dc40d2b27dd", "math": "c109ed7a56f25c977d1e0ca80337a40083fe4c8172eed6a2dc8ad081f91244b2", - "story": "ded163eb533ec4a7e47cd4351644d35c209852c8418503e060b988ddf6765414" + "story": "5549d66a2bfdd759b00954e048f79944245a07cdbf593cd92e99865ee5042602" }, "target_manifest_sha256": "007d9f3da8d54363a75e8af84cf0c9b762eafd6edf23679de48b08c62d584d55", - "target_proof_sha256": "1bd70995c42ce9d484c717d4169a000df36ce5181798c6f54b26138128daf296", + "target_proof_sha256": "824b1705373431eb0a79dd3774eaf1890a1849cc1de06375b1506a2f2073abd2", "tokens_compared": 3130 } diff --git a/gemma-4-12B-it/equivalence.json b/gemma-4-12B-it/equivalence.json index 5b115e389d20a5a48f7d84dd166767a0b8a5817b..4e240c3df8390919e3cf0d9b69e51e7db9885e94 100644 --- a/gemma-4-12B-it/equivalence.json +++ b/gemma-4-12B-it/equivalence.json @@ -1,5 +1,5 @@ { - "baseline_metadata_sha256": "903a232363321fef2ac2ba77417d1c0f7d86ae961136361f99cc614b97930c16", + "baseline_metadata_sha256": "cae1346eba41af2dff667ca25643bfcd998ddd038d91dc0c30202e735274c9a4", "exact_logits_equal": true, "manifest_sha256": "007d9f3da8d54363a75e8af84cf0c9b762eafd6edf23679de48b08c62d584d55", "precision_plan": { @@ -35,16 +35,18 @@ "code-stdlib-_pydecimal-1k", "long-book-2k" ], - "restored_metadata_sha256": "673a8647733567367ad5ec14f9a86493dddf1dbf12cb0d892adca329a9c6ed92", + "restored_metadata_sha256": "ce184930064503e5058292bb338332df1fc6a1575b3460d4ea9eb4542d01e52f", "runtime_env": { + "GEMMA4_DECODE_KV_ROWS": "1", + "GEMMA4_DECODE_MM_DN": "1", "GEMMA4_FUSE_GELU_MUL": "1", "GEMMA4_VERIFY_SDPA": "batched" }, "runtime_identity": { "binaries": { - "build/lib/_ttnncpp.so": "72186b132bf7be85bbb6e1b54bf6de125fbd7a557b12711c84efc6d3e2e08f8d", + "build/lib/_ttnncpp.so": "90680abdecf91fff303ea439313eef14e2358ba4b47ca68971171f63dbb5de33", "build/lib/libtt_metal.so": "0497308232f180a0973406e2afd290b9eafa9e310a16b09bf69acc4434eb0e78", - "ttnn/ttnn/_ttnn.so": "a6448405abd714d6ba310dbbc409c01842f42017b8c1d6bb1c30503fb194c7b2" + "ttnn/ttnn/_ttnn.so": "0d4a1051d2d284d14be8d72d0d818436cb0c736f17bdf39a8e82026dfab43300" }, "gemma4_tree": { "__init__.py": "501d5b563ca412fa21db82bf39117b5480ecf99bcaf11996b9c3d5f1093bb9a8", @@ -58,15 +60,15 @@ "tt/assistant/model.py": "0e3247c34a359a73bd4c8ecc7365e028372084554eeeba41682d1fd818c296af", "tt/attention/__init__.py": "df24b2bf4311ffaee781692ff8f465dea0f415e2323cf87fba4ce22227edb8ae", "tt/attention/config.py": "01b0f3e14c31aae5ed01104aee7f4cb03913007c0d0018f9d35d341d3fe049ef", - "tt/attention/decode.py": "0ec303ff75ded2a03f773c22aba27c719e2f055c71b9560fe18661d0272fd3db", + "tt/attention/decode.py": "df1f33dce05739d867ff0f4f11facabfa5625818118d8cf4d5fc8e97df36c4ac", "tt/attention/kv_cache.py": "3a1e9f520cdde61dd49550dbfca1072d009ee34878bafe5923c655821c09bff7", "tt/attention/kv_cache_hybrid.py": "9074a2d1de4b3a65d11271370f1fce39a78e5523b841752bcaa87dccb4b56470", - "tt/attention/operations.py": "f5cda2ade4031a66e631a732199f3b4b9a00b45ef2e252aee0d3138406bbcc43", + "tt/attention/operations.py": "8e87380345b46db028f882df1df2ea05b59a4e0278ed8516a1f15b9663384a65", "tt/attention/prefill.py": "32c6e1c9423b9c22f14558c33d76d9164cb8a27c939cb5d1ccb6da039b7c831f", "tt/attention/weights.py": "3ec94d60c6d1085bed0bf83fbc450aa9599a476341d3c6d77466b1b2e17c00d6", "tt/ccl.py": "3580f2e25950381465e55dd2886d2832c33d8d6fc618b5a2b94088d0b1a91105", "tt/common.py": "e4369b067c3c2e7930fb538a34ab04b237940048e771f2f8f0560776f03a578f", - "tt/decode_mm.py": "93ebedf2eb3b0c78ebbddfc00c630b13c726d9b13bb07fb53a073a207731bae4", + "tt/decode_mm.py": "85e5a625134c36925d559c4b7557b0139e824a77c4e724bf3b647015b77f7fac", "tt/experts/__init__.py": "567aef533a30a692e4c0b28695c15d13e2cfc024ee621e4147c836b63220d5c3", "tt/experts/config.py": "80bd7d18ffd0e0b6bd969d05a8e9972ff8473dc63631f1cb6de0bd619a50f33b", "tt/experts/decode.py": "112b9df4b88e8395053fdbf69315e3a43e9f3f4877e699ca5a48462ccd8e804b", @@ -98,8 +100,8 @@ } }, "runtime_run_sha256": { - "gemma4_tree_sha256": "ffae949a9a260c3a03a9b0a2ef891668bf79fb790df84f53e24e19048e9fe1fb", - "tool_sha256": "36485b68752f51c3642e3dd777f5da0cdb9de61ab819d5680f95402e1a24ede3" + "gemma4_tree_sha256": "caa9c51c76651f429cf7ea91c290ccac9d9d9d7541bb5afee0cbcf5df135dac4", + "tool_sha256": "505baaad9c88ccfb2e4f264f557477b30ae066673837b6c7fe56f0dc197f0934" }, "scope": "bit-exact full-vocabulary logits, original weights quantized at load vs native reload, same runtime; not a quality certification", "tokens_compared": 4093 diff --git a/launch.py b/launch.py index 2492f46ffd3a4c838e857de1f92c9ee40bb9b599..c3eab6b2d421ed5f6863a43a2c59932c0f1053d5 100644 --- a/launch.py +++ b/launch.py @@ -22,8 +22,9 @@ from urllib.request import Request, urlopen MODEL_REPO = "Lottolabs/gemma-4-12B-it-TT-BFP8-P150" -RELEASE_ID = "p3c-20260930" -IMAGE_TAG = "lottolabs/gemma4-12b-tt-p150:p3c" +RELEASE_ID = "p3d-20261001" +IMAGE_TAG = "lottolabs/gemma4-12b-tt-p150:p3d" +IMAGE_ID = "sha256:331e79b24fa2575c436158be994c6bc0e184d776d15305e82538d2e4d646f544" CHECKPOINT_SUBDIR = "gemma-4-12B-it" DRAFTER_SUBDIR = "gemma-4-12B-it-assistant" HF_BASE = "https://huggingface.co" @@ -82,8 +83,7 @@ def validate_release(release): image = release.get("image") image_id = release.get("image_id") image_archive = release.get("image_archive") - if (image != IMAGE_TAG or not isinstance(image_id, str) - or re.fullmatch(r"sha256:[0-9a-f]{64}", image_id) is None + if (image != IMAGE_TAG or image_id != IMAGE_ID or image_archive != "runtime-image.tar.gz"): raise ValueError("Release must identify the checksum-pinned bundled runtime image") files = file_records(release.get("files")) diff --git a/release-manifest.json b/release-manifest.json index 939fab89462bcbdc5ec132cc4cbe38647b23ed64..878870ecbde480a0c9bc00a5b669f613d5268611 100644 --- a/release-manifest.json +++ b/release-manifest.json @@ -7,125 +7,233 @@ "sha256": "cfc7749b96f63bd31c3c42b5c471bf756814053e847c10f3eb003417bc523d30" }, "README.md": { - "bytes": 10555, - "sha256": "98ed9bdcd1983193f71e4d7e327eebe72a3be34c6d19e8f378aa5b7e5e370a54" + "bytes": 12912, + "sha256": "fb17f137237deb11b801ab11ca09a446b02e6a35c9a278f294daee9a82faedfb" }, - "evidence/http-p3c/exact.json": { + "evidence/http-p3d/exact.json": { + "bytes": 3099, + "sha256": "2d2a52c8c9dab817a10c4bec4c9d483ee1214b3a22cc883fb754549dac8dbf0e" + }, + "evidence/http-p3d/long.json": { + "bytes": 5973, + "sha256": "4ae89585c639211abecfb7347a476c18d43f809385b0eb2a8077d3e994e80c8e" + }, + "evidence/http-p3d/tput-128/requests-128-c1.json": { + "bytes": 5998, + "sha256": "534ef3c7105f7f2d8fc7e0bad4836b6d5507139c0a18a0b97377d4a8051ba035" + }, + "evidence/http-p3d/tput-128/summary.json": { + "bytes": 876, + "sha256": "4d213e69b02e4b71062c358b3cda4a8db94a615b43c16f3dfd619c053788dac6" + }, + "evidence/http-p3d/tput-128/warmup-128.json": { + "bytes": 2966, + "sha256": "7a94e4f69a9e5ac6de2d96cd7eb9647fda3c50ba406cdaa2c68bba7aba5c74ad" + }, + "evidence/http-p3d/tput-131072/requests-131072-c1.json": { + "bytes": 6093, + "sha256": "2bbf64804564dfaf71c0f5b78db64db026eaf16ab5d515125e295f145c137548" + }, + "evidence/http-p3d/tput-131072/summary.json": { + "bytes": 880, + "sha256": "3b5fc67ff3e95a6ba8fc59f721fb79ba6df82993af27b8f6b1ef280cdd134ddf" + }, + "evidence/http-p3d/tput-131072/warmup-131072.json": { + "bytes": 3015, + "sha256": "3b0071bad4dda913079bd8d853d419220e5c6cb58b34618f3059d3e4c305d0a4" + }, + "evidence/http-p3d/tput-2048/requests-2048-c1.json": { + "bytes": 6117, + "sha256": "2f6176f102977e953ddd33a8d938fe0cbf1661c1466a287e406c2f1ac7215174" + }, + "evidence/http-p3d/tput-2048/summary.json": { + "bytes": 878, + "sha256": "4206f0885321e0cd121e2efbbb91daf2c87bb3343081e81709fb2f80b9d7fcb0" + }, + "evidence/http-p3d/tput-2048/warmup-2048.json": { + "bytes": 3025, + "sha256": "cc417c28c2a3952197e27e516b10c358665bd9792ef1cebb990de8037287292f" + }, + "evidence/http-p3d/tput-261632/requests-261632-c1.json": { + "bytes": 5942, + "sha256": "1cfdf00f237dccc9e9ddfed7d0f64f224f574b048d5b395edb0f81ced7a0dfdb" + }, + "evidence/http-p3d/tput-261632/summary.json": { + "bytes": 881, + "sha256": "bad5d6e144e98944b706ad12b62017d4355f4148b3ff2d6ee72543f3cefe72c8" + }, + "evidence/http-p3d/tput-261632/warmup-261632.json": { + "bytes": 2936, + "sha256": "c6a74eafe2f81e64b185d4327d79ca49bcc52f52a9572d353d55cf054656c196" + }, + "evidence/http-p3d/tput-32768/requests-32768-c1.json": { + "bytes": 6280, + "sha256": "91e11398058527c1d04f3911de6218b2f3f183dfadd3977280e738dd59907aa2" + }, + "evidence/http-p3d/tput-32768/summary.json": { + "bytes": 878, + "sha256": "89effe9fecb4560388df68a68cba91117c9861f68be7991c3f89773c1e515117" + }, + "evidence/http-p3d/tput-32768/warmup-32768.json": { + "bytes": 3106, + "sha256": "7de3bc8246a36dfcd42b9802f4312874fa5a4919f93c01552326727bcb14b3e3" + }, + "evidence/http-p3d/tput-8192/requests-8192-c1.json": { + "bytes": 6182, + "sha256": "7aa5cf9a79ccd791f80a4ddb1350d3da0ecc7a14901a4cee7ec81df8e2fc1fa3" + }, + "evidence/http-p3d/tput-8192/summary.json": { + "bytes": 877, + "sha256": "2657f8971531b9f9bb940856f83740aebf1fdd7f1c6f313356ad13c102b54ac3" + }, + "evidence/http-p3d/tput-8192/warmup-8192.json": { + "bytes": 3057, + "sha256": "5df2100d05536b268b8a811e814d66628320eca6016f7968d46f9747eeba74a1" + }, + "evidence/http-p3d/tput_vs_direct.json": { + "bytes": 235, + "sha256": "1c7778066e729cf8e2f302418c63f1640eba7d88a6b943b9681f3819c1579adb" + }, + "evidence/localmaxxing/p3d-2k/prompt-2k.txt": { + "bytes": 8155, + "sha256": "cd754ee70eca7fe4db4c243e6ca985e46e46dd89157735b8ebadd92dafcff4a1" + }, + "evidence/localmaxxing/p3d-2k/speed-test.json": { + "bytes": 15836, + "sha256": "e341a88dd41612504017df8c950b3670d8dc01b057887f697a8b761dd8a32106" + }, + "evidence/previous-p3c/http/exact.json": { "bytes": 3099, "sha256": "dcd509a51cd400fd48f042a8147c99df5381132386b53347b7b37a5f48b16094" }, - "evidence/http-p3c/long.json": { + "evidence/previous-p3c/http/long.json": { "bytes": 5976, "sha256": "67124b62e2fd13138863abda1b821e6e66e066541e572e5ec321a4b275b0e959" }, - "evidence/http-p3c/tput-128/requests-128-c1.json": { + "evidence/previous-p3c/http/tput-128/requests-128-c1.json": { "bytes": 5955, "sha256": "dfc6acee93d6a341f92e453592fd2ee5b52329a229be413fa4407f7bb681c7c0" }, - "evidence/http-p3c/tput-128/summary.json": { + "evidence/previous-p3c/http/tput-128/summary.json": { "bytes": 876, "sha256": "316fb623119a50661ae7e768c3bf8ad8733683ceb56953cdbdf2a2d963a5df9f" }, - "evidence/http-p3c/tput-128/warmup-128.json": { + "evidence/previous-p3c/http/tput-128/warmup-128.json": { "bytes": 2945, "sha256": "44bdcbb47481b62ab615f72a2bc29c2a4d2524a4ec318f8fd19159ef4ac00c5f" }, - "evidence/http-p3c/tput-131072/requests-131072-c1.json": { + "evidence/previous-p3c/http/tput-131072/requests-131072-c1.json": { "bytes": 5955, "sha256": "fdb5caa830280f0ba36b1045a905cfeb02bee98e54d72bfa8a0f54f712ea4c20" }, - "evidence/http-p3c/tput-131072/summary.json": { + "evidence/previous-p3c/http/tput-131072/summary.json": { "bytes": 882, "sha256": "ea4fa34dd27a8011fe1cdad7cf26f5acd3af67c0d7cab5c1dd8a3f7c680005c4" }, - "evidence/http-p3c/tput-131072/warmup-131072.json": { + "evidence/previous-p3c/http/tput-131072/warmup-131072.json": { "bytes": 2945, "sha256": "9aa37602c1c2ddfbe18be1c210ed50c953a20eda00f1a94b14fb3e2c249dedbd" }, - "evidence/http-p3c/tput-2048/requests-2048-c1.json": { + "evidence/previous-p3c/http/tput-2048/requests-2048-c1.json": { "bytes": 6177, "sha256": "d73bc532e74a494cb908c00f98f3745e465472e6fa66290b3de69ea0d673849a" }, - "evidence/http-p3c/tput-2048/summary.json": { + "evidence/previous-p3c/http/tput-2048/summary.json": { "bytes": 877, "sha256": "b0b79e75d5f23f31261ff91863a7737dc9e0849476eac663e7577b587daa7cf8" }, - "evidence/http-p3c/tput-2048/warmup-2048.json": { + "evidence/previous-p3c/http/tput-2048/warmup-2048.json": { "bytes": 3057, "sha256": "abc0819a5df996ee123cba52233fc2d81c16df49fffc0c2e1bd8e934eb28333c" }, - "evidence/http-p3c/tput-261632/requests-261632-c1.json": { + "evidence/previous-p3c/http/tput-261632/requests-261632-c1.json": { "bytes": 5914, "sha256": "e211831e0530c7b1cd1aa356a95beaec3a2215a7fe77d9908e858b9f3b453c05" }, - "evidence/http-p3c/tput-261632/summary.json": { + "evidence/previous-p3c/http/tput-261632/summary.json": { "bytes": 879, "sha256": "9ebdbacd8886c3faf8e2753fb010cd2a87d00d3da3a085051b9663265b6bcff5" }, - "evidence/http-p3c/tput-261632/warmup-261632.json": { + "evidence/previous-p3c/http/tput-261632/warmup-261632.json": { "bytes": 2927, "sha256": "94a93fe15d162528cb287df6f4ff7f95244e05ea8cfaae89076b4b518eaf39d4" }, - "evidence/http-p3c/tput-32768/requests-32768-c1.json": { + "evidence/previous-p3c/http/tput-32768/requests-32768-c1.json": { "bytes": 5959, "sha256": "37e1fc84216f03c5db76d35995054e367adc55c4fa6f2892dabb61ed6be1c12b" }, - "evidence/http-p3c/tput-32768/summary.json": { + "evidence/previous-p3c/http/tput-32768/summary.json": { "bytes": 878, "sha256": "6862ec9e00ddad920c0324a442ee5d28737a2e839de8e89c693c0cd809df1774" }, - "evidence/http-p3c/tput-32768/warmup-32768.json": { + "evidence/previous-p3c/http/tput-32768/warmup-32768.json": { "bytes": 2948, "sha256": "e4f3cc40e54e56410341d29bd26d9102f06dbd00e35901c4a05b2a05f085c3bd" }, - "evidence/http-p3c/tput-8192/requests-8192-c1.json": { + "evidence/previous-p3c/http/tput-8192/requests-8192-c1.json": { "bytes": 6022, "sha256": "e67e8fbe22e47880715b26cee62f0d999a6395388d36774117c880f6ff6e4842" }, - "evidence/http-p3c/tput-8192/summary.json": { + "evidence/previous-p3c/http/tput-8192/summary.json": { "bytes": 875, "sha256": "af278ee14869619894a8ff9a4c12c69f155d35c26f3b7868190d9559ec001981" }, - "evidence/http-p3c/tput-8192/warmup-8192.json": { + "evidence/previous-p3c/http/tput-8192/warmup-8192.json": { "bytes": 2978, "sha256": "957010a8c724f2a1f82ce7ffc5dcbdabfeb2e1420c5a9274fda0faa8c85c05bb" }, - "evidence/http-p3c/tput_vs_direct.json": { + "evidence/previous-p3c/http/tput_vs_direct.json": { "bytes": 235, "sha256": "1c7778066e729cf8e2f302418c63f1640eba7d88a6b943b9681f3819c1579adb" }, - "evidence/localmaxxing/speed-test.json": { + "evidence/previous-p3c/localmaxxing-speed-test.json": { "bytes": 6046, "sha256": "f03041863057e793e816ded6f048899034e04f11a3a4dd0fd102a04d40e7b5de" }, - "evidence/public-download-verification.json": { + "evidence/previous-p3c/public-download-verification.json": { "bytes": 1337, "sha256": "92c92f98fb43b89ff04bd970974f23a3647264a722da2e674b033cf0d8523a32" }, - "evidence/public-serving-smoke.json": { + "evidence/previous-p3c/public-serving-smoke.json": { "bytes": 1781, "sha256": "791d5b2dff0a130024200ab016e671c1d74495229b3461f83464bb35124478a9" }, - "evidence/public-serving-smoke/requests-2048-c1.json": { + "evidence/previous-p3c/public-serving-smoke/requests-2048-c1.json": { "bytes": 6177, "sha256": "07b8133a548b886bc30794b9f984f0a5dee9b9301f7c90ff671806cbb4ce766c" }, - "evidence/public-serving-smoke/summary.json": { + "evidence/previous-p3c/public-serving-smoke/summary.json": { "bytes": 875, "sha256": "a386aef5e382c304cc3c5315b8021166dfadec3d256b6af7887acf35383262e2" }, - "evidence/public-serving-smoke/warmup-2048.json": { + "evidence/previous-p3c/public-serving-smoke/warmup-2048.json": { "bytes": 3058, "sha256": "3cb6e6002631824144551034aeed233081c44182771b51d8c3ade48555ff7426" }, + "evidence/previous-p3c/runtime-gate-dtf-score.json": { + "bytes": 1130, + "sha256": "b118bf712e7d70058213f378920c0b129a2366b554f23d489c441ecbb2e5f2ba" + }, + "evidence/quality/p3d-gate/book-2048-score.json": { + "bytes": 371, + "sha256": "156ecb1746f24b3c6c7d6143e74034caf533b286270baab8f9bb7ad117fc8750" + }, + "evidence/quality/p3d-gate/dtf-score.json": { + "bytes": 1139, + "sha256": "51c580756dd62962bcad40d1cf53100319f8d9615902f2a5db24f78af42f510a" + }, + "evidence/quality/p3d-gate/gsm8k-100-exact-build.json": { + "bytes": 351, + "sha256": "9bb6af5e16abd148f7b7d528e045587bec96b94249a89be227352840257fb703" + }, + "evidence/quality/p3d-gate/gsm8k-100.json": { + "bytes": 349, + "sha256": "82c012ee2d4dd6f5986f7a7bbbec425589405c98e487b93747457b0180e220ab" + }, "evidence/quality/quant_table.json": { "bytes": 9571, "sha256": "32904d8e6be2daea2b797cc71b8d587e06605dc23421d96e1912c8593ed4b991" }, - "evidence/quality/runtime-gate-dtf-score.json": { - "bytes": 1130, - "sha256": "b118bf712e7d70058213f378920c0b129a2366b554f23d489c441ecbb2e5f2ba" - }, "gemma-4-12B-it-assistant/config.json": { "bytes": 2346, "sha256": "b6f19209588fcefe41f65b193fad6148446253c470d36e29441ecc5158a54e6d" @@ -143,8 +251,8 @@ "sha256": "3279c173daddd7186e79d652ad94022415736d3a1370625696c898429b06d6df" }, "gemma-4-12B-it-assistant/spec_equivalence.json": { - "bytes": 6567, - "sha256": "032cc468d6bf33bca7267be0e70d1d6932e7dcc4dfb97acb5a75e736f2ee8d0d" + "bytes": 6629, + "sha256": "1c09aae0f94d18ae7075800ccf55a7866769bec1fe41475ec781c054ad4ee489" }, "gemma-4-12B-it-assistant/tensors/assistant_tensor_cache_bfp8/final_norm/weight_dtype_BFLOAT16_layout_ROW_MAJOR.tensorbin": { "bytes": 2392, @@ -347,8 +455,8 @@ "sha256": "478c46e8d2c52d5c2d85bf67e3b3e8c90e7c9d91086cee27e3c267907e936bd9" }, "gemma-4-12B-it/equivalence.json": { - "bytes": 6617, - "sha256": "1bd70995c42ce9d484c717d4169a000df36ce5181798c6f54b26138128daf296" + "bytes": 6679, + "sha256": "824b1705373431eb0a79dd3774eaf1890a1849cc1de06375b1506a2f2073abd2" }, "gemma-4-12B-it/generation_config.json": { "bytes": 260, @@ -2499,8 +2607,8 @@ "sha256": "a62f4e85a47c0c136edaaa3a4f591fd6783717299a9def47e5ad03a49f6a5eb9" }, "launch.py": { - "bytes": 15438, - "sha256": "dd9ce6c7f2776f93acd7091dcd7c60eb0460a7d18e8c655e44e11846fc5d5df6" + "bytes": 15444, + "sha256": "22186ea9c0c9fb94823f666c740555459728d5750278315ac2ab364ccbb4cbfb" }, "native_checkpoint.py": { "bytes": 23890, @@ -2511,20 +2619,20 @@ "sha256": "3aa234a04e768e734a8ede9cdd9da1716366b06ea47b22a5ca466a49ad4fb56f" }, "reproduction.json": { - "bytes": 2066, - "sha256": "30578daa4ad6a0e529ac99470c5131c7738ed2a60e70fbaedb32ac0de5736d0a" + "bytes": 2701, + "sha256": "2f916687347fd9b47ddb7dc18585a80a7285aacf4bb727db4fe8b52218d203fa" }, "runtime-image.tar.gz": { - "bytes": 5383095342, - "sha256": "782cab35025b19b55fe2212a2619cd97dfd5f6da6c04a00e68b9c69a8f92d2dd" + "bytes": 5425816554, + "sha256": "59074c4431d616cd56e0a001684f320825cb3f5ade7d16d6e3e58667d127e6ce" }, "runtime-release.json": { "bytes": 136583, - "sha256": "48fdc1e0b7cb088fcd12c56de79300f90cb6a5a79a49e029f3ec6deec565b66a" + "sha256": "7d4c57f97b7ab4f2f353b4e4dffcd56899b24e73a459d5fc14578aaf48ae59e9" }, - "runtime/Dockerfile.fast1": { - "bytes": 267, - "sha256": "d9e47832fb3354959883d3069a9f1c7e2cc8be0a6e2b9d37f22c36c154749fdd" + "runtime/Dockerfile.fast2": { + "bytes": 326, + "sha256": "b99769e8cd55f06640377a142de2199d0a2cda7830081cab5a1974c15915d14d" }, "runtime/Dockerfile.p2b": { "bytes": 731, @@ -2534,9 +2642,13 @@ "bytes": 210, "sha256": "91fdd9a2cae0bffb6715422bfda322e47e9b05d40d99e3a6cc962aba89132b9e" }, - "runtime/Dockerfile.p3c": { + "runtime/Dockerfile.p3d": { "bytes": 254, - "sha256": "abaee1bb49548bf7b8abebe6519e68d9d7bc3fa9e86e3ef76f1ea9a86e64337c" + "sha256": "70c93a67b96223c1f0bcfdee039ba1ffb2e17ed1725678d9bad98fb1dd19ee18" + }, + "runtime/Dockerfile.ttnn-dn1": { + "bytes": 449, + "sha256": "005744b93e3b95864bd6152c3084226c448b48aa124311855fa4896dbfaa6fe6" }, "runtime/build.sh": { "bytes": 2293, @@ -2723,8 +2835,8 @@ "sha256": "01b0f3e14c31aae5ed01104aee7f4cb03913007c0d0018f9d35d341d3fe049ef" }, "runtime/gemma4/tt/attention/decode.py": { - "bytes": 50404, - "sha256": "0ec303ff75ded2a03f773c22aba27c719e2f055c71b9560fe18661d0272fd3db" + "bytes": 51136, + "sha256": "df1f33dce05739d867ff0f4f11facabfa5625818118d8cf4d5fc8e97df36c4ac" }, "runtime/gemma4/tt/attention/kv_cache.py": { "bytes": 5273, @@ -2735,8 +2847,8 @@ "sha256": "9074a2d1de4b3a65d11271370f1fce39a78e5523b841752bcaa87dccb4b56470" }, "runtime/gemma4/tt/attention/operations.py": { - "bytes": 25598, - "sha256": "f5cda2ade4031a66e631a732199f3b4b9a00b45ef2e252aee0d3138406bbcc43" + "bytes": 26179, + "sha256": "8e87380345b46db028f882df1df2ea05b59a4e0278ed8516a1f15b9663384a65" }, "runtime/gemma4/tt/attention/prefill.py": { "bytes": 24198, @@ -2755,8 +2867,8 @@ "sha256": "e4369b067c3c2e7930fb538a34ab04b237940048e771f2f8f0560776f03a578f" }, "runtime/gemma4/tt/decode_mm.py": { - "bytes": 4567, - "sha256": "93ebedf2eb3b0c78ebbddfc00c630b13c726d9b13bb07fb53a073a207731bae4" + "bytes": 5535, + "sha256": "85e5a625134c36925d559c4b7557b0139e824a77c4e724bf3b647015b77f7fac" }, "runtime/gemma4/tt/experts/__init__.py": { "bytes": 2988, @@ -2915,8 +3027,8 @@ "sha256": "e7e004902362d832cfea5f4e22eedcc530494c3e42ea1e720159712ae6d8e1d3" }, "runtime/source.json": { - "bytes": 2335, - "sha256": "03bfbc377872223a5aa72bee2148b941831969065e9241aa5dedd0d8ddb36c61" + "bytes": 3477, + "sha256": "dc0a865c919aecaee44bc113cced7f3be2d64172505f0a5c5fc166dcd8fb57f6" }, "runtime/tools/compare_gen.py": { "bytes": 1320, @@ -2927,12 +3039,20 @@ "sha256": "cc507fe91c3279813aa4bdf697e7a2a811d5f67ea25afe90206830ec4c3db9cd" }, "runtime/tools/tt_eval.py": { - "bytes": 39090, - "sha256": "36485b68752f51c3642e3dd777f5da0cdb9de61ab819d5680f95402e1a24ede3" + "bytes": 39367, + "sha256": "505baaad9c88ccfb2e4f264f557477b30ae066673837b6c7fe56f0dc197f0934" }, "runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/factory/matmul_multicore_reuse_mcast_1d_program_factory.cpp": { - "bytes": 283703, - "sha256": "1de00fc9d9d9ca537d7a90af22dd61b4723c54b6ecd236f7053124963789a62b" + "bytes": 297804, + "sha256": "fa967085d4bccc78e448454911c38468415c1403e3138e11dbbf9d176ffd3c67" + }, + "runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in0_sender_padding.cpp": { + "bytes": 23983, + "sha256": "a78857488315cac5257d5e8a93f03800bbf40acb528ba0edd0ad67660a2cbad1" + }, + "runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in1_sender_writer_padding.cpp": { + "bytes": 51192, + "sha256": "3620262683632718607446e2a8eba75ea74cd714eee00886d3b2a00ad4f08dc4" }, "serve_native.py": { "bytes": 7512, @@ -2941,13 +3061,19 @@ }, "integrity": "SHA256SUMS covers all payload files plus release-manifest.json; checksum file excludes itself", "model_repo": "Lottolabs/gemma-4-12B-it-TT-BFP8-P150", - "payload_bytes": 21277074233, - "release_id": "p3c-20260930", + "payload_bytes": 21319987546, + "previous_runtime": { + "image_id": "sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73", + "image_tag": "lottolabs/gemma4-12b-tt-p150:p3c", + "release_id": "p3c-20260930", + "repo_revision": "3d9001d5f52a5f3350b90e87703440a44255bd95" + }, + "release_id": "p3d-20261001", "reproduction": "reproduction.json", "runtime_archive": "runtime-image.tar.gz", - "runtime_archive_sha256": "782cab35025b19b55fe2212a2619cd97dfd5f6da6c04a00e68b9c69a8f92d2dd", - "runtime_image": "sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73", - "runtime_image_tag": "lottolabs/gemma4-12b-tt-p150:p3c", + "runtime_archive_sha256": "59074c4431d616cd56e0a001684f320825cb3f5ade7d16d6e3e58667d127e6ce", + "runtime_image": "sha256:331e79b24fa2575c436158be994c6bc0e184d776d15305e82538d2e4d646f544", + "runtime_image_tag": "lottolabs/gemma4-12b-tt-p150:p3d", "schema_version": 1, "scope": "Single-P150 TT-native all-BFP8 google/gemma-4-12B-it, text-only, single user, native 262,144-token context, greedy speculative decoding with the google/gemma-4-12B-it-assistant drafter (K=5)", "status": "complete", diff --git a/reproduction.json b/reproduction.json index 68a7015e2ae88fb2350b1c380efb79b8965953fb..7931598a1fa0624f3bac045fe359b10dce8180a7 100644 --- a/reproduction.json +++ b/reproduction.json @@ -1,6 +1,6 @@ { "schema_version": 1, - "release_id": "p3c-20260930", + "release_id": "p3d-20261001", "model_repo": "Lottolabs/gemma-4-12B-it-TT-BFP8-P150", "source_model": { "model_id": "google/gemma-4-12B-it", @@ -10,8 +10,8 @@ "model_id": "google/gemma-4-12B-it-assistant", "revision": "46d4c6f13f0ac0ad827b915669b8df9b81c64c51" }, - "runtime_image_tag": "lottolabs/gemma4-12b-tt-p150:p3c", - "runtime_image": "sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73", + "runtime_image_tag": "lottolabs/gemma4-12b-tt-p150:p3d", + "runtime_image": "sha256:331e79b24fa2575c436158be994c6bc0e184d776d15305e82538d2e4d646f544", "runtime_source": "runtime/source.json", "checkpoint_build": { "builder": "provenance/build_native_checkpoint.py", @@ -52,14 +52,28 @@ "drafter": "gemma-4-12B-it-assistant/spec_equivalence.json" }, "evidence": { - "quality": "evidence/quality/", - "http": "evidence/http-p3c/", - "localmaxxing": "evidence/localmaxxing/speed-test.json", + "quality_phase1": "evidence/quality/quant_table.json", + "quality_runtime_gate": "evidence/quality/p3d-gate/", + "http": "evidence/http-p3d/", + "localmaxxing": "evidence/localmaxxing/p3d-2k/speed-test.json", "public_download": "evidence/public-download-verification.json", - "public_serving": "evidence/public-serving-smoke.json" + "public_serving": "evidence/public-serving-smoke.json", + "previous_runtime": "evidence/previous-p3c/" }, "runtime_distribution": { "path": "runtime-image.tar.gz", "contract": "launch.py verifies the archive checksum, loads it with Docker, and requires the exact recorded image ID." + }, + "runtime_env": { + "GEMMA4_VERIFY_SDPA": "batched", + "GEMMA4_FUSE_GELU_MUL": "1", + "GEMMA4_DECODE_KV_ROWS": "1", + "GEMMA4_DECODE_MM_DN": "1", + "binding": "Set in the image (layer fast2) and recorded as runtime_env in both proofs; the loader rejects any other numerics environment." + }, + "previous_runtime": { + "release_id": "p3c-20260930", + "image_id": "sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73", + "repo_revision": "3d9001d5f52a5f3350b90e87703440a44255bd95" } } diff --git a/runtime-image.tar.gz b/runtime-image.tar.gz index bb0509b33ab64c9c073f0b6a4df2eadfb9c67ef7..568c16ce391d63517f1949080f29923570d10b81 100644 --- a/runtime-image.tar.gz +++ b/runtime-image.tar.gz @@ -1,3 +1,3 @@ version https://git-lfs.github.com/spec/v1 -oid sha256:782cab35025b19b55fe2212a2619cd97dfd5f6da6c04a00e68b9c69a8f92d2dd -size 5383095342 +oid sha256:59074c4431d616cd56e0a001684f320825cb3f5ade7d16d6e3e58667d127e6ce +size 5425816554 diff --git a/runtime-release.json b/runtime-release.json index 629f8a4fa8ac34c64de5a1842fe91587836af155..b248dded5dc0944925ccd4a6c06e8ba7cca41f88 100644 --- a/runtime-release.json +++ b/runtime-release.json @@ -25,8 +25,8 @@ "sha256": "3279c173daddd7186e79d652ad94022415736d3a1370625696c898429b06d6df" }, "gemma-4-12B-it-assistant/spec_equivalence.json": { - "bytes": 6567, - "sha256": "032cc468d6bf33bca7267be0e70d1d6932e7dcc4dfb97acb5a75e736f2ee8d0d" + "bytes": 6629, + "sha256": "1c09aae0f94d18ae7075800ccf55a7866769bec1fe41475ec781c054ad4ee489" }, "gemma-4-12B-it-assistant/tensors/assistant_tensor_cache_bfp8/final_norm/weight_dtype_BFLOAT16_layout_ROW_MAJOR.tensorbin": { "bytes": 2392, @@ -229,8 +229,8 @@ "sha256": "478c46e8d2c52d5c2d85bf67e3b3e8c90e7c9d91086cee27e3c267907e936bd9" }, "gemma-4-12B-it/equivalence.json": { - "bytes": 6617, - "sha256": "1bd70995c42ce9d484c717d4169a000df36ce5181798c6f54b26138128daf296" + "bytes": 6679, + "sha256": "824b1705373431eb0a79dd3774eaf1890a1849cc1de06375b1506a2f2073abd2" }, "gemma-4-12B-it/generation_config.json": { "bytes": 260, @@ -2385,19 +2385,19 @@ "sha256": "cc507fe91c3279813aa4bdf697e7a2a811d5f67ea25afe90206830ec4c3db9cd" }, "runtime-image.tar.gz": { - "bytes": 5383095342, - "sha256": "782cab35025b19b55fe2212a2619cd97dfd5f6da6c04a00e68b9c69a8f92d2dd" + "bytes": 5425816554, + "sha256": "59074c4431d616cd56e0a001684f320825cb3f5ade7d16d6e3e58667d127e6ce" }, "serve_native.py": { "bytes": 7512, "sha256": "a9ae81f0f28d1f5c7c093817485cd72d28df4072714a1a30672c43e0166a88d9" } }, - "image": "lottolabs/gemma4-12b-tt-p150:p3c", + "image": "lottolabs/gemma4-12b-tt-p150:p3d", "image_archive": "runtime-image.tar.gz", - "image_id": "sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73", + "image_id": "sha256:331e79b24fa2575c436158be994c6bc0e184d776d15305e82538d2e4d646f544", "model_repo": "Lottolabs/gemma-4-12B-it-TT-BFP8-P150", - "release_id": "p3c-20260930", + "release_id": "p3d-20261001", "schema_version": 1, "upstream": { "drafter": { diff --git a/runtime/Dockerfile.fast1 b/runtime/Dockerfile.fast2 similarity index 75% rename from runtime/Dockerfile.fast1 rename to runtime/Dockerfile.fast2 index 595c37dc3365d9846f3b98f5210e5c66ef365b28..e94fc43e1cc3b6525d54ec9aba21099d094ef0ce 100644 --- a/runtime/Dockerfile.fast1 +++ b/runtime/Dockerfile.fast2 @@ -1,7 +1,9 @@ -FROM gemma4-12b:p2c +FROM gemma4-12b:ttnn-dn1 USER root COPY --chown=1000:1000 gemma4/ /home/container_app_user/tt-metal/models/demos/gemma4/ COPY --chown=1000:1000 tools/ /home/container_app_user/gemma4-tools/ ENV GEMMA4_VERIFY_SDPA=batched ENV GEMMA4_FUSE_GELU_MUL=1 +ENV GEMMA4_DECODE_KV_ROWS=1 +ENV GEMMA4_DECODE_MM_DN=1 USER container_app_user diff --git a/runtime/Dockerfile.p3c b/runtime/Dockerfile.p3d similarity index 91% rename from runtime/Dockerfile.p3c rename to runtime/Dockerfile.p3d index ae01d07ba86170a8fcaca29c3276ee96bec6b1bd..5d500a8c15b2de22295c7ec1c07e8516f49b950d 100644 --- a/runtime/Dockerfile.p3c +++ b/runtime/Dockerfile.p3d @@ -1,4 +1,4 @@ -FROM gemma4-12b:fast1 +FROM gemma4-12b:fast2 USER root COPY --chown=1000:1000 gemma4/ /home/container_app_user/tt-metal/models/demos/gemma4/ diff --git a/runtime/Dockerfile.ttnn-dn1 b/runtime/Dockerfile.ttnn-dn1 new file mode 100644 index 0000000000000000000000000000000000000000..c24c2b0565eef904540d544c04c8ea9a544da66d --- /dev/null +++ b/runtime/Dockerfile.ttnn-dn1 @@ -0,0 +1,10 @@ +FROM gemma4-12b:p2c +USER root +COPY --chown=1000:1000 ttnn/ /home/container_app_user/tt-metal/ttnn/ +USER container_app_user +WORKDIR /home/container_app_user/tt-metal +RUN cmake --build build_Release --target ttnn --parallel 24 \ + && cp build_Release/ttnn/_ttnncpp.so build/lib/_ttnncpp.so \ + && cp build_Release/ttnn/_ttnn.so build/lib/_ttnn.so \ + && cp build_Release/ttnn/_ttnn.so ttnn/ttnn/_ttnn.so +WORKDIR /home/container_app_user/app/src diff --git a/runtime/gemma4/tt/attention/decode.py b/runtime/gemma4/tt/attention/decode.py index c22869375d3aca4518951948e6e922c665599e49..ffa4fe43889e736b8d8dfc7014f98ce6a56326d0 100644 --- a/runtime/gemma4/tt/attention/decode.py +++ b/runtime/gemma4/tt/attention/decode.py @@ -216,11 +216,24 @@ def decode_forward( paged_modulo_kwargs = _position_modulo_kwargs(config) if kv_cache is not None: k_cache, v_cache = kv_cache - if not is_kv_shared: + if ( + not is_kv_shared + and page_table is not None + and os.environ.get("GEMMA4_DECODE_KV_ROWS") == "1" + and tp == 1 + and not paged_modulo_kwargs + and tt_k.shape[1] == 1 + and position_idx is not None + and position_idx.dtype == ttnn.uint32 + and k_cache.dtype == ttnn.bfloat16 + ): + # One DMA-only launch writes this row's K and V (bit copy; same kernel as the verify rows) + # instead of two reshards + two paged_update_cache calls. + kv_rows_write(tt_k, tt_v, k_cache, v_cache, position_idx, page_table, 1, k_row_major=True, v_row_major=True) + elif not is_kv_shared: # After HF-style RoPE, tensors may be in DRAM. Move to HEIGHT_SHARDED for cache update. tt_k = ttnn.to_memory_config(tt_k, q_sharded_mem) tt_v = ttnn.to_memory_config(tt_v, q_sharded_mem) - if page_table is not None: # Per-device kv-head count of the layer's input view. When the cache # was allocated for a different layer type under HMA cross-group @@ -445,11 +458,11 @@ def verify_rows_forward( xqkv = apply_qkv_projection(hidden_states, weights) tt_q, tt_k, tt_v = split_qkv_heads_decode(xqkv, config, weights.is_global, kv_replicated=weights.kv_replicated) q = apply_per_head_norm(ttnn.to_memory_config(tt_q, ttnn.DRAM_MEMORY_CONFIG), weights.q_norm_weight, - config.rms_norm_eps, with_scale=True) + config.rms_norm_eps, with_scale=True, reshape=False) k = apply_per_head_norm(ttnn.to_memory_config(tt_k, ttnn.DRAM_MEMORY_CONFIG), weights.k_norm_weight, - config.rms_norm_eps, with_scale=True) + config.rms_norm_eps, with_scale=True, reshape=False) v = apply_per_head_norm(ttnn.to_memory_config(tt_v, ttnn.DRAM_MEMORY_CONFIG), None, config.rms_norm_eps, - with_scale=False) + with_scale=False, reshape=False) # [1, P, heads, D] -> [1, heads, P, D]: rows become the sequence of the fused RoPE. q_t = ttnn.transpose(q, 1, 2) k_t = ttnn.transpose(k, 1, 2) diff --git a/runtime/gemma4/tt/attention/operations.py b/runtime/gemma4/tt/attention/operations.py index 61b691789e214d35cfe9921971fa7fb4fde41e7e..f131f36e63b922e86d297ea79634a9550b87ada4 100644 --- a/runtime/gemma4/tt/attention/operations.py +++ b/runtime/gemma4/tt/attention/operations.py @@ -102,7 +102,7 @@ def split_qkv_heads_prefill( -def apply_per_head_norm(tensor, weight, eps, with_scale=True, memory_config=None): +def apply_per_head_norm(tensor, weight, eps, with_scale=True, memory_config=None, reshape=True): """ Apply RMSNorm per-head on the head_dim dimension. @@ -113,6 +113,13 @@ def apply_per_head_norm(tensor, weight, eps, with_scale=True, memory_config=None activation resident on L1 (packed-verify decode path). ``None`` keeps the op's default (follows the input's layout). """ + if not reshape: + # rms_norm reduces each row over head_dim, so the 4D tensor can be normed in place of its + # flattened view; skips the two reshape copies for multi-row decode ([1, P, heads, head_dim]). + if with_scale and weight is not None: + return ttnn.rms_norm(tensor, weight=weight, epsilon=eps, memory_config=memory_config, + compute_kernel_config=norm_compute_config()) + return ttnn.rms_norm(tensor, epsilon=eps, memory_config=memory_config, compute_kernel_config=norm_compute_config()) orig_shape = tensor.shape head_dim = orig_shape[-1] if len(orig_shape) == 4 and orig_shape[0] > 1: diff --git a/runtime/gemma4/tt/decode_mm.py b/runtime/gemma4/tt/decode_mm.py index 010a723fda323cea47ba439baba21a9b147943b2..1902d7915f045247c35a03e463af168547e1cf28 100644 --- a/runtime/gemma4/tt/decode_mm.py +++ b/runtime/gemma4/tt/decode_mm.py @@ -46,6 +46,28 @@ _TABLE = { (32, 288): (11, 3, 2, 9), # global qkv 1024 -> 9216 } +# GEMMA4_DECODE_MM_DN=1 (needs a TTNN build with the per-worker dual-NoC in1 reader, image ttnn-dn*; +# ttnn/overlay branch dn): grids spanning all 10 rows, per-grid NoC-1 column masks, 16-slot weight CB. +# Measured 431-472 GB/s per shape (tools/dn_bench.py) vs 372-398 with _TABLE. Interleaved weights. +_DN_TABLE = { + (120, 480): (10, 10, 2, 5), # MLP gate / up + (480, 120): (6, 10, 4, 2), # MLP down + (120, 256): (7, 10, 2, 4), # sliding fused QKV + (120, 288): (10, 10, 2, 3), # global fused QKV + (128, 120): (8, 8, 4, 2), # sliding o_proj + (256, 120): (7, 9, 4, 2), # global o_proj +} +_DN_ENV = { + "TT_MM1D_IN1_DUAL_NOC": "3", + "TT_MM1D_IN1_NOC_COLS": "10x10:0000011111,7x10:1000101,6x10:100100,7x9:1100000,8x8:10010000", + "TT_MM1D_IN1_CB_SLOTS": "16", + "TT_MM1D_IN0_CHUNK_TILES": "64", +} +if os.environ.get("GEMMA4_DECODE_MM_DN") == "1": + for _k, _v in _DN_ENV.items(): + os.environ[_k] = _v + _TABLE.update(_DN_TABLE) + def enabled(): return os.environ.get("GEMMA4_DECODE_MM", "1") == "1" diff --git a/runtime/source.json b/runtime/source.json index d819f1cfd2b455b040c910ddf6ff5fbd57797389..32f925dcfe67fc632d539c1b639abaab59c3b9b0 100644 --- a/runtime/source.json +++ b/runtime/source.json @@ -1,12 +1,12 @@ { "schema_version": 1, - "runtime_image": "sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73", - "runtime_image_tag": "lottolabs/gemma4-12b-tt-p150:p3c", + "runtime_image": "sha256:331e79b24fa2575c436158be994c6bc0e184d776d15305e82538d2e4d646f544", + "runtime_image_tag": "lottolabs/gemma4-12b-tt-p150:p3d", "gemma4_tree": { "path": "runtime/gemma4/", "installed_at": "/home/container_app_user/tt-metal/models/demos/gemma4/", "git_branch": "master", - "git_commit": "ce67960382fffed8550b38717cedf574e821da06", + "git_commit": "061e48ddf6a0eab5dac1f958731fcb70768231d1", "note": "Byte-identical to the tree in the runtime image (compared file by file against the image); .git and __pycache__ excluded." }, "tools": { @@ -21,7 +21,17 @@ "ttnn_overlay": { "path": "runtime/ttnn-overlay/", "installed_at": "/home/container_app_user/tt-metal/ttnn/", - "content": "in0-multicast 1D matmul program factory; TTNN host libraries rebuilt in layer p2b" + "git_branch": "dn", + "git_commit": "efc465f4b1a44025ae4a33e4dbc7d371b3f9266c", + "content": "1D matmul program factory with in0-multicast chunking and the dual-NoC in1 reader (TT_MM1D_IN1_* column masks), in0/in1 dataflow reader kernels; TTNN host libraries rebuilt in layer ttnn-dn1", + "note": "Byte-identical to the three files in the runtime image. Layer p2b built an earlier revision of the factory file (overlay master 5200a780d6feabdc387f7c106f25f43f4a85a0a8) that ttnn-dn1 overwrites." + }, + "runtime_env": { + "GEMMA4_VERIFY_SDPA": "batched", + "GEMMA4_FUSE_GELU_MUL": "1", + "GEMMA4_DECODE_KV_ROWS": "1", + "GEMMA4_DECODE_MM_DN": "1", + "binding": "Set in the image (layer fast2) and recorded as runtime_env in both proofs; the loader rejects any other numerics environment." }, "build_chain": [ { @@ -40,18 +50,29 @@ "image_id": "sha256:b91c7e92ec16c507c916229ba7d675093379fe18ab05c18732eedcd3b91156f2" }, { - "layer": "fast1", - "dockerfile": "runtime/Dockerfile.fast1", - "image_id": "sha256:aade01a9328d028c1979070df60aa17ade1ce729e936c74815b67802a086fff9" + "layer": "ttnn-dn1", + "dockerfile": "runtime/Dockerfile.ttnn-dn1", + "build_context": "runtime/ttnn-overlay/ as ttnn/", + "image_id": "sha256:2807668702d59ed5cfd914e76df10ab1f6276d3fce3e7e1d174040156ec5d8d0" + }, + { + "layer": "fast2", + "dockerfile": "runtime/Dockerfile.fast2", + "image_id": "sha256:1c23787a9da68bdbb289a190deb86647061c36c7e5fec709be8e358e269ae4b7" }, { - "layer": "p3c", - "dockerfile": "runtime/Dockerfile.p3c", - "image_id": "sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73" + "layer": "p3d", + "dockerfile": "runtime/Dockerfile.p3d", + "image_id": "sha256:331e79b24fa2575c436158be994c6bc0e184d776d15305e82538d2e4d646f544" } ], "build_script": "runtime/build.sh", - "rebuild_note": "Provenance only: intermediate layers copied earlier gemma4 trees that the p3c layer overwrites; the Dockerfiles reference local intermediate tags. Serve the checksum-pinned runtime-image.tar.gz.", + "rebuild_note": "Provenance only: intermediate layers copied earlier gemma4 trees that the p3d layer overwrites; the Dockerfiles reference local intermediate tags. Serve the checksum-pinned runtime-image.tar.gz.", + "previous_runtime": { + "tag": "lottolabs/gemma4-12b-tt-p150:p3c", + "image_id": "sha256:53428e7f60a705b75b51b5166f44a912b19c3032981013fb788dd1ce9757aa73", + "gemma4_git_commit": "ce67960382fffed8550b38717cedf574e821da06" + }, "licenses": { "tt-metal": "runtime/licenses/tt-metal/", "vllm": "runtime/licenses/vllm/" diff --git a/runtime/tools/tt_eval.py b/runtime/tools/tt_eval.py index 8f70107e19682c5bd466c42fa8893b68bd9253c5..5bc87a3cc84e052482588a75b9893e1b6857e9d9 100644 --- a/runtime/tools/tt_eval.py +++ b/runtime/tools/tt_eval.py @@ -661,8 +661,12 @@ def run_specparts(args, ctx): vx0 = ttnn.from_torch(torch.full((1, spec.P), 100, dtype=torch.int64), dtype=ttnn.uint32, layout=ttnn.ROW_MAJOR_LAYOUT, device=mesh) def verify(): - lg, hid = spec.model.ttnn_rows_forward(vx0, b["v_pu"], [b[f"v_pi{i}"] for i in range(spec.P)], spec.page_tables, - kv_cache=spec.kv_cache) + if getattr(spec, "batched", False): + lg, hid = spec.model.ttnn_rows_forward(vx0, b["v_pu"], None, spec.page_tables, kv_cache=spec.kv_cache, + batched=(b["v_pall"], spec.page_tables_rep)) + else: + lg, hid = spec.model.ttnn_rows_forward(vx0, b["v_pu"], [b[f"v_pi{i}"] for i in range(spec.P)], spec.page_tables, + kv_cache=spec.kv_cache) return [spec._argmax_rows(lg, spec.P), hid] dec = get_decoder(ctx) diff --git a/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/factory/matmul_multicore_reuse_mcast_1d_program_factory.cpp b/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/factory/matmul_multicore_reuse_mcast_1d_program_factory.cpp index 2862d1a8b10b11654cb7e8ba29cf01714635b331..0f435d8e5753928b6d733118f75526ed3195ac8a 100644 --- a/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/factory/matmul_multicore_reuse_mcast_1d_program_factory.cpp +++ b/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/factory/matmul_multicore_reuse_mcast_1d_program_factory.cpp @@ -50,6 +50,193 @@ static uint32_t env_u32(const char* name, uint32_t fallback) { return value == nullptr || value[0] == '\0' ? fallback : static_cast(std::strtoul(value, nullptr, 10)); } +// Weight (in1) reader of the in0-multicast 1D matmul for an interleaved DRAM weight and one output block per +// core (decode). Read when a program is created, so set before the first matmul of a process or clear the +// program cache. Data movement only: the K blocking and accumulation order are unchanged. Unset = default. +// TT_MM1D_IN1_DUAL_NOC=1..4 (Blackhole) read the weight over both NoCs; all data-movement kernels of the program +// run in dynamic NoC mode (the activation multicast becomes unlinked). 1: interleaved +// weight alternates tiles, packed weight splits every contiguous read into two +// DRAM-aligned halves; 2: interleaved weight gives each NoC half of a block's tiles +// (packed: as 1); 3: odd workers read on the second NoC; 4: odd K blocks do; +// 5: per DRAM bank by direction: a worker reads a bank on the NoC whose x route from +// the bank's DRAM column to the worker is short (banks in the leftmost DRAM column on +// the in1 NoC for workers west of the middle DRAM column, on the other NoC east of it; +// the reverse for middle-column banks); 6: the opposite choice (control). +// TT_MM1D_IN1_NOC_COLS=GXxGY:c0c1...[,...] (with DUAL_NOC=3) per program grid: the worker in logical column x +// reads on the second NoC iff c_x is 1. When set, programs on grids not listed keep one +// NoC (dedicated mode); unset, every eligible program splits odd/even workers. +// TT_MM1D_IN0_MCAST_FLUSH=1 (with DUAL_NOC, experimental) the unlinked activation multicast only flushes its data +// before the VALID flag instead of waiting for acknowledgements. +// TT_MM1D_IN1_CB_SLOTS=S the in1 CB holds S K blocks (default 2) } only in programs that use the dual-NoC +// TT_MM1D_IN1_OUTSTANDING=D keep up to D K-block reads in flight } or packed reader; D uses NoC transaction +// ids and needs S > D. +// TT_MM1D_PACKED_IN1=Kt:Nt:pn:G[:b0.b1...][,...] every [Kt, Nt]-tile weight is stored in the worker-bank +// packed layout for per_core_N = pn (G = 0: worker w's columns live in bank b_w, by +// default w % 8, each bank holding Nt / pn / 8 workers; G > 0: groups of G K rows +// rotate over the banks; layout in the in1 reader kernel). Only programs with +// per_core_N == pn can read such a weight; any other program on it is rejected. +struct PackedIn1 { + uint32_t pn = 0; // 0: not packed + uint32_t group = 0; + std::vector bank_map; // G == 0 only; empty: w % banks +}; + +struct In1ReaderOpts { + uint32_t dual_noc = 0; // 0..6 (G4_IN1_SPLIT) + uint32_t cb_slots = 0; // 0: default depth + uint32_t outstanding = 1; // K blocks in flight + std::string noc_cols; // DUAL_NOC=3: per logical column '0'/'1' (TT_MM1D_IN1_NOC_COLS), empty: odd workers + PackedIn1 packed; + bool active() const { return dual_noc != 0 || cb_slots != 0 || outstanding > 1 || packed.pn != 0; } +}; + +// TT_MM1D_IN1_NOC_COLS entry for a gx x gy program grid ("" when not listed). +static std::string noc_cols_entry(const char* value, uint32_t gx, uint32_t gy) { + const std::string key = std::to_string(gx) + "x" + std::to_string(gy) + ":"; + const std::string all(value); + for (size_t pos = 0; pos < all.size();) { + const size_t end = std::min(all.find(',', pos), all.size()); + const std::string entry = all.substr(pos, end - pos); + TT_FATAL( + entry.find(':') != std::string::npos && + entry.find_first_not_of("01", entry.find(':') + 1) == std::string::npos, + "TT_MM1D_IN1_NOC_COLS='{}': expected GXxGY:bits[,...]", + value); + if (entry.compare(0, key.size(), key) == 0) { + TT_FATAL(entry.size() - key.size() == gx, "TT_MM1D_IN1_NOC_COLS entry '{}' needs {} columns", entry, gx); + return entry.substr(key.size()); + } + pos = end + 1; + } + return ""; +} + +// TT_MM1D_PACKED_IN1 entry for a [Kt, Nt]-tile weight (pn == 0 when the weight is not packed). +static PackedIn1 packed_in1_entry(uint32_t Kt, uint32_t Nt) { + const char* value = std::getenv("TT_MM1D_PACKED_IN1"); + for (const char* p = value; p != nullptr && *p != '\0';) { + uint32_t f[4] = {0, 0, 0, 0}; + uint32_t n = 0; + while (n < 4) { + char* end = nullptr; + f[n++] = static_cast(std::strtoul(p, &end, 10)); + TT_FATAL(end != p, "TT_MM1D_PACKED_IN1='{}': expected Kt:Nt:pn:G[:b0.b1...],...", value); + p = end; + if (*p != ':' || n == 4) { + break; + } + ++p; + } + TT_FATAL(n == 4 && f[2] > 0, "TT_MM1D_PACKED_IN1='{}': expected Kt:Nt:pn:G[:b0.b1...],...", value); + std::vector bank_map; + if (*p == ':') { + do { + ++p; + char* end = nullptr; + bank_map.push_back(static_cast(std::strtoul(p, &end, 10))); + TT_FATAL(end != p, "TT_MM1D_PACKED_IN1='{}': bad bank map", value); + p = end; + } while (*p == '.'); + } + if (f[0] == Kt && f[1] == Nt) { + return {f[2], f[3], std::move(bank_map)}; + } + if (*p == ',') { + ++p; + } + } + return {}; +} + +static In1ReaderOpts in1_reader_opts( + const tt_metal::IDevice* device, + CoreCoord grid, + const MeshTensor& in1_tensor, + uint32_t Kt, + uint32_t Nt, + uint32_t per_core_N, + uint32_t out_block_w, + uint32_t per_core_M, + uint32_t out_block_h, + uint32_t B, + bool in0_is_sharded, + bool in1_is_sharded, + bool transpose_b, + bool fuse_op) { + In1ReaderOpts o; + const PackedIn1 packed = packed_in1_entry(Kt, Nt); + const bool eligible = !in1_is_sharded && in1_tensor.memory_config().is_dram() && !fuse_op; + if (packed.pn != 0) { + const uint32_t banks = device->num_dram_channels(); + const uint32_t pn = packed.pn; + const uint32_t group = packed.group; + TT_FATAL( + eligible && !transpose_b && B == 1 && per_core_N == pn && out_block_w == pn && per_core_M == out_block_h && + Nt % pn == 0 && + (group == 0 ? (Nt / pn) % banks == 0 : Kt % (group * banks) == 0 && packed.bank_map.empty()), + "TT_MM1D_PACKED_IN1 weight [{}, {}] tiles (pn {}, G {}): needs an unsharded DRAM weight read by a " + "non-transposed, unfused B=1 program with per_core_N == out_block_w == pn and one output block per core " + "(got per_core_N {}, out_block_w {}), Nt % pn == 0, and (G == 0) {} workers a multiple of {} banks or " + "(G > 0, no bank map) Kt a multiple of G * banks", + Kt, + Nt, + pn, + group, + per_core_N, + out_block_w, + Nt / pn, + banks); + if (!packed.bank_map.empty()) { + std::vector count(banks, 0); + for (const uint32_t b : packed.bank_map) { + TT_FATAL(b < banks, "TT_MM1D_PACKED_IN1 bank map: bank {} out of range", b); + ++count[b]; + } + TT_FATAL( + packed.bank_map.size() == Nt / pn && + std::all_of(count.begin(), count.end(), [&](uint32_t c) { return c == Nt / pn / banks; }), + "TT_MM1D_PACKED_IN1 bank map for [{}, {}] must list {} workers, {} per bank", + Kt, + Nt, + Nt / pn, + Nt / pn / banks); + } + o.packed = packed; + } + if (!eligible) { + return o; + } + // The block-sharded activation sender is not adapted to dynamic NoC mode. + o.dual_noc = device->arch() == tt::ARCH::BLACKHOLE && !in0_is_sharded ? env_u32("TT_MM1D_IN1_DUAL_NOC", 0) : 0; + TT_FATAL(o.dual_noc <= 6, "TT_MM1D_IN1_DUAL_NOC must be 0..6"); + if (const char* cols = std::getenv("TT_MM1D_IN1_NOC_COLS"); o.dual_noc == 3 && cols != nullptr && *cols != '\0') { + o.noc_cols = noc_cols_entry(cols, grid.x, grid.y); + if (o.noc_cols.empty() || transpose_b || per_core_N != out_block_w) { + o.dual_noc = 0; // unlisted grid (or no per-worker column mapping): one NoC + o.noc_cols.clear(); + } + } + // Deeper in1 buffering only for the programs using the dual-NoC or packed reader (keeps e.g. an unlisted LM + // head program's L1 footprint unchanged). + if (o.dual_noc == 0 && o.packed.pn == 0) { + return o; + } + o.cb_slots = env_u32("TT_MM1D_IN1_CB_SLOTS", 0); + o.outstanding = std::max(env_u32("TT_MM1D_IN1_OUTSTANDING", 1), 1); + const uint32_t slots = o.cb_slots != 0 ? o.cb_slots : operations::matmul::utilities::MCAST_INPUT_BUFFERING_DEPTH; + TT_FATAL( + o.cb_slots == 0 || (o.cb_slots >= 2 && o.cb_slots <= 16), "TT_MM1D_IN1_CB_SLOTS must be in [2, 16]"); + TT_FATAL( + o.outstanding == 1 || (o.outstanding < slots && o.outstanding < 15), + "TT_MM1D_IN1_OUTSTANDING={} needs more CB slots ({})", + o.outstanding, + slots); + if (o.outstanding > 1 && B != 1) { + o.outstanding = 1; + } + return o; +} + uint32_t get_preferred_noc( const ttnn::CoreCoord src, const ttnn::CoreCoord dst, @@ -3208,6 +3395,25 @@ static ProgramDescriptor create_program_mcast_in0_descriptor( uint32_t in1_shard_height_in_tiles = in1_tensor.shard_spec()->shape[0] / in1_tile.get_height(); in1_CB_tiles = per_core_N * in1_shard_height_in_tiles; } + // Gemma decode weight reader (TT_MM1D_IN1_*, TT_MM1D_PACKED_IN1; see In1ReaderOpts). + const In1ReaderOpts in1_opts = in1_reader_opts( + device, + compute_with_storage_grid_size, + in1_tensor, + K, + N, + per_core_N, + out_block_w, + per_core_M, + out_block_h, + std::max(in0_B, in1_B), + in0_is_sharded, + in1_is_sharded, + transpose_b, + fuse_op); + if (in1_opts.cb_slots != 0) { + in1_CB_tiles = in1_block_tiles * in1_opts.cb_slots; + } uint32_t in1_CB_size = in1_CB_tiles * in1_aligned_tile_size; @@ -3582,6 +3788,67 @@ static ProgramDescriptor create_program_mcast_in0_descriptor( } mm_kernel_in1_sender_writer_defines["SKIP_MCAST"] = "1"; + if (in1_opts.active()) { + auto& defines = mm_kernel_in1_sender_writer_defines; + defines["G4_IN1"] = "1"; + defines["G4_IN1_NOCS"] = in1_opts.dual_noc != 0 ? "2" : "1"; + defines["G4_IN1_SPLIT"] = std::to_string(in1_opts.dual_noc); + defines["G4_IN1_D"] = std::to_string(in1_opts.outstanding); + defines["G4_IN1_BANKS"] = std::to_string(device->num_dram_channels()); + if (!in1_opts.noc_cols.empty()) { + // Worker w (w-th core with work, row-major on the program grid) sits in logical column w % gx. + std::string map; + for (uint32_t w = 0; w < num_cores_with_work; ++w) { + map += (w == 0 ? "" : ",") + std::string(1, in1_opts.noc_cols[w % compute_with_storage_grid_size.x]); + } + defines["G4_NOC_MAP"] = map; + } + if (in1_opts.dual_noc >= 5) { + // DRAM-adjacent logical worker per bank (Blackhole: right of the bank's endpoint on the in1 NoC); banks + // whose worker sits in the leftmost column are "left", the others "middle". Workers at logical x >= + // G4_EAST_X (the middle column's adjacent worker) are east of the middle DRAM column. + const auto anchors = device->get_optimal_dram_bank_to_logical_worker_assignment( + tt::tt_metal::detail::preferred_noc_for_dram_read(device->arch())); + uint32_t left_x = UINT32_MAX; + uint32_t mid_x = 0; + for (const auto& a : anchors) { + left_x = std::min(left_x, a.x); + mid_x = std::max(mid_x, a.x); + } + uint32_t mid_banks = 0; + for (uint32_t b = 0; b < anchors.size(); ++b) { + mid_banks |= (anchors[b].x != left_x ? 1u : 0u) << b; + } + const uint32_t left_banks = ((1u << anchors.size()) - 1) & ~mid_banks; + // in1 NoC (NOC_0 on Blackhole) flows +x: west workers get left banks on it, middle banks on the other NoC. + const bool direct = in1_opts.dual_noc == 5; + defines["G4_EAST_X"] = std::to_string(mid_x); + defines["G4_NOC1_BANKS_WEST"] = std::to_string(direct ? mid_banks : left_banks); + defines["G4_NOC1_BANKS_EAST"] = std::to_string(direct ? left_banks : mid_banks); + } + if (in1_opts.packed.pn != 0) { + defines["G4_PACKED_IN1"] = "1"; + defines["G4_PACKED_KT"] = std::to_string(K); + defines["G4_PACKED_PN"] = std::to_string(in1_opts.packed.pn); + defines["G4_PACKED_G"] = std::to_string(in1_opts.packed.group); + defines["G4_PACKED_BANKS"] = std::to_string(device->num_dram_channels()); + if (!in1_opts.packed.bank_map.empty()) { + std::string map; + for (const uint32_t b : in1_opts.packed.bank_map) { + map += (map.empty() ? "" : ",") + std::to_string(b); + } + defines["G4_PACKED_BANK_MAP"] = map; + } + } + } + // Both data-movement processors must use dynamic NoC mode when the in1 reader uses both NoCs. + const auto dm_noc_mode = + in1_opts.dual_noc != 0 ? tt_metal::NOC_MODE::DM_DYNAMIC_NOC : tt_metal::NOC_MODE::DM_DEDICATED_NOC; + // TT_MM1D_IN0_MCAST_FLUSH=1 (experimental, dual NoC only): the unlinked activation multicast waits until its data + // has left L1 instead of for its acknowledgements before the VALID flag multicast (same NoC, VC and command buffer). + if (in1_opts.dual_noc != 0 && env_u32("TT_MM1D_IN0_MCAST_FLUSH", 0) != 0) { + mm_kernel_in0_sender_writer_defines["G4_IN0_MCAST_FLUSH"] = "1"; + } // Intermediate CB read @@ -3636,7 +3903,8 @@ static ProgramDescriptor create_program_mcast_in0_descriptor( }; } in0_sender_kernel_desc.config = - DataMovementConfigDescriptor{.processor = tt_metal::DataMovementProcessor::RISCV_1, .noc = in0_noc}; + DataMovementConfigDescriptor{ + .processor = tt_metal::DataMovementProcessor::RISCV_1, .noc = in0_noc, .noc_mode = dm_noc_mode}; if (in0_is_sharded) { if (in0_mcast_cores_without_work_and_in_receiver_grid.num_cores() > 0) { @@ -3657,7 +3925,8 @@ static ProgramDescriptor create_program_mcast_in0_descriptor( {"cb_l1_array", tt::CBIndex::c_6}, }; in0_no_work_in_receiver_kernel_desc.config = - DataMovementConfigDescriptor{.processor = tt_metal::DataMovementProcessor::RISCV_1, .noc = in0_noc}; + DataMovementConfigDescriptor{ + .processor = tt_metal::DataMovementProcessor::RISCV_1, .noc = in0_noc, .noc_mode = dm_noc_mode}; } if (in0_mcast_cores_without_work_and_not_in_receiver_grid.num_cores() > 0) { has_in0_no_work_not_in_receiver_kernel = true; @@ -3677,7 +3946,8 @@ static ProgramDescriptor create_program_mcast_in0_descriptor( {"cb_l1_array", tt::CBIndex::c_6}, }; in0_no_work_not_in_receiver_kernel_desc.config = - DataMovementConfigDescriptor{.processor = tt_metal::DataMovementProcessor::RISCV_1, .noc = in0_noc}; + DataMovementConfigDescriptor{ + .processor = tt_metal::DataMovementProcessor::RISCV_1, .noc = in0_noc, .noc_mode = dm_noc_mode}; } } @@ -3692,7 +3962,8 @@ static ProgramDescriptor create_program_mcast_in0_descriptor( {"cb_in0", tt::CBIndex::c_0}, }; in0_receiver_kernel_desc.config = - DataMovementConfigDescriptor{.processor = tt_metal::DataMovementProcessor::RISCV_1, .noc = in0_noc}; + DataMovementConfigDescriptor{ + .processor = tt_metal::DataMovementProcessor::RISCV_1, .noc = in0_noc, .noc_mode = dm_noc_mode}; } in1_sender_writer_kernel_desc.kernel_source = @@ -3708,7 +3979,8 @@ static ProgramDescriptor create_program_mcast_in0_descriptor( {"cb_sparsity", tt::CBIndex::c_7}, }; in1_sender_writer_kernel_desc.config = - DataMovementConfigDescriptor{.processor = tt_metal::DataMovementProcessor::RISCV_0, .noc = in1_noc}; + DataMovementConfigDescriptor{ + .processor = tt_metal::DataMovementProcessor::RISCV_0, .noc = in1_noc, .noc_mode = dm_noc_mode}; // Compute kernel compile time args uint32_t in0_subblock_num_tiles = out_subblock_h * in0_block_w; @@ -4225,6 +4497,11 @@ static ProgramDescriptor create_program_mcast_in1_descriptor( bool fuse_op = false; + TT_FATAL( + packed_in1_entry(K, N).pn == 0, + "TT_MM1D_PACKED_IN1 lists this [{}, {}]-tile weight as packed; the in1-multicast program cannot read it", + K, + N); uint32_t num_blocks = K / in0_block_w; // Only enable packer l1 accumulation when there are num_blocks > 2, otherwise // unnecessary overhead for reconfigs are added. Last iteration of l1 accumulation diff --git a/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in0_sender_padding.cpp b/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in0_sender_padding.cpp new file mode 100644 index 0000000000000000000000000000000000000000..69833d05ddc7729bad0a20556c579f6ff007f4f2 --- /dev/null +++ b/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in0_sender_padding.cpp @@ -0,0 +1,453 @@ +// SPDX-FileCopyrightText: © 2025 Tenstorrent USA, Inc. +// +// SPDX-License-Identifier: Apache-2.0 + +#include + +#include "api/dataflow/dataflow_api.h" +#include "api/debug/assert.h" +#include "hostdevcommon/common_values.hpp" +#include "ttnn/operations/ccl/kernel_common/worker_sync_utils.hpp" +#include "ttnn/operations/kernel_helper_functions/pad_tile.hpp" +#include "ckernel.h" +#include "ckernel_defs.h" +#include "api/dataflow/noc.h" +#include "api/dataflow/dataflow_buffer.h" +#include "api/dataflow/noc_semaphore.h" +#include "api/tensor/noc_traits.h" +#include "api/dataflow/endpoints.h" +#include "api/core_local_mem.h" +void kernel_main() { + uint32_t rt_args_idx = 0; + // in0 tensor args + const uint32_t in0_tensor_addr = get_arg_val(rt_args_idx++); + uint32_t in0_tensor_start_tile_id = get_arg_val(rt_args_idx++); + // in0 mcast args + const uint32_t in0_mcast_dest_noc_start_x = get_arg_val(rt_args_idx++); + const uint32_t in0_mcast_dest_noc_start_y = get_arg_val(rt_args_idx++); + const uint32_t in0_mcast_dest_noc_end_x = get_arg_val(rt_args_idx++); + const uint32_t in0_mcast_dest_noc_end_y = get_arg_val(rt_args_idx++); + + // padding args + const uint32_t last_block_h = get_arg_val(rt_args_idx++); + // sparsity args + const uint32_t sparsity_addr = get_arg_val(rt_args_idx++); + + // COMPILE TIME ARGS + // in0 tensor args + constexpr uint32_t in0_tensor_stride_w = get_compile_time_arg_val(0); + constexpr uint32_t in0_tensor_stride_h = get_compile_time_arg_val(1); + constexpr uint32_t in0_tensor_next_inner_dim_block_stride = get_compile_time_arg_val(2); + constexpr uint32_t in0_tensor_next_h_dim_block_stride = get_compile_time_arg_val(3); + // in0 block args + constexpr uint32_t in0_block_w = get_compile_time_arg_val(4); + constexpr uint32_t in0_block_h = get_compile_time_arg_val(5); + constexpr uint32_t in0_block_num_tiles = get_compile_time_arg_val(6); + constexpr uint32_t in0_last_ktile_w = get_compile_time_arg_val(7); + constexpr uint32_t in0_last_ktile_h = get_compile_time_arg_val(8); + + constexpr bool extract_shard_sub_blocks = (bool)get_compile_time_arg_val(9); + constexpr uint32_t shard_width_in_tiles = get_compile_time_arg_val(10); + constexpr uint32_t shard_height_in_tiles = get_compile_time_arg_val(11); + // in0/in1 common args + constexpr uint32_t num_blocks_inner_dim = get_compile_time_arg_val(12); + constexpr uint32_t num_blocks_w_dim = get_compile_time_arg_val(13); + constexpr uint32_t num_blocks_h_dim = get_compile_time_arg_val(14); + // in0 mcast args + constexpr uint32_t in0_mcast_num_dests = get_compile_time_arg_val(17); + constexpr uint32_t in0_mcast_num_cores = get_compile_time_arg_val(18); + // batch args + constexpr uint32_t MtKt = get_compile_time_arg_val(19); // if 0 + constexpr uint32_t in0_B = get_compile_time_arg_val(20); + constexpr uint32_t in1_B = get_compile_time_arg_val(21); + constexpr uint32_t in0_reuse_in_CB = get_compile_time_arg_val(22); + + // sparsity args + + constexpr uint32_t batchB = get_compile_time_arg_val(23); + constexpr uint32_t sparsity_pagesize = get_compile_time_arg_val(24); + // Boolean that is set when input A is sparse. If set, both input A and B are assumed to be sparse. + // Based on the sparsity tensor, the corresponding batch in input A and B are skipped. + constexpr bool bcast_A = (bool)get_compile_time_arg_val(25); + // This boolean is set when the number of batches is only known at runtime, typically based on a sparsity tensor. + constexpr bool get_batch_from_reader = (bool)get_compile_time_arg_val(26); + + constexpr bool fuse_op = (bool)get_compile_time_arg_val(27); + + constexpr auto in0_args = TensorAccessorArgs<28>(); + + constexpr auto sparsity_args = TensorAccessorArgs(); + + // Number of valid (non-zero sparsity) batches the receiver and compute kernels are configured to + // process. When nnz is supplied (get_batch_from_reader == false), those kernels loop exactly + // num_batch_compute times, while this sender only multicasts once per non-zero sparsity entry, i.e. + // count_nonzero(sparsity) times. The op silently requires count_nonzero(sparsity) == num_batch_compute; + // if they disagree, the receivers wait on multicasts that never come (or the sender waits on receivers + // that already finished) and the device deadlocks. count_nonzero(sparsity) is data-dependent and only + // known here at runtime, so we validate it on-device by counting the multicasts we actually issue and + // asserting the contract holds -- surfacing a loud assert (under watcher) instead of a silent hang. + // See https://github.com/tenstorrent/tt-metal/issues/45943. + [[maybe_unused]] constexpr uint32_t num_batch_compute = + get_compile_time_arg_val(sparsity_args.next_compile_time_args_offset()); + + // 0 is used to specify "INVALID" state, i.e. when the multicasted data has not been received by the receiver. + // 0x1 is used to specify "VALID" state, i.e. when the batch is valid. + // 0x2 is used to specify "IGNORE_BATCH" state, i.e. when the batch is not valid. + constexpr uint32_t IGNORE_BATCH = 0x2; + + // When sparsity is disabled, we just loop once + constexpr uint32_t batchB_lim = batchB == 0 ? 1u : batchB; + + MatmulOpReceiver fused_op_receiver; + if constexpr (fuse_op) { + fused_op_receiver = MatmulOpReceiver( + true, /* wait_for_op_signal */ + rt_args_idx, + num_blocks_inner_dim, + in0_block_w /* tiles_per_block (in the same dimension as tensor slice) */ + ); + } + + constexpr uint32_t dfb_id_in0 = get_named_compile_time_arg_val("cb_in0"); + constexpr uint32_t in0_single_tile_size_bytes = get_tile_size(dfb_id_in0); + // Tiles whose size is not a multiple of the DRAM alignment are padded to it in DRAM, and the + // interleaved in0 CB pages are sized to match (see the program factory). The NOC reads the + // unpadded tile of data into each padded slot, and tiles are laid out / multicast at the padded + // stride. No-op when already aligned. The sharded path keeps the natural (unpadded) stride. + constexpr uint32_t in0_aligned_tile_size_bytes = + (in0_single_tile_size_bytes + (DRAM_ALIGNMENT - 1)) & ~(DRAM_ALIGNMENT - 1); +#ifdef IN0_SHARDED + constexpr uint32_t in0_block_size_bytes = in0_block_num_tiles * in0_single_tile_size_bytes; +#else + constexpr uint32_t in0_block_size_bytes = in0_block_num_tiles * in0_aligned_tile_size_bytes; +#endif + + Noc noc; + DataflowBuffer dfb_in0(dfb_id_in0); + Semaphore<> sender_sem(get_compile_time_arg_val(15)); + Semaphore<> receiver_sem(get_compile_time_arg_val(16)); + +#ifdef IN0_SHARDED + // In case we need to send multiple blocks per shard, in0 sharded cb is cb2 and we extract the sub-blocks to cb0 + constexpr uint32_t shard_read_stride = shard_width_in_tiles * in0_single_tile_size_bytes; + constexpr uint32_t shard_read_width = in0_single_tile_size_bytes * in0_block_w; + constexpr uint32_t shard_num_tiles = shard_width_in_tiles * shard_height_in_tiles; + constexpr uint32_t in0_tensor_next_h_dim_block_stride_bytes = + in0_tensor_next_h_dim_block_stride * in0_single_tile_size_bytes; + + uint32_t noc_shard_read_start_addr = 0; + if constexpr (extract_shard_sub_blocks) { + constexpr uint32_t dfb_id_in2 = + get_named_compile_time_arg_val("cb_in0_sharded"); // in0 sharded cb if extract_shard_sub_blocks + DataflowBuffer dfb_in2(dfb_id_in2); + noc_shard_read_start_addr = dfb_in2.get_read_ptr(); + } + +#else + const auto s0 = TensorAccessor(in0_args, in0_tensor_addr); +#endif // IN0_SHARDED + + // sparsity accessor + constexpr uint32_t dfb_id_sparsity = get_named_compile_time_arg_val("cb_sparsity"); + DataflowBuffer dfb_sparsity(dfb_id_sparsity); + const auto s_sparsity = TensorAccessor(sparsity_args, sparsity_addr); + +#ifndef SKIP_MCAST + // Set ur local VALID value, to be mcasted to destinations flag address after the data has been mcasted + receiver_sem.set(VALID); + // local address that will be atomically incremented by mcast receivers, to know when all receivers are ready + // to receive the mcast + +#ifdef IN0_SHARDED + uint32_t in0_start_address = dfb_in0.get_write_ptr(); +#endif // IN0_SHARDED +#endif // SKIP_MCAST + + uint32_t l1_write_addr_sparsity = 0; + if constexpr (batchB > 0) { + dfb_sparsity.reserve_back(1); + l1_write_addr_sparsity = dfb_sparsity.get_write_ptr(); + } + + // Counts the in0 multicasts actually issued (one per non-zero sparsity entry). Used to validate + // count_nonzero(sparsity) == num_batch_compute when nnz is supplied (see num_batch_compute above). + [[maybe_unused]] uint32_t num_valid_batches = 0; + + for (uint32_t b = 0; b < in0_B; ++b) { + if constexpr (batchB > 0) { + noc.async_read(s_sparsity, dfb_sparsity, sparsity_pagesize, {.page_id = b}, {.offset_bytes = 0}); + noc.async_read_barrier(); + } + + for (uint32_t bB = 0; bB < batchB_lim; ++bB) { + if constexpr (batchB > 0) { + volatile auto is_batch_valid = + ((reinterpret_cast(l1_write_addr_sparsity))[bB]) != 0; + + if constexpr (get_batch_from_reader) { +#ifndef SKIP_MCAST + // First broadcast this to other cores + sender_sem.wait(in0_mcast_num_dests); + sender_sem.set(0); + receiver_sem.set(is_batch_valid ? VALID : IGNORE_BATCH); + receiver_sem.set_multicast( + noc, + in0_mcast_dest_noc_start_x, + in0_mcast_dest_noc_start_y, + in0_mcast_dest_noc_end_x, + in0_mcast_dest_noc_end_y, + in0_mcast_num_cores); + noc.async_writes_flushed(); + // Reset the semaphore value to VALID + receiver_sem.set(VALID); +#endif // SKIP_MCAST + + // We need to pass the value to compute cores regardless of the value of is_batch_valid + ckernel::mailbox_write(ckernel::ThreadId::UnpackThreadId, is_batch_valid); + ckernel::mailbox_write(ckernel::ThreadId::MathThreadId, is_batch_valid); + ckernel::mailbox_write(ckernel::ThreadId::PackThreadId, is_batch_valid); + } + + if (!is_batch_valid) { + if constexpr (!bcast_A) { + in0_tensor_start_tile_id += MtKt; + } + continue; + } + + // This is a valid (non-zero) batch that we are about to multicast. When nnz was supplied, + // catch count_nonzero(sparsity) > num_batch_compute here, before the sender blocks below + // waiting on receivers that have already finished their num_batch_compute iterations. + if constexpr (!get_batch_from_reader) { + ++num_valid_batches; + ASSERT(num_valid_batches <= num_batch_compute); + } + } + +#ifdef IN0_SHARDED + uint32_t in0_tensor_current_h_dim_block_start_addr = noc_shard_read_start_addr; +#endif // IN0_SHARDED + uint32_t in0_tensor_current_h_dim_block_tile_id = in0_tensor_start_tile_id; + for (uint32_t bh = 0; bh < num_blocks_h_dim; ++bh) { + for (uint32_t bw = 0; bw < num_blocks_w_dim; ++bw) { +#ifdef IN0_SHARDED + uint32_t in0_tensor_current_inner_dim_block_start_addr = in0_tensor_current_h_dim_block_start_addr; +#endif // IN0_SHARDED + uint32_t in0_tensor_current_inner_dim_block_start_tile_id = in0_tensor_current_h_dim_block_tile_id; + for (uint32_t block = 0; block < num_blocks_inner_dim; ++block) { + if constexpr (fuse_op) { + fused_op_receiver.update_current_block_start_tile_id( + block, in0_tensor_current_inner_dim_block_start_tile_id, in0_tensor_start_tile_id); + } + + // Operand 0 + // Common for sharded and interleaved paths + dfb_in0.reserve_back(in0_block_num_tiles); +#ifndef IN0_SHARDED + + uint32_t in0_write_offset = 0; + +#ifndef SKIP_MCAST + uint32_t in0_start_address = + dfb_in0.get_write_ptr(); // copy start address of block, to be used for mcasting +#endif // SKIP_MCAST + + // Copy in0 block into CB, as the default kernel + uint32_t in0_tensor_row_start_tile_id = in0_tensor_current_inner_dim_block_start_tile_id; + for (uint32_t h = 0; h < in0_block_h; ++h) { + uint32_t in0_tensor_tile_id = in0_tensor_row_start_tile_id; + for (uint32_t w = 0; w < in0_block_w; ++w) { + if (bh < num_blocks_h_dim - 1 || h < last_block_h) { + noc.async_read( + s0, + dfb_in0, + in0_single_tile_size_bytes, + {.page_id = in0_tensor_tile_id}, + {.offset_bytes = in0_write_offset}); + } + + // Zero out padded regions for the very last tile + if constexpr (in0_last_ktile_w > 0) { + if ((block == num_blocks_inner_dim - 1) && (w == in0_block_w - 1)) { + noc.async_read_barrier(); + constexpr DataFormat in0_data_format = get_dataformat(dfb_id_in0); + pad_last_ktile( + dfb_in0.get_write_ptr() + in0_write_offset); + } + } + if constexpr (in0_last_ktile_h > 0) { + if ((block == num_blocks_inner_dim - 1) && (w == in0_block_w - 1)) { + noc.async_read_barrier(); + constexpr DataFormat in0_data_format = get_dataformat(dfb_id_in0); + pad_last_transposed_ktile( + dfb_in0.get_write_ptr() + in0_write_offset); + } + } + + in0_write_offset += in0_aligned_tile_size_bytes; + in0_tensor_tile_id += in0_tensor_stride_w; + } + in0_tensor_row_start_tile_id += in0_tensor_stride_h; + } + in0_tensor_current_inner_dim_block_start_tile_id += in0_tensor_next_inner_dim_block_stride; + + // Barrier! make sure the reads are done + noc.async_read_barrier(); +#else + if constexpr (extract_shard_sub_blocks) { + uint32_t l1_write_addr_in0 = dfb_in0.get_write_ptr(); + +#ifndef SKIP_MCAST + in0_start_address = + l1_write_addr_in0; // copy start address of block, to be used for mcasting +#endif // SKIP_MCAST + + UnicastEndpoint self_ep; + uint32_t noc_shard_read_l1_addr = in0_tensor_current_inner_dim_block_start_addr; + + for (uint32_t i = 0; i < in0_block_h; i++) { + noc.async_read( + self_ep, + CoreLocalMem(l1_write_addr_in0), + shard_read_width, + {.noc_x = my_x[0], .noc_y = my_y[0], .addr = noc_shard_read_l1_addr}, + {}); + + l1_write_addr_in0 += shard_read_width; + noc_shard_read_l1_addr += shard_read_stride; + } + + in0_tensor_current_inner_dim_block_start_addr += shard_read_width; + noc.async_read_barrier(); + } + + { + constexpr DataFormat in0_data_format = get_dataformat(dfb_id_in0); + uint32_t in0_pad_base_addr = dfb_in0.get_write_ptr(); + if constexpr (in0_last_ktile_w > 0) { + if ((block == num_blocks_inner_dim - 1)) { + for (uint32_t h = 0; h < in0_block_h; ++h) { + auto ptr = in0_pad_base_addr + + (h * in0_block_w + in0_block_w - 1) * in0_single_tile_size_bytes; + pad_last_ktile(ptr); + } + } + } + if constexpr (in0_last_ktile_h > 0) { + if ((block == num_blocks_inner_dim - 1)) { + for (uint32_t w = 0; w < in0_block_w; ++w) { + auto ptr = in0_pad_base_addr + + ((in0_block_h - 1) * in0_block_w + w) * in0_single_tile_size_bytes; + pad_last_transposed_ktile(ptr); + } + } + } + } +#endif // IN0_SHARDED + +#ifndef SKIP_MCAST + // wait until all in0 mcast destinations have atomically incremented the in0 semaphore_addr + // (i.e. its value should be in0_mcast_num_dests), then reset the semaphore_addr value back to + // zero for the next block + sender_sem.wait(in0_mcast_num_dests); + sender_sem.set(0); + + // Now we have the block in the CB address, we can mcast to dests! + MulticastEndpoint mcast_dst; + // num_dests must not include source, since we are NOT really doing a local copy! + noc.async_write_multicast( + CoreLocalMem(in0_start_address), + mcast_dst, + in0_block_size_bytes, + in0_mcast_num_cores, + {}, + {.noc_x_start = in0_mcast_dest_noc_start_x, + .noc_y_start = in0_mcast_dest_noc_start_y, + .noc_x_end = in0_mcast_dest_noc_end_x, + .noc_y_end = in0_mcast_dest_noc_end_y, + .addr = in0_start_address}, + noc_mode == DM_DEDICATED_NOC); + + if constexpr (noc_mode == DM_DYNAMIC_NOC) { + // Dynamic NoC mode (TT_MM1D_IN1_DUAL_NOC: the in1 reader also issues unicast reads on + // this NoC): a linked multicast forbids unicasts on any command buffer of the NoC, so + // the data multicast is unlinked and completes before its VALID flag is published. + // G4_IN0_MCAST_FLUSH (TT_MM1D_IN0_MCAST_FLUSH=1, experimental) only waits until the data + // has left L1 and relies on same-NoC/VC/command-buffer ordering of the two multicasts. +#ifdef G4_IN0_MCAST_FLUSH + noc.async_writes_flushed(); +#else + noc.async_write_barrier(); +#endif + } else { + // Note: no need for write barrier, since these two multicasts are done on the same noc + // id, same vc, same cmd_buf Also, this only works because we are setting VCs statically + // (using NOC_CMD_STATIC_VC). +#ifdef ARCH_BLACKHOLE + // On Blackhole the flush is needed because NoC latency is higher than L1 <-> RISCV + // latency which means data could be changed before write is issued. + noc.async_writes_flushed(); +#endif // ARCH_BLACKHOLE + } + + // We should also multicast the flag to destinations + // num_dests must not include source, since we are NOT really doing a local copy! + receiver_sem.set_multicast( + noc, + in0_mcast_dest_noc_start_x, + in0_mcast_dest_noc_start_y, + in0_mcast_dest_noc_end_x, + in0_mcast_dest_noc_end_y, + in0_mcast_num_cores); +#endif // SKIP_MCAST + + // Common for sharded and interleaved paths + dfb_in0.push_back(in0_block_num_tiles); + } + } +#ifdef IN0_SHARDED + in0_tensor_current_h_dim_block_start_addr += in0_tensor_next_h_dim_block_stride_bytes; +#endif // IN0_SHARDED + in0_tensor_current_h_dim_block_tile_id += in0_tensor_next_h_dim_block_stride; + } + + if constexpr (!bcast_A) { + in0_tensor_start_tile_id += MtKt; + } + } + + if constexpr (bcast_A) { + in0_tensor_start_tile_id += MtKt; + } + + // this is an optimization for the case when in0 is [1, 1, M, K] and in1 is [1, H, K, N], i.e. when in0_B == + // 1 and in1_B > 1 in this case we originally had to replicate the in0 block for each batch, but with this + // optimization we just read the tensor slice once into each core's L1 and keep it there for all weight + // batches since the needed in0 data is already in L1 after batch 0, we can just move read pointer for this + // CB so compute kernel thinks it has new data + if (in0_reuse_in_CB) { + for (uint32_t fake_batch = 0; fake_batch < in1_B - in0_B; ++fake_batch) { + for (uint32_t blk = 0; blk < num_blocks_inner_dim; ++blk) { + dfb_in0.reserve_back(in0_block_num_tiles); + dfb_in0.push_back(in0_block_num_tiles); + } + } + } + } + noc.async_write_barrier(); + + // When nnz was supplied, the receiver and compute kernels loop exactly num_batch_compute times. + // If we issued fewer multicasts than that (count_nonzero(sparsity) < num_batch_compute), those + // kernels are now waiting on multicasts that will never arrive and the device would deadlock. + // Fail loudly instead. See https://github.com/tenstorrent/tt-metal/issues/45943. + if constexpr (!get_batch_from_reader && batchB > 0) { + ASSERT(num_valid_batches == num_batch_compute); + } + + // For completeness, we empty the sparsity CB if it was reserved earlier + if constexpr (batchB > 0) { + dfb_sparsity.push_back(1); + dfb_sparsity.wait_front(1); + dfb_sparsity.pop_front(1); + } +} diff --git a/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in1_sender_writer_padding.cpp b/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in1_sender_writer_padding.cpp new file mode 100644 index 0000000000000000000000000000000000000000..b9034fa2581e9471f608054a1862a3f7725487ea --- /dev/null +++ b/runtime/ttnn-overlay/cpp/ttnn/operations/matmul/device/kernels/dataflow/reader_bmm_tile_layout_in1_sender_writer_padding.cpp @@ -0,0 +1,922 @@ +// SPDX-FileCopyrightText: © 2023 Tenstorrent USA, Inc. +// +// SPDX-License-Identifier: Apache-2.0 + +#include + +#include "api/dataflow/dataflow_api.h" +#include "hostdevcommon/common_values.hpp" +#include "ttnn/operations/ccl/kernel_common/worker_sync_utils.hpp" +#include "api/dataflow/noc.h" +#include "api/dataflow/dataflow_buffer.h" +#include "api/dataflow/noc_semaphore.h" +#include "api/remote_circular_buffer.h" +#include "api/tensor/noc_traits.h" +#include "api/dataflow/endpoints.h" +#include "api/core_local_mem.h" + +// Gemma decode weight reader (G4_IN1, set by the 1D in0-multicast factory from TT_MM1D_IN1_* and +// TT_MM1D_PACKED_IN1). Interleaved DRAM weight, no in1 multicast. Per K block: +// - G4_IN1_NOCS == 2: the in1 reads use both NoCs (dynamic NoC mode). G4_IN1_SPLIT 1: interleaved weight +// alternates tiles, packed weight splits every contiguous read into two DRAM-aligned byte halves; 2: like 1, +// but the interleaved weight gives the first half of each block's tiles to one NoC and the rest to the other; +// 3: odd workers read everything on the second NoC; 4: odd K blocks are read on the second NoC. +// - G4_IN1_D > 1: up to G4_IN1_D blocks in flight, one NoC transaction id per block; blocks are published +// to compute in K order as their reads complete (the CB holds more than G4_IN1_D blocks). +// - G4_PACKED_IN1: the weight is stored in a worker-bank packed layout for per_core_N = G4_PACKED_PN +// (W = Nt / pn workers, worker w owns logical columns [w*pn, (w+1)*pn), B = G4_PACKED_BANKS). Logical +// tile (k, n), w = n / pn, c = n % pn, lives in interleaved page bank + B * slot with +// G4_PACKED_G == 0: bank = map[w], slot = rank(w) * Kt * pn + k * pn + c +// G4_PACKED_G == G: bank = (w + k / G) % B, slot = w * (Kt / B) * pn + (k / G / B) * G * pn +// + (k % G) * pn + c +// map = G4_PACKED_BANK_MAP (one bank per worker, W / B workers per bank) or w % B by default, and rank(w) +// = #{w' < w : map[w'] == map[w]}. A worker's rows of one bank are contiguous and each K block is read with +// one read per bank run. +// Tiles land in the CB in the same order as on the default path, so compute is unchanged. +#if defined(G4_IN1) && (!defined(SKIP_MCAST) || defined(IN1_SHARDED) || defined(ENABLE_GLOBAL_CB) || \ + defined(IN1_DRAM_WIDTH_SHARDED) || defined(IN1_DRAM_HEIGHT_SHARDED)) +#error "G4_IN1 requires the interleaved, non-multicast in1 reader" +#endif +void kernel_main() { + // READER + uint32_t rt_args_idx = 0; + // in1 tensor args + const uint32_t in1_tensor_addr = get_arg_val(rt_args_idx++); + uint32_t in1_tensor_start_tile_id = get_arg_val(rt_args_idx++); + // in1 mcast args + const uint32_t in1_mcast_dest_noc_start_x = get_arg_val(rt_args_idx++); + const uint32_t in1_mcast_dest_noc_start_y = get_arg_val(rt_args_idx++); + const uint32_t in1_mcast_dest_noc_end_x = get_arg_val(rt_args_idx++); + const uint32_t in1_mcast_dest_noc_end_y = get_arg_val(rt_args_idx++); + + // sparsity args + const uint32_t sparsity_addr = get_arg_val(rt_args_idx++); + + // WRITER + // out tensor args + const uint32_t out_tensor_addr = get_arg_val(rt_args_idx++); + uint32_t out_tensor_start_tile_id = get_arg_val(rt_args_idx++); + + // padding args (READER) + const uint32_t last_block_w = get_arg_val(rt_args_idx++); + // padding args (WRITER) + const uint32_t out_num_nonzero_subblocks_h = get_arg_val(rt_args_idx++); + const uint32_t out_last_subblock_h = get_arg_val(rt_args_idx++); + const uint32_t padded_block_tiles_h_skip = get_arg_val(rt_args_idx++); + const uint32_t out_num_nonzero_subblocks_w = get_arg_val(rt_args_idx++); + const uint32_t out_last_num_nonzero_subblocks_w = get_arg_val(rt_args_idx++); + const uint32_t out_last_subblock_w = get_arg_val(rt_args_idx++); + const uint32_t padded_subblock_tiles_addr_skip = get_arg_val(rt_args_idx++); + const uint32_t padded_block_tiles_w_skip = get_arg_val(rt_args_idx++); + + // COMPILE TIME ARGS + // READER + // in1 tensor args + constexpr uint32_t in1_tensor_stride_w = get_compile_time_arg_val(0); + constexpr uint32_t in1_tensor_stride_h = get_compile_time_arg_val(1); + constexpr uint32_t in1_tensor_next_block_stride = get_compile_time_arg_val(2); + constexpr uint32_t in1_tensor_next_w_dim_block_stride = get_compile_time_arg_val(3); + // in1 block args + constexpr uint32_t in1_block_w = get_compile_time_arg_val(4); + constexpr uint32_t in1_block_h = get_compile_time_arg_val(5); + constexpr uint32_t in1_block_num_tiles = get_compile_time_arg_val(6); + // in0/in1 common args + constexpr uint32_t num_blocks_inner_dim = get_compile_time_arg_val(7); + constexpr uint32_t num_blocks_w_dim = get_compile_time_arg_val(8); + constexpr uint32_t num_blocks_h_dim = get_compile_time_arg_val(9); + + // in1 mcast args + constexpr uint32_t in1_mcast_num_dests = get_compile_time_arg_val(12); + constexpr uint32_t in1_mcast_num_cores = get_compile_time_arg_val(13); + // batch args + constexpr uint32_t KtNt = get_compile_time_arg_val(14); + constexpr uint32_t batch = get_compile_time_arg_val(15); + constexpr uint32_t bcast_B = get_compile_time_arg_val(16); + // sparsity args + constexpr uint32_t batchB = get_compile_time_arg_val(17); + constexpr uint32_t sparsity_pagesize = get_compile_time_arg_val(18); + + // WRITER + // out tensor args + constexpr uint32_t out_tensor_stride_w = get_compile_time_arg_val(19); + constexpr uint32_t out_tensor_stride_h = get_compile_time_arg_val(20); + constexpr uint32_t out_tensor_next_subblock_stride_w = get_compile_time_arg_val(21); + constexpr uint32_t out_tensor_next_subblock_stride_h = get_compile_time_arg_val(22); + constexpr uint32_t out_tensor_next_w_dim_block_stride = get_compile_time_arg_val(23); + constexpr uint32_t out_tensor_next_h_dim_block_stride = get_compile_time_arg_val(24); + // out subblock args + constexpr uint32_t out_subblock_w = get_compile_time_arg_val(25); + constexpr uint32_t out_subblock_h = get_compile_time_arg_val(26); + constexpr uint32_t out_subblock_tile_count = get_compile_time_arg_val(27); + // batch args + constexpr uint32_t MtNt = get_compile_time_arg_val(28); // if 0 + // Don't need batch; same as batch from READER args + + // When sparsity is disabled, we just loop once + constexpr uint32_t batchB_lim = batchB == 0 ? 1u : batchB; + +#ifdef FUSE_BIAS + // in3 mcast args + const uint32_t in3_tensor_addr = get_arg_val(rt_args_idx++); + const uint32_t in3_tensor_start_tile_id = get_arg_val(rt_args_idx++); + + constexpr uint32_t in3_tensor_stride_w = get_compile_time_arg_val(29); + + constexpr uint32_t dfb_id_in3 = get_named_compile_time_arg_val("cb_bias"); + // Use the CB page size (padded to the DRAM alignment by the factory) for DRAM reads + // and L1 write strides, NOT the raw tile size. On Blackhole, the DRAM read alignment + // is 64B, so a sub-64B tile (e.g. 32B for a (1,16) bf16 bias tile) cannot be read + // directly from DRAM, and 32B-strided L1 writes land at non-64B-aligned addresses + // that disagree with the 64B-aligned DRAM source. The factory pads the CB page to + // 64B; the unpacker still reads the actual 32B tile from the padded page via tile + // dims. For tiles already >= dram_alignment (e.g. 32x32 bf16 = 2048B), the page size + // equals the tile size, so this is a no-op. Mirrors how in0/in1 readers walk at the + // aligned stride. + const uint32_t bias_single_tile_size_bytes = get_local_cb_interface(dfb_id_in3).fifo_page_size; + constexpr const uint32_t in3_tile_hw = get_tile_hw(dfb_id_in3); + +#ifndef BIAS_SHARDED + uint32_t l1_write_addr_in3; + // Bias accessor will be defined later after TensorAccessor args +#endif // BIAS_SHARDED +#else + rt_args_idx += 2; // Skip over placeholders +#endif // FUSE_BIAS +#ifndef OUT_SHARDED + const uint32_t last_num_blocks_w_dim = get_arg_val(rt_args_idx++); +#endif // OUT_SHARDED + + constexpr bool fuse_op_all_gather = (bool)get_compile_time_arg_val(30); + constexpr bool fuse_op_reduce_scatter = (bool)get_compile_time_arg_val(31); + + MatmulOpReceiver fused_op_receiver; + OpSignaler op_signaler; + if constexpr (fuse_op_all_gather) { + fused_op_receiver = MatmulOpReceiver( + false, /* wait_for_op_signal */ + rt_args_idx, + num_blocks_inner_dim, + in1_block_h /* tiles_per_block (in the same dimension */ + ); + } else if constexpr (fuse_op_reduce_scatter) { + op_signaler = OpSignaler(rt_args_idx); + } + + constexpr auto in1_args = TensorAccessorArgs<32>(); + constexpr auto sparsity_args = TensorAccessorArgs(); + constexpr auto out_args = TensorAccessorArgs(); +#ifdef FUSE_BIAS + constexpr auto bias_args = TensorAccessorArgs(); + constexpr auto after_bias_offset = bias_args.next_compile_time_args_offset(); +#else + constexpr auto after_bias_offset = out_args.next_compile_time_args_offset(); +#endif // FUSE_BIAS + +// RT and COMPILE TIME ARGS for DRAM sharded weights +#ifdef IN1_DRAM_WIDTH_SHARDED + const uint32_t vc = get_arg_val(rt_args_idx++); + const uint32_t num_dram_shards_to_read = get_arg_val(rt_args_idx++); + const uint32_t dram_tensor_start_offset = get_arg_val(rt_args_idx++); + tt_l1_ptr uint32_t* in1_block_w_dram_stride_bytes = (tt_l1_ptr uint32_t*)get_arg_addr(rt_args_idx++); + tt_l1_ptr uint32_t* current_dram_bank_id = (tt_l1_ptr uint32_t*)get_arg_addr(rt_args_idx++); + + constexpr uint32_t in1_dram_block_num_tiles = get_compile_time_arg_val(after_bias_offset); + constexpr uint32_t in1_block_w_dram_bytes = get_compile_time_arg_val(after_bias_offset + 1); +#endif // IN1_DRAM_WIDTH_SHARDED + +#ifdef IN1_DRAM_HEIGHT_SHARDED + constexpr uint32_t in1_KtNt_per_batch = get_compile_time_arg_val(after_bias_offset); // K*N tiles per batch + constexpr uint32_t in1_batches_per_bank = get_compile_time_arg_val(after_bias_offset + 1); // batches per DRAM bank +#endif // IN1_DRAM_HEIGHT_SHARDED + +#ifdef FUSE_BIAS +#ifndef BIAS_SHARDED + const auto s3 = TensorAccessor(bias_args, in3_tensor_addr); +#endif // BIAS_SHARDED +#endif // FUSE_BIAS + + constexpr uint32_t dfb_id_in1 = get_named_compile_time_arg_val("cb_in1"); + constexpr uint32_t in1_single_tile_size_bytes = get_tile_size(dfb_id_in1); + constexpr const uint32_t in1_tile_hw = get_tile_hw(dfb_id_in1); + // Tiles whose size is not a multiple of the DRAM alignment are padded to it in DRAM, and the + // interleaved in1 CB pages are sized to match (see the program factory). On the plain interleaved + // path the NOC reads the unpadded tile of data into each padded slot and tiles are laid out / + // multicast at the padded stride. No-op when already aligned. The sharded / DRAM-sharded paths + // keep their natural (unpadded) stride. + constexpr uint32_t in1_aligned_tile_size_bytes = + (in1_single_tile_size_bytes + (DRAM_ALIGNMENT - 1)) & ~(DRAM_ALIGNMENT - 1); +#if !defined(IN1_SHARDED) && !defined(IN1_DRAM_WIDTH_SHARDED) && !defined(IN1_DRAM_HEIGHT_SHARDED) && \ + !defined(ENABLE_GLOBAL_CB) + constexpr uint32_t in1_block_size_bytes = in1_block_num_tiles * in1_aligned_tile_size_bytes; +#else + constexpr uint32_t in1_block_size_bytes = in1_block_num_tiles * in1_single_tile_size_bytes; +#endif + + constexpr uint32_t dfb_id_out0 = get_named_compile_time_arg_val("cb_out"); + constexpr uint32_t output_single_tile_size_bytes = get_tile_size(dfb_id_out0); + constexpr const uint32_t output_tile_hw = get_tile_hw(dfb_id_out0); + + Noc noc; + DataflowBuffer dfb_in1(dfb_id_in1); + DataflowBuffer dfb_out(dfb_id_out0); + Semaphore<> sender_sem(get_compile_time_arg_val(10)); + Semaphore<> receiver_sem(get_compile_time_arg_val(11)); +#ifdef FUSE_BIAS + DataflowBuffer dfb_in3(dfb_id_in3); +#endif + +// READER +#ifdef IN1_SHARDED + dfb_in1.reserve_back(in1_block_num_tiles * num_blocks_inner_dim); + dfb_in1.push_back(in1_block_num_tiles * num_blocks_inner_dim); +#elif !defined(ENABLE_GLOBAL_CB) + uint32_t l1_write_addr_in1; + + [[maybe_unused]] const auto s1 = TensorAccessor(in1_args, in1_tensor_addr); +#endif // IN1_SHARDED / ENABLE_GLOBAL_CB + +#ifdef ENABLE_GLOBAL_CB + constexpr uint32_t remote_cb_id = tt::CBIndex::c_31; + const uint32_t in1_fifo_tiles = get_local_cb_interface(dfb_id_in1).fifo_num_pages; +#endif + +#ifdef G4_IN1 + constexpr uint32_t g4_d = G4_IN1_D; + constexpr uint32_t g4_nocs = G4_IN1_NOCS; + static_assert(g4_d >= 1 && g4_d < NOC_MAX_TRANSACTION_ID && (g4_nocs == 1 || g4_nocs == 2)); + const Noc g4_noc0(noc.get_noc_id()); + const Noc g4_noc1(noc.get_noc_id() ^ 1); + // Whole-read NoC choice for G4_IN1_SPLIT 3 (per worker: odd workers, or G4_NOC_MAP[worker] when given; K block + // parity for 4 is applied per block). +#ifdef G4_NOC_MAP + constexpr uint8_t g4_noc_map[] = {G4_NOC_MAP}; + const bool g4_odd_worker = g4_noc_map[in1_tensor_start_tile_id / in1_block_w] != 0; +#else + const bool g4_odd_worker = ((in1_tensor_start_tile_id / in1_block_w) & 1) != 0; +#endif +#if G4_IN1_SPLIT >= 5 + // G4_IN1_SPLIT 5/6: per DRAM bank, bit b of the mask selects the second NoC for reads from bank b. The factory + // gives workers west / east of logical column G4_EAST_X the masks that keep each bank's data on the NoC whose + // route from the bank's column to the worker is short. + const uint32_t g4_noc1_banks = + get_absolute_logical_x() < G4_EAST_X ? G4_NOC1_BANKS_WEST : G4_NOC1_BANKS_EAST; +#else + constexpr uint32_t g4_noc1_banks = 0; +#endif + const uint32_t g4_fifo_limit = get_local_cb_interface(dfb_id_in1).fifo_limit; + const uint32_t g4_fifo_size = get_local_cb_interface(dfb_id_in1).fifo_size; + uint32_t g4_pending = 0; // blocks issued but not yet pushed + uint32_t g4_issue_trid = 1; + uint32_t g4_wait_trid = 1; +#ifdef G4_PACKED_IN1 + static_assert(in1_block_w == G4_PACKED_PN); + // This core's worker index follows from its first output column (non-transposed weight). + const uint32_t g4_worker = in1_tensor_start_tile_id / G4_PACKED_PN; +#if G4_PACKED_G == 0 +#ifdef G4_PACKED_BANK_MAP + constexpr uint8_t g4_bank_map[] = {G4_PACKED_BANK_MAP}; + const uint32_t g4_fixed_bank = g4_bank_map[g4_worker]; + uint32_t g4_rank = 0; + for (uint32_t i = 0; i < g4_worker; ++i) { + g4_rank += g4_bank_map[i] == g4_fixed_bank ? 1 : 0; + } +#else + const uint32_t g4_fixed_bank = g4_worker % G4_PACKED_BANKS; + const uint32_t g4_rank = g4_worker / G4_PACKED_BANKS; +#endif +#endif + AllocatorBank g4_dram; + auto g4_read_run = [&](const Noc& n, uint32_t bank, uint32_t src, uint32_t dst, uint32_t bytes) { + if constexpr (g4_d > 1) { + n.async_read( + g4_dram, CoreLocalMem(dst), bytes, {.bank_id = bank, .addr = src}, {}, + NocOptVals{.trid = g4_issue_trid}); + } else { + n.async_read(g4_dram, CoreLocalMem(dst), bytes, {.bank_id = bank, .addr = src}, {}); + } + }; +#else + auto g4_read_page = [&](const Noc& n, uint32_t page, uint32_t dst) { + if constexpr (g4_d > 1) { + n.async_read( + s1, CoreLocalMem(dst), in1_single_tile_size_bytes, {.page_id = page}, {}, + NocOptVals{.trid = g4_issue_trid}); + } else { + n.async_read( + s1, CoreLocalMem(dst), in1_single_tile_size_bytes, {.page_id = page}, {}); + } + }; +#endif + auto g4_wait = [&]() { + if constexpr (g4_d > 1) { + g4_noc0.async_read_barrier({.trid = g4_wait_trid}); + if constexpr (g4_nocs == 2) { + g4_noc1.async_read_barrier({.trid = g4_wait_trid}); + } + } else { + g4_noc0.async_read_barrier(); + if constexpr (g4_nocs == 2) { + g4_noc1.async_read_barrier(); + } + } + }; +#endif // G4_IN1 + + // WRITER + const auto s = TensorAccessor(out_args, out_tensor_addr); + // `s` is only consumed inside the `#ifndef OUT_SHARDED` write path below; mark it used so + // sharded builds don't warn (-Wunused-but-set-variable). + (void)s; + + // sparsity accessor + constexpr uint32_t dfb_id_sparsity = get_named_compile_time_arg_val("cb_sparsity"); + DataflowBuffer dfb_sparsity(dfb_id_sparsity); + const auto s_sparsity = TensorAccessor(sparsity_args, sparsity_addr); + +#ifndef SKIP_MCAST + // Set ur local VALID value, to be mcasted to destinations flag address after the data has been mcasted + receiver_sem.set(VALID); + // local address that will be atomically incremented by mcast receivers, to know when all receivers are ready + // to receive the mcast + +#ifdef IN1_SHARDED + uint64_t in1_start_address = dfb_in1.get_write_ptr(); +#endif // IN1_SHARDED +#endif // SKIP_MCAST + + uint32_t l1_write_addr_sparsity = 0; + if constexpr (batchB > 0) { + dfb_sparsity.reserve_back(1); + l1_write_addr_sparsity = dfb_sparsity.get_write_ptr(); + } + +#ifdef IN1_DRAM_WIDTH_SHARDED + constexpr uint32_t in1_dram_block_size_bytes = in1_dram_block_num_tiles * in1_single_tile_size_bytes; + uint32_t in1_block_w_bytes = in1_block_w * in1_single_tile_size_bytes; +#endif // IN1_DRAM_WIDTH_SHARDED + +#ifdef IN1_DRAM_HEIGHT_SHARDED + constexpr uint32_t in1_batch_stride_bytes = in1_KtNt_per_batch * in1_single_tile_size_bytes; +#endif // IN1_DRAM_HEIGHT_SHARDED + + for (uint32_t b = 0; b < batch; ++b) { + uint32_t in1_batch_tile_id = in1_tensor_start_tile_id; + +#ifdef IN1_DRAM_HEIGHT_SHARDED + // Compute DRAM bank and offset for this batch + uint32_t in1_dram_bank_id = b / in1_batches_per_bank; + uint32_t in1_batch_in_shard = b % in1_batches_per_bank; + AllocatorBank dram_src; + uint32_t in1_dram_batch_offset = in1_batch_in_shard * in1_batch_stride_bytes; +#endif // IN1_DRAM_HEIGHT_SHARDED + + if constexpr (batchB > 0) { + noc.async_read(s_sparsity, dfb_sparsity, sparsity_pagesize, {.page_id = b}, {.offset_bytes = 0}); + noc.async_read_barrier(); + } + + for (uint32_t bB = 0; bB < batchB_lim; ++bB) { + if constexpr (batchB > 0) { + if (reinterpret_cast(l1_write_addr_sparsity)[bB] == 0) { + out_tensor_start_tile_id += MtNt; + in1_batch_tile_id += KtNt; + continue; + } + } + + uint32_t in1_tensor_current_h_dim_block_tile_id = in1_batch_tile_id; + uint32_t out_tensor_current_h_dim_block_tile_id = out_tensor_start_tile_id; + for (uint32_t bh = 0; bh < num_blocks_h_dim; ++bh) { + uint32_t in1_tensor_current_w_dim_block_tile_id = in1_tensor_current_h_dim_block_tile_id; + uint32_t out_tensor_current_w_dim_block_tile_id = out_tensor_current_h_dim_block_tile_id; +#ifdef FUSE_BIAS + uint32_t in3_tensor_current_w_dim_block_tile_id = in3_tensor_start_tile_id; +#endif // FUSE_BIAS + for (uint32_t bw = 0; bw < num_blocks_w_dim; ++bw) { + uint32_t in1_tensor_current_inner_dim_block_start_tile_id = in1_tensor_current_w_dim_block_tile_id; +#ifdef IN1_DRAM_WIDTH_SHARDED + // Reset DRAM read offset for each bh block — the inner dim loop + // advances through K, and each output row block re-reads the same + // in1 columns from K=0. (bw is always 1 for DRAM-sharded senders.) + uint32_t l1_read_addr_in1_offset = 0; +#endif // IN1_DRAM_WIDTH_SHARDED + + for (uint32_t block = 0; block < num_blocks_inner_dim; ++block) { + if constexpr (fuse_op_all_gather) { + fused_op_receiver.update_current_block_start_tile_id( + block, in1_tensor_current_inner_dim_block_start_tile_id, in1_batch_tile_id); + } +#if defined(ENABLE_GLOBAL_CB) + // The tensor prefetcher pushes this receiver's K-blocks in natural order. + // Keep one block of lookahead: publish the current block to compute, then + // wait for the unpack engine to drain the previous block before returning + // its remote-CB credit to the prefetcher. + dfb_in1.reserve_back(in1_block_num_tiles); + experimental::remote_cb_wait_front(remote_cb_id, block == 0 ? 1u : 2u); +#elif defined(IN1_DRAM_WIDTH_SHARDED) + // Operand 1 - DRAM width sharded + dfb_in1.reserve_back(in1_block_num_tiles); + + uint64_t in1_start_address = + dfb_in1.get_write_ptr(); // copy start address of block, to be used for mcasting + + uint32_t l1_write_addr_in1_offset = 0; + uint32_t next_bank_id_and_dram_stride_index = 0; + + AllocatorBank dram_bank; + for (uint32_t i = 0; i < num_dram_shards_to_read; ++i) { + uint32_t shard_bank_id = current_dram_bank_id[next_bank_id_and_dram_stride_index]; + uint32_t shard_base_addr = in1_tensor_addr; + if (i == 0) { + shard_base_addr += dram_tensor_start_offset; + } + noc.set_async_read_state( + dram_bank, + in1_single_tile_size_bytes, + {.bank_id = shard_bank_id, .addr = shard_base_addr}, + NocOptVals{.vc = vc}); + + uint32_t l1_read_addr_in1 = l1_read_addr_in1_offset; + uint32_t l1_write_addr_in1 = dfb_in1.get_write_ptr() + l1_write_addr_in1_offset; + uint32_t in1_block_w_dram = + in1_block_w_dram_stride_bytes[next_bank_id_and_dram_stride_index] / + in1_single_tile_size_bytes; + + for (uint32_t m = 0; m < in1_block_h; ++m) { + uint32_t l1_read_addr_in1_temp = l1_read_addr_in1; + uint32_t l1_write_addr_in1_temp = l1_write_addr_in1; + for (uint32_t w = 0; w < in1_block_w_dram; ++w) { + noc.async_read_with_state( + dram_bank, + CoreLocalMem(l1_write_addr_in1_temp), + in1_single_tile_size_bytes, + {.bank_id = shard_bank_id, .addr = shard_base_addr + l1_read_addr_in1_temp}, + {}, + NocOptVals{.vc = vc}); + l1_read_addr_in1_temp += in1_single_tile_size_bytes; + l1_write_addr_in1_temp += in1_single_tile_size_bytes; + } + l1_read_addr_in1 += in1_block_w_dram_bytes; + l1_write_addr_in1 += in1_block_w_bytes; + } + l1_write_addr_in1_offset += + in1_block_w_dram_stride_bytes[next_bank_id_and_dram_stride_index]; + next_bank_id_and_dram_stride_index += 2; + } + l1_read_addr_in1_offset += in1_dram_block_size_bytes; + noc.async_read_barrier(); +#elif defined(IN1_DRAM_HEIGHT_SHARDED) + // Operand 1 - DRAM height sharded (batched) + // Each DRAM bank holds batches_per_bank complete [K, N] matrices + // Bank and offset computed at start of batch loop + dfb_in1.reserve_back(in1_block_num_tiles); + + l1_write_addr_in1 = dfb_in1.get_write_ptr(); + uint64_t in1_start_address = + l1_write_addr_in1; // copy start address of block, to be used for mcasting + + // Read in1 block from the correct DRAM bank + // Tile layout within a batch: row-major [K, N], same strides as interleaved + uint32_t in1_tensor_row_start_tile_id = in1_tensor_current_inner_dim_block_start_tile_id; + for (uint32_t h = 0; h < in1_block_h; ++h) { + uint32_t in1_tensor_tile_id = in1_tensor_row_start_tile_id; + for (uint32_t w = 0; w < in1_block_w; ++w) { + if (bw < num_blocks_w_dim - 1 || w < last_block_w) { + uint32_t tile_byte_offset = + in1_dram_batch_offset + in1_tensor_tile_id * in1_single_tile_size_bytes; + noc.async_read( + dram_src, + CoreLocalMem(l1_write_addr_in1), + in1_single_tile_size_bytes, + {.bank_id = in1_dram_bank_id, .addr = in1_tensor_addr + tile_byte_offset}, + {}); + } + l1_write_addr_in1 += in1_single_tile_size_bytes; + in1_tensor_tile_id += in1_tensor_stride_w; + } + in1_tensor_row_start_tile_id += in1_tensor_stride_h; + } + in1_tensor_current_inner_dim_block_start_tile_id += in1_tensor_next_block_stride; + + // Barrier! make sure the reads are done + noc.async_read_barrier(); +#elif defined(G4_IN1) + // Operand 1 - Gemma decode reader (see G4_IN1 above): reserve room for every unpublished + // block plus this one, issue this block's reads and publish the oldest block once g4_d + // blocks are in flight. + dfb_in1.reserve_back(in1_block_num_tiles * (g4_pending + 1)); + { + uint32_t g4_dst = dfb_in1.get_write_ptr() + g4_pending * in1_block_size_bytes; + if (g4_dst >= g4_fifo_limit) { + g4_dst -= g4_fifo_size; + } + const bool g4_block_other = g4_nocs == 2 && (G4_IN1_SPLIT == 3 ? g4_odd_worker + : G4_IN1_SPLIT == 4 ? (block & 1) != 0 + : false); +#ifdef G4_PACKED_IN1 + constexpr uint32_t pn = G4_PACKED_PN; + constexpr uint32_t banks = G4_PACKED_BANKS; + uint32_t g4_k = block * in1_block_h; + for (uint32_t h = 0; h < in1_block_h;) { +#if G4_PACKED_G == 0 + const uint32_t g4_bank = g4_fixed_bank; + const uint32_t g4_slot = g4_rank * G4_PACKED_KT * pn + g4_k * pn; + const uint32_t g4_rows = in1_block_h - h; +#else + constexpr uint32_t group = G4_PACKED_G; + const uint32_t g4_group = g4_k / group; + const uint32_t g4_row = g4_k - g4_group * group; + const uint32_t g4_bank = (g4_worker + g4_group) % banks; + const uint32_t g4_slot = g4_worker * (G4_PACKED_KT / banks) * pn + + (g4_group / banks) * group * pn + g4_row * pn; + const uint32_t g4_rows = + in1_block_h - h < group - g4_row ? in1_block_h - h : group - g4_row; +#endif + const uint32_t g4_bytes = g4_rows * pn * in1_aligned_tile_size_bytes; + const uint32_t g4_src = in1_tensor_addr + g4_slot * in1_aligned_tile_size_bytes; + if constexpr (g4_nocs == 2 && G4_IN1_SPLIT <= 2) { + const uint32_t first = (g4_bytes / 2) & ~(DRAM_ALIGNMENT - 1); + g4_read_run(g4_noc0, g4_bank, g4_src, g4_dst, first); + g4_read_run(g4_noc1, g4_bank, g4_src + first, g4_dst + first, g4_bytes - first); + } else { + const bool g4_other = G4_IN1_SPLIT >= 5 ? ((g4_noc1_banks >> g4_bank) & 1) != 0 + : g4_block_other; + g4_read_run(g4_other ? g4_noc1 : g4_noc0, g4_bank, g4_src, g4_dst, g4_bytes); + } + g4_dst += g4_bytes; + h += g4_rows; + g4_k += g4_rows; + } +#else + uint32_t in1_tensor_row_start_tile_id = in1_tensor_current_inner_dim_block_start_tile_id; + uint32_t g4_t = 0; + for (uint32_t h = 0; h < in1_block_h; ++h) { + uint32_t in1_tensor_tile_id = in1_tensor_row_start_tile_id; + for (uint32_t w = 0; w < in1_block_w; ++w, ++g4_t) { + if (bw < num_blocks_w_dim - 1 || w < last_block_w) { + const bool other = + g4_nocs == 2 && + (G4_IN1_SPLIT == 1 ? (g4_t & 1) != 0 + : G4_IN1_SPLIT == 2 ? g4_t >= in1_block_num_tiles / 2 + : G4_IN1_SPLIT >= 5 + ? ((g4_noc1_banks >> (in1_tensor_tile_id % G4_IN1_BANKS)) & 1) != 0 + : g4_block_other); + g4_read_page(other ? g4_noc1 : g4_noc0, in1_tensor_tile_id, g4_dst); + } + g4_dst += in1_aligned_tile_size_bytes; + in1_tensor_tile_id += in1_tensor_stride_w; + } + in1_tensor_row_start_tile_id += in1_tensor_stride_h; + } +#endif + } + in1_tensor_current_inner_dim_block_start_tile_id += in1_tensor_next_block_stride; + g4_issue_trid = g4_issue_trid == g4_d ? 1 : g4_issue_trid + 1; + if (++g4_pending == g4_d) { + g4_wait(); + dfb_in1.push_back(in1_block_num_tiles); + g4_wait_trid = g4_wait_trid == g4_d ? 1 : g4_wait_trid + 1; + --g4_pending; + } +#elif !defined(IN1_SHARDED) + // Operand 1 - interleaved + dfb_in1.reserve_back(in1_block_num_tiles); + uint32_t in1_write_offset = 0; + uint64_t in1_start_address = + dfb_in1.get_write_ptr(); // copy start address of block, to be used for mcasting + + // Copy in1 block into CB, as the default kernel + uint32_t in1_tensor_row_start_tile_id = in1_tensor_current_inner_dim_block_start_tile_id; + for (uint32_t h = 0; h < in1_block_h; ++h) { + uint32_t in1_tensor_tile_id = in1_tensor_row_start_tile_id; + for (uint32_t w = 0; w < in1_block_w; ++w) { + if (bw < num_blocks_w_dim - 1 || w < last_block_w) { + noc.async_read( + s1, + dfb_in1, + in1_single_tile_size_bytes, + {.page_id = in1_tensor_tile_id}, + {.offset_bytes = in1_write_offset}); + } + in1_write_offset += in1_aligned_tile_size_bytes; + in1_tensor_tile_id += in1_tensor_stride_w; + } + in1_tensor_row_start_tile_id += in1_tensor_stride_h; + } + in1_tensor_current_inner_dim_block_start_tile_id += in1_tensor_next_block_stride; + + // Barrier! make sure the reads are done + noc.async_read_barrier(); +#endif // IN1_DRAM_WIDTH_SHARDED / IN1_DRAM_HEIGHT_SHARDED / IN1_SHARDED + +#ifndef SKIP_MCAST + // wait until all in1 mcast destinations have atomically incremented the in1 semaphore_addr + // (i.e. its value should be in0_mcast_num_dests), then reset the semaphore_addr value back to + // zero for the next block + sender_sem.wait(in1_mcast_num_dests); + sender_sem.set(0); + + // Now we have the block in the CB address, we can mcast to dests! + MulticastEndpoint mcast_dst; + // num_dests must not include source, since we are NOT really doing a local copy! + noc.async_write_multicast( + CoreLocalMem(static_cast(in1_start_address)), + mcast_dst, + in1_block_size_bytes, + in1_mcast_num_cores, + {}, + {.noc_x_start = in1_mcast_dest_noc_start_x, + .noc_y_start = in1_mcast_dest_noc_start_y, + .noc_x_end = in1_mcast_dest_noc_end_x, + .noc_y_end = in1_mcast_dest_noc_end_y, + .addr = static_cast(in1_start_address)}, + true); + + // Note: no need for write barrier, since these two multicasts are done on the same noc id and + // same vc even though cmd bufs are different Also, this only works because we are setting VCs + // statically (using NOC_CMD_STATIC_VC). +#ifdef ARCH_BLACKHOLE + // On Blackhole the flush is needed because NoC latency is higher than L1 <-> RISCV latency + // which means data could be changed before + // write is issued. + noc.async_writes_flushed(); +#endif // ARCH_BLACKHOLE + + // We should also multicast the flag to destinations + // num_dests must not include source, since we are NOT really doing a local copy! + receiver_sem.set_multicast( + noc, + in1_mcast_dest_noc_start_x, + in1_mcast_dest_noc_start_y, + in1_mcast_dest_noc_end_x, + in1_mcast_dest_noc_end_y, + in1_mcast_num_cores); +#endif // SKIP_MCAST + +#if !defined(IN1_SHARDED) && !defined(G4_IN1) + dfb_in1.push_back(in1_block_num_tiles); +#endif // IN1_SHARDED +#ifdef ENABLE_GLOBAL_CB + if (block >= 1) { + while (!dfb_in1.pages_reservable_at_back(in1_fifo_tiles - in1_block_num_tiles)) { + invalidate_l1_cache(); + } + experimental::remote_cb_pop_front(remote_cb_id, 1); + } +#endif + } +#ifdef G4_IN1 + // Publish the blocks still in flight, oldest first. + for (; g4_pending > 0; --g4_pending) { + g4_wait(); + dfb_in1.push_back(in1_block_num_tiles); + g4_wait_trid = g4_wait_trid == g4_d ? 1 : g4_wait_trid + 1; + } + if constexpr (g4_d > 1) { + // The packet tag register is sticky: leave later plain reads untagged. + noc_async_read_set_trid(0, g4_noc0.get_noc_id()); + if constexpr (g4_nocs == 2) { + noc_async_read_set_trid(0, g4_noc1.get_noc_id()); + } + } +#endif +#ifdef ENABLE_GLOBAL_CB + if (num_blocks_inner_dim > 0) { + while (!dfb_in1.pages_reservable_at_back(in1_fifo_tiles)) { + invalidate_l1_cache(); + } + experimental::remote_cb_pop_front(remote_cb_id, 1); + } +#endif +#ifdef FUSE_BIAS + // Only read bias on first batch, or we have multiple output blocks + if ((b == 0 && bh == 0) || num_blocks_w_dim > 1) { + // Operand 1 +#ifndef BIAS_SHARDED + dfb_in3.reserve_back(in1_block_w); + uint32_t in3_write_offset = 0; + + uint64_t in3_start_address = + dfb_in3.get_write_ptr(); // copy start address of block, to be used for mcasting + uint32_t in3_block_size_bytes = 0; // can be optimized later, pass it to kernel + +#ifdef IN1_DRAM_WIDTH_SHARDED + uint32_t l1_write_addr_in3_offset = 0; + uint32_t next_bank_id_and_dram_stride_index = 0; + + AllocatorBank bias_dram_bank; + for (uint32_t i = 0; i < num_dram_shards_to_read; ++i) { + uint32_t bias_shard_bank_id = current_dram_bank_id[next_bank_id_and_dram_stride_index]; + uint32_t bias_shard_base_addr = in3_tensor_addr; + if (i == 0) { + // dram_tensor_start_offset is in in1 tile bytes; convert to + // bias tile bytes since bias_dtype may differ from in1_dtype. + bias_shard_base_addr += (dram_tensor_start_offset / in1_single_tile_size_bytes) * + bias_single_tile_size_bytes; + } + + noc.set_async_read_state( + bias_dram_bank, + bias_single_tile_size_bytes, + {.bank_id = bias_shard_bank_id, .addr = bias_shard_base_addr}, + NocOptVals{.vc = vc}); + + uint32_t l1_read_addr_in3 = 0; + l1_write_addr_in3 = dfb_in3.get_write_ptr() + l1_write_addr_in3_offset; + // in1_block_w_dram_stride_bytes is in in1 tile bytes, so divide + // by in1_single_tile_size_bytes (not bias) to get the tile count. + uint32_t in3_block_w_dram = + in1_block_w_dram_stride_bytes[next_bank_id_and_dram_stride_index] / + in1_single_tile_size_bytes; + + for (uint32_t w = 0; w < in3_block_w_dram; ++w) { + noc.async_read_with_state( + bias_dram_bank, + CoreLocalMem(l1_write_addr_in3), + bias_single_tile_size_bytes, + {.bank_id = bias_shard_bank_id, .addr = bias_shard_base_addr + l1_read_addr_in3}, + {}, + NocOptVals{.vc = vc}); + l1_read_addr_in3 += bias_single_tile_size_bytes; + l1_write_addr_in3 += bias_single_tile_size_bytes; + in3_block_size_bytes += bias_single_tile_size_bytes; + } + // Advance L1 offset in bias tile bytes, not in1 stride bytes. + l1_write_addr_in3_offset += in3_block_w_dram * bias_single_tile_size_bytes; + next_bank_id_and_dram_stride_index += 2; + } + noc.async_read_barrier(); +#else + // Copy in1 block into CB, as the default kernel + uint32_t in3_tensor_tile_id = in3_tensor_current_w_dim_block_tile_id; + for (uint32_t w = 0; w < in1_block_w; ++w) { + if (bw < num_blocks_w_dim - 1 || w < last_block_w) { + noc.async_read( + s3, + dfb_in3, + bias_single_tile_size_bytes, + {.page_id = in3_tensor_tile_id}, + {.offset_bytes = in3_write_offset}); + } + in3_write_offset += bias_single_tile_size_bytes; + in3_tensor_tile_id += in3_tensor_stride_w; + in3_block_size_bytes += bias_single_tile_size_bytes; + } + // Barrier! make sure the reads are done + noc.async_read_barrier(); +#endif // IN1_DRAM_WIDTH_SHARDED + +#ifndef SKIP_MCAST + + // wait until all in1 mcast destinations have atomically incremented the in1 semaphore_addr + // (i.e. its value should be in0_mcast_num_dests), then reset the semaphore_addr value back to + // zero for the next block + sender_sem.wait(in1_mcast_num_dests); + sender_sem.set(0); + + // Now we have the block in the CB address, we can mcast to dests! + MulticastEndpoint mcast_dst; + // num_dests must not include source, since we are NOT really doing a local copy! + noc.async_write_multicast( + CoreLocalMem(static_cast(in3_start_address)), + mcast_dst, + in3_block_size_bytes, + in1_mcast_num_cores, + {}, + {.noc_x_start = in1_mcast_dest_noc_start_x, + .noc_y_start = in1_mcast_dest_noc_start_y, + .noc_x_end = in1_mcast_dest_noc_end_x, + .noc_y_end = in1_mcast_dest_noc_end_y, + .addr = static_cast(in3_start_address)}, + true); + // Note: no need for write barrier, since these two multicasts are done on the same noc id, same + // vc, same cmd_buf Also, this only works because we are setting VCs statically (using + // NOC_CMD_STATIC_VC). +#ifdef ARCH_BLACKHOLE + // On Blackhole the flush is needed because NoC latency is higherthan L1 <-> RISCV + // latency which means data could be changed before write is issued. + noc.async_writes_flushed(); +#endif // ARCH_BLACKHOLE + + // We should also multicast the flag to destinations + // num_dests must not include source, since we are NOT really doing a local copy! + receiver_sem.set_multicast( + noc, + in1_mcast_dest_noc_start_x, + in1_mcast_dest_noc_start_y, + in1_mcast_dest_noc_end_x, + in1_mcast_dest_noc_end_y, + in1_mcast_num_cores); +#endif // SKIP_MCAST + + dfb_in3.push_back(in1_block_w); +#else + dfb_in3.reserve_back(in1_block_w); + dfb_in3.push_back(in1_block_w); +#endif // BIAS_SHARDED + } +#endif // FUSE_BIAS + +#ifndef OUT_SHARDED + // WRITER + uint32_t num_blocks_w_dim_ = + bw >= last_num_blocks_w_dim - 1 ? last_num_blocks_w_dim : num_blocks_w_dim; + uint32_t out_num_nonzero_subblocks_h_ = out_num_nonzero_subblocks_h; + uint32_t out_num_nonzero_subblocks_w_ = out_num_nonzero_subblocks_w; + if (bw == num_blocks_w_dim_ - 1) { + out_num_nonzero_subblocks_w_ = out_last_num_nonzero_subblocks_w; + } + uint32_t out_tensor_sbh_start_tile_id = out_tensor_current_w_dim_block_tile_id; + for (uint32_t sbh = 0; sbh < out_num_nonzero_subblocks_h_; ++sbh) { + uint32_t out_tensor_sbw_start_tile_id = out_tensor_sbh_start_tile_id; + for (uint32_t sbw = 0; sbw < out_num_nonzero_subblocks_w_; ++sbw) { + uint32_t out_tensor_sb_row_start_tile_id = out_tensor_sbw_start_tile_id; + + uint32_t out_subblock_h_ = out_subblock_h; + uint32_t out_subblock_w_ = out_subblock_w; + uint32_t subblock_tiles_addr_skip = 0; + if (bh == num_blocks_h_dim - 1 && sbh == out_num_nonzero_subblocks_h - 1) { + out_subblock_h_ = out_last_subblock_h; + } + if (bw == num_blocks_w_dim_ - 1 && sbw == out_num_nonzero_subblocks_w_ - 1) { + out_subblock_w_ = out_last_subblock_w; + subblock_tiles_addr_skip = padded_subblock_tiles_addr_skip; + } + + dfb_out.wait_front(out_subblock_tile_count); + uint32_t out_read_offset = 0; + + for (uint32_t h = 0; h < out_subblock_h_; ++h) { + uint32_t out_tensor_tile_id = out_tensor_sb_row_start_tile_id; + for (uint32_t w = 0; w < out_subblock_w_; ++w) { + if (bw < num_blocks_w_dim_) { + noc.async_write( + dfb_out, + s, + output_single_tile_size_bytes, + {.offset_bytes = out_read_offset}, + {.page_id = out_tensor_tile_id}); + } + + out_read_offset += output_single_tile_size_bytes; + + out_tensor_tile_id += out_tensor_stride_w; + } + // Skip padded tiles in subblock along row + out_read_offset += subblock_tiles_addr_skip; + out_tensor_sb_row_start_tile_id += out_tensor_stride_h; + } + + noc.async_write_barrier(); + dfb_out.pop_front(out_subblock_tile_count); + out_tensor_sbw_start_tile_id += out_tensor_next_subblock_stride_w; + } + // Pop fully padded subblocks along the row + if (bw == num_blocks_w_dim_ - 1) { + dfb_out.wait_front(padded_block_tiles_w_skip); + dfb_out.pop_front(padded_block_tiles_w_skip); + } + out_tensor_sbh_start_tile_id += out_tensor_next_subblock_stride_h; + } + // Pop row(s) of fully padded subblocks + if (bh == num_blocks_h_dim - 1) { + dfb_out.wait_front(padded_block_tiles_h_skip); + dfb_out.pop_front(padded_block_tiles_h_skip); + } + +#endif + in1_tensor_current_w_dim_block_tile_id += in1_tensor_next_w_dim_block_stride; + out_tensor_current_w_dim_block_tile_id += out_tensor_next_w_dim_block_stride; +#ifdef FUSE_BIAS + in3_tensor_current_w_dim_block_tile_id += in1_block_w; +#endif + } + out_tensor_current_h_dim_block_tile_id += out_tensor_next_h_dim_block_stride; + } + out_tensor_start_tile_id += MtNt; + in1_batch_tile_id += KtNt; + } + if constexpr (bcast_B == 0) { +#ifndef IN1_DRAM_HEIGHT_SHARDED + // For height-sharded DRAM, tile IDs are relative within a batch; + // batch offset is handled by switching DRAM banks + in1_tensor_start_tile_id += KtNt; +#endif + } + + if (fuse_op_reduce_scatter) { + // Signal reduce_scatter to go + op_signaler.synchronize_workers_and_signal_op(0); + } + } + +#if OUT_SHARDED + dfb_out.wait_front( + batch * out_num_nonzero_subblocks_h * out_num_nonzero_subblocks_w * out_subblock_w * out_subblock_h); +#endif +#ifdef ENABLE_GLOBAL_CB + experimental::update_remote_cb_config_in_l1(remote_cb_id); + noc.async_atomic_barrier(); +#endif + noc.async_write_barrier(); +}