CUDA backend by rafaelha · Pull Request #219 · QuEraComputing/tsim

rafaelha · 2026-10-27T04:23:44Z

With strategy="active_statevector" (the default), the samplers now run on an NVIDIA GPU whenever CuPy sees one and every statevector component has at most 11 active qubits. This follows the stabilizer-rank strategies, which already use the GPU through JAX when it is installed. The noise sampler and every statevector component run as CUDA kernels that CuPy compiles with NVRTC, in float32.

Stacked on #284, which holds the host code that this backend and the Metal backend (#290) share: the tape generator (statevector/gpu/tape.py), the noise tables (gpu/noise.py), the row packing (gpu/rows.py), shared.common_prefix and the deferred-measurement plans. This PR now holds only what is CUDA's. Its history was squashed into two commits on top of #284; the previous head was 20b3eef.

The public API is main's: this PR adds no functions, methods or arguments. With a postselection_mask, sample returns only the shots that pass, as on main, and only those shots are copied off the GPU.

What's in it

statevector/cuda/codegen.py: the CUDA dialect of gpu.tape (__shfl_xor_sync among the shot's lanes, goto to the end of a pruned shot) and the CUDA prelude. The outcome draw is in float32, with a double-precision fallback near the decision boundary.
statevector/cuda/noise.py: the source of the noise and noise_kept kernels, which read gpu.noise's tables with a 46 KB shared-memory budget.
statevector/cuda/runtime.py: loads CuPy with a pip CUDA 13 NVRTC preloaded, and detects forks and out-of-memory errors.
statevector/cuda/sampler.py: CudaShotSampler. It runs chunks sized from the estimated pass rate on two streams, packs the kept rows on the GPU and copies them to pinned memory while the next chunk's kernels run. Noise that postselection cannot see is drawn on the kept shots only. Its private sample_accepted(..., expectations=True) returns observable expectations, which the tests use.
sampler.py: _gpu_shots chooses the GPU, with the same shape as the Metal PR's:
- TSIM_DEVICE=cpu (or CUDA_VISIBLE_DEVICES="") keeps the samplers on the CPU.
- sampler.stats.device and the repr (samples on cuda) say where a sampler samples. The new tsim-dev bench --device reads both.
- These cases stay on the CPU, with a warning when a GPU is present:
  - a component wider than 11;
  - a forked child;
  - noise beyond the GPU tables' limits;
  - rows too wide for the kernels' shared memory;
  - out of GPU memory, after the call has retried in smaller batches.
- A program without shot words stays on the CPU.
Install: the new cuda extra (cupy-cuda13x[ctk], CUDA 13, driver 580 or newer). The JAX extras are now jax-cuda12 and jax-cuda13; Tsim 1.0's cuda12 and cuda13 still install them. Docs (install.mdx, index.mdx, README), AGENTS.md, the agent skill and the changelog say so.
Seeds: on the GPU, a seed's samples differ from the CPU's and depend on the batch size. On the CPU they are unchanged: the GPU draws from a fourth child of the seed's SeedSequence.

Changes since the last revision

Rebuilt on feat(statevector): add the host code that GPU backends share #284. With the CUDA dialect, gpu.tape.tape_function writes byte-identical source to the old cuda/codegen.py for 120 tapes: the test circuits and random circuits, unrolled and looping, at 1, 2 and 32 threads per shot, with and without masks.
New fallbacks:
- Programs without shot words used to divide by zero. An example is a detector sampler of a circuit without detectors.
- Noise beyond the tables' limits used to raise a ValueError. An example is distillation_15to1_T_d9, whose 232 shot words do not fit in shared memory.
- Both now sample on the CPU.
Review fixes:
- Rows of 97 to 176 output words overflowed the gather kernels' 48 KB of shared memory (CUDA_ERROR_INVALID_VALUE). They now go to the CPU.
- Calls on one sampler from several threads are serialized.
- An empty mask no longer shares the module of no mask, which returned no rows.
- A sampler compiled before a fork samples on the CPU in the child.
- GPU samplers pickle and continue their streams.

Benchmark

uv run tsim-dev bench --suite full --side .:device=gpu --side .:device=cpu --thorough, tsim-bench eb4a6d11, at d11d4d6. It runs on an RTX 5090, with bit-packed rows, the postselection that tsim-bench records, and the copy of the kept rows to the host included. The reference is this branch's CPU executor on 1 thread of a 24-core machine; all cores would be roughly 10 to 20 times faster than 1 thread. kept is the share of shots that pass postselection (1 without). Compile times include NVRTC with CuPy's disk cache warm. The first 19 rows come from a rerun on an idle machine, because the second full run overlapped other jobs; they agree with the first full run to within 2%.

circuit	k	kept	GPU (shots/s)	CPU, 1 thread (shots/s)	GPU / CPU	compile s (GPU / CPU)
inject_cultivate_d3_p0.001	4	0.687	1.27e+10	3.22e+07	394x	0.24 / 0.02
inject_cultivate_d5_p0.001	6	0.144	4.37e+09	1.46e+07	299x	0.36 / 0.14
end2end_d3_dsurface15_p0.001	4	0.539	2.45e+08	1.06e+06	231x	2.34 / 1.83
hook_inject_T_d15_pd5_p0.001	1	0.667	1.36e+08	1.10e+06	124x	2.72 / 2.08
distillation_color_d5_Z	3	0.0798	1.46e+10	1.28e+07	1140x	0.23 / 0.03
coherent_d3_r3	4	1	8.67e+09	1.96e+07	442x	0.22 / 0.02
surface_d7_r7	0	1	1.06e+09	1.32e+07	81x	0.37 / 0.11
qv_q10	10	1	1.94e+08	1.11e+07	18x	0.38 / 0.18
end2end_d5_dsurface15_p0.001	6	0.0518	4.88e+08	8.59e+05	568x	3.34 / 2.75
cult_growing_T_d3_to_d7_p0.001	4	0.643	1.74e+09	7.86e+06	221x	0.50 / 0.17
unrotated_H_injection-ghz_d1=3_p=0.001	9	0.548	9.88e+07	4.06e+05	243x	0.33 / 0.09
expanded-rotated-injection-double-ghz-HXY_d1=3_d2=11_p=0.001	8	0.626	2.08e+08	2.18e+06	95x	1.38 / 0.94
rp_3_T_cult	4	0.514	8.44e+09	2.04e+07	414x	0.25 / 0.04
rp_3_rp_5_T_cult	10	0.0945	2.74e+08	6.66e+06	41x	0.52 / 0.21
fold-d5_native_T_p0.001	8	0.16	3.02e+08	4.31e+06	70x	0.66 / 0.34
inject_cultivate_T_d3_p0.001	4	0.687	1.44e+10	3.23e+07	446x	0.23 / 0.02
gauging_cultivation_d5_T_p0.001	6	0.0401	2.22e+09	8.05e+06	275x	0.61 / 0.33
d5a19_inject_cultivate_T_p0.001	6	0.144	3.11e+09	1.46e+07	213x	0.36 / 0.14
lattice-surgery-complementary-gap-3-to-5-to-7_T	4	0.574	3.01e+09	1.47e+07	204x	0.65 / 0.35
s_injection_T_d5_p0.001_pinj0.1	1	0.912	2.88e+09	3.47e+07	83x	0.26 / 0.02
bb144_cfg369_inj1_T_p0.001	1	0.0703	2.88e+09	6.63e+06	433x	0.66 / 0.32
star_d=5	2	0.237	3.66e+09	1.59e+07	231x	0.29 / 0.07
distillation_steane_d3_Z	3	0.353	1.68e+10	2.10e+07	797x	0.26 / 0.01
distillation_15to1_T_d9_p0.001	4	1	3.94e+05 (CPU fallback)	3.95e+05	1.00x	6.60 / 6.40
unfolded_surface_no_bitflip_pz1e-3_rounds4	4	1	1.31e+09	1.13e+07	116x	0.37 / 0.12
distillation_15to1_p0.001	4	0.985	1.56e+10	6.07e+07	258x	0.21 / 0.01
bh10-T	3	0.963	1.37e+10	1.98e+07	692x	0.22 / 0.02
coherent_d5_r5	11	1	2.51e+06	1.87e+03	1346x	0.54 / 0.31
stability_d3_X_t20_alpha0.5_p0.001	3	1	1.20e+09	3.49e+05	3433x	1.17 / 0.86
tricycle_6_3_2	9	0.92	1.35e+08	4.46e+06	30x	0.24 / 0.03
tricycle_9_3_2	11	0.867	1.59e+07	1.96e+05	81x	0.31 / 0.05
tricycle_12_3_2	16	0.806	5.75e+02 (CPU fallback)	5.72e+02	1.01x	0.25 / 0.06

30 of 32 circuits run on the GPU, 18x (qv_q10) to 3400x faster than one CPU thread. The median is about 240x.
Two stay on the CPU, at the CPU's rate:
- tricycle_12_3_2 needs 16 active qubits.
- distillation_15to1_T_d9 has 232 shot words (14832 detectors), more than CUDA's noise kernel keeps in shared memory. Metal keeps such rows in device memory (shared_words=False); CUDA could do the same.
Compiling takes 0.2 to 3.3 s, 0.1 to 0.6 s longer than on the CPU, for NVRTC.
tsim-dev bench --ref origin/feat/gpu-shared (CPU, quick suite): 0 of 8 circuits slower. The run was noisy, since other processes were using the CPU.

Testing

uv run tsim-dev check on the RTX 5090: lint passes and 2731 tests pass, the CUDA tests included. Since then, the CUDA tests pass at d11d4d6: 82, five of them new for the review fixes, which fail without them.
test/integration/test_cuda_sampler.py skips without CuPy and a GPU. It compares detector, measurement and postselected samples with the stabilizer-rank strategy and the CPU executor. It also covers:
- main's column layouts;
- postselection that keeps columns, bit for bit;
- expectations and their fallback under later postselection;
- widths 8 to 11 and looping tapes;
- tiny noise and outcome probabilities, drawn at their rates;
- every CPU fallback;
- out-of-memory retries;
- fork, pickle and threads.
CI has no GPU, so the CUDA tests only run locally.
Statistical cross-checks at scale (earlier revisions, same kernels):
- d=5 end-to-end cultivation reproduces Clifft's curve at every compared gap threshold.
- Tsim GPU, Tsim CPU and Stim agree on Hetényi's d=3 gauging circuit at 1e10 shots each.

Known limits, not fixed here

Memory:
- Each call's result rows come from CuPy's pinned-memory pool, which reuses a block only for the same size. A long loop of postselected calls can therefore grow page-locked host memory.
- A call holds every chunk's packed device rows until it returns.
- Each sampler keeps its device buffers for its lifetime.
NVRTC failures: if NVRTC cannot build the kernels (for example an old CuPy or system NVRTC), the error is raised instead of falling back to the CPU.
Booleans: boolean and rearranged outputs are unpacked on the host, single-threaded. Metal unpacks booleans on the GPU.
CPU-faster cases: Clifford-only programs and booleans can be slower on the GPU than on the CPU.
stats.device: it still says cuda after a call falls back to the CPU.
CERTAIN: outcomes of weight at most 1e-11 of the total count as impossible.

🤖 Generated with Claude Code

github-actions · 2026-10-27T04:24:59Z

PR Preview Action v1.8.2
🚀 View preview at https://queracomputing.github.io/tsim/pr-preview/pr-219/
Built to branch `gh-pages` at 2026-11-05 23:15 UTC. Preview will be ready when the GitHub Pages deployment is complete.

github-actions · 2026-10-27T04:30:17Z

☂️ Code Coverage

current status: ✅

Overall Coverage

Statements	Covered	Coverage	Threshold	Status
6865	5388	78%	0%	🟢

New Files

File	Coverage	Status
src/tsim/statevector/cuda/_init_.py	34%	🟢
src/tsim/statevector/cuda/codegen.py	100%	🟢
src/tsim/statevector/cuda/kernels.py	77%	🟢
src/tsim/statevector/cuda/noise.py	70%	🟢
src/tsim/statevector/cuda/runtime.py	57%	🟢
src/tsim/statevector/cuda/sampler.py	23%	🟢
TOTAL	60%	🟢

Modified Files

File	Coverage	Status
src/tsim/circuit.py	95%	🟢
src/tsim/core/types.py	100%	🟢
src/tsim/sampler.py	95%	🟢
src/tsim/statevector/gpu/tape.py	98%	🟢
src/tsim/statevector/metal/_init_.py	37%	🟢
TOTAL	85%	🟢

updated for commit: 6f99566 by action🐍

rafaelha

Review of the CUDA backend. The first three comments change sampled statistics relative to device="cpu"; the next two are memory-scaling failures; the last is test coverage.

rafaelha · 2026-10-28T04:23:03Z

+        while (true) {{
+            state += GOLDEN;
+            u64 r = mix64(state);
+            float skip = lk < -3.0e38f ? 0.f : floorf(__logf(unit(r)) / lk);


Geometric skip is biased for small fire probabilities. unit(r) is a 24-bit uniform and __logf is the fast approximation, so the per-segment fire probability has a floor of 2^-24 (u = 1.0 always gives skip 0). The CPU next_skip is exact down to 2^-64.

Measured on an RTX 5090 over all 2^24 unit() values, 256-trial segment, one channel per block:

p P(fire) / exact

1e-8 1.02x

1e-10 4.7x

1e-12 233x

A circuit with many negligible-but-nonzero channels (e.g. 1000 idle channels at p=1e-12) gets ~230x the CPU fault rate, with no warning. __logf is never positive for any unit() output, so there is no negative-skip / out-of-bounds issue, only the bias. Suggest a 53-bit uniform with log/log1p in double, or at least a 32-bit uniform on (0,1).

rafaelha · 2026-10-28T04:23:03Z

+            self._buffer("values", (shots, no), np.float32) if expectations else None
+        )
+        counts.fill(0)
+        seeds = cp.asarray(


Samples depend on batch_size. Component seeds and the noise seed are drawn from self._rng once per device batch, and the noise stream is keyed by blockIdx within the batch. So compile_sampler(seed=1, device="cuda").sample(2**23) and .sample(2**23, batch_size=2**22) return different samples.

This contradicts the active_statevector docstring in sampler.py ("the sampled outcomes of a seed do not depend on the batch size") and ShotSampler.sample's batch_size doc. Either key seeds/noise by the global shot index, or document that device="cuda" is exempt.

rafaelha · 2026-10-28T04:23:03Z

+                )
+        self._num_detectors = num_detectors
+        self._observables = np.arange(num_detectors, layout.num_outputs)
+        self._rng = np.random.default_rng([seed, 2])


RNG stream collides with the reference sample. shots.py uses default_rng([seed, 2]) for _reference_seeds. Here the same stream is used, and both the host reference() and the first _run_batch start with integers(0, 2**64, size=num_components), so they get identical component seeds. draw(seed, 0, ...) then matches the CPU _draw for shot 0.

With use_detector_reference_sample=True, shot 0 of the first batch reproduces the noiseless reference whenever noise misses the component, so its detectors are deterministically 0 after the XOR. Use a distinct stream, e.g. [seed, 3].

rafaelha · 2026-10-28T04:23:03Z

+        """Return the shot words of ``shots`` shots and the largest norm deviation."""
+        module = self._module(mask)
+        layout = self.layout
+        out = np.empty((shots, layout.num_words), dtype=np.uint64)


Whole-call host allocation. _sample_words allocates (shots, num_words) uint64 up front, including key words, and copies every word back per batch. ShotSampler.sample on CPU splits into _part_shots() parts bounded by memory_budget() and keeps only output columns.

E.g. compile_detector_sampler(device="cuda").sample(10**9) on a 12-word layout tries to allocate 96 GB and raises MemoryError before any kernel runs; the same call on CPU succeeds. The gather kernel (compiled but never launched) looks like it was meant for this: gather only output_words per batch and size the host array accordingly.

rafaelha · 2026-10-28T04:23:03Z

+MAX_WIDTH = 11
+"""Widest component the CUDA executor runs: 32 threads of 64 amplitudes."""
+
+BATCH_SHOTS = 1 << 22


Batch size ignores device memory and word count. A layout with 150 shot words (the noise kernel admits up to ~176) needs 2^22 × 150 × 8 B ≈ 5 GB for words alone, plus queues and the float32 values buffer. On an 8–12 GB GPU this hits OutOfMemoryError unless the user guesses a smaller batch_size. Suggest deriving the default from cp.cuda.Device().mem_info and layout.num_words, analogous to memory_budget() on CPU.

rafaelha · 2026-10-28T04:23:03Z

+from tsim.statevector.cuda import available
+from tsim.statevector.program import StatevectorComponent
+
+pytestmark = pytest.mark.skipif(not available(), reason="needs CuPy and a CUDA GPU")


Module-wide skip hides device-independent tests in CI. test_sample_postselected_returns_only_passing_shots[cpu], test_expectations_need_the_cuda_device, and the device="tpu" / strategy="cat5" ValueError checks in test_cuda_needs_the_statevector_strategy don't need CuPy, but CI (no GPU) skips them. A regression in the CPU branch of sample_postselected would go unnoticed. Suggest moving these to test/unit/test_sampler.py or applying the skip only to the CUDA cases.

…223) Stacked on #222 (stack #245). This PR's diff is against that branch. ## Summary This PR changes no behaviour. It is cleanup from the statevector review, made of one new test and three refactors. Samples are bit-identical after every commit. No feature or fix depends on it, so it can wait behind #222 and #219. - **Brute-force test for `_Relations`** (`test_pbc.py`). `_Relations` is the GF(2) elimination that turns classical spiders into output parities. Until now it was only covered through whole-circuit distributions. The test builds 300 random small systems and sums the sign over the eliminated variables by enumeration. It then checks the reduced rows, and that the reduction raises exactly when no such rows exist. - **One symbol encoding** (`program.py`, `shared.py`, `shots.py`; 13 lines added, 32 removed). The tape used to pack symbol sets into uint32 words, which `lane_program` unpacked again before packing them into uint64 lane words. The tape now keeps a boolean column per symbol. `_pack`, `unpack_symbols` and `StatevectorComponent.num_words` are gone, and `num_symbols` replaces `num_words`. The tape grows slightly: its symbol rows take 1.3 MB on d=5 cultivation, up from about 0.16 MB. - **No graves in `run_lanes`** (`shared.py`; 161 lines added, 100 removed). A pruning now finishes the dying lanes' shots right away, with their outputs, coins and hidden bits, and drops them from `live`. Before, they waited in a grave buffer until the next random measurement or the end of the run. This removes `graves`, `grave_at`, `dead_shot`, `dead_grave`, the per-entry hidden masks and the `-1 - grave` codes. The per-shot work moves into `_finish_shot`, and `lane_bytes` counts two rows of lane symbols per lane instead of four. - **One sampling entry point** (`sampler.py`). `CompiledMeasurementSampler.sample` and `CompiledDetectorSampler.sample` both call `_sample_outputs(shots, batch_size, columns, bit_packed, mask, referenced)`. That method picks the shot-word or stabilizer-rank path once. The stabilizer-rank path's mask screening and reference XOR move into `_sample_stabilizer_rank_outputs`, so `CompiledDetectorSampler.sample` no longer duplicates them. ## Open before merge - [x] **Single-thread slowdown on the shared, unpruned path.** Fixed in `ad8d042`: the end-of-run loop is inline again and has no `hidden` pass, and `_finish_shot` serves only pruned shots. On one thread, a width-7 component with 12 detectors, 2M shots and sharing on takes 0.031 s, against 0.033 s on #222 and 0.045 s before this fix. Sample hashes are identical. - [x] **The `_Relations` test checks the reduced echelon form** (`ad8d042`, `04b7ae3`): pivots are unique and in `kept` order, every row term is in `kept`, and no other row contains a pivot. Dropping the back-substitution and reversing the pivot priority both fail it now. - [x] **Docstring nit in `_sample_outputs`**, fixed in `ad8d042`. - [x] **Stacked on #222 in stack #245.** Both branches are rebased linearly onto `main` at `6ee9e1b` (#240), and the stack merges them together. - [ ] **Overlap with #219 (CUDA).** It merges cleanly with this branch. Its new `sample_postselected` builds `referenced`/`flip` the same way `_sample_outputs` now does, so whichever PR lands second can reuse that code. ## Verification - After the rebase onto `main` at `6ee9e1b`, the full suite passes locally (2367 passed) and pyright is clean. `CX_d1=5_cult`, with every detector postselected, gives the same hash as #222 with sharing on and off. Before the rebase, CI was green on `c221fea`. - **Executor commits.** A snapshot hashed exact samples and lane statistics for 48 circuits, with sharing on and off and three masks, plus two prune-heavy workloads. The hash is identical before and after. - **Sampler commit.** A second snapshot covered both strategies and every combination of mask, both reference flags, `bit_packed` and observable layout. The hash is identical before and after. - **Review harness.** 74,304 `active_statevector` cases and 4,788 `cat5` cases produced identical hashes against the base branch. So did lane statistics on 62 more circuits. - **Timing.** The table below uses cultivation, with every detector postselected, best of 5. Sample hashes are identical. | Workload | Base | This PR | | --- | --- | --- | | `CX_d1=5_cult`, 2M shots, 1 thread, sharing on | 1.38 s | 1.31 s | | `CX_d1=5_cult`, 2M shots, 1 thread, sharing off | 1.87 s | 1.82 s | | `CX_d1=5_cult`, 2M shots, all threads, sharing on | 0.18 s | 0.19 s | | `CX_d1=5_cult`, 2M shots, all threads, sharing off | 0.25 s | 0.26 s | | Width-7 component, 12 detectors with half postselected, 400k shots, sharing on | 0.05 s | 0.05 s | Pruning now scans the live shots each time a lane dies, instead of waiting for the next random measurement. That cost doesn't show up even on this prune-heavy workload. The only slowdown measured was the single-thread case without pruning, now fixed. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

Defines what tsim 1.0.0 commits to under semantic versioning. ## Public API - `tsim.__all__ = ["Circuit", "CompiledDetectorSampler", "CompiledMeasurementSampler"]`. These are the only objects semver covers. `tsim.__init__`'s docstring and the docs index ("API stability") say that every other module is internal, even where it is importable. - New `tsim.experimental` re-exports `CompiledStateProbs`, `SteaneEncoder` and `ColorEncoder5`, which the tutorials use. They may change in any release. The old import paths still work. - `docs/scripts/gen_ref_nav.py` no longer generates a page per module. The API reference has two pages, "Tsim API" (the three public objects) and "Experimental". - `overview.ipynb` and `magic_state_distillation.ipynb` import from `tsim.experimental`. ## Sampling strategies Only `strategy="active_statevector"` is fully stable. The stabilizer rank strategies (`"cat5"`, `"bss"`, `"cutting"`) are provisional: they may be removed in a minor release, after at least one minor release in which they emit a `DeprecationWarning`. The docstrings say only "provisional"; the policy is in the "API stability" section of `docs/index.md` and in the CHANGELOG. ## Seed reproducibility The `seed` docstrings of `compile_sampler`, `compile_detector_sampler`, the sampler classes and `CompiledStateProbs` now say that a seed reproduces its samples only with the same Tsim version and backend (CPU or GPU). Without this, changing a random stream (as #263 did), or adding a float32 GPU executor (#219), would technically be a breaking change. ## Test plan - [x] `uv run pytest test/unit`: 2072 passed. Includes a new `test/unit/test_public_api.py`. - [x] pre-commit (black, isort, ruff, pyright, codespell) - [x] `mkdocs build` from a clean checkout. The tutorials run, and the reference contains only `tsim.Circuit`, the two samplers and the `tsim.experimental` objects. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

The CUDA and Metal backends of the statevector executor share their host side: `statevector/gpu/` holds the noise kernels' tables (`noise`), a tape generator that writes a component's tape in either language through a small `Dialect` (`tape`), and the packing of returned rows (`rows`). Two changes to the statevector plans come with it. `shared.common_prefix` returns the leading actions of a tape that every shot runs alike, which a GPU backend runs once. `plan_component` also plans each component with its measurements deferred (`schedule(defer=True)`) and keeps that plan when it is no wider and leaves less shot-dependent work. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…rk model - noise_tables raises when an outcome flips more than 255 words (the record's 8-bit count would wrap) and when the events alone do not fit shared_bytes without shared words (MIN_EVENTS is now a named constant). - Dialect rejects an amplitude_chunk that is not a positive power of two, which would declare fewer chunk arrays than the tape indexes. - _shot_work ends the shared prefix where shared.common_prefix does: an unread dormant measurement is a fair coin and no longer counts as shot work. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>

Below p ~ 2^-24, a kernel's 24-bit geometric skip reaches only some trials of a segment, so a group of many such channels fired them at biased per-channel rates (0.93 to 1.17 of their rates at p = 1e-9 on 3200 channels; the total rate was right). Such groups are now cut into segments of fewer than 2^-13 expected faults, which the kernels draw from 48 bits (their fine path). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Rebuild #219's CUDA backend on #284's tsim.statevector.gpu: the tape comes from gpu.tape with a CUDA dialect (byte-identical to #219's codegen), the noise tables from gpu.noise and the row packing from gpu.rows. cuda/ keeps the CUDA prelude, the noise kernels' source, the runtime and the sampler. The dispatch follows the Metal backend's shape (_gpu_shots): TSIM_DEVICE=cpu keeps the samplers on the CPU, sampler.stats.device and the repr say where a sampler samples, programs without shot words stay on the CPU, and noise beyond the GPU tables' limits falls back to the CPU with a warning. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Both dialects are internal constants, so a ValueError and its tests guard nothing a user can reach. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Both backends return every output, in order, so each packed word is a whole shot word and the piece tables only encoded the identity. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Both GPU backends built the read mask of a postselected batch, chose the deferred channels and estimated the screen's pass rate with the same code; gpu.noise now has split, screen_pass_rate and words_fit (the shared-memory fit test that Metal repeated with literals). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Both GPU backends wrote the same component calls, seeds, word stores and postselection hits around the tape functions, and repeated MAX_WIDTH, the threads per shot, the deviation tolerance and the layout check. gpu.tape now has tape_kernel, threads_per_shot, deviation and supports. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Rebuild #219's CUDA backend on #284's tsim.statevector.gpu: the tape comes from gpu.tape with a CUDA dialect (byte-identical to #219's codegen), the noise tables from gpu.noise and the row packing from gpu.rows. cuda/ keeps the CUDA prelude, the noise kernels' source, the runtime and the sampler. The dispatch follows the Metal backend's shape (_gpu_shots): TSIM_DEVICE=cpu keeps the samplers on the CPU, sampler.stats.device and the repr say where a sampler samples, programs without shot words stay on the CPU, and noise beyond the GPU tables' limits falls back to the CPU with a warning. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

- Rows of 97 to 176 output words overflowed the shared memory of the gather and noise_kept kernels (CUDA_ERROR_INVALID_VALUE); such circuits now sample on the CPU with a warning. - Calls on one sampler from several threads are serialized, as on Metal. - An empty postselection mask no longer shares the module of no mask, which returned no rows. - A sampler compiled before a fork samples on the CPU in the child, with a warning, instead of failing to initialize CUDA. - GPU samplers pickle: the inputs and the stream state are kept, and the GPU state is rebuilt on load (the CPU sampler where CUDA is missing). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The sampler always returns every output in order, so pack_rows copies output word k and clears the bits past the last output, instead of reading a piece table per word (gpu.rows.pieces is gone). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The tape kernel's glue, the noise split, the screen's pass rate, the deviation, the width limit and the fork warning now come from tsim.statevector.gpu; the samples of a seed are unchanged. codegen's tape_function and hex64 pass-throughs, the duplicate NE argument and a misplaced docstring go. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

tsim/sampler.py knew CUDA's width limit, its fork state, its out-of-memory errors and the no-words rule; cuda.shot_sampler now decides, and sampler.py keeps only the TSIM_DEVICE check and one call. ShotSampler says device = "cpu" itself. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The kernels' __launch_bounds__ and their launches spelled out 128, 256 and 1024 separately, and the counts buffer was sliced by bare indices. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The kernel template and the source assembly leave sampler.py, so _module only plans, compiles and looks up the kernels. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The detector, measurement and looping-tape comparisons repeated one body; the same 24 cases run. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The changelog entry described the kernels, and three docstrings repeated a six-line paragraph; the install page keeps the details. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

## Problem `tsim-dev bench --device gpu` reported GPU throughput about 3.7x too low. The tsim-bench worker capped one sampling call at `2^28 / outputs` shots, as if every shot came back as unpacked booleans. `tsim-dev bench` asks for bit-packed rows, though, and a postselected call only returns the rows it keeps. On `end2end_d5_dsurface15_p0.001` (2745 outputs, 5.2% kept) the cap allowed 97,790 shots per call, and the CUDA sampler's fixed ~0.2 ms per call dominated. The served (non-`--thorough`) rounds had a second problem: the first side's worker calibrated one call and every side timed it. With a GPU side against a CPU side, the CPU then had to run a GPU-sized call, which takes minutes. ## Fix - Re-pin tsim-bench to rafaelha/tsim-bench#2 (`19804f26`). Its worker bounds a call by the bytes it actually returns: `ceil(outputs/8)` per bit-packed row, times the acceptance measured during calibration. The returned array stays under 256 MB, and a call is also capped at 2^27 shots. The timed size gets two warm calls before timing, because the first calls at a new size are slow on a GPU (pinned host buffers). - `_served_rounds` calibrates once per device, so sides on different devices each time a call of their own size. There is a unit test for this. - The benchmarking skill notes both changes. The new pin also brings in tsim-bench `9bc81ad` (`qv_q10`: one observable per measured qubit). That changes what `qv_q10` measures, so its before/after rates below are not like-for-like. ## Numbers (RTX 5090, 1 CPU thread) GPU side: the CUDA branch `simplify/cuda@70e55b8` (#219), since main has no GPU backend. `tsim-dev bench end2end_d5_dsurface15_p0.001 --side <cuda>:device=gpu --side <cuda>:device=cpu --thorough`: | tsim-bench | GPU shots/s | shots per call | CPU shots/s | GPU vs CPU | |---|---|---|---|---| | `eb4a6d11` (old pin) | 512M | 97,790 | 0.86M | 595x | | `19804f26` (new pin) | 1.89G | 15.1M | 0.86M | 2196x | The default served mode now gives the same result: 1.89G against 0.86M, 2197x. GPU quick suite (`--side <cuda>:device=gpu`), old pin → new pin: inject_cultivate_d3 12.6G → 14.8G, inject_cultivate_d5 4.36G → 6.89G, end2end_d3 242M → 349M, hook_inject_T_d15 139M → 178M, distillation_color_d5_Z 14.7G → 19.6G, coherent_d3_r3 8.77G → 10.0G, surface_d7_r7 1.08G → 1.23G, qv_q10 197M → 195M. CPU quick suite (`--side .:device=cpu`, Tsim main), old pin / new pin (CPU calls are bounded by time, not by the cap, so the rates should not move): | circuit | old | new | |---|---|---| | inject_cultivate_d3_p0.001 | 28.5M | 28.7M | | inject_cultivate_d5_p0.001 | 14.8M | 14.8M | | end2end_d3_dsurface15_p0.001 | 1.07M | 1.07M | | hook_inject_T_d15_pd5_p0.001 | 1.11M | 1.10M | | distillation_color_d5_Z | 12.9M | 13.0M | | coherent_d3_r3 | 19.6M | 19.7M | | surface_d7_r7 | 13.3M | 13.3M | | qv_q10 | 10.7M | 10.7M | `uv run tsim-dev check` passes (lint, 2618 tests), and `python -m tsim_bench check` reports 0 failed. ## Merge order rafaelha/tsim-bench#2 must merge first. `REV` pins its branch commit `19804f26`, which `_export` can fetch because it lives on a pushed branch. If that PR is squash-merged, re-pin `REV` here to the merged commit on tsim-bench `main` before merging. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

codecov · 2026-11-11T19:55:47Z

Codecov Report

❌ Patch coverage is 32.06612% with 411 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
src/tsim/statevector/cuda/sampler.py	22.76%	346 Missing ⚠️
src/tsim/statevector/cuda/runtime.py	56.66%	26 Missing ⚠️
src/tsim/statevector/cuda/__init__.py	34.48%	19 Missing ⚠️
src/tsim/statevector/gpu/tape.py	27.27%	8 Missing ⚠️
src/tsim/statevector/cuda/noise.py	70.00%	6 Missing ⚠️
src/tsim/statevector/cuda/kernels.py	77.27%	5 Missing ⚠️
src/tsim/sampler.py	90.00%	1 Missing ⚠️

📢 Thoughts on this report? Let us know!

First of three PRs for the statevector executor on GPUs. It holds the host-side code that the CUDA backend (#219) and a Metal backend for Apple GPUs share, so that either can land first and the other rebases onto it. Nothing on `main` uses `statevector/gpu/` until a backend does. ## Changes - `statevector/gpu/noise.py`: the noise kernels' host tables (channel groups cut into segments, 32-byte outcome records with alias thresholds, overflow flips), `deferred_channels` and `pass_rate`. This is the host part of #219's `cuda/noise.py`. The fast-memory budget is now a parameter (CUDA has 46 KB of shared memory, Metal 14 KB is fastest). New: `shared_words=False` keeps a block's words in device memory and gives the events the whole budget. Rows of more than about 47 words (over 3000 outputs) need that on Metal. A layout with no shot words raises a `ValueError` instead of dividing by zero. - `statevector/gpu/tape.py`: #219's tape generator, with a frozen `Dialect` for what differs between CUDA C++ and the Metal Shading Language: - spelling: qualifiers, address spaces, `make_float2`/`float2`, literal suffixes; - the shuffle templates and the lane mask; - control flow: `goto done` or a `do { … break; } while (false)` block, since Metal has no `goto`; - two options that Metal needs and CUDA keeps off: splitting an unrolled tape's amplitudes into arrays of 32, and projecting partner pairs together. - `statevector/gpu/rows.py`: `pieces` and `select` from #219's sampler. - `shared.common_prefix` and the deferred-measurement plans (`schedule(defer=True)`, `plan_component`, `_shot_work`), unchanged from #219. ## Checks - **CUDA source identical:** with the CUDA dialect, `gpu.tape.tape_function` writes the same source, byte for byte, as #219's `cuda/codegen.py` at 20b3eef. 202 tapes compared: the test circuits, random circuits up to width 11, cultivation d3 and tricycle 6-3-2, unrolled and looping, with and without postselection masks, at 1, 2 and 32 threads per shot. #219's CUDA tests therefore stay valid once it imports this module. I have no NVIDIA GPU, so nothing was run on CUDA. - **New CPU tests** (`test/unit/statevector/test_gpu.py`): - a NumPy model of the noise kernel draws from the tables, against `ChannelSampler.sample_words`: per-bit rates, and a four-word correlated fault, with words in fast and in device memory; - `pass_rate` against Monte Carlo, and `deferred_channels` against the flip tables; - `pieces` and `select` against direct bit selection; - the generator's `do`-block, amplitude chunks and paired projections. - **`common_prefix` test** in `test_shared.py`: against dense `reference_rotate` on the test circuits and a noiseless 6-qubit circuit with a prefix of 20 or more actions. - `uv run tsim-dev check` on an Apple M4 Pro, macOS 15.7: 2630 passed, 2 failed. The two failures are `test_sampler.py::test_rotated_surface_code_single_noise_channel_matches_stim[*-before_round_data_depolarization]` (relative difference 0.0067 > 0.005, for `cat5` as well). They fail the same way on a clean `origin/main` worktree. - Lint passes except pyright's `os.sched_getaffinity` error in `dev/tsim_dev/common.py`, which also fails on `main` on macOS. The commit skipped that one hook. ## Benchmark `uv run tsim-dev bench --ref origin/main`, default suite, Apple M4 Pro, 1 thread, tolerance 5%: 8 circuits, 0 slower, 0 faster. | circuit | k | this branch (shots/s) | main (shots/s) | rate ratio | compile ms (branch / main) | |---|---|---|---|---|---| | inject_cultivate_d3 | 4 | 3.43e7 | 3.42e7 | 1.00 | 24 / 21 | | inject_cultivate_d5 | 6 | 1.89e7 | 1.89e7 | 1.00 | 141 / 130 | | end2end_d3_dsurface15 | 4 | 1.37e6 | 1.37e6 | 0.99 | 1732 / 1644 | | hook_inject_T_d15_pd5 | 1 | 1.15e6 | 1.16e6 | 0.99 | 1952 / 1932 | | distillation_color_d5_Z | 3 | 2.42e7 | 2.41e7 | 1.00 | 24 / 23 | | coherent_d3_r3 | 4 | 2.41e7 | 2.41e7 | 1.00 | 17 / 13 | | surface_d7_r7 | 0 | 1.40e7 | 1.40e7 | 1.00 | 102 / 102 | | qv_q10 | 10 | 1.27e7 | 1.25e7 | 1.00 | 186 / 149 | Sampling is unchanged. Compile time goes up by up to 1.3× (coherent_d3_r3 +4 ms, qv_q10 +37 ms), which the harness reports as "same". The reason is that `plan_component` now schedules and plans every component twice, to try the deferred order. Only a GPU backend gains from the deferred plan. If the CPU cost matters, the deferred plan can be computed only when a GPU sampler is built. ## Not verified - Nothing ran on a CUDA GPU or, in this PR, on Metal. - `--suite full` was not run. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

## Problem On an Apple GPU, detector samples that return a subset or reordering of the outputs split each call's rows on the host: - **Bit-packed:** columns that do not lead the rows (the observables of `separate_observables=True`, every column of `prepend_observables=True`) went through `gpu.rows.select`, which unpacks every row to booleans, indexes the columns and repacks them. - **Booleans:** the GPU unpacked every output, then the host copied the requested columns out of them (`_columns`). The GPU idled meanwhile. On distance-3 cultivation postselected on every detector, `sample(2**26, separate_observables=True, bit_packed=True)` spent 0.48 s of its 0.53 s in `rows.select`, and the GPU work took 0.04 s. The most Stim-like call, `detectors, observables = sampler.sample(..., separate_observables=True)`, was 10 to 15 times slower than appended observables. On end-to-end cultivation d=5 (2745 outputs), several layouts sampled 0 to 3 M shots/s. ## Fix A new Metal kernel, `select_rows`, writes each requested column list into its own array, bit-packed (zero bits past the columns) or as one boolean per output. It runs one thread per output byte, and the column lists of a call share one command buffer. - All outputs in order are still the rows themselves (bit-packed) or `unpack_rows` (booleans), and leading whole bytes stay a view of the rows. - Calls past 32-bit indices keep the host path. - Boolean arrays stay C-contiguous. `_columns` is gone, and `MetalShotSampler.sample_accepted` takes the column lists (`columns=`) and returns one array per list. ## Numbers Apple M4 Pro, 20-core GPU, Metal. The work is on the GPU, so the numba thread count does not change these. `main` (`7ec136e6`) → this branch, in M shots/s, postselected as tsim-bench records: | cultivation d3 (21 outputs) | bit-packed | booleans | |---|---|---| | separate observables | 147 → **1316** | 97 → **1104** | | prepended observables | 55 → **1508** | 89 → **1249** | | detectors only | 1312 → 1468 | 101 → **1241** | | appended observables | 1835 → 1863 | 1424 → 1391 | | end2end d5 (2745 outputs) | bit-packed | booleans | |---|---|---| | separate observables | 59 → **291** | 3 → **156** | | prepended observables | 1 → **221** | 3 → **171** | | detectors only | 305 → 306 | 3 → **170** | | appended observables | 308 → 307 | 214 → 213 | Without postselection the same layouts gain alike: cultivation goes from 21–102 to 888–1187, and end2end from 0–7 to 14–65. Every layout of end2end is then bound near 68 M shots/s by its 344-byte rows. These grids come from a timing script, because `tsim-dev bench` only times appended observables. That bench is unchanged: `tsim-dev bench --device gpu --ref origin/main`, M4 Pro, 1 numba thread, quick suite: ``` circuit k device shots/s kept compile shots/s vs ref (min, max) compile vs ref inject_cultivate_d3_p0.001 4 metal 1.84G 0.687 0.18 0.99x [0.98, 0.99] same 1.14x same inject_cultivate_d5_p0.001 6 metal 722M 0.144 0.30 1.02x [1.00, 1.06] same 1.09x same end2end_d3_dsurface15_p0.001 4 metal 166M 0.539 2.13 1.00x [0.99, 1.00] same 1.02x same hook_inject_T_d15_pd5_p0.001 1 metal 52.7M 0.667 2.47 0.98x [0.94, 1.01] same 0.99x same distillation_color_d5_Z 3 metal 2.38G 0.0798 0.18 1.00x [0.99, 1.01] same 0.81x same coherent_d3_r3 4 metal 1.09G - 0.17 1.00x [0.99, 1.00] same 1.16x same surface_d7_r7 0 metal 727M - 0.28 1.00x [1.00, 1.00] same 1.14x same qv_q10 10 metal 11M - 0.37 1.00x [1.00, 1.00] same 1.12x same SUMMARY: 8 circuits, 0 errors, 0 slower, 0 faster, 0 noisy (tolerance 5%). ``` ## Tests - `test_metal_selects_every_layout_of_columns` checks separate, prepended and detectors-only output, bit-packed and as booleans, with and without postselection, against appended booleans (the unchanged `unpack_rows` path). Its five shapes have observables that start inside a byte, end inside one, span two bytes, or align with bytes. It also checks that booleans are C-contiguous. - `test_metal_samples_no_shots` now also checks `separate_observables=True` with zero shots, which an early version of this change returned as one array instead of two. - I also checked every layout against appended booleans for six shapes, three postselection masks and the measurement sampler (not committed). ## Verification - `uv run tsim-dev check` passes, apart from `test_dev_skills_have_valid_frontmatter`. That one fails only because my checkout has untracked local skills in `.claude/skills`, so CI will not see it. - The Metal tests run only on Apple silicon, so CI does not run the new kernel. - Not verified: CUDA (#219) has its own sampler and is unaffected; the shared `gpu/rows.py` is unchanged. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

Takes main's shared GPU code (40-bit keep thresholds, RESOLVABLE), the device argument, and Metal's docs; the CUDA backend is wired back in by the next commits.

CUDA now behaves as the Metal backend does on main: - device="cuda" samples on an NVIDIA GPU or raises with the reason, and device="auto" tries Metal, then CUDA, then the CPU, without warnings. Kernel compile failures (runtime.CompileError) move a sampler to the CPU; out of memory, forked processes and pickles follow 'required'. - The noise kernels draw outcomes from 40-bit keep thresholds, hash each block's stream, and keep rows too wide for shared memory in device memory; gather_bytes packs rows too wide to stage, so no output count falls back. - Requested output columns (separate, prepended, detectors only) and booleans are selected on the GPU (select_rows) before the copy to the host. - The reason a layout cannot run on a GPU is shared (gpu.tape.unsupported). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…t of line Bit-packed outputs in order that end on a byte boundary (detectors only, the detectors of separate_observables) are a view of the copied rows again, as on Metal, and the 40-bit keep threshold's rare tie check is __noinline__. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

TSIM_DEVICE is gone, as on main; the rotation-angle message moves into the shared gpu.tape.unsupported, which CUDA uses too; CUDA's CHANGELOG entry goes under Unreleased. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

rafaelha force-pushed the fast-amplitude-kernels branch from e1d3453 to 73656a2 Compare October 27, 2026 04:30

rafaelha force-pushed the fast-amplitude-kernels branch from 73656a2 to ab4a19d Compare October 27, 2026 05:13

Base automatically changed from fast-amplitude-kernels to main October 27, 2026 17:33

rafaelha commented Oct 28, 2026

View reviewed changes

rafaelha force-pushed the cuda-statevector-executor branch from 211fed3 to 840be38 Compare October 28, 2026 04:46

rafaelha mentioned this pull request Nov 2, 2026

Simplify the statevector executor and the sampler's backend dispatch #223

Merged

4 of 5 tasks

This was referenced Nov 3, 2026

Per-shot observable expectations and Pauli-target OBSERVABLE_INCLUDE (like Clifft's EXP_VAL) #249

Open

Define the public API and qualify seed reproducibility for 1.0.0 #268

Merged

rafaelha force-pushed the cuda-statevector-executor branch from 2092c79 to a2af975 Compare November 4, 2026 13:16

rafaelha changed the title ~~Run the statevector sampler on CUDA GPUs (device="cuda")~~ Sample on CUDA GPUs automatically with the statevector strategy Nov 4, 2026

rafaelha force-pushed the cuda-statevector-executor branch 2 times, most recently from a99fb38 to 20b3eef Compare November 5, 2026 23:14

This was referenced Nov 10, 2026

chore: add a development harness for agents working on Tsim #281

Merged

feat(statevector): add the host code that GPU backends share #284

Merged

rafaelha and others added 3 commits November 10, 2026 20:17

rafaelha force-pushed the cuda-statevector-executor branch from 20b3eef to d11d4d6 Compare November 10, 2026 21:26

rafaelha changed the base branch from main to feat/gpu-shared November 10, 2026 21:26

rafaelha and others added 6 commits November 11, 2026 03:56

refactor(statevector): assert the tape dialects' chunk size

0bff456

Both dialects are internal constants, so a ValueError and its tests guard nothing a user can reach. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

refactor(statevector): drop rows.pieces

710ddf2

Both backends return every output, in order, so each packed word is a whole shot word and the piece tables only encoded the identity. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

refactor(statevector): share the GPU samplers' fork warning

f25fcfc

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

rafaelha and others added 8 commits November 11, 2026 06:14

refactor(cuda): name the thread counts and the counts words

23841cd

The kernels' __launch_bounds__ and their launches spelled out 128, 256 and 1024 separately, and the counts buffer was sliced by bare indices. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

refactor(cuda): move the kernels' CUDA C++ to cuda/kernels.py

1d3b8ee

The kernel template and the source assembly leave sampler.py, so _module only plans, compiles and looks up the kernels. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

test(cuda): one parametrized comparison with cat5

e211b59

The detector, measurement and looping-tape comparisons repeated one body; the same 24 cases run. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

docs(cuda): shorten the changelog entry and the strategy docstrings

70e55b8

The changelog entry described the kernels, and three docstrings repeated a six-line paragraph; the install page keeps the details. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

rafaelha force-pushed the cuda-statevector-executor branch from d11d4d6 to 70e55b8 Compare November 11, 2026 13:33

rafaelha mentioned this pull request Nov 11, 2026

chore(dev): pin tsim-bench with call sizes that fit a GPU #292

Merged

rafaelha added 2 commits November 11, 2026 12:13

Merge remote-tracking branch 'origin/main' into simplify/gpu-shared

8a7f4cf

Merge remote-tracking branch 'origin/feat/gpu-shared' into simplify/cuda

27b9806

Base automatically changed from feat/gpu-shared to main November 12, 2026 04:42

This was referenced Nov 12, 2026

feat(sampler): choose the sampling device with a device argument #293

Merged

perf(metal): select the requested output columns on the GPU #294

Merged

rafaelha changed the title ~~Sample on CUDA GPUs automatically with the statevector strategy~~ CUDA backend Nov 12, 2026

rafaelha and others added 5 commits November 12, 2026 14:55

Merge origin/main into the CUDA backend

f9d8591

Takes main's shared GPU code (40-bit keep thresholds, RESOLVABLE), the device argument, and Metal's docs; the CUDA backend is wired back in by the next commits.

docs: document the CUDA backend beside Metal

9371e02

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Merge origin/main (1.1.0, #296) into the CUDA backend

6f99566

TSIM_DEVICE is gone, as on main; the rotation-angle message moves into the shared gpu.tape.unsupported, which CUDA uses too; CUDA's CHANGELOG entry goes under Unreleased. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Navigation Menu

Repository navigation

CUDA backend - #219

CUDA backend#219
rafaelha wants to merge 24 commits into
mainfrom
cuda-statevector-executor

rafaelha commented Oct 27, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Oct 27, 2026 •

edited

Loading

Built to branch `gh-pages` at 2026-11-05 23:15 UTC.
Preview will be ready when the GitHub Pages deployment is complete.

Uh oh!

github-actions Bot commented Oct 27, 2026 •

edited

Loading

Uh oh!

rafaelha left a comment

Uh oh!

rafaelha Oct 28, 2026

Uh oh!

rafaelha Oct 28, 2026

Uh oh!

rafaelha Oct 28, 2026

Uh oh!

rafaelha Oct 28, 2026

Uh oh!

rafaelha Oct 28, 2026

Uh oh!

rafaelha Oct 28, 2026

Uh oh!

codecov Bot commented Nov 11, 2026 •

edited

Loading

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

rafaelha commented Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What's in it

Changes since the last revision

Benchmark

Testing

Known limits, not fixed here

Uh oh!

github-actions Bot commented Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Built to branch gh-pages at 2026-11-05 23:15 UTC. Preview will be ready when the GitHub Pages deployment is complete.

Uh oh!

github-actions Bot commented Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

☂️ Code Coverage

Overall Coverage

New Files

Modified Files

Uh oh!

rafaelha left a comment

Choose a reason for hiding this comment

Uh oh!

rafaelha Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

rafaelha Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

rafaelha Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

rafaelha Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

rafaelha Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

rafaelha Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

codecov Bot commented Nov 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

rafaelha commented Oct 27, 2026 •

edited

Loading

github-actions Bot commented Oct 27, 2026 •

edited

Loading

Built to branch `gh-pages` at 2026-11-05 23:15 UTC.
Preview will be ready when the GitHub Pages deployment is complete.

github-actions Bot commented Oct 27, 2026 •

edited

Loading

codecov Bot commented Nov 11, 2026 •

edited

Loading