From 9e1202e5b62f222acf0c3ed8a40961df02cf00b2 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Sat, 26 Sep 2026 13:27:38 +0800 Subject: [PATCH 1/2] example(warp_reduce_redux_bench): measure both warp_reduce forms #1305 selects the reduction form at compile time from the target floor, so the comparison has to be two builds of one source rather than two kernels in one binary. This adds that harness, plus the control the comparison needs. - `api_chain` / `api_independent` drive `warp_reduce::`: one is a dependency chain (reduction latency), the other four unrelated reductions per iteration (throughput). - `shuffle_chain` / `shuffle_independent` spell the butterfly out with `WarpCollective::shfl_xor`, so at `sm_80` the pair isolates the reduction from the rest of the build: same target, same compiler, same binary. - 128 threads, 40,000 iterations, 15 interleaved samples and 3 warm-up rounds per arm, medians reported per reduction, and the first iterations checked against a host simulation so a folded-away loop cannot pass as a timing. Measured on an A100-SXM4-40GB (cc 8.0, driver 595.91.07). At `sm_80`, where the API takes `redux.sync.add.s32` and the control keeps its five `shfl.sync.bfly.b32` + `add.s32`: | workload | redux ns/reduction | butterfly ns/reduction | ratio | |---|---|---|---| | dependent chain | 8.4 - 10.8 | 27.2 - 35.0 | 3.24x | | four independent reductions | 2.9 - 3.7 | 9.4 - 12.1 | 3.24 - 3.27x | Three batches agree on the ratio to the printed digit. Built at `sm_75` the two arms are the same code and time identically, which is the control for the harness itself. Needs #1305 for the `sm_80` arm to be a `redux.sync` build; without it both arms are the butterfly. Signed-off-by: 0z5a --- .../warp_reduce_redux_bench/Cargo.lock | 723 ++++++++++++++++++ .../warp_reduce_redux_bench/Cargo.toml | 17 + .../warp_reduce_redux_bench/src/main.rs | 367 +++++++++ 3 files changed, 1107 insertions(+) create mode 100644 crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/Cargo.lock create mode 100644 crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/Cargo.toml create mode 100644 crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/src/main.rs diff --git a/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/Cargo.lock b/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/Cargo.lock new file mode 100644 index 0000000000..abe1001049 --- /dev/null +++ b/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/Cargo.lock @@ -0,0 +1,723 @@ +# This file is automatically @generated by Cargo. +# It is not intended for manual editing. +version = 4 + +[[package]] +name = "aho-corasick" +version = "1.1.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ddd31a130427c27518df266943a5308ed92d4b226cc639f5a8f1002816174301" +dependencies = [ + "memchr", +] + +[[package]] +name = "anyhow" +version = "1.0.102" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7f202df86484c868dbad7eaa557ef785d5c66295e41b460ef922eca0723b842c" + +[[package]] +name = "bindgen" +version = "0.69.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "271383c67ccabffb7381723dea0672a673f292304fcb45c01cc648c7a8d58088" +dependencies = [ + "bitflags", + "cexpr", + "clang-sys", + "itertools", + "lazy_static", + "lazycell", + "log", + "prettyplease", + "proc-macro2", + "quote", + "regex", + "rustc-hash", + "shlex", + "syn", + "which", +] + +[[package]] +name = "bitflags" +version = "2.13.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b4388bee8683e3d04af747c73422af53102d2bd24d9eadb6cbc100baef4b43f8" + +[[package]] +name = "block-buffer" +version = "0.10.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3078c7629b62d3f0439517fa394996acacc5cbc91c5a20d8c658e77abd503a71" +dependencies = [ + "generic-array", +] + +[[package]] +name = "cexpr" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6fac387a98bb7c37292057cffc56d62ecb629900026402633ae9160df93a8766" +dependencies = [ + "nom", +] + +[[package]] +name = "cfg-if" +version = "1.0.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801" + +[[package]] +name = "clang-sys" +version = "1.8.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0b023947811758c97c59bf9d1c188fd619ad4718dcaa767947df1cadb14f39f4" +dependencies = [ + "glob", + "libc", + "libloading", +] + +[[package]] +name = "cpufeatures" +version = "0.2.17" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "59ed5838eebb26a2bb2e58f6d5b5316989ae9d08bab10e0e6d103e656d1b0280" +dependencies = [ + "libc", +] + +[[package]] +name = "crc32fast" +version = "1.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9481c1c90cbf2ac953f07c8d4a58aa3945c425b7185c9154d67a65e4230da511" +dependencies = [ + "cfg-if", +] + +[[package]] +name = "crunchy" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "460fbee9c2c2f33933d720630a6a0bac33ba7053db5344fac858d4b8952d77d5" + +[[package]] +name = "crypto-common" +version = "0.1.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "78c8292055d1c1df0cce5d180393dc8cce0abec0a7102adb6c7b1eef6016d60a" +dependencies = [ + "generic-array", + "typenum", +] + +[[package]] +name = "cuda-artifact-finalizer" +version = "0.2.1" +dependencies = [ + "libnvvm-sys", + "nvjitlink-sys", + "serde", + "sha2", + "thiserror", +] + +[[package]] +name = "cuda-bindings" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0f42e2c99c56bf2eef9d569435bfd72a478dfaea3f178b8e5bae891875bbbe2e" +dependencies = [ + "bindgen", + "libloading", + "prettyplease", + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "cuda-core" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "84bdb592a2844148fb13bdfbbab158879f56916bde1d1a1b6580ca27f23850cd" +dependencies = [ + "anyhow", + "cuda-bindings", + "cuda-core-derive", + "half", + "oxide-artifacts", +] + +[[package]] +name = "cuda-core-derive" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8a48da7d10c1078f7fbb87d7176e9476affadea148babf31946e348073872d44" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "cuda-device" +version = "0.2.1" +dependencies = [ + "cuda-macros", + "cuda-target-spec", +] + +[[package]] +name = "cuda-host" +version = "0.2.1" +dependencies = [ + "cuda-artifact-finalizer", + "cuda-core", + "cuda-macros", + "half", + "ptx-parse", + "sha2", + "thiserror", +] + +[[package]] +name = "cuda-macros" +version = "0.2.1" +dependencies = [ + "proc-macro2", + "quote", + "reserved-oxide-symbols", + "syn", +] + +[[package]] +name = "cuda-target-spec" +version = "0.1.0" + +[[package]] +name = "digest" +version = "0.10.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9ed9a281f7bc9b7576e61468ba615a66a5c8cfdff42420a70aa82701a3b1e292" +dependencies = [ + "block-buffer", + "crypto-common", +] + +[[package]] +name = "either" +version = "1.16.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "91622ff5e7162018101f2fea40d6ebf4a78bbe5a49736a2020649edf9693679e" + +[[package]] +name = "equivalent" +version = "1.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "877a4ace8713b0bcf2a4e7eec82529c029f1d0619886d18145fea96c3ffe5c0f" + +[[package]] +name = "errno" +version = "0.3.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "39cab71617ae0d63f51a36d69f866391735b51691dbda63cf6f96d042b63efeb" +dependencies = [ + "libc", + "windows-sys 0.61.2", +] + +[[package]] +name = "foldhash" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d9c4f5dac5e15c24eb999c26181a6ca40b39fe946cbe4c263c7209467bc83af2" + +[[package]] +name = "generic-array" +version = "0.14.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "85649ca51fd72272d7821adaf274ad91c288277713d9c18820d8499a7ff69e9a" +dependencies = [ + "typenum", + "version_check", +] + +[[package]] +name = "glob" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0cc23270f6e1808e30a928bdc84dea0b9b4136a8bc82338574f23baf47bbd280" + +[[package]] +name = "half" +version = "2.7.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6ea2d84b969582b4b1864a92dc5d27cd2b77b622a8d79306834f1be5ba20d84b" +dependencies = [ + "cfg-if", + "crunchy", + "zerocopy", +] + +[[package]] +name = "hashbrown" +version = "0.15.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9229cfe53dfd69f0609a49f65461bd93001ea1ef889cd5529dd176593f5338a1" +dependencies = [ + "foldhash", +] + +[[package]] +name = "hashbrown" +version = "0.17.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ed5909b6e89a2db4456e54cd5f673791d7eca6732202bbf2a9cc504fe2f9b84a" + +[[package]] +name = "home" +version = "0.5.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cc627f471c528ff0c4a49e1d5e60450c8f6461dd6d10ba9dcd3a61d3dff7728d" +dependencies = [ + "windows-sys 0.61.2", +] + +[[package]] +name = "indexmap" +version = "2.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d466e9454f08e4a911e14806c24e16fba1b4c121d1ea474396f396069cf949d9" +dependencies = [ + "equivalent", + "hashbrown 0.17.1", +] + +[[package]] +name = "itertools" +version = "0.12.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ba291022dbbd398a455acf126c1e341954079855bc60dfdda641363bd6922569" +dependencies = [ + "either", +] + +[[package]] +name = "lazy_static" +version = "1.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bbd2bcb4c963f2ddae06a2efc7e9f3591312473c50c6685e1f298068316e66fe" + +[[package]] +name = "lazycell" +version = "1.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55" + +[[package]] +name = "libc" +version = "0.2.186" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68ab91017fe16c622486840e4c83c9a37afeff978bd239b5293d61ece587de66" + +[[package]] +name = "libloading" +version = "0.8.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d7c4b02199fee7c5d21a5ae7d8cfa79a6ef5bb2fc834d6e9058e89c825efdc55" +dependencies = [ + "cfg-if", + "windows-link", +] + +[[package]] +name = "libnvvm-sys" +version = "0.2.1" +dependencies = [ + "cuda-target-spec", + "libloading", + "thiserror", +] + +[[package]] +name = "linux-raw-sys" +version = "0.4.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d26c52dbd32dccf2d10cac7725f8eae5296885fb5703b261f7d0a0739ec807ab" + +[[package]] +name = "log" +version = "0.4.29" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5e5032e24019045c762d3c0f28f5b6b8bbf38563a65908389bf7978758920897" + +[[package]] +name = "memchr" +version = "2.8.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98" + +[[package]] +name = "minimal-lexical" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a" + +[[package]] +name = "nom" +version = "7.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d273983c5a657a70a3e8f2a01329822f3b8c8172b73826411a55751e404a0a4a" +dependencies = [ + "memchr", + "minimal-lexical", +] + +[[package]] +name = "nvjitlink-sys" +version = "0.2.1" +dependencies = [ + "libloading", + "thiserror", +] + +[[package]] +name = "object" +version = "0.36.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "62948e14d923ea95ea2c7c86c71013138b66525b86bdc08d2dcc262bdb497b87" +dependencies = [ + "crc32fast", + "hashbrown 0.15.5", + "indexmap", + "memchr", +] + +[[package]] +name = "once_cell" +version = "1.21.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9f7c3e4beb33f85d45ae3e3a1792185706c8e16d043238c593331cc7cd313b50" + +[[package]] +name = "oxide-artifacts" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "74d9a941c0c43d2dceaf25f5c9ca4c9f9cceb28dfa5e4e3cfd14e476f003e7c9" +dependencies = [ + "object", +] + +[[package]] +name = "prettyplease" +version = "0.2.37" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "479ca8adacdd7ce8f1fb39ce9ecccbfe93a3f1344b3d0d97f20bc0196208f62b" +dependencies = [ + "proc-macro2", + "syn", +] + +[[package]] +name = "proc-macro2" +version = "1.0.106" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8fd00f0bb2e90d81d1044c2b32617f68fcb9fa3bb7640c23e9c748e53fb30934" +dependencies = [ + "unicode-ident", +] + +[[package]] +name = "ptx-parse" +version = "0.1.0" + +[[package]] +name = "quote" +version = "1.0.45" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "41f2619966050689382d2b44f664f4bc593e129785a36d6ee376ddf37259b924" +dependencies = [ + "proc-macro2", +] + +[[package]] +name = "regex" +version = "1.12.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e10754a14b9137dd7b1e3e5b0493cc9171fdd105e0ab477f51b72e7f3ac0e276" +dependencies = [ + "aho-corasick", + "memchr", + "regex-automata", + "regex-syntax", +] + +[[package]] +name = "regex-automata" +version = "0.4.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6e1dd4122fc1595e8162618945476892eefca7b88c52820e74af6262213cae8f" +dependencies = [ + "aho-corasick", + "memchr", + "regex-syntax", +] + +[[package]] +name = "regex-syntax" +version = "0.8.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4" + +[[package]] +name = "reserved-oxide-symbols" +version = "0.2.1" + +[[package]] +name = "rustc-hash" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2" + +[[package]] +name = "rustix" +version = "0.38.44" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fdb5bc1ae2baa591800df16c9ca78619bf65c0488b41b96ccec5d11220d8c154" +dependencies = [ + "bitflags", + "errno", + "libc", + "linux-raw-sys", + "windows-sys 0.59.0", +] + +[[package]] +name = "serde" +version = "1.0.228" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9a8e94ea7f378bd32cbbd37198a4a91436180c5bb472411e48b5ec2e2124ae9e" +dependencies = [ + "serde_core", + "serde_derive", +] + +[[package]] +name = "serde_core" +version = "1.0.228" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "41d385c7d4ca58e59fc732af25c3983b67ac852c1a25000afe1175de458b67ad" +dependencies = [ + "serde_derive", +] + +[[package]] +name = "serde_derive" +version = "1.0.228" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d540f220d3187173da220f885ab66608367b6574e925011a9353e4badda91d79" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "sha2" +version = "0.10.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a7507d819769d01a365ab707794a4084392c824f54a7a6a7862f8c3d0892b283" +dependencies = [ + "cfg-if", + "cpufeatures", + "digest", +] + +[[package]] +name = "shlex" +version = "1.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0fda2ff0d084019ba4d7c6f371c95d8fd75ce3524c3cb8fb653a3023f6323e64" + +[[package]] +name = "syn" +version = "2.0.117" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e665b8803e7b1d2a727f4023456bbbbe74da67099c585258af0ad9c5013b9b99" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "thiserror" +version = "2.0.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4288b5bcbc7920c07a1149a35cf9590a2aa808e0bc1eafaade0b80947865fbc4" +dependencies = [ + "thiserror-impl", +] + +[[package]] +name = "thiserror-impl" +version = "2.0.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ebc4ee7f67670e9b64d05fa4253e753e016c6c95ff35b89b7941d6b856dec1d5" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "typenum" +version = "1.20.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b6f5e870be6c3b371b77fe0ee0bafb859fa4964b4404c27de1d380043c4dda20" + +[[package]] +name = "unicode-ident" +version = "1.0.24" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75" + +[[package]] +name = "vecadd" +version = "0.1.0" +dependencies = [ + "cuda-core", + "cuda-device", + "cuda-host", +] + +[[package]] +name = "version_check" +version = "0.9.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" + +[[package]] +name = "which" +version = "4.4.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "87ba24419a2078cd2b0f2ede2691b6c66d8e47836da3b6db8265ebad47afbfc7" +dependencies = [ + "either", + "home", + "once_cell", + "rustix", +] + +[[package]] +name = "windows-link" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5" + +[[package]] +name = "windows-sys" +version = "0.59.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1e38bc4d79ed67fd075bcc251a1c39b32a1776bbe92e5bef1f0bf1f8c531853b" +dependencies = [ + "windows-targets", +] + +[[package]] +name = "windows-sys" +version = "0.61.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ae137229bcbd6cdf0f7b80a31df61766145077ddf49416a728b02cb3921ff3fc" +dependencies = [ + "windows-link", +] + +[[package]] +name = "windows-targets" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9b724f72796e036ab90c1021d4780d4d3d648aca59e491e6b98e725b84e99973" +dependencies = [ + "windows_aarch64_gnullvm", + "windows_aarch64_msvc", + "windows_i686_gnu", + "windows_i686_gnullvm", + "windows_i686_msvc", + "windows_x86_64_gnu", + "windows_x86_64_gnullvm", + "windows_x86_64_msvc", +] + +[[package]] +name = "windows_aarch64_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "32a4622180e7a0ec044bb555404c800bc9fd9ec262ec147edd5989ccd0c02cd3" + +[[package]] +name = "windows_aarch64_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "09ec2a7bb152e2252b53fa7803150007879548bc709c039df7627cabbd05d469" + +[[package]] +name = "windows_i686_gnu" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8e9b5ad5ab802e97eb8e295ac6720e509ee4c243f69d781394014ebfe8bbfa0b" + +[[package]] +name = "windows_i686_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0eee52d38c090b3caa76c563b86c3a4bd71ef1a819287c19d586d7334ae8ed66" + +[[package]] +name = "windows_i686_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "240948bc05c5e7c6dabba28bf89d89ffce3e303022809e73deaefe4f6ec56c66" + +[[package]] +name = "windows_x86_64_gnu" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "147a5c80aabfbf0c7d901cb5895d1de30ef2907eb21fbbab29ca94c5b08b1a78" + +[[package]] +name = "windows_x86_64_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "24d5b23dc417412679681396f2b49f3de8c1473deb516bd34410872eff51ed0d" + +[[package]] +name = "windows_x86_64_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "589f6da84c646204747d1270a2a5661ea66ed1cced2631d546fdfb155959f9ec" + +[[package]] +name = "zerocopy" +version = "0.8.52" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ce1022995ff5ff5d841ad7d994facc23098cd40152f2c1d11cd607c6f530653f" +dependencies = [ + "zerocopy-derive", +] + +[[package]] +name = "zerocopy-derive" +version = "0.8.52" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1ae7f38b72ec2a254e2b87ef277cf2cd4fb97cbebf944faa6f33354da0867930" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] diff --git a/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/Cargo.toml b/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/Cargo.toml new file mode 100644 index 0000000000..6ed4df29e2 --- /dev/null +++ b/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/Cargo.toml @@ -0,0 +1,17 @@ +[package] +name = "warp_reduce_redux_bench" +version = "0.1.0" +edition = "2024" +license = "Apache-2.0" + +# Mark as standalone crate (not part of parent workspace) +[workspace] + +# Paired measurement of the two `warp_reduce` forms from one source: +# cargo oxide run warp_reduce_redux_bench --arch sm_80 redux.sync +# cargo oxide run warp_reduce_redux_bench --arch sm_75 shuffle butterfly + +[dependencies] +cuda-device = { path = "../../../cuda-device" } +cuda-host = { path = "../../../cuda-host" } +cuda-core = "0.3.1" diff --git a/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/src/main.rs b/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/src/main.rs new file mode 100644 index 0000000000..c80d9d356b --- /dev/null +++ b/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench/src/main.rs @@ -0,0 +1,367 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +//! Paired measurement of `warp_reduce`'s two forms (#811 / #1305). +//! +//! One source, two builds: the fast path is selected by the +//! `cuda_oxide_sm_at_least` cfg the device build script derives from the target, +//! so there is no run-time switch to hold constant. That cfg is what #1305 +//! adds; on a tree without it both builds take the butterfly and the harness +//! still runs, it just has nothing to compare. +//! +//! ```text +//! cargo oxide run warp_reduce_redux_bench --arch sm_80 redux.sync +//! cargo oxide run warp_reduce_redux_bench --arch sm_75 shuffle butterfly +//! ``` +//! +//! Two timings per build, because they answer different questions: +//! +//! * `*_chain` is a dependency chain — each reduction consumes the previous +//! one — so it measures the latency of one reduction, which is what a +//! serial scan or a scan-like recurrence pays. +//! * `*_independent` runs four unrelated reductions per iteration, so the +//! issue slots and shuffle port have something to overlap. A form that only +//! wins in the chain is a latency win, not a throughput win. +//! +//! `shuffle_*` is the same butterfly `warp_reduce` uses when the fast path does +//! not apply, written out with `WarpCollective::shfl_xor`. Building at sm_80 and +//! comparing `api_chain` against `shuffle_chain` isolates the reduction from the +//! rest of the build, because both walk the same target and the same compiler. +//! +//! Results are per reduction, amortised over a long kernel, so launch and host +//! overhead are not in the number. The first iterations are checked against a +//! host simulation: a kernel the compiler could fold away is not a timing. + +use cuda_device::cooperative_groups::{ + WarpCollective, WarpTile, ops::Sum, this_thread_block, warp_reduce, +}; +use cuda_device::{DisjointSlice, kernel, thread}; +use cuda_host::cuda_module; + +/// Iterations of the timed loop. Long enough that the kernel dominates its own +/// launch; short enough that the whole example stays interactive. +const ITERATIONS: u32 = 40_000; + +/// Iterations used for the correctness pass, where the host can follow along. +const CHECK_ITERATIONS: u32 = 3; + +const SAMPLES: usize = 15; +const WARMUPS: usize = 3; + +// ============================================================================= +// KERNELS +// ============================================================================= +#[cuda_module] +mod kernels { + use super::*; + + /// The shuffle form of an unsigned sum, written out so it can be timed next + /// to the API on the same target. + #[inline(always)] + fn butterfly_sum(warp: &WarpTile<32>, value: u32) -> u32 { + let mut acc = value; + let mut delta: u32 = 16; + while delta > 0 { + acc = acc.wrapping_add(warp.shfl_xor(acc, delta)); + delta >>= 1; + } + acc + } + + #[inline(always)] + fn butterfly_xor(warp: &WarpTile<32>, value: u32) -> u32 { + let mut acc = value; + let mut delta: u32 = 16; + while delta > 0 { + acc ^= warp.shfl_xor(acc, delta); + delta >>= 1; + } + acc + } + + #[inline(always)] + fn butterfly_max(warp: &WarpTile<32>, value: u32) -> u32 { + let mut acc = value; + let mut delta: u32 = 16; + while delta > 0 { + let other = warp.shfl_xor(acc, delta); + if other > acc { + acc = other; + } + delta >>= 1; + } + acc + } + + #[inline(always)] + fn butterfly_min(warp: &WarpTile<32>, value: u32) -> u32 { + let mut acc = value; + let mut delta: u32 = 16; + while delta > 0 { + let other = warp.shfl_xor(acc, delta); + if other < acc { + acc = other; + } + delta >>= 1; + } + acc + } + + /// Full-warp sum through the public API, chained. + #[kernel] + pub fn api_chain(input: &[u32], iterations: u32, mut out: DisjointSlice) { + let warp = this_thread_block().tiled_partition::<32>(); + let i = thread::index_1d().get() as usize; + let mut v = input[i]; + let mut n = 0; + while n < iterations { + v = warp_reduce::(&warp, v); + n += 1; + } + unsafe { + *out.get_unchecked_mut(i) = v; + } + } + + /// The same chain with the butterfly spelled out. + #[kernel] + pub fn shuffle_chain(input: &[u32], iterations: u32, mut out: DisjointSlice) { + let warp = this_thread_block().tiled_partition::<32>(); + let i = thread::index_1d().get() as usize; + let mut v = input[i]; + let mut n = 0; + while n < iterations { + v = butterfly_sum(&warp, v); + n += 1; + } + unsafe { + *out.get_unchecked_mut(i) = v; + } + } + + /// Four unrelated reductions per iteration, so latency can be hidden. + #[kernel] + pub fn api_independent(input: &[u32], iterations: u32, mut out: DisjointSlice) { + let warp = this_thread_block().tiled_partition::<32>(); + let i = thread::index_1d().get() as usize; + let seed = input[i]; + let (mut a, mut b, mut c, mut d) = ( + seed, + seed ^ 0x9e37_79b9, + seed.wrapping_mul(3), + seed.wrapping_add(0x85eb_ca6b), + ); + let mut n = 0; + while n < iterations { + a = warp_reduce::(&warp, a); + b = warp_reduce::(&warp, b); + c = warp_reduce::(&warp, c); + d = warp_reduce::(&warp, d); + n += 1; + } + unsafe { + *out.get_unchecked_mut(i) = a ^ b ^ c ^ d; + } + } + + /// The same four chains over the butterfly. + #[kernel] + pub fn shuffle_independent(input: &[u32], iterations: u32, mut out: DisjointSlice) { + let warp = this_thread_block().tiled_partition::<32>(); + let i = thread::index_1d().get() as usize; + let seed = input[i]; + let (mut a, mut b, mut c, mut d) = ( + seed, + seed ^ 0x9e37_79b9, + seed.wrapping_mul(3), + seed.wrapping_add(0x85eb_ca6b), + ); + let mut n = 0; + while n < iterations { + a = butterfly_sum(&warp, a); + b = butterfly_xor(&warp, b); + c = butterfly_max(&warp, c); + d = butterfly_min(&warp, d); + n += 1; + } + unsafe { + *out.get_unchecked_mut(i) = a ^ b ^ c ^ d; + } + } +} + +// ============================================================================= +// HOST +// ============================================================================= + +const BLOCK: u32 = 128; +const WARPS: u64 = (BLOCK / 32) as u64; + +fn lcg(state: &mut u32) -> u32 { + *state = state.wrapping_mul(1_664_525).wrapping_add(1_013_904_223); + *state +} + +fn input(len: usize) -> Vec { + let mut state = 0x2468_ace0; + (0..len).map(|_| lcg(&mut state)).collect() +} + +/// Follows `api_chain` / `shuffle_chain`: every lane ends an iteration holding +/// its warp's wrapping sum, so the next iteration reduces 32 copies of it. +fn simulate_chain(values: &[u32], iterations: u32) -> Vec { + let mut v = values.to_vec(); + for _ in 0..iterations { + let mut next = v.clone(); + for warp in 0..v.len() / 32 { + let base = warp * 32; + let sum = v[base..base + 32] + .iter() + .fold(0u32, |acc, x| acc.wrapping_add(*x)); + next[base..base + 32].fill(sum); + } + v = next; + } + v +} + +fn median_ms(mut samples: Vec) -> f64 { + samples.sort_by(|a, b| a.partial_cmp(b).unwrap()); + samples[samples.len() / 2] +} + +fn main() { + use cuda_core::simt::LaunchConfig; + use cuda_core::{CudaContext, DeviceBuffer}; + use std::time::Instant; + + let ctx = CudaContext::new(0).expect("Failed to create CUDA context"); + let stream = ctx.default_stream(); + let module = kernels::load(&ctx).expect("Failed to load module"); + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (BLOCK, 1, 1), + shared_mem_bytes: 0, + }; + + let len = BLOCK as usize; + let host_input = input(len); + let input_dev = DeviceBuffer::from_host(&stream, &host_input).unwrap(); + let mut out_dev = DeviceBuffer::::zeroed(&stream, len).unwrap(); + + // ---- correctness: the timed loop, run short enough to simulate ---- + let mut failed = false; + for (name, launch) in [ + ("api_chain", 0u8), + ("shuffle_chain", 1u8), + ("api_independent", 2u8), + ("shuffle_independent", 3u8), + ] { + // SAFETY: launch shape matches the buffers; every lane stays in range. + let result = unsafe { + match launch { + 0 => module.api_chain(stream.as_ref(), cfg, &input_dev, CHECK_ITERATIONS, &mut out_dev), + 1 => module.shuffle_chain(stream.as_ref(), cfg, &input_dev, CHECK_ITERATIONS, &mut out_dev), + 2 => module.api_independent(stream.as_ref(), cfg, &input_dev, CHECK_ITERATIONS, &mut out_dev), + _ => module.shuffle_independent(stream.as_ref(), cfg, &input_dev, CHECK_ITERATIONS, &mut out_dev), + } + }; + result.expect("launch failed"); + let got = out_dev.to_host_vec(&stream).unwrap(); + let paired = if launch < 2 { + let want = simulate_chain(&host_input, CHECK_ITERATIONS); + if got == want { + "matches host".to_string() + } else { + failed = true; + format!("MISMATCH first lane got {:#x} want {:#x}", got[0], want[0]) + } + } else if got.iter().any(|v| *v != 0) { + "non-zero".to_string() + } else { + failed = true; + "ALL ZERO".to_string() + }; + println!("correctness {name}: {paired}"); + } + if failed { + std::process::exit(1); + } + + // ---- timing ---- + // + // Every sample of every arm is taken in one round-robin pass, so a clock or + // power drift over the run lands on all four arms instead of on whichever + // arm happened to be measured last. + let arms: [(&str, u8, u64); 4] = [ + ("api_chain", 0, 1), + ("shuffle_chain", 1, 1), + ("api_independent", 2, 4), + ("shuffle_independent", 3, 4), + ]; + + let mut launch_arm = |arm: u8| { + // SAFETY: every arm's buffers cover its own indices; `out_dev` is + // written by lane `i` only. + unsafe { + match arm { + 0 => module.api_chain(stream.as_ref(), cfg, &input_dev, ITERATIONS, &mut out_dev), + 1 => module.shuffle_chain( + stream.as_ref(), + cfg, + &input_dev, + ITERATIONS, + &mut out_dev, + ), + 2 => module.api_independent( + stream.as_ref(), + cfg, + &input_dev, + ITERATIONS, + &mut out_dev, + ), + _ => module.shuffle_independent( + stream.as_ref(), + cfg, + &input_dev, + ITERATIONS, + &mut out_dev, + ), + } + } + .expect("launch failed"); + }; + + for _ in 0..WARMUPS { + for (_, arm, _) in arms { + launch_arm(arm); + } + } + stream.synchronize().unwrap(); + + let mut samples: Vec> = vec![Vec::with_capacity(SAMPLES); arms.len()]; + for _ in 0..SAMPLES { + for (index, (_, arm, _)) in arms.iter().enumerate() { + let start = Instant::now(); + launch_arm(*arm); + stream.synchronize().unwrap(); + samples[index].push(start.elapsed().as_secs_f64() * 1e3); + } + } + + println!(); + println!("| kernel | median ms | ns per reduction |"); + println!("|---|---|---|"); + for (index, (name, _, factors)) in arms.iter().enumerate() { + let reductions = u64::from(ITERATIONS) * WARPS * factors; + let median = median_ms(samples[index].clone()); + let ns = median * 1e6 / reductions as f64; + println!("| {name} | {median:.3} | {ns:.1} |"); + } + + println!(); + println!("SUCCESS: warp_reduce_redux_bench finished"); +} From fe221e4e1544eb5190bab658e5102bae9ebf0df2 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Sat, 26 Sep 2026 13:39:59 +0800 Subject: [PATCH 2/2] docs(evidence): record the A100 validation and the paired measurement Keeps the measurement next to the change it supports, so the numbers in the description can be checked against the transcript that produced them. Signed-off-by: 0z5a --- evidence/a100-20260926/A2/a2-ptx-sm75.txt | 296 ++++++++++++ evidence/a100-20260926/A2/a2-ptx-sm80.txt | 216 +++++++++ evidence/a100-20260926/A2/a2.log | 135 ++++++ .../a100-20260926/A2/a2bench2-batches.log | 73 +++ evidence/a100-20260926/A2/a2bench2.log | 73 +++ evidence/a100-20260926/A2/bench-ptx-sm80.txt | 454 ++++++++++++++++++ evidence/a100-20260926/A2/conclusion.md | 174 +++++++ 7 files changed, 1421 insertions(+) create mode 100644 evidence/a100-20260926/A2/a2-ptx-sm75.txt create mode 100644 evidence/a100-20260926/A2/a2-ptx-sm80.txt create mode 100644 evidence/a100-20260926/A2/a2.log create mode 100644 evidence/a100-20260926/A2/a2bench2-batches.log create mode 100644 evidence/a100-20260926/A2/a2bench2.log create mode 100644 evidence/a100-20260926/A2/bench-ptx-sm80.txt create mode 100644 evidence/a100-20260926/A2/conclusion.md diff --git a/evidence/a100-20260926/A2/a2-ptx-sm75.txt b/evidence/a100-20260926/A2/a2-ptx-sm75.txt new file mode 100644 index 0000000000..cf192cb5d4 --- /dev/null +++ b/evidence/a100-20260926/A2/a2-ptx-sm75.txt @@ -0,0 +1,296 @@ + Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.08s + Running `target/debug/cargo-oxide inspect warp_reduce_redux --arch sm_75` +Building rustc-codegen-cuda backend... + Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.08s +✓ Backend built: /workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/target/x86_64-unknown-linux-gnu/debug/librustc_codegen_cuda.so +========================================= +RUSTC-CODEGEN-CUDA BUILD: warp_reduce_redux +========================================= + +Building warp_reduce_redux... + + Compiling cuda-device v0.2.1 (/workspace/0z5a/oxide-1305/crates/cuda-device) + Compiling warp_reduce_redux v0.1.0 (/workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/examples/warp_reduce_redux) + Finished `release` profile [optimized] target(s) in 2.00s + +========================================= +PTX (warp_reduce_redux.ptx) +========================================= +// +// Generated by LLVM NVPTX Back-End +// + +.version 6.3 +.target sm_75 +.address_size 64 + + // .globl reduce_i32_signed_ops // -- Begin function reduce_i32_signed_ops + // @reduce_i32_signed_ops +.visible .entry reduce_i32_signed_ops( + .param .u64 .ptr .align 1 reduce_i32_signed_ops_param_0, + .param .u64 reduce_i32_signed_ops_param_1 +) +{ + .reg .pred %p<2>; + .reg .b32 %r<33>; + .reg .b64 %rd<3>; + +// %bb.0: // %entry + ld.param.b64 %rd2, [reduce_i32_signed_ops_param_0]; + cvta.to.global.u64 %rd1, %rd2; + mov.u32 %r4, %laneid; + add.s32 %r5, %r4, -16; + shfl.sync.bfly.b32 %r6, %r5, 16, 31, -1; + add.s32 %r7, %r6, %r5; + shfl.sync.bfly.b32 %r8, %r7, 8, 31, -1; + add.s32 %r9, %r8, %r7; + shfl.sync.bfly.b32 %r10, %r9, 4, 31, -1; + add.s32 %r11, %r10, %r9; + shfl.sync.bfly.b32 %r12, %r11, 2, 31, -1; + add.s32 %r13, %r12, %r11; + shfl.sync.bfly.b32 %r14, %r13, 1, 31, -1; + shfl.sync.bfly.b32 %r15, %r5, 16, 31, -1; + min.s32 %r16, %r5, %r15; + shfl.sync.bfly.b32 %r17, %r16, 8, 31, -1; + min.s32 %r18, %r16, %r17; + shfl.sync.bfly.b32 %r19, %r18, 4, 31, -1; + min.s32 %r20, %r18, %r19; + shfl.sync.bfly.b32 %r21, %r20, 2, 31, -1; + min.s32 %r22, %r20, %r21; + shfl.sync.bfly.b32 %r23, %r22, 1, 31, -1; + shfl.sync.bfly.b32 %r24, %r5, 16, 31, -1; + max.s32 %r25, %r5, %r24; + shfl.sync.bfly.b32 %r26, %r25, 8, 31, -1; + max.s32 %r27, %r25, %r26; + shfl.sync.bfly.b32 %r28, %r27, 4, 31, -1; + max.s32 %r29, %r27, %r28; + shfl.sync.bfly.b32 %r30, %r29, 2, 31, -1; + max.s32 %r31, %r29, %r30; + shfl.sync.bfly.b32 %r32, %r31, 1, 31, -1; + setp.ne.b32 %p1, %r4, 0; + @%p1 bra $L__BB0_2; + bra.uni $L__BB0_1; +$L__BB0_1: // %bb4 + add.s32 %r1, %r14, %r13; + min.s32 %r2, %r22, %r23; + max.s32 %r3, %r31, %r32; + st.global.b32 [%rd1], %r1; + st.global.b32 [%rd1+4], %r2; + st.global.b32 [%rd1+8], %r3; +$L__BB0_2: // %bb5 + ret; + // -- End function +} + // .globl reduce_f32_control // -- Begin function reduce_f32_control +.visible .entry reduce_f32_control( + .param .u64 .ptr .align 1 reduce_f32_control_param_0, + .param .u64 reduce_f32_control_param_1 +) // @reduce_f32_control +{ + .reg .pred %p<12>; + .reg .b32 %r<33>; + .reg .b64 %rd<3>; + +// %bb.0: // %entry + ld.param.b64 %rd2, [reduce_f32_control_param_0]; + cvta.to.global.u64 %rd1, %rd2; + mov.u32 %r4, %laneid; + cvt.rn.f32.u32 %r5, %r4; + shfl.sync.bfly.b32 %r6, %r5, 16, 31, -1; + add.f32 %r7, %r6, %r5; + shfl.sync.bfly.b32 %r8, %r7, 8, 31, -1; + add.f32 %r9, %r7, %r8; + shfl.sync.bfly.b32 %r10, %r9, 4, 31, -1; + add.f32 %r11, %r9, %r10; + shfl.sync.bfly.b32 %r12, %r11, 2, 31, -1; + add.f32 %r13, %r11, %r12; + shfl.sync.bfly.b32 %r14, %r13, 1, 31, -1; + shfl.sync.bfly.b32 %r15, %r5, 16, 31, -1; + setp.gt.f32 %p1, %r15, %r5; + selp.f32 %r16, %r5, %r15, %p1; + shfl.sync.bfly.b32 %r17, %r16, 8, 31, -1; + setp.lt.f32 %p2, %r16, %r17; + selp.f32 %r18, %r16, %r17, %p2; + shfl.sync.bfly.b32 %r19, %r18, 4, 31, -1; + setp.lt.f32 %p3, %r18, %r19; + selp.f32 %r20, %r18, %r19, %p3; + shfl.sync.bfly.b32 %r21, %r20, 2, 31, -1; + setp.lt.f32 %p4, %r20, %r21; + selp.f32 %r22, %r20, %r21, %p4; + shfl.sync.bfly.b32 %r23, %r22, 1, 31, -1; + shfl.sync.bfly.b32 %r24, %r5, 16, 31, -1; + setp.lt.f32 %p5, %r24, %r5; + selp.f32 %r25, %r5, %r24, %p5; + shfl.sync.bfly.b32 %r26, %r25, 8, 31, -1; + setp.gt.f32 %p6, %r25, %r26; + selp.f32 %r27, %r25, %r26, %p6; + shfl.sync.bfly.b32 %r28, %r27, 4, 31, -1; + setp.gt.f32 %p7, %r27, %r28; + selp.f32 %r29, %r27, %r28, %p7; + shfl.sync.bfly.b32 %r30, %r29, 2, 31, -1; + setp.gt.f32 %p8, %r29, %r30; + selp.f32 %r31, %r29, %r30, %p8; + shfl.sync.bfly.b32 %r32, %r31, 1, 31, -1; + setp.ne.b32 %p9, %r4, 0; + @%p9 bra $L__BB1_2; + bra.uni $L__BB1_1; +$L__BB1_1: // %bb4 + add.f32 %r1, %r13, %r14; + setp.lt.f32 %p10, %r22, %r23; + selp.f32 %r2, %r22, %r23, %p10; + setp.gt.f32 %p11, %r31, %r32; + selp.f32 %r3, %r31, %r32, %p11; + st.global.b32 [%rd1], %r1; + st.global.b32 [%rd1+4], %r2; + st.global.b32 [%rd1+8], %r3; +$L__BB1_2: // %bb5 + ret; + // -- End function +} + // .globl reduce_subwarp_control // -- Begin function reduce_subwarp_control +.visible .entry reduce_subwarp_control( + .param .u64 .ptr .align 1 reduce_subwarp_control_param_0, + .param .u64 reduce_subwarp_control_param_1 +) // @reduce_subwarp_control +{ + .reg .pred %p<6>; + .reg .b32 %r<21>; + .reg .b64 %rd<13>; + +// %bb.0: // %entry + ld.param.b64 %rd2, [reduce_subwarp_control_param_0]; + cvta.to.global.u64 %rd1, %rd2; + mov.u32 %r10, %laneid; + and.b32 %r11, %r10, 15; + and.b32 %r12, %r10, 16; + mov.b32 %r13, 65535; + shl.b32 %r14, %r13, %r12; + shfl.sync.bfly.b32 %r15, %r11, 8, 31, %r14; + add.s32 %r16, %r15, %r11; + shfl.sync.bfly.b32 %r17, %r16, 4, 31, %r14; + add.s32 %r18, %r17, %r16; + shfl.sync.bfly.b32 %r19, %r18, 2, 31, %r14; + add.s32 %r1, %r19, %r18; + shfl.sync.bfly.b32 %r2, %r1, 1, 31, %r14; + setp.ne.b32 %p1, %r11, 0; + @%p1 bra $L__BB2_2; + bra.uni $L__BB2_1; +$L__BB2_1: // %bb3 + mov.u32 %r3, %ctaid.x; + mov.u32 %r4, %ntid.x; + mov.u32 %r5, %tid.x; + mov.u32 %r6, %ntid.y; + mov.u32 %r7, %nctaid.y; + mov.u32 %r8, %ntid.z; + mov.u32 %r9, %nctaid.z; + setp.eq.b32 %p2, %r6, 1; + setp.eq.b32 %p3, %r7, 1; + setp.eq.b32 %p4, %r8, 1; + setp.eq.b32 %p5, %r9, 1; + mul.wide.u32 %rd3, %r3, %r4; + cvt.u64.u32 %rd4, %r5; + add.s64 %rd5, %rd3, %rd4; + add.s32 %r20, %r2, %r1; + shr.u64 %rd6, %rd5, 4; + selp.b64 %rd7, %rd6, 1152921504606846975, %p5; + selp.b64 %rd8, %rd7, 1152921504606846975, %p4; + selp.b64 %rd9, %rd8, 1152921504606846975, %p3; + selp.b64 %rd10, %rd9, 1152921504606846975, %p2; + shl.b64 %rd11, %rd10, 2; + add.s64 %rd12, %rd1, %rd11; + st.global.b32 [%rd12], %r20; +$L__BB2_2: // %bb4 + ret; + // -- End function +} + // .globl reduce_u32_all_ops // -- Begin function reduce_u32_all_ops +.visible .entry reduce_u32_all_ops( + .param .u64 .ptr .align 1 reduce_u32_all_ops_param_0, + .param .u64 reduce_u32_all_ops_param_1 +) // @reduce_u32_all_ops +{ + .reg .pred %p<2>; + .reg .b32 %r<65>; + .reg .b64 %rd<3>; + +// %bb.0: // %entry + ld.param.b64 %rd2, [reduce_u32_all_ops_param_0]; + cvta.to.global.u64 %rd1, %rd2; + mov.u32 %r13, %laneid; + shfl.sync.bfly.b32 %r14, %r13, 16, 31, -1; + add.s32 %r15, %r14, %r13; + shfl.sync.bfly.b32 %r16, %r15, 8, 31, -1; + add.s32 %r17, %r16, %r15; + shfl.sync.bfly.b32 %r18, %r17, 4, 31, -1; + add.s32 %r19, %r18, %r17; + shfl.sync.bfly.b32 %r20, %r19, 2, 31, -1; + add.s32 %r1, %r20, %r19; + shfl.sync.bfly.b32 %r2, %r1, 1, 31, -1; + shfl.sync.bfly.b32 %r21, %r13, 16, 31, -1; + min.u32 %r22, %r13, %r21; + shfl.sync.bfly.b32 %r23, %r22, 8, 31, -1; + min.u32 %r24, %r22, %r23; + shfl.sync.bfly.b32 %r25, %r24, 4, 31, -1; + min.u32 %r26, %r24, %r25; + shfl.sync.bfly.b32 %r27, %r26, 2, 31, -1; + min.u32 %r3, %r26, %r27; + shfl.sync.bfly.b32 %r4, %r3, 1, 31, -1; + shfl.sync.bfly.b32 %r28, %r13, 16, 31, -1; + max.u32 %r29, %r13, %r28; + shfl.sync.bfly.b32 %r30, %r29, 8, 31, -1; + max.u32 %r31, %r29, %r30; + shfl.sync.bfly.b32 %r32, %r31, 4, 31, -1; + max.u32 %r33, %r31, %r32; + shfl.sync.bfly.b32 %r34, %r33, 2, 31, -1; + max.u32 %r5, %r33, %r34; + shfl.sync.bfly.b32 %r6, %r5, 1, 31, -1; + mov.b32 %r35, 1; + shl.b32 %r36, %r35, %r13; + not.b32 %r37, %r36; + shfl.sync.bfly.b32 %r38, %r37, 16, 31, -1; + and.b32 %r39, %r38, %r37; + shfl.sync.bfly.b32 %r40, %r39, 8, 31, -1; + and.b32 %r41, %r40, %r39; + shfl.sync.bfly.b32 %r42, %r41, 4, 31, -1; + and.b32 %r43, %r42, %r41; + shfl.sync.bfly.b32 %r44, %r43, 2, 31, -1; + and.b32 %r7, %r44, %r43; + shfl.sync.bfly.b32 %r8, %r7, 1, 31, -1; + shfl.sync.bfly.b32 %r45, %r36, 16, 31, -1; + or.b32 %r46, %r45, %r36; + shfl.sync.bfly.b32 %r47, %r46, 8, 31, -1; + or.b32 %r48, %r47, %r46; + shfl.sync.bfly.b32 %r49, %r48, 4, 31, -1; + or.b32 %r50, %r49, %r48; + shfl.sync.bfly.b32 %r51, %r50, 2, 31, -1; + or.b32 %r9, %r51, %r50; + shfl.sync.bfly.b32 %r10, %r9, 1, 31, -1; + shfl.sync.bfly.b32 %r52, %r36, 16, 31, -1; + xor.b32 %r53, %r52, %r36; + shfl.sync.bfly.b32 %r54, %r53, 8, 31, -1; + xor.b32 %r55, %r54, %r53; + shfl.sync.bfly.b32 %r56, %r55, 4, 31, -1; + xor.b32 %r57, %r56, %r55; + shfl.sync.bfly.b32 %r58, %r57, 2, 31, -1; + xor.b32 %r11, %r58, %r57; + shfl.sync.bfly.b32 %r12, %r11, 1, 31, -1; + setp.ne.b32 %p1, %r13, 0; + @%p1 bra $L__BB3_2; + bra.uni $L__BB3_1; +$L__BB3_1: // %bb7 + add.s32 %r59, %r2, %r1; + min.u32 %r60, %r3, %r4; + max.u32 %r61, %r5, %r6; + and.b32 %r62, %r8, %r7; + or.b32 %r63, %r10, %r9; + xor.b32 %r64, %r12, %r11; + st.global.b32 [%rd1], %r59; + st.global.b32 [%rd1+4], %r60; + st.global.b32 [%rd1+8], %r61; + st.global.b32 [%rd1+12], %r62; + st.global.b32 [%rd1+16], %r63; + st.global.b32 [%rd1+20], %r64; +$L__BB3_2: // %bb8 + ret; + // -- End function +} diff --git a/evidence/a100-20260926/A2/a2-ptx-sm80.txt b/evidence/a100-20260926/A2/a2-ptx-sm80.txt new file mode 100644 index 0000000000..903049d8c5 --- /dev/null +++ b/evidence/a100-20260926/A2/a2-ptx-sm80.txt @@ -0,0 +1,216 @@ + Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.08s + Running `target/debug/cargo-oxide inspect warp_reduce_redux --arch sm_80` +Building rustc-codegen-cuda backend... + Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.08s +✓ Backend built: /workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/target/x86_64-unknown-linux-gnu/debug/librustc_codegen_cuda.so +========================================= +RUSTC-CODEGEN-CUDA BUILD: warp_reduce_redux +========================================= + +Building warp_reduce_redux... + + Compiling warp_reduce_redux v0.1.0 (/workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/examples/warp_reduce_redux) + Finished `release` profile [optimized] target(s) in 1.17s + +========================================= +PTX (warp_reduce_redux.ptx) +========================================= +// +// Generated by LLVM NVPTX Back-End +// + +.version 7.0 +.target sm_80 +.address_size 64 + + // .globl reduce_i32_signed_ops // -- Begin function reduce_i32_signed_ops + // @reduce_i32_signed_ops +.visible .entry reduce_i32_signed_ops( + .param .u64 .ptr .align 1 reduce_i32_signed_ops_param_0, + .param .u64 reduce_i32_signed_ops_param_1 +) +{ + .reg .pred %p<2>; + .reg .b32 %r<7>; + .reg .b64 %rd<3>; + +// %bb.0: // %entry + ld.param.b64 %rd2, [reduce_i32_signed_ops_param_0]; + cvta.to.global.u64 %rd1, %rd2; + mov.u32 %r4, %laneid; + add.s32 %r5, %r4, -16; + mov.b32 %r6, -1; + redux.sync.add.s32 %r1, %r5, %r6; + redux.sync.min.s32 %r2, %r5, %r6; + redux.sync.max.s32 %r3, %r5, %r6; + setp.ne.b32 %p1, %r4, 0; + @%p1 bra $L__BB0_2; + bra.uni $L__BB0_1; +$L__BB0_1: // %bb4 + st.global.b32 [%rd1], %r1; + st.global.b32 [%rd1+4], %r2; + st.global.b32 [%rd1+8], %r3; +$L__BB0_2: // %bb5 + ret; + // -- End function +} + // .globl reduce_f32_control // -- Begin function reduce_f32_control +.visible .entry reduce_f32_control( + .param .u64 .ptr .align 1 reduce_f32_control_param_0, + .param .u64 reduce_f32_control_param_1 +) // @reduce_f32_control +{ + .reg .pred %p<12>; + .reg .b32 %r<33>; + .reg .b64 %rd<3>; + +// %bb.0: // %entry + ld.param.b64 %rd2, [reduce_f32_control_param_0]; + cvta.to.global.u64 %rd1, %rd2; + mov.u32 %r4, %laneid; + cvt.rn.f32.u32 %r5, %r4; + shfl.sync.bfly.b32 %r6, %r5, 16, 31, -1; + add.f32 %r7, %r6, %r5; + shfl.sync.bfly.b32 %r8, %r7, 8, 31, -1; + add.f32 %r9, %r7, %r8; + shfl.sync.bfly.b32 %r10, %r9, 4, 31, -1; + add.f32 %r11, %r9, %r10; + shfl.sync.bfly.b32 %r12, %r11, 2, 31, -1; + add.f32 %r13, %r11, %r12; + shfl.sync.bfly.b32 %r14, %r13, 1, 31, -1; + shfl.sync.bfly.b32 %r15, %r5, 16, 31, -1; + setp.gt.f32 %p1, %r15, %r5; + selp.f32 %r16, %r5, %r15, %p1; + shfl.sync.bfly.b32 %r17, %r16, 8, 31, -1; + setp.lt.f32 %p2, %r16, %r17; + selp.f32 %r18, %r16, %r17, %p2; + shfl.sync.bfly.b32 %r19, %r18, 4, 31, -1; + setp.lt.f32 %p3, %r18, %r19; + selp.f32 %r20, %r18, %r19, %p3; + shfl.sync.bfly.b32 %r21, %r20, 2, 31, -1; + setp.lt.f32 %p4, %r20, %r21; + selp.f32 %r22, %r20, %r21, %p4; + shfl.sync.bfly.b32 %r23, %r22, 1, 31, -1; + shfl.sync.bfly.b32 %r24, %r5, 16, 31, -1; + setp.lt.f32 %p5, %r24, %r5; + selp.f32 %r25, %r5, %r24, %p5; + shfl.sync.bfly.b32 %r26, %r25, 8, 31, -1; + setp.gt.f32 %p6, %r25, %r26; + selp.f32 %r27, %r25, %r26, %p6; + shfl.sync.bfly.b32 %r28, %r27, 4, 31, -1; + setp.gt.f32 %p7, %r27, %r28; + selp.f32 %r29, %r27, %r28, %p7; + shfl.sync.bfly.b32 %r30, %r29, 2, 31, -1; + setp.gt.f32 %p8, %r29, %r30; + selp.f32 %r31, %r29, %r30, %p8; + shfl.sync.bfly.b32 %r32, %r31, 1, 31, -1; + setp.ne.b32 %p9, %r4, 0; + @%p9 bra $L__BB1_2; + bra.uni $L__BB1_1; +$L__BB1_1: // %bb4 + add.f32 %r1, %r13, %r14; + setp.lt.f32 %p10, %r22, %r23; + selp.f32 %r2, %r22, %r23, %p10; + setp.gt.f32 %p11, %r31, %r32; + selp.f32 %r3, %r31, %r32, %p11; + st.global.b32 [%rd1], %r1; + st.global.b32 [%rd1+4], %r2; + st.global.b32 [%rd1+8], %r3; +$L__BB1_2: // %bb5 + ret; + // -- End function +} + // .globl reduce_subwarp_control // -- Begin function reduce_subwarp_control +.visible .entry reduce_subwarp_control( + .param .u64 .ptr .align 1 reduce_subwarp_control_param_0, + .param .u64 reduce_subwarp_control_param_1 +) // @reduce_subwarp_control +{ + .reg .pred %p<6>; + .reg .b32 %r<21>; + .reg .b64 %rd<13>; + +// %bb.0: // %entry + ld.param.b64 %rd2, [reduce_subwarp_control_param_0]; + cvta.to.global.u64 %rd1, %rd2; + mov.u32 %r10, %laneid; + and.b32 %r11, %r10, 15; + and.b32 %r12, %r10, 16; + mov.b32 %r13, 65535; + shl.b32 %r14, %r13, %r12; + shfl.sync.bfly.b32 %r15, %r11, 8, 31, %r14; + add.s32 %r16, %r15, %r11; + shfl.sync.bfly.b32 %r17, %r16, 4, 31, %r14; + add.s32 %r18, %r17, %r16; + shfl.sync.bfly.b32 %r19, %r18, 2, 31, %r14; + add.s32 %r1, %r19, %r18; + shfl.sync.bfly.b32 %r2, %r1, 1, 31, %r14; + setp.ne.b32 %p1, %r11, 0; + @%p1 bra $L__BB2_2; + bra.uni $L__BB2_1; +$L__BB2_1: // %bb3 + mov.u32 %r3, %ctaid.x; + mov.u32 %r4, %ntid.x; + mov.u32 %r5, %tid.x; + mov.u32 %r6, %ntid.y; + mov.u32 %r7, %nctaid.y; + mov.u32 %r8, %ntid.z; + mov.u32 %r9, %nctaid.z; + setp.eq.b32 %p2, %r6, 1; + setp.eq.b32 %p3, %r7, 1; + setp.eq.b32 %p4, %r8, 1; + setp.eq.b32 %p5, %r9, 1; + mul.wide.u32 %rd3, %r3, %r4; + cvt.u64.u32 %rd4, %r5; + add.s64 %rd5, %rd3, %rd4; + add.s32 %r20, %r2, %r1; + shr.u64 %rd6, %rd5, 4; + selp.b64 %rd7, %rd6, 1152921504606846975, %p5; + selp.b64 %rd8, %rd7, 1152921504606846975, %p4; + selp.b64 %rd9, %rd8, 1152921504606846975, %p3; + selp.b64 %rd10, %rd9, 1152921504606846975, %p2; + shl.b64 %rd11, %rd10, 2; + add.s64 %rd12, %rd1, %rd11; + st.global.b32 [%rd12], %r20; +$L__BB2_2: // %bb4 + ret; + // -- End function +} + // .globl reduce_u32_all_ops // -- Begin function reduce_u32_all_ops +.visible .entry reduce_u32_all_ops( + .param .u64 .ptr .align 1 reduce_u32_all_ops_param_0, + .param .u64 reduce_u32_all_ops_param_1 +) // @reduce_u32_all_ops +{ + .reg .pred %p<2>; + .reg .b32 %r<12>; + .reg .b64 %rd<3>; + +// %bb.0: // %entry + ld.param.b64 %rd2, [reduce_u32_all_ops_param_0]; + cvta.to.global.u64 %rd1, %rd2; + mov.u32 %r7, %laneid; + mov.b32 %r8, 1; + shl.b32 %r9, %r8, %r7; + not.b32 %r10, %r9; + mov.b32 %r11, -1; + redux.sync.add.s32 %r1, %r7, %r11; + redux.sync.min.u32 %r2, %r7, %r11; + redux.sync.max.u32 %r3, %r7, %r11; + redux.sync.and.b32 %r4, %r10, %r11; + redux.sync.or.b32 %r5, %r9, %r11; + redux.sync.xor.b32 %r6, %r9, %r11; + setp.ne.b32 %p1, %r7, 0; + @%p1 bra $L__BB3_2; + bra.uni $L__BB3_1; +$L__BB3_1: // %bb7 + st.global.b32 [%rd1], %r1; + st.global.b32 [%rd1+4], %r2; + st.global.b32 [%rd1+8], %r3; + st.global.b32 [%rd1+12], %r4; + st.global.b32 [%rd1+16], %r5; + st.global.b32 [%rd1+20], %r6; +$L__BB3_2: // %bb8 + ret; + // -- End function +} diff --git a/evidence/a100-20260926/A2/a2.log b/evidence/a100-20260926/A2/a2.log new file mode 100644 index 0000000000..58bbfe6086 --- /dev/null +++ b/evidence/a100-20260926/A2/a2.log @@ -0,0 +1,135 @@ +### head +cf4c55c38615d973cebf059211ab0d4624881bda +### list + ], + "schema_version": 1 +} +### run sm_80 (redux fast path) + Compiling cuda-core v0.3.1 + Compiling reserved-oxide-symbols v0.2.1 (/workspace/0z5a/oxide-1305/crates/reserved-oxide-symbols) + Compiling ptx-parse v0.1.0 (/workspace/0z5a/oxide-1305/crates/ptx-parse) + Compiling libloading v0.8.9 + Compiling itertools v0.12.1 + Compiling hashbrown v0.15.5 + Compiling cuda-device v0.2.1 (/workspace/0z5a/oxide-1305/crates/cuda-device) + Compiling generic-array v0.14.7 + Compiling nom v7.1.3 + Compiling clang-sys v1.8.1 + Compiling indexmap v2.14.0 + Compiling regex-automata v0.4.14 + Compiling syn v2.0.117 + Compiling block-buffer v0.10.4 + Compiling crypto-common v0.1.7 + Compiling digest v0.10.7 + Compiling sha2 v0.10.9 + Compiling cexpr v0.6.0 + Compiling which v4.4.2 + Compiling regex v1.12.3 + Compiling oxide-artifacts v0.2.1 + Compiling thiserror-impl v2.0.18 + Compiling zerocopy-derive v0.8.52 + Compiling serde_derive v1.0.228 + Compiling cuda-macros v0.2.1 (/workspace/0z5a/oxide-1305/crates/cuda-macros) + Compiling cuda-core-derive v0.3.1 + Compiling nvjitlink-sys v0.2.1 (/workspace/0z5a/oxide-1305/crates/nvjitlink-sys) + Compiling libnvvm-sys v0.2.1 (/workspace/0z5a/oxide-1305/crates/libnvvm-sys) + Compiling cuda-artifact-finalizer v0.2.1 (/workspace/0z5a/oxide-1305/crates/cuda-artifact-finalizer) + Compiling cuda-bindings v0.3.1 + Compiling half v2.7.1 + Compiling cuda-host v0.2.1 (/workspace/0z5a/oxide-1305/crates/cuda-host) + Compiling warp_reduce_redux v0.1.0 (/workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/examples/warp_reduce_redux) + Finished `release` profile [optimized] target(s) in 21.70s + Running `target/release/warp_reduce_redux` +=== warp_reduce: single-instruction form and its controls === + +GPU Compute Capability: sm_80 + +--- Test 1: warp_reduce over u32, all six ops --- +[sum, min, max, and, or, xor] = [496, 0, 31, 0, 4294967295, 4294967295] +expected = [496, 0, 31, 0, 4294967295, 4294967295] +✓ every u32 reduction correct + +--- Test 2: warp_reduce over i32, values -16..=15 --- +[sum, min, max] = [-16, -16, 15] +expected = [-16, -16, 15] (unsigned min/max would read 0 / 4294967295) +✓ signed reductions correct + +--- Test 3: control, f32 keeps the butterfly --- +[sum, min, max] = [496.0, 0.0, 31.0] +expected = [496.0, 0.0, 31.0] +✓ f32 reductions correct + +--- Test 4: control, WarpTile<16> reduces only its own lanes --- +per-tile sums = [120, 120] +expected = [120, 120] (a whole-warp reduction would read 496) +✓ sub-warp tiles reduce independently + +SUCCESS: warp_reduce is correct in both forms +### run sm_75 build (shuffle path, compiled only) + Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.08s + Running `target/debug/cargo-oxide build warp_reduce_redux --arch sm_75` +Building rustc-codegen-cuda backend... + Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.08s +✓ Backend built: /workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/target/x86_64-unknown-linux-gnu/debug/librustc_codegen_cuda.so +========================================= +RUSTC-CODEGEN-CUDA BUILD: warp_reduce_redux +========================================= + +Building warp_reduce_redux... + + Compiling cuda-device v0.2.1 (/workspace/0z5a/oxide-1305/crates/cuda-device) + Compiling warp_reduce_redux v0.1.0 (/workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/examples/warp_reduce_redux) + Finished `release` profile [optimized] target(s) in 2.02s +### sanitize memcheck +--- Test 1: warp_reduce over u32, all six ops --- +[sum, min, max, and, or, xor] = [496, 0, 31, 0, 4294967295, 4294967295] +expected = [496, 0, 31, 0, 4294967295, 4294967295] +✓ every u32 reduction correct + +--- Test 2: warp_reduce over i32, values -16..=15 --- +[sum, min, max] = [-16, -16, 15] +expected = [-16, -16, 15] (unsigned min/max would read 0 / 4294967295) +✓ signed reductions correct + +--- Test 3: control, f32 keeps the butterfly --- +[sum, min, max] = [496.0, 0.0, 31.0] +expected = [496.0, 0.0, 31.0] +✓ f32 reductions correct + +--- Test 4: control, WarpTile<16> reduces only its own lanes --- +per-tile sums = [120, 120] +expected = [120, 120] (a whole-warp reduction would read 496) +✓ sub-warp tiles reduce independently + +SUCCESS: warp_reduce is correct in both forms +========= ERROR SUMMARY: 0 errors + +Compute Sanitizer completed with exit code 0. +Inspect the sanitizer report above; exit status alone is not a clean-report assertion. +### sanitize synccheck +--- Test 1: warp_reduce over u32, all six ops --- +[sum, min, max, and, or, xor] = [496, 0, 31, 0, 4294967295, 4294967295] +expected = [496, 0, 31, 0, 4294967295, 4294967295] +✓ every u32 reduction correct + +--- Test 2: warp_reduce over i32, values -16..=15 --- +[sum, min, max] = [-16, -16, 15] +expected = [-16, -16, 15] (unsigned min/max would read 0 / 4294967295) +✓ signed reductions correct + +--- Test 3: control, f32 keeps the butterfly --- +[sum, min, max] = [496.0, 0.0, 31.0] +expected = [496.0, 0.0, 31.0] +✓ f32 reductions correct + +--- Test 4: control, WarpTile<16> reduces only its own lanes --- +per-tile sums = [120, 120] +expected = [120, 120] (a whole-warp reduction would read 496) +✓ sub-warp tiles reduce independently + +SUCCESS: warp_reduce is correct in both forms +========= ERROR SUMMARY: 0 errors + +Compute Sanitizer completed with exit code 0. +Inspect the sanitizer report above; exit status alone is not a clean-report assertion. +=== A2_DONE === diff --git a/evidence/a100-20260926/A2/a2bench2-batches.log b/evidence/a100-20260926/A2/a2bench2-batches.log new file mode 100644 index 0000000000..4bd99312f5 --- /dev/null +++ b/evidence/a100-20260926/A2/a2bench2-batches.log @@ -0,0 +1,73 @@ +===== batch 1 : sm_80 (redux) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 1.731 | 10.8 | +| shuffle_chain | 5.599 | 35.0 | +| api_independent | 2.384 | 3.7 | +| shuffle_independent | 7.718 | 12.1 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 1 : sm_75 (butterfly) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 4.350 | 27.2 | +| shuffle_chain | 4.350 | 27.2 | +| api_independent | 5.995 | 9.4 | +| shuffle_independent | 5.996 | 9.4 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 2 : sm_80 (redux) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 1.346 | 8.4 | +| shuffle_chain | 4.349 | 27.2 | +| api_independent | 1.853 | 2.9 | +| shuffle_independent | 5.995 | 9.4 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 2 : sm_75 (butterfly) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 4.349 | 27.2 | +| shuffle_chain | 4.349 | 27.2 | +| api_independent | 5.995 | 9.4 | +| shuffle_independent | 5.995 | 9.4 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 3 : sm_80 (redux) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 1.732 | 10.8 | +| shuffle_chain | 5.599 | 35.0 | +| api_independent | 2.384 | 3.7 | +| shuffle_independent | 7.718 | 12.1 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 3 : sm_75 (butterfly) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 4.350 | 27.2 | +| shuffle_chain | 4.350 | 27.2 | +| api_independent | 5.996 | 9.4 | +| shuffle_independent | 5.996 | 9.4 | +SUCCESS: warp_reduce_redux_bench finished +=== A2BENCH2_DONE === diff --git a/evidence/a100-20260926/A2/a2bench2.log b/evidence/a100-20260926/A2/a2bench2.log new file mode 100644 index 0000000000..4bd99312f5 --- /dev/null +++ b/evidence/a100-20260926/A2/a2bench2.log @@ -0,0 +1,73 @@ +===== batch 1 : sm_80 (redux) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 1.731 | 10.8 | +| shuffle_chain | 5.599 | 35.0 | +| api_independent | 2.384 | 3.7 | +| shuffle_independent | 7.718 | 12.1 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 1 : sm_75 (butterfly) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 4.350 | 27.2 | +| shuffle_chain | 4.350 | 27.2 | +| api_independent | 5.995 | 9.4 | +| shuffle_independent | 5.996 | 9.4 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 2 : sm_80 (redux) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 1.346 | 8.4 | +| shuffle_chain | 4.349 | 27.2 | +| api_independent | 1.853 | 2.9 | +| shuffle_independent | 5.995 | 9.4 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 2 : sm_75 (butterfly) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 4.349 | 27.2 | +| shuffle_chain | 4.349 | 27.2 | +| api_independent | 5.995 | 9.4 | +| shuffle_independent | 5.995 | 9.4 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 3 : sm_80 (redux) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 1.732 | 10.8 | +| shuffle_chain | 5.599 | 35.0 | +| api_independent | 2.384 | 3.7 | +| shuffle_independent | 7.718 | 12.1 | +SUCCESS: warp_reduce_redux_bench finished +===== batch 3 : sm_75 (butterfly) ===== +correctness api_chain: matches host +correctness shuffle_chain: matches host +correctness api_independent: non-zero +correctness shuffle_independent: non-zero +| kernel | median ms | ns per reduction | +|---|---|---| +| api_chain | 4.350 | 27.2 | +| shuffle_chain | 4.350 | 27.2 | +| api_independent | 5.996 | 9.4 | +| shuffle_independent | 5.996 | 9.4 | +SUCCESS: warp_reduce_redux_bench finished +=== A2BENCH2_DONE === diff --git a/evidence/a100-20260926/A2/bench-ptx-sm80.txt b/evidence/a100-20260926/A2/bench-ptx-sm80.txt new file mode 100644 index 0000000000..94367c8449 --- /dev/null +++ b/evidence/a100-20260926/A2/bench-ptx-sm80.txt @@ -0,0 +1,454 @@ + Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.08s + Running `target/debug/cargo-oxide inspect warp_reduce_redux_bench --arch sm_80` +Building rustc-codegen-cuda backend... + Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.08s +✓ Backend built: /workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/target/x86_64-unknown-linux-gnu/debug/librustc_codegen_cuda.so +========================================= +RUSTC-CODEGEN-CUDA BUILD: warp_reduce_redux_bench +========================================= + +Building warp_reduce_redux_bench... + + Compiling cuda-device v0.2.1 (/workspace/0z5a/oxide-1305/crates/cuda-device) + Compiling warp_reduce_redux_bench v0.1.0 (/workspace/0z5a/oxide-1305/crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench) +warning: unused import: `ThreadGroup` + --> src/main.rs:36:5 + | +36 | ThreadGroup, WarpCollective, WarpTile, ops::Sum, this_thread_block, warp_reduce, + | ^^^^^^^^^^^ + | + = note: `#[warn(unused_imports)]` (part of `#[warn(unused)]`) on by default + +warning: `warp_reduce_redux_bench` (bin "warp_reduce_redux_bench") generated 1 warning (run `cargo fix --bin "warp_reduce_redux_bench" -p warp_reduce_redux_bench` to apply 1 suggestion) + Finished `release` profile [optimized] target(s) in 1.97s + +========================================= +PTX (warp_reduce_redux_bench.ptx) +========================================= +// +// Generated by LLVM NVPTX Back-End +// + +.version 7.0 +.target sm_80 +.address_size 64 + + // .globl shuffle_chain // -- Begin function shuffle_chain + // @shuffle_chain +.visible .entry shuffle_chain( + .param .u64 .ptr .align 4 shuffle_chain_param_0, + .param .u64 shuffle_chain_param_1, + .param .u32 shuffle_chain_param_2, + .param .u64 .ptr .align 1 shuffle_chain_param_3, + .param .u64 shuffle_chain_param_4 +) +{ + .reg .pred %p<10>; + .reg .b32 %r<40>; + .reg .b64 %rd<16>; + +// %bb.0: // %entry + ld.param.b64 %rd4, [shuffle_chain_param_0]; + cvta.to.global.u64 %rd1, %rd4; + ld.param.b64 %rd5, [shuffle_chain_param_1]; + ld.param.b64 %rd6, [shuffle_chain_param_3]; + cvta.to.global.u64 %rd2, %rd6; + mov.u32 %r3, %ctaid.x; + mov.u32 %r4, %ntid.x; + mov.u32 %r5, %tid.x; + mul.wide.u32 %rd7, %r3, %r4; + cvt.u64.u32 %rd8, %r5; + add.s64 %rd9, %rd7, %rd8; + mov.u32 %r6, %ntid.y; + mov.u32 %r7, %nctaid.y; + setp.eq.b32 %p1, %r6, 1; + setp.eq.b32 %p2, %r7, 1; + mov.u32 %r8, %ntid.z; + setp.eq.b32 %p3, %r8, 1; + mov.u32 %r9, %nctaid.z; + setp.eq.b32 %p4, %r9, 1; + selp.b64 %rd10, %rd9, -1, %p4; + selp.b64 %rd11, %rd10, -1, %p3; + selp.b64 %rd12, %rd11, -1, %p2; + selp.b64 %rd3, %rd12, -1, %p1; + setp.ge.u64 %p5, %rd3, %rd5; + @%p5 bra $L__BB0_8; + bra.uni $L__BB0_1; +$L__BB0_1: // %bb3 + ld.param.b32 %r2, [shuffle_chain_param_2]; + shl.b64 %rd13, %rd3, 2; + add.s64 %rd14, %rd1, %rd13; + ld.global.b32 %r39, [%rd14]; + setp.eq.b32 %p6, %r2, 0; + @%p6 bra $L__BB0_7; + bra.uni $L__BB0_2; +$L__BB0_2: // %bb7.preheader.preheader + setp.eq.b32 %p7, %r2, 1; + @%p7 bra $L__BB0_6; + bra.uni $L__BB0_3; +$L__BB0_3: // %bb7.preheader.preheader.new + and.b32 %r1, %r2, 1; + and.b32 %r38, %r2, -2; +$L__BB0_4: // %bb7.preheader + // =>This Inner Loop Header: Depth=1 + shfl.sync.bfly.b32 %r10, %r39, 16, 31, -1; + add.s32 %r11, %r10, %r39; + shfl.sync.bfly.b32 %r12, %r11, 8, 31, -1; + add.s32 %r13, %r12, %r11; + shfl.sync.bfly.b32 %r14, %r13, 4, 31, -1; + add.s32 %r15, %r14, %r13; + shfl.sync.bfly.b32 %r16, %r15, 2, 31, -1; + add.s32 %r17, %r16, %r15; + shfl.sync.bfly.b32 %r18, %r17, 1, 31, -1; + add.s32 %r19, %r18, %r17; + shfl.sync.bfly.b32 %r20, %r19, 16, 31, -1; + add.s32 %r21, %r20, %r19; + shfl.sync.bfly.b32 %r22, %r21, 8, 31, -1; + add.s32 %r23, %r22, %r21; + shfl.sync.bfly.b32 %r24, %r23, 4, 31, -1; + add.s32 %r25, %r24, %r23; + shfl.sync.bfly.b32 %r26, %r25, 2, 31, -1; + add.s32 %r27, %r26, %r25; + shfl.sync.bfly.b32 %r28, %r27, 1, 31, -1; + add.s32 %r39, %r28, %r27; + add.s32 %r38, %r38, -2; + setp.ne.b32 %p8, %r38, 0; + @%p8 bra $L__BB0_4; + bra.uni $L__BB0_5; +$L__BB0_5: // %bb6.loopexit.unr-lcssa + setp.eq.b32 %p9, %r1, 0; + @%p9 bra $L__BB0_7; + bra.uni $L__BB0_6; +$L__BB0_6: // %bb7.preheader.epil.preheader + shfl.sync.bfly.b32 %r29, %r39, 16, 31, -1; + add.s32 %r30, %r29, %r39; + shfl.sync.bfly.b32 %r31, %r30, 8, 31, -1; + add.s32 %r32, %r31, %r30; + shfl.sync.bfly.b32 %r33, %r32, 4, 31, -1; + add.s32 %r34, %r33, %r32; + shfl.sync.bfly.b32 %r35, %r34, 2, 31, -1; + add.s32 %r36, %r35, %r34; + shfl.sync.bfly.b32 %r37, %r36, 1, 31, -1; + add.s32 %r39, %r37, %r36; +$L__BB0_7: // %bb6 + add.s64 %rd15, %rd2, %rd13; + st.global.b32 [%rd15], %r39; + ret; +$L__BB0_8: // %bb11 + trap; exit; + // -- End function +} + // .globl api_chain // -- Begin function api_chain +.visible .entry api_chain( + .param .u64 .ptr .align 4 api_chain_param_0, + .param .u64 api_chain_param_1, + .param .u32 api_chain_param_2, + .param .u64 .ptr .align 1 api_chain_param_3, + .param .u64 api_chain_param_4 +) // @api_chain +{ + .reg .pred %p<11>; + .reg .b32 %r<21>; + .reg .b64 %rd<16>; + +// %bb.0: // %entry + ld.param.b64 %rd4, [api_chain_param_0]; + cvta.to.global.u64 %rd1, %rd4; + ld.param.b64 %rd5, [api_chain_param_1]; + ld.param.b64 %rd6, [api_chain_param_3]; + cvta.to.global.u64 %rd2, %rd6; + mov.u32 %r2, %ctaid.x; + mov.u32 %r3, %ntid.x; + mov.u32 %r4, %tid.x; + mul.wide.u32 %rd7, %r2, %r3; + cvt.u64.u32 %rd8, %r4; + add.s64 %rd9, %rd7, %rd8; + mov.u32 %r5, %ntid.y; + mov.u32 %r6, %nctaid.y; + setp.eq.b32 %p1, %r5, 1; + setp.eq.b32 %p2, %r6, 1; + mov.u32 %r7, %ntid.z; + setp.eq.b32 %p3, %r7, 1; + mov.u32 %r8, %nctaid.z; + setp.eq.b32 %p4, %r8, 1; + selp.b64 %rd10, %rd9, -1, %p4; + selp.b64 %rd11, %rd10, -1, %p3; + selp.b64 %rd12, %rd11, -1, %p2; + selp.b64 %rd3, %rd12, -1, %p1; + setp.ge.u64 %p5, %rd3, %rd5; + @%p5 bra $L__BB1_9; + bra.uni $L__BB1_1; +$L__BB1_1: // %bb3 + ld.param.b32 %r1, [api_chain_param_2]; + shl.b64 %rd13, %rd3, 2; + add.s64 %rd14, %rd1, %rd13; + ld.global.b32 %r20, [%rd14]; + setp.eq.b32 %p6, %r1, 0; + @%p6 bra $L__BB1_8; + bra.uni $L__BB1_2; +$L__BB1_2: // %bb5.preheader + and.b32 %r19, %r1, 7; + setp.lt.u32 %p7, %r1, 8; + @%p7 bra $L__BB1_6; + bra.uni $L__BB1_3; +$L__BB1_3: // %bb5.preheader.new + and.b32 %r18, %r1, -8; +$L__BB1_4: // %bb5 + // =>This Inner Loop Header: Depth=1 + mov.b32 %r9, -1; + redux.sync.add.s32 %r10, %r20, %r9; + redux.sync.add.s32 %r11, %r10, %r9; + redux.sync.add.s32 %r12, %r11, %r9; + redux.sync.add.s32 %r13, %r12, %r9; + redux.sync.add.s32 %r14, %r13, %r9; + redux.sync.add.s32 %r15, %r14, %r9; + redux.sync.add.s32 %r16, %r15, %r9; + redux.sync.add.s32 %r20, %r16, %r9; + add.s32 %r18, %r18, -8; + setp.ne.b32 %p8, %r18, 0; + @%p8 bra $L__BB1_4; + bra.uni $L__BB1_5; +$L__BB1_5: // %bb7.loopexit.unr-lcssa + setp.eq.b32 %p9, %r19, 0; + @%p9 bra $L__BB1_8; + bra.uni $L__BB1_6; +$L__BB1_6: // %bb5.epil.preheader +$L__BB1_7: // %bb5.epil + // =>This Inner Loop Header: Depth=1 + .pragma "nounroll"; + mov.b32 %r17, -1; + redux.sync.add.s32 %r20, %r20, %r17; + add.s32 %r19, %r19, -1; + setp.ne.b32 %p10, %r19, 0; + @%p10 bra $L__BB1_7; + bra.uni $L__BB1_8; +$L__BB1_8: // %bb7 + add.s64 %rd15, %rd2, %rd13; + st.global.b32 [%rd15], %r20; + ret; +$L__BB1_9: // %bb8 + trap; exit; + // -- End function +} + // .globl api_independent // -- Begin function api_independent +.visible .entry api_independent( + .param .u64 .ptr .align 4 api_independent_param_0, + .param .u64 api_independent_param_1, + .param .u32 api_independent_param_2, + .param .u64 .ptr .align 1 api_independent_param_3, + .param .u64 api_independent_param_4 +) // @api_independent +{ + .reg .pred %p<11>; + .reg .b32 %r<32>; + .reg .b64 %rd<16>; + +// %bb.0: // %entry + ld.param.b64 %rd4, [api_independent_param_0]; + cvta.to.global.u64 %rd1, %rd4; + ld.param.b64 %rd5, [api_independent_param_1]; + ld.param.b64 %rd6, [api_independent_param_3]; + cvta.to.global.u64 %rd2, %rd6; + mov.u32 %r2, %ctaid.x; + mov.u32 %r3, %ntid.x; + mov.u32 %r4, %tid.x; + mul.wide.u32 %rd7, %r2, %r3; + cvt.u64.u32 %rd8, %r4; + add.s64 %rd9, %rd7, %rd8; + mov.u32 %r5, %ntid.y; + mov.u32 %r6, %nctaid.y; + setp.eq.b32 %p1, %r5, 1; + setp.eq.b32 %p2, %r6, 1; + mov.u32 %r7, %ntid.z; + setp.eq.b32 %p3, %r7, 1; + mov.u32 %r8, %nctaid.z; + setp.eq.b32 %p4, %r8, 1; + selp.b64 %rd10, %rd9, -1, %p4; + selp.b64 %rd11, %rd10, -1, %p3; + selp.b64 %rd12, %rd11, -1, %p2; + selp.b64 %rd3, %rd12, -1, %p1; + setp.ge.u64 %p5, %rd3, %rd5; + @%p5 bra $L__BB2_10; + bra.uni $L__BB2_1; +$L__BB2_1: // %bb3 + ld.param.b32 %r1, [api_independent_param_2]; + shl.b64 %rd13, %rd3, 2; + add.s64 %rd14, %rd1, %rd13; + ld.global.b32 %r29, [%rd14]; + mul.lo.s32 %r27, %r29, 3; + add.s32 %r26, %r29, -2048144789; + setp.eq.b32 %p6, %r1, 0; + mov.b32 %r31, -1640531527; + @%p6 bra $L__BB2_9; + bra.uni $L__BB2_2; +$L__BB2_2: // %bb5.preheader + xor.b32 %r30, %r29, -1640531527; + and.b32 %r28, %r1, 3; + setp.lt.u32 %p7, %r1, 4; + @%p7 bra $L__BB2_6; + bra.uni $L__BB2_3; +$L__BB2_3: // %bb5.preheader.new + and.b32 %r25, %r1, -4; +$L__BB2_4: // %bb5 + // =>This Inner Loop Header: Depth=1 + mov.b32 %r9, -1; + redux.sync.add.s32 %r10, %r29, %r9; + redux.sync.xor.b32 %r11, %r30, %r9; + redux.sync.max.u32 %r12, %r27, %r9; + redux.sync.min.u32 %r13, %r26, %r9; + redux.sync.add.s32 %r14, %r10, %r9; + redux.sync.xor.b32 %r15, %r11, %r9; + redux.sync.max.u32 %r16, %r12, %r9; + redux.sync.min.u32 %r17, %r13, %r9; + redux.sync.add.s32 %r18, %r14, %r9; + redux.sync.xor.b32 %r19, %r15, %r9; + redux.sync.max.u32 %r20, %r16, %r9; + redux.sync.min.u32 %r21, %r17, %r9; + redux.sync.add.s32 %r29, %r18, %r9; + redux.sync.xor.b32 %r30, %r19, %r9; + redux.sync.max.u32 %r27, %r20, %r9; + redux.sync.min.u32 %r26, %r21, %r9; + add.s32 %r25, %r25, -4; + setp.ne.b32 %p8, %r25, 0; + @%p8 bra $L__BB2_4; + bra.uni $L__BB2_5; +$L__BB2_5: // %bb10.loopexit.unr-lcssa + setp.eq.b32 %p9, %r28, 0; + @%p9 bra $L__BB2_8; + bra.uni $L__BB2_6; +$L__BB2_6: // %bb5.epil.preheader +$L__BB2_7: // %bb5.epil + // =>This Inner Loop Header: Depth=1 + .pragma "nounroll"; + mov.b32 %r22, -1; + redux.sync.add.s32 %r29, %r29, %r22; + redux.sync.xor.b32 %r30, %r30, %r22; + redux.sync.max.u32 %r27, %r27, %r22; + redux.sync.min.u32 %r26, %r26, %r22; + add.s32 %r28, %r28, -1; + setp.ne.b32 %p10, %r28, 0; + @%p10 bra $L__BB2_7; + bra.uni $L__BB2_8; +$L__BB2_8: // %bb10.loopexit + xor.b32 %r31, %r30, %r29; +$L__BB2_9: // %bb10 + xor.b32 %r23, %r31, %r27; + add.s64 %rd15, %rd2, %rd13; + xor.b32 %r24, %r23, %r26; + st.global.b32 [%rd15], %r24; + ret; +$L__BB2_10: // %bb11 + trap; exit; + // -- End function +} + // .globl shuffle_independent // -- Begin function shuffle_independent +.visible .entry shuffle_independent( + .param .u64 .ptr .align 4 shuffle_independent_param_0, + .param .u64 shuffle_independent_param_1, + .param .u32 shuffle_independent_param_2, + .param .u64 .ptr .align 1 shuffle_independent_param_3, + .param .u64 shuffle_independent_param_4 +) // @shuffle_independent +{ + .reg .pred %p<8>; + .reg .b32 %r<52>; + .reg .b64 %rd<16>; + +// %bb.0: // %entry + ld.param.b64 %rd4, [shuffle_independent_param_0]; + cvta.to.global.u64 %rd1, %rd4; + ld.param.b64 %rd5, [shuffle_independent_param_1]; + ld.param.b64 %rd6, [shuffle_independent_param_3]; + cvta.to.global.u64 %rd2, %rd6; + mov.u32 %r1, %ctaid.x; + mov.u32 %r2, %ntid.x; + mov.u32 %r3, %tid.x; + mul.wide.u32 %rd7, %r1, %r2; + cvt.u64.u32 %rd8, %r3; + add.s64 %rd9, %rd7, %rd8; + mov.u32 %r4, %ntid.y; + mov.u32 %r5, %nctaid.y; + setp.eq.b32 %p1, %r4, 1; + setp.eq.b32 %p2, %r5, 1; + mov.u32 %r6, %ntid.z; + setp.eq.b32 %p3, %r6, 1; + mov.u32 %r7, %nctaid.z; + setp.eq.b32 %p4, %r7, 1; + selp.b64 %rd10, %rd9, -1, %p4; + selp.b64 %rd11, %rd10, -1, %p3; + selp.b64 %rd12, %rd11, -1, %p2; + selp.b64 %rd3, %rd12, -1, %p1; + setp.ge.u64 %p5, %rd3, %rd5; + @%p5 bra $L__BB3_6; + bra.uni $L__BB3_1; +$L__BB3_1: // %bb3 + ld.param.b32 %r46, [shuffle_independent_param_2]; + shl.b64 %rd13, %rd3, 2; + add.s64 %rd14, %rd1, %rd13; + ld.global.b32 %r50, [%rd14]; + mul.lo.s32 %r48, %r50, 3; + add.s32 %r47, %r50, -2048144789; + setp.eq.b32 %p6, %r46, 0; + mov.b32 %r51, -1640531527; + @%p6 bra $L__BB3_5; + bra.uni $L__BB3_2; +$L__BB3_2: // %bb7.preheader.preheader + xor.b32 %r49, %r50, -1640531527; +$L__BB3_3: // %bb7.preheader + // =>This Inner Loop Header: Depth=1 + shfl.sync.bfly.b32 %r8, %r50, 16, 31, -1; + add.s32 %r9, %r8, %r50; + shfl.sync.bfly.b32 %r10, %r9, 8, 31, -1; + add.s32 %r11, %r10, %r9; + shfl.sync.bfly.b32 %r12, %r11, 4, 31, -1; + add.s32 %r13, %r12, %r11; + shfl.sync.bfly.b32 %r14, %r13, 2, 31, -1; + add.s32 %r15, %r14, %r13; + shfl.sync.bfly.b32 %r16, %r15, 1, 31, -1; + shfl.sync.bfly.b32 %r17, %r49, 16, 31, -1; + xor.b32 %r18, %r17, %r49; + shfl.sync.bfly.b32 %r19, %r18, 8, 31, -1; + xor.b32 %r20, %r19, %r18; + shfl.sync.bfly.b32 %r21, %r20, 4, 31, -1; + xor.b32 %r22, %r21, %r20; + shfl.sync.bfly.b32 %r23, %r22, 2, 31, -1; + xor.b32 %r24, %r23, %r22; + shfl.sync.bfly.b32 %r25, %r24, 1, 31, -1; + shfl.sync.bfly.b32 %r26, %r48, 16, 31, -1; + max.u32 %r27, %r26, %r48; + shfl.sync.bfly.b32 %r28, %r27, 8, 31, -1; + max.u32 %r29, %r28, %r27; + shfl.sync.bfly.b32 %r30, %r29, 4, 31, -1; + max.u32 %r31, %r30, %r29; + shfl.sync.bfly.b32 %r32, %r31, 2, 31, -1; + max.u32 %r33, %r32, %r31; + shfl.sync.bfly.b32 %r34, %r33, 1, 31, -1; + shfl.sync.bfly.b32 %r35, %r47, 16, 31, -1; + min.u32 %r36, %r35, %r47; + shfl.sync.bfly.b32 %r37, %r36, 8, 31, -1; + min.u32 %r38, %r37, %r36; + shfl.sync.bfly.b32 %r39, %r38, 4, 31, -1; + min.u32 %r40, %r39, %r38; + shfl.sync.bfly.b32 %r41, %r40, 2, 31, -1; + min.u32 %r42, %r41, %r40; + shfl.sync.bfly.b32 %r43, %r42, 1, 31, -1; + min.u32 %r47, %r43, %r42; + max.u32 %r48, %r34, %r33; + xor.b32 %r49, %r25, %r24; + add.s32 %r50, %r16, %r15; + add.s32 %r46, %r46, -1; + setp.ne.b32 %p7, %r46, 0; + @%p7 bra $L__BB3_3; + bra.uni $L__BB3_4; +$L__BB3_4: // %bb6.loopexit + xor.b32 %r51, %r49, %r50; +$L__BB3_5: // %bb6 + xor.b32 %r44, %r51, %r48; + add.s64 %rd15, %rd2, %rd13; + xor.b32 %r45, %r44, %r47; + st.global.b32 [%rd15], %r45; + ret; +$L__BB3_6: // %bb29 + trap; exit; + // -- End function +} diff --git a/evidence/a100-20260926/A2/conclusion.md b/evidence/a100-20260926/A2/conclusion.md new file mode 100644 index 0000000000..c67e3e696e --- /dev/null +++ b/evidence/a100-20260926/A2/conclusion.md @@ -0,0 +1,174 @@ +# A2 · #811 / #1305 — native A100 validation + +**Status: `#1305` head validated on real A100/SM80. Correctness, sanitizers, +target gating and a paired measurement are all green. No new implementation +was written; the vendored implementation and the #1304 dependency are unchanged +and remain with their author.** + +| item | value | +|---|---| +| PR | [#1305](https://github.com/NVlabs/cuda-oxide/pull/1305) `feat(cuda-device): derive an SM floor cfg and use redux.sync in warp_reduce` | +| validated HEAD | `cf4c55c38615d973cebf059211ab0d4624881bda` | +| dependency | [#1304](https://github.com/NVlabs/cuda-oxide/pull/1304) `6b4ac7d0147129c8a9c7ab86185e61b6b616c3fd` | +| merge base | `b9847e9515ed3a23096f22567d3eaf0a6e3e440c` | +| stack | 5 commits, `b9847e95..cf4c55c3` | +| hardware | 1× `NVIDIA A100-SXM4-40GB`, cc 8.0, driver `595.91.07`, no MIG | +| toolchain | rustc `nightly-2026-08-28`, `llc` LLVM `23.1.0-rust-1.100.0-nightly`, CUDA `13.0.88`, `libNVVM 2.0` | + +## Scope audit — what was already covered + +The author's body reports 1,395 CPU tests at `96ea0b56`, strict Clippy and fmt, +and B200 GPU runs at `db5c23f8` for `sm_80` and `sm_75`, including memcheck and +synccheck on the `sm_80` run. It also states plainly: *"GPU tests were not +repeated after the scanner rebase; no new speedup measurement was made."* + +| area | author | this run | +|---|---|---| +| CPU selector / scanner / cfg regressions | covered | not repeated | +| device oracle, 9 integer pairs + f32 + sub-warp controls | A10G/B200 | **A100, this report** | +| memcheck / synccheck | B200 | **A100, this report** | +| `sm_75` compile + PTX fallback | B200 | **A100 host, this report** | +| paired measurement of the two forms | not made | **A100, this report** | +| malformed-hint regression | covered | not repeated | + +Nothing here re-implements the architecture gating. The one artefact added is a +measurement harness (below); it does not change device behaviour. + +## Correctness on A100 + +`crates/rustc-codegen-cuda/examples/warp_reduce_redux` (the PR's own oracle, +unmodified): + +``` +cargo oxide run warp_reduce_redux --arch sm_80 -> SUCCESS +cargo oxide sanitize warp_reduce_redux --arch sm_80 --tool memcheck -> 0 errors +cargo oxide sanitize warp_reduce_redux --arch sm_80 --tool synccheck -> 0 errors +cargo oxide build warp_reduce_redux --arch sm_75 -> ok +``` + +| case | expected | observed | +|---|---|---| +| `u32` full warp: sum / min / max / and / or / xor | `[496, 0, 31, 0, 0xffffffff, 0xffffffff]` | equal | +| `i32` full warp: sum / min / max over `-16..=15` | `[-16, -16, 15]` | equal | +| control `f32` full warp: sum / min / max | `[496.0, 0.0, 31.0]` | equal | +| control `WarpTile<16>`: per-tile sums | `[120, 120]` | equal | + +The `i32` case is the one that matters for a single-instruction form: unsigned +min/max over the same bits would read `0 / 4294967295`, and the oracle rejects +that. The `f32` case is the negative control — there is no float `redux` below +`sm_100`, so a hook that claimed one would name an instruction that cannot +assemble. `WarpTile<16>` is the sub-warp control: a broken `N == 32` gate would +make both tiles answer 496. + +## Emitted code, per target + +`cargo oxide inspect warp_reduce_redux --arch sm_80` and `--arch sm_75`: + +| kernel | sm_80 | sm_75 | +|---|---|---| +| `u32` / `i32` full warp | `redux.sync.*` — the reduction is one instruction | 5× `shfl.sync.bfly.b32` + 5× `add.s32` | +| `f32` full warp | `shfl.sync.bfly.b32` + `add.f32` (fallback) | same | +| `WarpTile<16>` | `shfl.sync.bfly.b32` pairs (fallback) | same | + +`sfl`/`shfl` counts in the two module dumps are 19 vs 64, i.e. the `sm_80` +module keeps shuffles only where the fast path does not apply. This is the +"expected path vs observed path" check: the fast path is not merely requested by +a flag, it is what the emitted module contains. + +For the measurement harness the same check was done per kernel, which is the +sharper form: + +```text +.visible .entry shuffle_chain(...) # 5.599 ms arm + shfl.sync.bfly.b32 %r10, %r39, 16, 31, -1; + add.s32 %r11, %r10, %r39; + ... five rounds, no redux in this kernel ... +.visible .entry api_chain(...) # 1.731 ms arm + redux.sync.add.s32 %r10, %r20, %r9; + redux.sync.add.s32 %r11, %r10, %r9; + ... no shfl in this kernel ... +``` + +## Paired measurement + +Harness: `crates/rustc-codegen-cuda/examples/warp_reduce_redux_bench` (new, +A100-only evidence; not part of the PR's device code). One source, two builds — +the fast path is chosen by the compile-time cfg, so there is no run-time switch +to hold constant. + +* `api_chain` — `warp_reduce::` fed its own result: a dependency + chain, so it exposes reduction latency. +* `shuffle_chain` — the same butterfly, written out with + `WarpCollective::shfl_xor`. At `sm_80` this isolates the reduction from the + rest of the build: same target, same compiler, same binary. +* `*_independent` — four unrelated reductions per iteration, so issue slots and + the shuffle port have something to overlap. + +128 threads (4 warps), 40,000 iterations, 15 samples per arm, 3 warm-up rounds, +all four arms sampled **round-robin within one pass** so clock drift lands on +every arm. Values are amortised per reduction over a multi-millisecond kernel, +so launch and host overhead are not in the number. The first iterations are +compared against a host simulation. + +Each cell is `median ms / ns per reduction`; three independent batches: + +| build | batch | api_chain | shuffle_chain | api_independent | shuffle_independent | +|---|---|---|---|---|---| +| `sm_80` (redux) | 1 | 1.731 / 10.8 | 5.599 / 35.0 | 2.384 / 3.7 | 7.718 / 12.1 | +| `sm_80` (redux) | 2 | 1.346 / 8.4 | 4.349 / 27.2 | 1.853 / 2.9 | 5.995 / 9.4 | +| `sm_80` (redux) | 3 | 1.732 / 10.8 | 5.599 / 35.0 | 2.384 / 3.7 | 7.718 / 12.1 | +| `sm_75` (butterfly) | 1 | 4.350 / 27.2 | 4.350 / 27.2 | 5.995 / 9.4 | 5.996 / 9.4 | +| `sm_75` (butterfly) | 2 | 4.349 / 27.2 | 4.349 / 27.2 | 5.995 / 9.4 | 5.995 / 9.4 | +| `sm_75` (butterfly) | 3 | 4.350 / 27.2 | 4.350 / 27.2 | 5.996 / 9.4 | 5.996 / 9.4 | + +**Speedup, `sm_80`, paired inside one binary and one interleaved pass:** + +| workload | redux (ns/reduction) | shuffle (ns/reduction) | ratio | +|---|---|---|---| +| dependent chain | 8.4 – 10.8 | 27.2 – 35.0 | **3.24×** (all three batches) | +| four independent reductions | 2.9 – 3.7 | 9.4 – 12.1 | **3.24 – 3.27×** | + +**Cross-build, the PR's own framing** (`--arch sm_80` against `--arch sm_75`): +`api_chain` 8.4–10.8 ns against 27.2 ns, i.e. **2.5×–3.2×**. + +The control is exact: at `sm_75` the API path and the hand-written butterfly +agree to the printed digit on all six cells (27.2 = 27.2, 9.4 = 9.4), which is +what a build where `HAS_FULL_WARP_FORM` is false must do. That is the evidence +that the harness measures the reduction and not something else. + +This is a primitive measurement. The reduction is 1 instruction against 10, so +the ratio is bounded by Amdahl at the call site: a kernel that reduces once +after a memory-bound pass sees far less, and nothing here is an end-to-end +model or serving number. + +## Residual issues with #1304 that this run did not close + +The PR body lists two known scanner gaps: valid whitespace/comment-separated +opcode modifiers are still missed, and a quoted `.file` path is still treated as +an instruction. These are false negatives on a *float* `redux` in hand-written +module text, so the failure direction is "an sm_100a-only float redux could be +mistaken for an integer one" rather than a false positive on integer redux. They +live in the #1304 dependency and are not re-fixed here; the integer +false-positive case the PR was filed for is covered by the PR's own regression +and was not re-tested by this run. + +## Evidence index + +| file | content | +|---|---| +| `a2.log` | PR oracle: run `sm_80`, build `sm_75`, memcheck, synccheck | +| `a2bench2-full.log` | three interleaved batches, both builds | +| `bench-ptx-sm80.txt` | per-kernel PTX for the harness at `sm_80` | +| `a2-ptx-sm80.txt`, `a2-ptx-sm75.txt` | module dumps for the PR oracle | + +## Limitations + +- One A100 SKU, one MIG mode (`none`), one driver. Not A10G, RTX 5090 or B200. +- The measurement is a microbenchmark of the primitive. Registers, occupancy and + SASS were **not** captured: `cargo oxide inspect` stops at PTX, the driver JITs + the PTX at load, and no `cuobjdump`-visible cubin was extracted. This run + therefore reports an executed-PTX path, not executed SASS. +- The `sm_75` build is compile/PTX coverage on an A100 host. It is not a Turing + GPU run and is not reported as one. +- No CI claim: `just check`, the full example matrix and the libNVVM route were + not re-run for this head.