loopyt 5e2838aa5e feat(triton): reproduce cuRAND's Philox stream, so the sampler can be ported ... The last blocker for `fused_sample` was the random stream: `T.rng_init`/`T.rng_rand` lower to `curand_init`/`curand` with `curandStatePhilox4_32_10_t`, and the seeding is not documented. Its mixing function is public but `curand_init` folds the seed, the subsequence and the offset into the counter in a way that has to be measured rather than derived. Measured (`scripts/h20_rng_dump.py` dumps the real stream), the layout is: key = (seed & 0xFFFFFFFF, seed >> 32) ctr = (offset & 0xFFFFFFFF, offset >> 32, subsequence & 0xFFFFFFFF, subsequence >> 32) with draws taking the block's four words in order before `ctr.x` advances. The subsequence sits in the third and fourth counter words, not the first -- a cuRAND subsequence spans 2^64 blocks. That detail is invisible at subsequence 0 and wrong everywhere else, which is why the check covers both. `scripts/h20_philox_check.py` compares our kernel against the dumped reference over five seeds (including 2^64-1) and four subsequences (0, 1, 2048, 4097 -- the last two are the actual `sample_positions` from a run): 20/20 streams bit-identical. One trap worth noting: the multipliers and the second Weyl constant all exceed int32, so referencing them as Python ints inside the kernel silently wraps them negative. They are materialised as uint32 tiles instead. This does not yet change any output -- `fused_sample` is still the TileLang kernel. It removes the reason it had to be, which is the point.		2026-10-22 08:37:08 +08:00
..
tilelang_ops	feat(triton): reproduce cuRAND's Philox stream, so the sampler can be ported	2026-10-22 08:37:08 +08:00
__init__.py	initial commit	2026-10-20 12:09:42 +08:00
attention_backend.py	initial commit	2026-10-20 12:09:42 +08:00
layernorm.py	initial commit	2026-10-20 12:09:42 +08:00
linear.py	feat: add SM90 (Hopper) support for DeepSeek-V4.1 inference	2026-10-21 15:18:56 +08:00
mtp_info_pool.py	initial commit	2026-10-20 12:09:42 +08:00
rotary_embedding.py	initial commit	2026-10-20 12:09:42 +08:00
sampler_utils.py	initial commit	2026-10-20 12:09:42 +08:00
sampler.py	fix(debug): move debug_dump out of h20_ops so other platforms can import it	2026-10-21 22:37:49 +08:00
vocab_parallel_embedding.py	initial commit	2026-10-20 12:09:42 +08:00