tilelang.contrib.cutedsl.rand¶

Philox4x32-10 random numbers for the CuTeDSL backend.

This is a port of the curandStatePhilox4_32_10_t part of curand_kernel.h. The generator state lives in a 16-word register array because it has to survive across conditional code, which is how curand keeps it in a local struct.

Integers and uniform floats reproduce curand bit for bit. The normals run the same Box-Muller transform with the same cached sibling and the same fast sin/cos, but that approximation is not bit-identical to curand’s __sincosf for every input, and the double pair uses sincospi, which CuTeDSL does not expose; both agree to a few ulps and the generator position always matches.

Functions¶

`rng_state`(*[, loc, ip])	Allocate the generator state as a register array.
`rng_init`(state, seed, seq, off[, generator, loc, ip])	Initialize the generator, like `curand_init(seed, seq, off, &state)`.
`rng_rand`(state, *[, loc, ip])	Return the next 32-bit unsigned value, like `curand(&state)`.
`rng_uniform2_double`(state, out, *[, loc, ip])	Fill two doubles with `curand_uniform2_double`, which draws one block.
`rng_normal2`(state, out, *[, loc, ip])	Two float32 normals, like `curand_normal2`.
`rng_normal4`(state, out, *[, loc, ip])	Four float32 normals from one block, like `curand_normal4`.
`rng_normal2_double`(state, out, *[, loc, ip])	Two float64 normals from one block, like `curand_normal2_double`.
`rng_rand_float`(state[, bit, dist, loc, ip])	Return the next random float, like `curand_uniform` / `curand_normal`.

Module Contents¶

tilelang.contrib.cutedsl.rand.rng_state(*, loc=None, ip=None)¶: Allocate the generator state as a register array.

tilelang.contrib.cutedsl.rand.rng_init(state, seed, seq, off, generator='curandStatePhilox4_32_10_t', *, loc=None, ip=None)¶

Initialize the generator, like curand_init(seed, seq, off, &state).

Return type:: None

tilelang.contrib.cutedsl.rand.rng_rand(state, *, loc=None, ip=None)¶: Return the next 32-bit unsigned value, like curand(&state).

tilelang.contrib.cutedsl.rand.rng_uniform2_double(state, out, *, loc=None, ip=None)¶

Fill two doubles with curand_uniform2_double, which draws one block.

The 64-bit uniform store is the one place where curand’s vector helper is not the elementwise composition of its scalar form: it takes 53-bit samples from a whole Philox block instead of one 32-bit word per result.

Return type:: None

tilelang.contrib.cutedsl.rand.rng_normal2(state, out, *, loc=None, ip=None)¶

Two float32 normals, like curand_normal2.

curand’s vector helpers generate a fresh pair and keep the scalar Box-Muller cache as it is, so a later scalar draw still returns its sibling.

Return type:: None

tilelang.contrib.cutedsl.rand.rng_normal4(state, out, *, loc=None, ip=None)¶

Four float32 normals from one block, like curand_normal4.

Return type:: None

tilelang.contrib.cutedsl.rand.rng_normal2_double(state, out, *, loc=None, ip=None)¶

Two float64 normals from one block, like curand_normal2_double.

Return type:: None

tilelang.contrib.cutedsl.rand.rng_rand_float(state, bit=32, dist='uniform', *, loc=None, ip=None)¶: Return the next random float, like curand_uniform / curand_normal.