Integrate verified v1.3 CPU CUDA contact fix and snowmobile center of mass

.gitattributes

@@ -1,2 +1,3 @@
 bin/* filter=lfs diff=lfs merge=lfs -text
 native/selective-freeze-source.tar.gz filter=lfs diff=lfs merge=lfs -text
+native/physics-source.tar.gz filter=lfs diff=lfs merge=lfs -text

BINARY_PROVENANCE.json

@@ -51,5 +51,19 @@
   "source_base_revision": "b90b8af06ffc4ea4e4fd703becd1107a48b8e5e0",
   "source_repository": "https://github.com/Zorro909/driving_sim_rs",
   "source_snapshot_path": "native/selective-freeze-source.tar.gz",
-  "source_snapshot_sha256": "297d2ce34569994acd2a477e64c5315cb2081699b4333b8abdabb3ff7451d8db"
+  "source_snapshot_sha256": "297d2ce34569994acd2a477e64c5315cb2081699b4333b8abdabb3ff7451d8db",
+  "physics_update": {
+    "upstream_version": "1.3.0",
+    "upstream_revision": "ab3ec67",
+    "local_additions_preserved": true,
+    "source_receipt": "native/physics-source-identities.json",
+    "status": "candidate; requires deployment and hosted app acceptance",
+    "artifacts": {
+      "bin/altd-sim": "792139f30bcfecf25f2d78595d2611234d30b0a254b5ece0ec587f1d27a80df2",
+      "bin/libaltd_gpu.so": "e5e9a61d863bc3b429af69a13cd8d50f63cad8984b3e6a9001ae342fc5981e05"
+    },
+    "source_archive": "native/physics-source.tar.gz",
+    "source_archive_sha256": "6abefb234de18dee58d69758c6cc8900f259993663ac37f90ce39aa114a30718",
+    "cuda_original_cases_passed": 57
+  }
 }

RELEASE_MANIFEST.json

@@ -3,7 +3,7 @@
   "files": {
     ".gitattributes": "3aa7a5a6b603279d7a0826d015985988192b67e3cedb38a59785ef43c04fc77b",
     "BETA_TESTING.md": "829dae67ec37dac2a4e7b178ab5ee4fc5f63896af132fbff457a5e35141fe94d",
-    "BINARY_PROVENANCE.json": "e05f4c3a4a2a0da39b213708c532f3e514109c5530250b402347a8d751105052",
+    "BINARY_PROVENANCE.json": "99c24677a039993b6100ce24551a01753c0187980f6c9c54be865b3dc7a65d00",
     "BUCKET_BUNDLE.md": "cb7a0f023d5a1f93f0288c6a5646b64f8f3e9d36150c0730aea94590e5b80260",
     "CUDA_REFERENCE.md": "f7453c3254d244f8b0de41fde5851cb025159f5cea4a4d1a5d8208a955550c6c",
     "DATA_POLICY.md": "fc5978528abe1c21c14882417ac8a706e04a0831bbd3fdb3de8c64442c1a7b3d",
@@ -22,15 +22,14 @@
     "app_release.py": "7482559c667a665fd3cb6aa8af6c78422a822b515e14589db8306f7fd4860a8b",
     "assets/formula_template.json": "80f745df82bb7be45af52785fd582460bcddc63c5a4e1d8375d85b3122db4c6b",
     "assets/rally_template.json": "7350a3fe28d9b4da7a0965ae6ea7ca3472fb7885303e342310d63a0dce80851c",
-    "assets/snowmobile_template.json": "12d30545876011fd6fccdd3e721f9c465938c9c58ac99338205481d5e7ac6742",
+    "assets/snowmobile_template.json": "8232f4d54829445c44de9292e12557bfe66b3d43e0363e07d2ae0b4118721c08",
     "assets/track-atlas.json": "d97c8900e3e5fb17734b5bd61d1f4f4b196d9be30af9848e2e273ae7e4079ae4",
     "assets/truck_template.json": "a83bffea4d5e818c135ab3aba68a47aa288d6ad0e7b8bd9c95ed0dea9e9df2f7",
-    "bin/altd-sim": "ce31f754f6b6f3afaf4dcc44c1e6fda3e6414eb2924fa167b8778523bcaf2070",
+    "bin/altd-sim": "792139f30bcfecf25f2d78595d2611234d30b0a254b5ece0ec587f1d27a80df2",
     "bin/extract_saved_track": "3a7733c45b4301a786fcbfe92660a1c484749272da5f608d62f34878bd4ff31a",
-    "bin/libaltd_gpu.so": "3b1b88bccc06a21acc33ac4400af016d55a7e3d8727136a0c43b541f3c7887ca",
+    "bin/libaltd_gpu.so": "e5e9a61d863bc3b429af69a13cd8d50f63cad8984b3e6a9001ae342fc5981e05",
     "bucket_storage.py": "f6132ba5f95d106a667742ec092f2799de00e4baf1bf33600da4f0daaa46012a",
     "build_info.py": "2a57f73abd9e062cf029b65b09b26b8d12e301bac2f7efec7edb4b85b47671c3",
-    "builtin_tracks.py": "52cb4a285a6079cec7f7610de8af80570f1644a02cf9979e47d2d49736e24cb5",
     "builtin_tracks/catalog.json": "4a7bba4385713e985d1427a0da35dcc32d5ecfd73d93cad7520563deb8685670",
     "builtin_tracks/tracks/Replica A01 Beginner.track": "ca5749537450430794e2eade2e0ac72fa141a4f225e13e2f8ef842abb948ff15",
     "builtin_tracks/tracks/Replica A02 Sharp Turns.track": "097f1ea8edf041112b690d0638075630716ddc03c7d0b1980ddcbeaf5071c8d4",
@@ -60,6 +59,7 @@
     "builtin_tracks/tracks/Replica R08 Demolition Derby.track": "6f11a9ffbcb477c9d19964566b1b612d2a94675605246e93e21f27bbbdb357ef",
     "builtin_tracks/tracks/Replica R09 Eau Noir.track": "3ec0a78478cca1a4f0d2d98ac81f40783928adf271b70f0c32d295a6da2cdaa9",
     "builtin_tracks/tracks/Replica The Champions Trial.track": "a3bfb76422affc774b44c739b0636d08acb901fa7615b9adcb52c069f79834e1",
+    "builtin_tracks.py": "52cb4a285a6079cec7f7610de8af80570f1644a02cf9979e47d2d49736e24cb5",
     "client.js": "cbb1dd26a8fe4256abf827903df94749dc8e4f99912d29dba7f74ba6e477b851",
     "delivery.py": "6bdba2e87c3e19a41440cb16103695af27105f8bafbae38948bbb1b936ed0c1b",
     "gpu_tickets.py": "20928fcf1ca24648809ad8712c929377bdcc8b5e6b4387bfff3503ec5711928b",
@@ -69,9 +69,14 @@
     "index.html": "7bdaf7ca11a455525c22432298e120ca2fa5257283b9815467e005e09da07a26",
     "native/analysis.patch": "42f79de67cb0734cffb2a7582f5479cab42ab82675884b8b42b9f158d46a98e8",
     "native/build-receipt.json": "559a7e4e382f9f7caa1c5ceb566a5a0bbefad2ed3a2e2cf604e31c9c3890c348",
+    "native/physics-parity.json": "006cc6b455e8f49af395ed60f3c5c88dbb2865ee6cad51d8cfdae040293d9f9e",
+    "native/physics-source-identities.json": "e3414e24144d517f4a0309e1f9d0dbab8309b4c0b70bd61ccb2c0f40b8f14998",
+    "native/physics-source.tar.gz": "6abefb234de18dee58d69758c6cc8900f259993663ac37f90ce39aa114a30718",
+    "native/physics.patch": "bae8eeb9342cad1e52ad77dc8026b8bd2ff510a1b836165fc753c043a04859b3",
     "native/selective-freeze-build-receipt.json": "ed804b1cd49a9ba34abc6f7cadaf4fb97807e27f3d14af117f8292a2d3f2d172",
     "native/selective-freeze-source.tar.gz": "297d2ce34569994acd2a477e64c5315cb2081699b4333b8abdabb3ff7451d8db",
     "native/source-receipt.json": "36feba861501baf8672884476bbeed5e375254cbb4feffb9f4f90aa61a415a93",
+    "native/upstream-v1.3.patch": "a9810c9abd40ddcbac5fcf6914320d5fc04fcde23849df9437ec293c83be15ed",
     "portable_bundle.py": "58df5b87956745ea909c41c4b12a23cf9ab8c40a7df47d1529f47ae4de9bf51f",
     "requirements.txt": "fc7e0703ad2289fd3926a36efbdce8c61f0edc00f7575a87e2b565e46588c5a4",
     "server.py": "1f82c9046324661d7bd4acf1cf37eef3445f985351159d88acdcaf89d1edf859",
@@ -89,5 +94,6 @@
     "worker.py": "d18053363baf8147425808e9fd939f44f00af23781662d353196ade974144974"
   },
   "interface": "single-engine-runner-with-map-winners-analysis-v1",
-  "upstream_revision": "b90b8af06ffc4ea4e4fd703becd1107a48b8e5e0"
+  "upstream_revision": "b90b8af06ffc4ea4e4fd703becd1107a48b8e5e0",
+  "physics_revision": "20261017-v1.3-candidate"
 }

assets/snowmobile_template.json

@@ -98,6 +98,10 @@
         "brake_power": 1.5,
         "handbrake_power": 0
       }
+    ],
+    "center_of_mass": [
+      -0.00011444099800428376,
+      0.49998602271080017
     ]
   },
   "physics": {

bin/altd-sim

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:ce31f754f6b6f3afaf4dcc44c1e6fda3e6414eb2924fa167b8778523bcaf2070
-size 3647808
+oid sha256:792139f30bcfecf25f2d78595d2611234d30b0a254b5ece0ec587f1d27a80df2
+size 3646480

bin/libaltd_gpu.so

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:3b1b88bccc06a21acc33ac4400af016d55a7e3d8727136a0c43b541f3c7887ca
-size 4711704
+oid sha256:e5e9a61d863bc3b429af69a13cd8d50f63cad8984b3e6a9001ae342fc5981e05
+size 4707512

native/physics-parity.json

@@ -0,0 +1,292 @@
+{
+  "cases": 57,
+  "passed": 57,
+  "library_sha256": "e5e9a61d863bc3b429af69a13cd8d50f63cad8984b3e6a9001ae342fc5981e05",
+  "results": [
+    {
+      "case": "B01-formula-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-formula-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-formula-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-formula-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-formula-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-rally-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-rally-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-rally-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-rally-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-rally-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-truck-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-truck-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-truck-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-truck-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-truck-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-snowmobile-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-snowmobile-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-snowmobile-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "B01-snowmobile-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-formula-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-formula-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-formula-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-formula-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-formula-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-rally-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-rally-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-rally-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-rally-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-rally-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-truck-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-truck-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-truck-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-truck-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-truck-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-snowmobile-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-snowmobile-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-snowmobile-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A04-snowmobile-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-formula-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-formula-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-formula-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-formula-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-formula-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-rally-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-rally-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-rally-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-rally-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-rally-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-truck-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-truck-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-truck-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-truck-turn-boost",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-truck-turn-handbrake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-snowmobile-straight",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-snowmobile-turn",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-snowmobile-turn-brake",
+      "passed": true,
+      "differences": []
+    },
+    {
+      "case": "A06-snowmobile-turn-boost",
+      "passed": true,
+      "differences": []
+    }
+  ]
+}

native/physics-source-identities.json

@@ -0,0 +1,50 @@
+{
+  "local/Cargo.toml": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/web-dashboard/reviews/native-upstream-20261014/Cargo.toml",
+    "sha256": "f9bf44c197883bc105a7d9ac4248c8d9ef450fa99b6b6d04924f7612aaff8875"
+  },
+  "local/Cargo.lock": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/web-dashboard/reviews/native-upstream-20261014/Cargo.lock",
+    "sha256": "3486e24f81b254c6e9d7fdbb2be8d645df69340b96f2401b6d4d8dcd9fa3ba85"
+  },
+  "local/src/physics/collision.rs": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/web-dashboard/reviews/native-upstream-20261014/src/physics/collision.rs",
+    "sha256": "c5226d6e10637725fa374b184e258a89932c239186f88ae69f4815a538ff4b82"
+  },
+  "local/gpu/sim/step.h": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/web-dashboard/reviews/native-upstream-20261014/gpu/sim/step.h",
+    "sha256": "9cd8847b53ad289dfc94288db7630537a5b5c4abf6a0a7561bf2430dcaacd1d0"
+  },
+  "local/gpu/sim/altd_gpu.hip": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/web-dashboard/reviews/native-upstream-20261014/gpu/sim/altd_gpu.hip",
+    "sha256": "cd4e95792ab5f70ec1a698fc327576abac4e9b19322fc37136f48099d56e0055"
+  },
+  "local/examples/parity_controls.rs": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/web-dashboard/reviews/native-upstream-20261014/examples/parity_controls.rs",
+    "sha256": "28e23df53c9db9e636893eb1dd1bb400f1344d68b100880c222e5a75bd561332"
+  },
+  "hf/Cargo.toml": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/hf-winners-evaluation-v1.2-20261013/build/source/Cargo.toml",
+    "sha256": "f9bf44c197883bc105a7d9ac4248c8d9ef450fa99b6b6d04924f7612aaff8875"
+  },
+  "hf/Cargo.lock": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/hf-winners-evaluation-v1.2-20261013/build/source/Cargo.lock",
+    "sha256": "3486e24f81b254c6e9d7fdbb2be8d645df69340b96f2401b6d4d8dcd9fa3ba85"
+  },
+  "hf/src/physics/collision.rs": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/hf-winners-evaluation-v1.2-20261013/build/source/src/physics/collision.rs",
+    "sha256": "443b9700229f6973e2d22163f8ef31c90aefe674a21ec0a3fd5bd91186cb4c14"
+  },
+  "hf/gpu/sim/step.h": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/hf-winners-evaluation-v1.2-20261013/build/source/gpu/sim/step.h",
+    "sha256": "dfbe41cc3891e94ba7ea15d5277661431fd78be08ee16dc2b300c79571ccf6ee"
+  },
+  "hf/gpu/sim/altd_gpu.hip": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/hf-winners-evaluation-v1.2-20261013/build/source/gpu/sim/altd_gpu.hip",
+    "sha256": "ce7a97e0b2f43df6f3b9d1cb70e284a8242a7386088ba0a72ef0b4bbe32a42ae"
+  },
+  "hf/examples/parity_controls.rs": {
+    "path": "/Volumes/PATRIOT/ALD-TAO-lab/hf-winners-evaluation-v1.2-20261013/build/source/examples/parity_controls.rs",
+    "sha256": "28e23df53c9db9e636893eb1dd1bb400f1344d68b100880c222e5a75bd561332"
+  }
+}

native/physics-source.tar.gz

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:6abefb234de18dee58d69758c6cc8900f259993663ac37f90ce39aa114a30718
+size 917770

native/physics.patch

@@ -0,0 +1,59 @@
+--- a/src/physics/collision.rs
++++ b/src/physics/collision.rs
+@@ -329,6 +329,7 @@
+     contacts: &mut Vec<Contact>,
+ ) {
+     contacts.clear();
++    scratch.collided_shapes.clear();
+     let (sx, sy) = cfg.shape_basis;
+     let basis_x = V2::from(body_basis.0 * sx.x + body_basis.1 * sx.y);
+     let basis_y = V2::from(body_basis.0 * sy.x + body_basis.1 * sy.y);
+@@ -483,6 +484,8 @@
+         if separated || push == V2::ZERO {
+             continue;
+         }
++        // SAT overlap keeps cached contacts eligible even without a new contact.
++        scratch.collided_shapes.push(shape_index);
+         let support_car = rectangle_supports(&car, -push, basis_x, basis_y);
+         let support_wall = polygon_supports(wall, shape, push);
+         let mut pairs = [(V2::ZERO, V2::ZERO); 2];
+@@ -601,11 +604,7 @@
+     }
+     find_contacts_with_basis(cfg, track, position, rotation, body_basis, scratch, contacts);
+     previous.retain(|c| scratch.broad_pairs.contains(&c.shape_index));
+-    scratch.collided_shapes.clear();
+     for new in contacts.iter() {
+-        if !scratch.collided_shapes.contains(&new.shape_index) {
+-            scratch.collided_shapes.push(new.shape_index);
+-        }
+         let radius2 = (cfg.contact_recycle_radius as f32).powi(2);
+         if let Some(old) = previous.iter_mut().find(|old| {
+             let a = F2::from(old.local_point) - F2::from(new.local_point);
+--- a/gpu/sim/step.h
++++ b/gpu/sim/step.h
+@@ -471,7 +471,8 @@
+ }
+ 
+ // The rest of `solve_wall_contacts` once this tick's contacts are merged;
+-// `made_contacts` has a bit per broad pair that made contacts.
++// `made_contacts` has a bit per SAT-overlapping broad pair, including
++// overlaps with no new points but a valid retained contact.
+ __device__ inline SolveResult solve_contacts(const World& w, const VehicleDesc& v, Car& car, uint32_t made_contacts, Vec2 velocity,
+                                              float angular, Vec2 initial_velocity, float initial_angular) {
+     STEP_MARK(ts);
+@@ -682,12 +683,13 @@
+             f = car_frame(v, car);
+             framed = true;
+         }
+-        uint32_t made = pair_contacts(w, v, car, f, ci, push, [&](const Contact& c) {
++        pair_contacts(w, v, car, f, ci, push, [&](const Contact& c) {
+             STEP_MARK(tm);
+             merge_contact(v, car, c);
+             STEP_SECTION(11, tm);
+         });
+-        made_contacts |= made ? 1u << ci : 0u;
++        // Keep cached contacts eligible when SAT overlaps but produces no new point.
++        made_contacts |= 1u << ci;
+     }
+     STEP_SECTION(2, tc);
+     SolveResult r = solve_contacts(w, v, car, made_contacts, k.velocity, k.angular, k.previous_velocity, k.previous_angular);

native/upstream-v1.3.patch

@@ -0,0 +1,466 @@
+diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml
+index 2c0e811..a7540f4 100644
+--- a/.github/workflows/release.yml
++++ b/.github/workflows/release.yml
+@@ -30,6 +30,8 @@ concurrency:
+ env:
+   CARGO_TERM_COLOR: always
+   CARGO_INCREMENTAL: 0
++  # Turing, Ampere, Ada and Blackwell consumer GPUs, as in ci.yml; newer GPUs JIT the sm_120 PTX.
++  CUDA_ARCH: sm_75,sm_86,sm_89,sm_120
+ 
+ jobs:
+   version:
+@@ -75,7 +77,9 @@ jobs:
+           echo "publish=$publish" >> "$GITHUB_OUTPUT"
+           echo "Building $version (publish: $publish)" >> "$GITHUB_STEP_SUMMARY"
+ 
+-  gpu-library:
++  # Each GPU library job uploads gpu-<system>-<vendor>, which the package job
++  # puts beside the binary in the matching altd-sim-*-<vendor> archive.
++  gpu-linux-amd:
+     runs-on: ubuntu-latest
+     timeout-minutes: 30
+     container: rocm/dev-ubuntu-24.04:7.1.1
+@@ -91,13 +95,117 @@ jobs:
+           sort gpu/release-targets | diff -u - embedded
+       - uses: actions/upload-artifact@v4
+         with:
+-          name: gpu-library
++          name: gpu-linux-amd
+           path: dist/libaltd_gpu.so
+           if-no-files-found: error
+ 
++  gpu-linux-nvidia:
++    runs-on: ubuntu-latest
++    timeout-minutes: 30
++    # Keep the CUDA version in step with ci.yml. Ubuntu 22.04 matches the
++    # package job, so the library needs no newer glibc than the binary.
++    container: nvidia/cuda:13.4.2-devel-ubuntu22.04
++    steps:
++      - name: Install the tools the CUDA image lacks
++        run: apt-get update && apt-get install -y --no-install-recommends python3
++      - uses: actions/checkout@v4
++      - name: Build for every consumer GPU target
++        run: OUT=dist gpu/build-cuda.sh
++      - uses: actions/upload-artifact@v4
++        with:
++          name: gpu-linux-nvidia
++          path: dist/libaltd_gpu.so
++          if-no-files-found: error
++
++  gpu-windows-amd:
++    # The HIP SDK 7.1 clang headers clash with the MSVC 14.5x <cmath> of Visual
++    # Studio 2026 (windows-latest): its isgreater and friends are already
++    # __host__ __device__. Visual Studio 2022 builds.
++    runs-on: windows-2022
++    timeout-minutes: 60
++    env:
++      RUSTFLAGS: -C target-cpu=x86-64
++      HIP_SDK: C:\hip-sdk
++    steps:
++      - run: git config --global core.autocrlf false
++      - uses: actions/checkout@v4
++      - name: Install the HIP SDK
++        # HIP SDK for Windows 7.1.1, in step with the ROCm 7.1.1 Linux builds. The
++        # bundle also installs a driver, so only its SDK core package is installed.
++        run: |
++          $installer = Join-Path $env:RUNNER_TEMP "hip-sdk.exe"
++          curl.exe -fsSL -o $installer https://download.amd.com/developer/eula/rocm-hub/AMD-Software-PRO-Edition-26.Q1-Win11-For-HIP.exe
++          if ($LASTEXITCODE -ne 0) { throw "cannot download the HIP SDK" }
++          $unpacked = Join-Path $env:RUNNER_TEMP "hip-sdk"
++          7z x $installer "-o$unpacked" | Out-Null
++          if ($LASTEXITCODE -ne 0) { throw "cannot unpack the HIP SDK" }
++          $msi = Get-ChildItem $unpacked -Recurse -Filter ROCm_SDK_Core.msi | Select-Object -First 1
++          if (-not $msi) { throw "the HIP SDK bundle has no ROCm_SDK_Core.msi" }
++          $log = Join-Path $env:RUNNER_TEMP "hip-sdk.log"
++          $install = Start-Process msiexec -Wait -PassThru -ArgumentList "/i", "`"$($msi.FullName)`"", "INSTALLDIR=$env:HIP_SDK", "/qn", "/norestart", "/l*v", "`"$log`""
++          if ($install.ExitCode -ne 0) { Get-Content $log -Tail 50; throw "HIP SDK install failed with $($install.ExitCode)" }
++          $hipcc = Get-ChildItem $env:HIP_SDK -Recurse -Filter hipcc.exe | Select-Object -First 1
++          if (-not $hipcc) { throw "no hipcc.exe under $env:HIP_SDK" }
++          & $hipcc.FullName --version
++          "HIP_PATH=$(Split-Path $hipcc.DirectoryName)" >> $env:GITHUB_ENV
++          # gpu-info below loads the HIP runtime (amdhip64_7.dll) from here.
++          $hipcc.DirectoryName >> $env:GITHUB_PATH
++      - name: Build for every consumer GPU target
++        run: ./gpu/build.ps1 -Out dist -Arch (Get-Content gpu/release-targets -Raw)
++      - uses: dtolnay/rust-toolchain@1.94.0
++      - uses: Swatinem/rust-cache@v2
++        with:
++          key: gpu-windows-x86-64
++      - name: Check the library against the Rust layouts
++        # No GPU is needed: gpu-info checks the device struct layouts and the exports on the host.
++        run: cargo run --locked --bin altd-sim -- gpu-info --library dist/altd_gpu.dll
++      - uses: actions/upload-artifact@v4
++        with:
++          name: gpu-windows-amd
++          path: dist/altd_gpu.dll
++          if-no-files-found: error
++
++  gpu-windows-nvidia:
++    runs-on: windows-latest
++    timeout-minutes: 30
++    env:
++      CUDA_VERSION: 13.4.2
++    steps:
++      - run: git config --global core.autocrlf false
++      - uses: actions/checkout@v4
++      - name: Install the CUDA compiler
++        # Only the redistributable parts nvcc needs, not the whole toolkit installer.
++        run: |
++          $base = "https://developer.download.nvidia.com/compute/cuda/redist"
++          $manifest = Invoke-RestMethod "$base/redistrib_$env:CUDA_VERSION.json"
++          $parts = Join-Path $env:RUNNER_TEMP "cuda-parts"
++          $cuda = Join-Path $env:RUNNER_TEMP "cuda"
++          foreach ($name in "cuda_nvcc", "cuda_crt", "cuda_cudart", "libnvvm", "cccl") {
++            $zip = Join-Path $env:RUNNER_TEMP "$name.zip"
++            curl.exe -fsSL -o $zip "$base/$($manifest.$name.'windows-x86_64'.relative_path)"
++            if ($LASTEXITCODE -ne 0) { throw "cannot download $name" }
++            7z x $zip "-o$parts" | Out-Null
++            if ($LASTEXITCODE -ne 0) { throw "cannot unpack $name" }
++          }
++          # Each archive unpacks into its own directory; merge them into one toolkit.
++          foreach ($dir in Get-ChildItem $parts -Directory) {
++            robocopy $dir.FullName $cuda /E /NFL /NDL /NJH /NJS /NP | Out-Null
++            if ($LASTEXITCODE -ge 8) { throw "cannot merge $($dir.Name)" }
++          }
++          "CUDA_PATH=$cuda" >> $env:GITHUB_ENV
++          "$cuda\bin" >> $env:GITHUB_PATH
++          & "$cuda\bin\nvcc.exe" --version
++      - name: Build for every consumer GPU target
++        run: ./gpu/build-cuda.ps1 -Out dist
++      - uses: actions/upload-artifact@v4
++        with:
++          name: gpu-windows-nvidia
++          path: dist/altd_gpu.dll
++          if-no-files-found: error
++
+   package:
+     name: package (${{ matrix.target }})
+-    needs: [version, gpu-library]
++    needs: [version, gpu-linux-amd, gpu-linux-nvidia, gpu-windows-amd, gpu-windows-nvidia]
+     runs-on: ${{ matrix.os }}
+     timeout-minutes: 45
+     strategy:
+@@ -108,18 +216,29 @@ jobs:
+           - os: ubuntu-22.04
+             target: x86_64-unknown-linux-gnu
+             rustflags: -C target-cpu=x86-64
+-            gpu: true
++            system: linux
++            library: libaltd_gpu.so
++            variants: cpu amd nvidia
+           - os: windows-latest
+             target: x86_64-pc-windows-msvc
+             rustflags: -C target-cpu=x86-64
+             exe: .exe
++            system: windows
++            library: altd_gpu.dll
++            variants: cpu amd nvidia
+           - os: macos-latest
+             target: aarch64-apple-darwin
++            system: macos
++            variants: cpu
+     env:
+       # Overrides target-cpu=native from .cargo/config.toml on x86 hosts.
+       RUSTFLAGS: ${{ matrix.rustflags }}
+       ALTD_BUILD_VERSION: ${{ needs.version.outputs.version }}
++      # Each variant becomes the archive $PACKAGE-<variant>: cpu has no GPU
++      # library, amd and nvidia carry that vendor's library beside the binary.
+       PACKAGE: altd-sim-${{ needs.version.outputs.label }}-${{ matrix.target }}
++      VARIANTS: ${{ matrix.variants }}
++      LIBRARY: ${{ matrix.library }}
+     defaults:
+       run:
+         shell: bash
+@@ -134,48 +253,72 @@ jobs:
+         with:
+           key: release-${{ matrix.target }}
+       - run: cargo build --locked --release --bin altd-sim --target ${{ matrix.target }}
+-      - name: Assemble the package
+-        run: |
+-          mkdir -p "dist/$PACKAGE"
+-          cp "target/${{ matrix.target }}/release/altd-sim${{ matrix.exe }}" README.md LICENSE GODOT_LICENSE "dist/$PACKAGE/"
+-          cp -R assets docs "dist/$PACKAGE/"
+-          mkdir -p "dist/$PACKAGE/math"
+-          cp math/README.md "dist/$PACKAGE/math/"
+-      - if: matrix.gpu
++      - if: matrix.library
+         uses: actions/download-artifact@v4
+         with:
+-          name: gpu-library
+-          path: dist/${{ env.PACKAGE }}
++          pattern: gpu-${{ matrix.system }}-*
++          path: gpu-libraries
++      - name: Assemble the packages
++        run: |
++          for variant in $VARIANTS; do
++            name="$PACKAGE-$variant"
++            mkdir -p "dist/$name/math"
++            cp "target/${{ matrix.target }}/release/altd-sim${{ matrix.exe }}" README.md LICENSE GODOT_LICENSE "dist/$name/"
++            cp -R assets docs "dist/$name/"
++            cp math/README.md "dist/$name/math/"
++            if [ "$variant" != cpu ]; then
++              cp "gpu-libraries/gpu-${{ matrix.system }}-$variant/$LIBRARY" "dist/$name/"
++            fi
++          done
+       - name: Archive
+         working-directory: dist
+         run: |
+-          if [ "$RUNNER_OS" = Windows ]; then
+-            7z a -tzip "$PACKAGE.zip" "$PACKAGE" > /dev/null
+-          else
+-            tar -czf "$PACKAGE.tar.gz" "$PACKAGE"
+-          fi
+-          rm -rf "$PACKAGE"
+-      - name: Smoke test the unpacked archive
++          for variant in $VARIANTS; do
++            name="$PACKAGE-$variant"
++            if [ "$RUNNER_OS" = Windows ]; then
++              7z a -tzip "$name.zip" "$name" > /dev/null
++            else
++              tar -czf "$name.tar.gz" "$name"
++            fi
++            rm -rf "$name"
++          done
++      - name: Smoke test the unpacked archives
+         run: |
+-          mkdir smoke
+-          cd smoke
+-          if [ "$RUNNER_OS" = Windows ]; then 7z x "../dist/$PACKAGE.zip" > /dev/null; else tar -xzf "../dist/$PACKAGE.tar.gz"; fi
+-          cd "$PACKAGE"
+-          ./altd-sim --version | tee version.txt
+-          grep -qxF "altd-sim $ALTD_BUILD_VERSION" version.txt
+-          ./altd-sim --threads 2 train-scratch \
+-            --scene assets/scenes/rally_template.json --track-mode random \
+-            --random-track-settings assets/random_track_settings.json \
+-            --population 16 --generations 1 --ticks 60 --seed 1 --out-dir smoke-run
+-          test -s smoke-run/run.json
+-          if [ "${{ matrix.gpu }}" = true ]; then
+-            # Hosted runners lack ROCm: the packaged library must be found and
+-            # its missing HIP runtime reported. gpu-verify loads it with ROCm.
+-            if ./altd-sim gpu-info 2> gpu-info.txt; then exit 1; fi
+-            cat gpu-info.txt
+-            grep -q "$PWD/libaltd_gpu.so" gpu-info.txt
+-            grep -q "ROCm 7.x runtime" gpu-info.txt
+-          fi
++          for variant in $VARIANTS; do
++            name="$PACKAGE-$variant"
++            mkdir -p "smoke/$variant"
++            cd "smoke/$variant"
++            if [ "$RUNNER_OS" = Windows ]; then 7z x "../../dist/$name.zip" > /dev/null; else tar -xzf "../../dist/$name.tar.gz"; fi
++            cd "$name"
++            ./altd-sim --version | tee version.txt
++            grep -qxF "altd-sim $ALTD_BUILD_VERSION" version.txt
++            ./altd-sim --threads 2 train-scratch \
++              --scene assets/scenes/rally_template.json --track-mode random \
++              --random-track-settings assets/random_track_settings.json \
++              --population 16 --generations 1 --ticks 60 --seed 1 --out-dir smoke-run
++            test -s smoke-run/run.json
++            # The path the binary reports for the library beside it.
++            if [ "$RUNNER_OS" = Windows ]; then library=$(cygpath -w "$PWD/$LIBRARY"); else library="$PWD/$LIBRARY"; fi
++            case "$variant" in
++              amd)
++                # Hosted runners lack the HIP runtime: the packaged library must be
++                # found and fail to load. gpu-verify and gpu-windows-amd load it with HIP.
++                if ./altd-sim gpu-info 2> gpu-info.txt; then exit 1; fi
++                cat gpu-info.txt
++                grep -qF "$library" gpu-info.txt
++                if [ "$RUNNER_OS" = Linux ]; then grep -q "ROCm 7.x runtime" gpu-info.txt; fi
++                ;;
++              nvidia)
++                # The CUDA runtime is linked in and finds no driver, so the library
++                # loads and gpu-info checks its layouts without a GPU.
++                ./altd-sim gpu-info | tee gpu-info.json
++                grep -q '"layout": "ok"' gpu-info.json
++                grep -q '"platform": "CUDA"' gpu-info.json
++                grep -qF "\"${library//\\/\\\\}\"" gpu-info.json
++                ;;
++            esac
++            cd "$GITHUB_WORKSPACE"
++          done
+       - uses: actions/upload-artifact@v4
+         with:
+           name: ${{ env.PACKAGE }}
+@@ -195,8 +338,8 @@ jobs:
+           name: ${{ env.PACKAGE }}
+       - name: Load the packaged HIP library with the ROCm runtime
+         run: |
+-          tar -xzf "$PACKAGE.tar.gz"
+-          cd "$PACKAGE"
++          tar -xzf "$PACKAGE-amd.tar.gz"
++          cd "$PACKAGE-amd"
+           # No GPU is needed: gpu-info checks the device struct layouts on the host.
+           ./altd-sim gpu-info | tee gpu-info.json
+           grep -q '"layout": "ok"' gpu-info.json
+@@ -242,3 +385,7 @@ jobs:
+             gh release create nightly --prerelease --latest=false --verify-tag --title "$title" --notes "$notes"
+           fi
+           gh release upload nightly dist/* --clobber
++          # Drop assets this build no longer produces, such as renamed archives.
++          gh release view nightly --json assets --jq '.assets[].name' | while read -r asset; do
++            if [ ! -e "dist/$asset" ]; then gh release delete-asset nightly "$asset" --yes; fi
++          done
+diff --git a/Cargo.lock b/Cargo.lock
+index 26d9448..dd79258 100644
+--- a/Cargo.lock
++++ b/Cargo.lock
+@@ -4,7 +4,7 @@ version = 4
+ 
+ [[package]]
+ name = "altd_sim"
+-version = "1.2.0"
++version = "1.3.0"
+ dependencies = [
+  "clap",
+  "console_error_panic_hook",
+diff --git a/Cargo.toml b/Cargo.toml
+index 5d466cd..2fbcad8 100644
+--- a/Cargo.toml
++++ b/Cargo.toml
+@@ -1,6 +1,6 @@
+ [package]
+ name = "altd_sim"
+-version = "1.2.0"
++version = "1.3.0"
+ edition = "2021"
+ description = "CPU game-fidelity simulator and trainer for AI Learns To Drive"
+ license = "MIT"
+diff --git a/README.md b/README.md
+index 9e9cdc3..d8384d3 100644
+--- a/README.md
++++ b/README.md
+@@ -16,15 +16,19 @@ Download the archive for your system from [Releases](https://github.com/Zorro909
+ 
+ | System | Archive | Training on |
+ | --- | --- | --- |
+-| Linux x86_64 | `altd-sim-<version>-x86_64-unknown-linux-gnu.tar.gz` | CPU, or AMD GPU through HIP |
+-| Windows x86_64 | `altd-sim-<version>-x86_64-pc-windows-msvc.zip` | CPU |
+-| macOS, Apple silicon | `altd-sim-<version>-aarch64-apple-darwin.tar.gz` | CPU |
++| Linux x86_64 | `altd-sim-<version>-x86_64-unknown-linux-gnu-cpu.tar.gz` | CPU |
++| | `altd-sim-<version>-x86_64-unknown-linux-gnu-amd.tar.gz` | CPU, or AMD GPU through HIP |
++| | `altd-sim-<version>-x86_64-unknown-linux-gnu-nvidia.tar.gz` | CPU, or NVIDIA GPU through CUDA |
++| Windows x86_64 | `altd-sim-<version>-x86_64-pc-windows-msvc-cpu.zip` | CPU |
++| | `altd-sim-<version>-x86_64-pc-windows-msvc-amd.zip` | CPU, or AMD GPU through HIP |
++| | `altd-sim-<version>-x86_64-pc-windows-msvc-nvidia.zip` | CPU, or NVIDIA GPU through CUDA |
++| macOS, Apple silicon | `altd-sim-<version>-aarch64-apple-darwin-cpu.tar.gz` | CPU |
+ 
+-Unpack it and open a terminal in the unpacked folder:
++The `amd` and `nvidia` archives add that vendor's GPU library beside the same `altd-sim` binary. Unpack one and open a terminal in the unpacked folder:
+ 
+ ```sh
+-tar -xzf altd-sim-nightly-x86_64-unknown-linux-gnu.tar.gz
+-cd altd-sim-nightly-x86_64-unknown-linux-gnu
++tar -xzf altd-sim-nightly-x86_64-unknown-linux-gnu-cpu.tar.gz
++cd altd-sim-nightly-x86_64-unknown-linux-gnu-cpu
+ ./altd-sim --version
+ ```
+ 
+@@ -68,13 +72,13 @@ Each browser connection gets its own session. Only one session at a time can use
+ 
+ All builds include all three math profiles. Windows 11 24H2 and newer default to `win11-fma3`; earlier Windows releases use `win10-fma3`; Linux and macOS use `proton`. To train on Linux for a game running on current Windows, for example, pass `--math-profile win11-fma3`. Checkpoints retain their profile. [Math profiles](math/README.md) describes the kernels and captured reference data.
+ 
+-### Training on an AMD GPU
++### Training on a GPU
+ 
+-The Linux archive contains `libaltd_gpu.so` next to `altd-sim`. It needs a ROCm 7.x runtime and one of the [covered GPU architectures](gpu/README.md). `./altd-sim gpu-info` checks that the library loads and prints `"layout": "ok"`; with a GPU present it also runs one kernel and prints the device, and fails if the library has no code for that GPU; the server's `HIP:` line shows the device or why HIP is unavailable.
++The `amd` and `nvidia` archives contain the GPU library next to `altd-sim`: `libaltd_gpu.so` on Linux, `altd_gpu.dll` on Windows. AMD libraries need a ROCm 7.x runtime that provides `libamdhip64.so.7` on Linux, or an AMD driver or HIP SDK 7.1 that provides `amdhip64_7.dll` on Windows, and one of the [covered GPU architectures](gpu/README.md). NVIDIA libraries need only an NVIDIA driver that supports CUDA 13 (580 or newer) and a consumer GPU from RTX 20 (Turing) to RTX 50 (Blackwell) or newer; datacenter GPUs such as A100 and H100 need a library built for their target, see the [GPU guide](gpu/README.md#nvidia-cuda). `./altd-sim gpu-info` checks that the library loads and prints `"layout": "ok"`; with a GPU present it also runs one kernel and prints the device, and fails if the library has no code for that GPU; the server's `HIP:` line shows the device or why HIP is unavailable.
+ 
+ A HIP session drives its first 12 ticks on both the GPU and the CPU and compares them. If they differ, or the network or track needs something HIP does not support, the session trains on CPU and Drive Lab shows the reason. Native CPU and HIP produce the same results. Browser training can differ from both in the last digits of a few values ([issue #3](https://github.com/Zorro909/driving_sim_rs/issues/3)), so a run moved between the browser and the server can develop differently from then on.
+ 
+-NVIDIA GPUs use the same backend: build the library with `gpu/build-cuda.sh` (Linux) or `gpu/build-cuda.ps1` (Windows) as described in the [GPU guide](gpu/README.md#nvidia-cuda). The release archives do not include it.
++NVIDIA GPUs use the same backend, so HIP in Drive Lab and the server output also covers CUDA. To build a library yourself, see the [GPU guide](gpu/README.md).
+ 
+ ### Troubleshooting
+ 
+@@ -178,6 +182,6 @@ cargo build --offline --target wasm32-unknown-unknown --no-default-features --fe
+ 
+ Tests retain self-contained native captures and label generated golden data as current-code regressions. Browser integration checks live in `wasm/test/`; HIP equivalence checks use `gpu_check`. The exhaustive trigonometry scan and hardware-dependent tests are opt-in. See [performance](docs/performance.md), [GPU internals](docs/gpu.md) and [WASM fidelity](docs/wasm.md).
+ 
+-GitHub Actions checks formatting, native and WASM Clippy, native release and debug tests, Windows and macOS Clippy and release tests, the HIP library build and layout, rustdoc, Rust 1.93 compatibility, and WASM packages with CPU browser comparisons. The release workflow builds and smoke-tests the archives, then publishes tags and nightly builds. Its native jobs override `target-cpu=native` with `x86-64` and include the flags in their cache keys, so cached builds can move between runner CPUs. Serial/threaded browser tests retain the documented expanded native differences. HIP equivalence and hardware WebGPU checks run separately because the hosted runners have no GPU.
++GitHub Actions checks formatting, native and WASM Clippy, native release and debug tests, Windows and macOS Clippy and release tests, the HIP library build and layout, rustdoc, Rust 1.93 compatibility, and WASM packages with CPU browser comparisons. The release workflow builds the AMD and NVIDIA libraries for Linux and Windows, smoke-tests the CPU, AMD and NVIDIA archives, then publishes tags and nightly builds. Its native jobs override `target-cpu=native` with `x86-64` and include the flags in their cache keys, so cached builds can move between runner CPUs. Serial/threaded browser tests retain the documented expanded native differences. HIP equivalence and hardware WebGPU checks run separately because the hosted runners have no GPU.
+ 
+ The project uses the [MIT license](LICENSE). Godot-derived code retains the [Godot notice](GODOT_LICENSE), and numerical modules retain their source notices.
+diff --git a/docs/server.md b/docs/server.md
+index 3c7c41d..937fb1e 100644
+--- a/docs/server.md
++++ b/docs/server.md
+@@ -39,7 +39,7 @@ There is no LAN binding, TLS or authentication token. These checks restrict brow
+ The first frame is a JSON text greeting:
+ 
+ ```json
+-{"type":"hello","protocol":1,"version":"1.2.0","threads":8,"carStateStride":14,"carStateFields":["x","y","..."],"hip":{"available":true,"device":"AMD Radeon RX 7900 XTX","platform":"HIP"},"mathProfile":"proton","mathProfiles":["proton","win10-fma3","win11-fma3"]}
++{"type":"hello","protocol":1,"version":"1.3.0","threads":8,"carStateStride":14,"carStateFields":["x","y","..."],"hip":{"available":true,"device":"AMD Radeon RX 7900 XTX","platform":"HIP"},"mathProfile":"proton","mathProfiles":["proton","win10-fma3","win11-fma3"]}
+ ```
+ 
+ Read `carStateStride` and `carStateFields` from the greeting rather than hard-coding them. The example's fields are abbreviated. `platform` is `"HIP"` for the AMD library and `"CUDA"` for the NVIDIA one. When HIP is unavailable, `hip` instead contains `{"available":false,"reason":"..."}`. Availability means the process opened the library and device; a particular session can still fall back to CPU. `mathProfile` is the default math profile and `mathProfiles` lists every profile a session can choose; [math profiles](../math/README.md) documents their runtime variants.
+diff --git a/gpu/README.md b/gpu/README.md
+index 4a51bb6..5fbc237 100644
+--- a/gpu/README.md
++++ b/gpu/README.md
+@@ -4,7 +4,7 @@ The GPU library runs whole driving generations on an AMD GPU (HIP) or an NVIDIA
+ 
+ ## Build and use
+ 
+-Linux release archives include `libaltd_gpu.so` beside `altd-sim`, built with ROCm 7.1 for the targets in `release-targets`. They need a ROCm 7.x runtime that provides `libamdhip64.so.7`, but no compiler. `altd-sim gpu-info` reports the loaded library. To build the library yourself, install ROCm with `hipcc`, then build from the repository root:
++The `amd` release archives include the library beside `altd-sim`, built for the targets in `release-targets`: `libaltd_gpu.so` with ROCm 7.1 on Linux, `altd_gpu.dll` with the HIP SDK 7.1 on Windows. They need a ROCm 7.x runtime that provides `libamdhip64.so.7` (Linux), or an AMD driver or HIP SDK 7.1 that provides `amdhip64_7.dll` (Windows), but no compiler. `altd-sim gpu-info` reports the loaded library. To build the library yourself, install ROCm with `hipcc` (or the HIP SDK on Windows, then use `gpu/build.ps1`), and build from the repository root:
+ 
+ ```sh
+ gpu/build.sh
+@@ -46,6 +46,8 @@ gpu/build-cuda.sh                  # Linux, writes target/gpu/libaltd_gpu.so
+ gpu/build-cuda.ps1                 # Windows, writes target/gpu/altd_gpu.dll
+ ```
+ 
++The `nvidia` release archives include a library built with CUDA 13.4 for `sm_75,sm_86,sm_89,sm_120` plus `sm_120` PTX: RTX 20 to RTX 50 GPUs, and later GPUs through the PTX. Other targets, such as `sm_80` (A100) and `sm_90` (H100), need their own build. It links the CUDA runtime statically, so it needs only an NVIDIA driver that supports CUDA 13 (580 or newer).
++
+ `OUT` overrides the output directory. `CUDA_ARCH` defaults to `native`, which builds for the GPUs of the build machine. Without a GPU, nvcc 13.4 warns and builds `sm_75` code and PTX instead, which newer GPUs compile at load time. To build for other GPUs, or several, list `sm_` targets separated by commas or spaces; `nvidia-smi --query-gpu=compute_cap --format=csv` reports the number, so 8.6 is `sm_86`. A list also embeds PTX for its newest target, which the driver compiles at load time for a newer GPU. CUDA 13 builds for Turing (`sm_75`, RTX 20 series) and newer only; older GPUs need a CUDA 12 toolkit, which is untested. Both scripts check the double tables with `gpu/gen_tables.py --check` (`python3` on Linux, the `py` launcher or `python` on Windows). The binary loads `ALTD_GPU_LIB`, then the library beside the executable, then `target/gpu`, as on AMD. `altd-sim gpu-info` runs one kernel when a GPU is present, so a library built for another GPU fails there with HIP error 209.
+ 
+ nvcc rejects a Visual Studio newer than the toolkit supports. `gpu/build-cuda.ps1 -AllowUnsupportedCompiler` (or `CUDA_ALLOW_UNSUPPORTED_COMPILER=1`) builds anyway with a warning; NVIDIA has not validated that combination, so run the checks below before trusting the library.
+diff --git a/gpu/build.ps1 b/gpu/build.ps1
+new file mode 100644
+index 0000000..499450f
+--- /dev/null
++++ b/gpu/build.ps1
+@@ -0,0 +1,28 @@
++# Build altd_gpu.dll (the GPU simulator parts) for AMD GPUs with the HIP SDK for Windows into $OUT.
++# GPU_ARCH lists one or more offload targets separated by commas or spaces.
++# -ffp-contract=off: hipcc otherwise fuses a*b+c into FMA, which rounds once.
++# HIP_PATH is the HIP SDK root, which its installer sets.
++param(
++    [string]$Out = $(if ($env:OUT) { $env:OUT } else { Join-Path $PSScriptRoot "..\target\gpu" }),
++    [string]$Arch = $(if ($env:GPU_ARCH) { $env:GPU_ARCH } else { "gfx1100" })
++)
++$ErrorActionPreference = "Stop"
++New-Item -ItemType Directory -Force $Out | Out-Null
++$Out = (Resolve-Path $Out).Path
++
++# The py launcher first: python3 on Windows is often the Microsoft Store placeholder.
++$tables = Join-Path $PSScriptRoot "gen_tables.py"
++if (Get-Command py -ErrorAction SilentlyContinue) { py -3 $tables --check } else { python $tables --check }
++if ($LASTEXITCODE -ne 0) { throw "gen_tables.py --check failed" }
++
++$targets = @($Arch -split '[,\s]+' | Where-Object { $_ })
++if (-not $targets) { throw "GPU_ARCH lists no targets" }
++if (-not $env:HIP_PATH) { throw "HIP_PATH is not set; install the HIP SDK for Windows" }
++$hipcc = Join-Path $env:HIP_PATH "bin\hipcc.exe"
++
++$src = Join-Path $PSScriptRoot "sim\altd_gpu.hip"
++$dll = Join-Path $Out "altd_gpu.dll"
++& $hipcc -O3 -ffp-contract=off @($targets | ForEach-Object { "--offload-arch=$_" }) -shared -std=c++20 `
++    -fhip-fp32-correctly-rounded-divide-sqrt -o $dll $src -L (Join-Path $env:HIP_PATH "lib") -lamdhip64
++if ($LASTEXITCODE -ne 0) { throw "hipcc failed" }
++Write-Host "built $dll for $($targets -join ' ')"
+diff --git a/gpu/sim/altd_gpu.hip b/gpu/sim/altd_gpu.hip
+index b0962e3..fafb3ee 100644
+--- a/gpu/sim/altd_gpu.hip
++++ b/gpu/sim/altd_gpu.hip
+@@ -502,7 +502,8 @@ double thread_cpu_seconds() {
+ // runtime's single-dispatch limit, so by default every launch allocates and
+ // frees it (0.2-1.7 ms per launch). Keeping the scratch avoids that. Runs at
+ // load, before the first HIP call initialises HSA; an explicit setting wins.
+-#if !defined(__CUDACC__)
++// HIP on Windows runs on PAL instead of HSA, and its C runtime has no setenv.
++#if !defined(__CUDACC__) && !defined(_WIN32)
+ __attribute__((constructor)) void keep_scratch() { setenv("HSA_NO_SCRATCH_RECLAIM", "1", 0); }
+ #endif
+