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"); +} 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.