Device ops

The building blocks the built-in generators are made of, public for plugin generators; Writing a Generator shows how they fit together. The ops modules import only Warp and each other, so importing track_gen.ops loads no generator module and no torch. A device op takes a LaunchContext first and launches on its stream; it never allocates, never synchronizes and never reads device data back (the conventions are listed in the module docstring below). grid_cells, hash32_host, check_array and the scratch classes’ alloc run on the host.

sample_grid_points and normalize_extent take a keyword-only size=: an [E] float32 per-env size multiplier (the built-in generators’ size, Size: the sizing rule), each value clamped at use to SIZE_BOUNDS. size=None (the default) runs the unchanged kernels, and a size-1 array gives the same bytes; normalize_extent refuses size= with count= (ValueError). A plugin generator declares its own size with SIZE_BOUNDS and passes b.param("size") to the op that applies its scale (A per-env size).

Public device primitives for generator authors (track_gen.ops).

These are the building blocks the built-in generators are made of, so a plugin generator can be assembled from the same documented pieces instead of private track_gen._src helpers. Conventions shared by every op:

  • op(ctx, ...): ctx is a track_gen.exec.LaunchContext (anything with .launch and .device works; the device is canonicalised, so "cuda" matches arrays on the current CUDA device). Every launch goes through ctx.launch, so the work lands on the context’s stream and can be recorded into a caller’s CUDA graph. An op never calls ctx.sync(), never reads device data back to the host and never allocates device memory: working buffers are passed in (e.g. an ArcScratch allocated once in alloc_scratch).

  • Flat per-env layout: env e owns [e*stride, (e+1)*stride) of a flat array of E*stride elements; per-env scalars are [E] arrays.

  • Argument names: src is the input point array and stride its points per env; num is an exact number of output points per env and out_stride an output row capacity (NaN-padded past the env’s count); outputs are out (plus count_out/truncated_out for per-env results).

  • E is src.shape[0] // stride (which must divide exactly) for every op that takes src, and seeds.shape[0] for the samplers; every other per-env array must then be exactly [E].

  • Carried channels: tables= (a ResampleTables) makes a resample also record where each output point lies on its source loop, and resample_channels() carries per-point channels through those tables.

  • Count semantics: count is always keyword-only. Where it is given, only the first count[e] points of a row are real (a count above the stride is clamped to it); slots past it are written (nan, nan) by ops that produce points. Masks are the keyword-only active (active[e] > 0 selects the env).

  • Validation is host-side and happens before any launch: data arrays must match their exact shape, dtype and ctx.device; scratch arrays need the dtype, the device and at least the stated length. Mismatches raise ValueError naming the argument, except that fill raises TypeError for an unsupported dtype and any_set / all_set for an argument that is not an int32 wp.array.

Device functions for your own kernels live in track_gen.ops.funcs; shared numeric conventions in track_gen.ops.constants.

Resampling

class ArcScratch[source]

Bases: object

Working buffers for the resample ops (wraps existing arrays; never allocates).

For E envs and an input stride stride (points per env of the op’s src):

  • real: [>= E*stride] vec2f – compacted real (finite) points (arc_resample() only).

  • seg: [>= E*stride] float32 – per-segment lengths.

  • s: [>= E*(stride+1)] float32 – cumulative arc length.

  • count_real: [>= E] int32 – real-point count per env (arc_resample() only; readable after the call).

resample_uniform() and resample_constant_spacing() only use seg and s; real/count_real may be None for them.

Build one per generator in alloc_scratch with alloc(), or wrap buffers you already own (views/slices included) with the constructor.

__init__(real, seg, s, count_real)[source]
Return type:

None

real
seg
s
count_real
classmethod alloc(E, stride, device, *, compact=True)[source]

Allocate scratch for E envs of stride input points each (construction time only: this allocates).

Parameters:

  • E (int) – number of envs.

  • stride (int) – input points per env (the stride of the op’s src).

  • device – Warp device.

  • compact (bool) – also allocate real and count_real (needed by arc_resample()); False leaves them None.

Return type:

ArcScratch

arc_resample(ctx, src, stride, num, scratch, out, count_out, *, active=None)[source]

NaN-aware arc-length resample of closed loops: stride points in, num out.

Per env, the points whose x AND y are finite are compacted in order (interior NaN rows are dropped too) into R real points. With R >= 2 the env’s output row is num points uniformly spaced by arc length along the CLOSED loop through them (point 0 is the first real point; the last segment wraps back to it) and count_out[e] = num. With R < 2 the row is all NaN and count_out[e] = 0. scratch.count_real[e] receives R.

With active, only envs with active[e] > 0 are resampled; every other env keeps its out row bit for bit and gets count_out[e] = scratch.count_real[e] = 0.

Two launches (E scan threads, E*num lookup threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f input loops (NaN rows allowed); E is src.shape[0] // stride.

  • stride (int) – input points per env (>= 1).

  • num (int) – output points per env (>= 1).

  • scratch (ArcScratch) – ArcScratch with real >= E*stride, seg >= E*stride, s >= E*(stride+1), count_real >= E.

  • out (wp.array) – [E*num] vec2f output; must not alias a scratch buffer (it may alias src: the lookup reads only the scratch).

  • count_out (wp.array) – [E] int32 output.

  • active (wp.array | None) – optional [E] int32 env mask.

Raises:

ValueError – a wrong type, shape, dtype or device, stride/num below 1, or out aliasing the scratch (host-side, before any launch).

Return type:

None

resample_uniform(ctx, src, stride, scratch, out, *, count, tables=None)[source]

Count-aware arc-length-uniform resample of closed loops (stride in and out).

Env e’s first count[e] points (which must be finite; no NaN compaction) are treated as a closed loop and resampled to count[e] points uniformly spaced by arc length, starting at point 0. Slots >= count[e] of the row are written NaN; count[e] <= 0 gives an all-NaN row and a count above stride is clamped to stride.

Two launches (E scan threads, E*stride lookup threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f input loops; E is src.shape[0] // stride.

  • stride (int) – points per env of src and out (>= 1).

  • scratch (ArcScratch) – ArcScratch with seg >= E*stride and s >= E*(stride+1) (real/count_real unused, may be None).

  • out (wp.array) – [E*stride] vec2f output; must not alias src or the scratch.

  • count (wp.array) – [E] int32 real points per env (keyword-only, required).

  • tables (ResampleTables | None) – optional ResampleTables of E*stride points over n_src = E*stride, written too (the lookup’s table-writing variant: the same out bytes); carry channels through them with resample_channels().

Raises:

ValueError – a wrong type, shape, dtype or device, stride < 1, or out aliasing src or the scratch (host-side, before any launch).

Return type:

None

resample_constant_spacing(ctx, src, stride, spacing, out_stride, scratch, out, count_out, truncated_out, *, tables=None)[source]

Resample fully-real closed loops of stride points to a constant spacing.

Per env: total = closed-loop perimeter; count = floor(total/spacing) + 1 trimmed so (count-1)*spacing < total, then capped at out_stride. Point k sits at arc length k*spacing from point 0. truncated_out[e] = 1 iff the uncapped count exceeded out_stride (the loop’s tail was dropped). A loop whose perimeter is not > 0 (zero or NaN) gets count_out[e] = 0 and truncated_out[e] = 0. Slots >= count_out[e] are written NaN.

Two launches; no allocation, no host sync (the truncation flag stays on device).

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f input loops (every point real); E is src.shape[0] // stride.

  • stride (int) – input points per env (>= 1).

  • spacing (float) – target spacing (finite, > 0).

  • out_stride (int) – output row capacity per env (>= 1).

  • scratch (ArcScratch) – ArcScratch with seg >= E*stride and s >= E*(stride+1).

  • out (wp.array) – [E*out_stride] vec2f output; must not alias src or the scratch.

  • count_out (wp.array) – [E] int32 output.

  • truncated_out (wp.array) – [E] int32 output (0/1).

  • tables (ResampleTables | None) – optional ResampleTables of E*out_stride points over n_src = E*stride, written too (the lookup’s table-writing variant: the same out bytes); carry channels of the source loop through them with resample_channels().

Raises:

ValueError – a wrong type, shape, dtype or device, a bad scalar, or out aliasing src or the scratch (host-side, before any launch).

Return type:

None

Carrying per-point channels through a resample (tables= above): the tables record where each output point lies on its source loop, and resample_channels interpolates any channel of the source through them; see Carried channels.

class ResampleTables[source]

Bases: object

Where each output point of a resample lies on its source loop (wraps existing arrays; never allocates except alloc()).

For an output of n points over a source of n_src points:

  • i0, i1: [n] int32 – the GLOBAL source indices of the segment ends (e * stride + idx and e * stride + (idx + 1) % count: the tables describe CLOSED loops);

  • frac: [n] float32 – the point’s clamped fraction along the segment (never NaN at a real point);

  • n_src: the source channel length the indices point into.

A padding slot (past the env’s count, or an env that produced no point) holds the sentinel i0 = i1 = -1, frac = NaN: every table-writing op writes it on every exit path, and resample_channels() writes the channel’s padding there (NaN for floats and every vector component, -1 for int32).

__init__(i0, i1, frac, n_src)[source]
Parameters:

n_src (int)

Return type:

None

i0
i1
frac
n_src
classmethod alloc(n_out, n_src, device)[source]

Allocate tables for n_out output points over n_src source points, initialised to the sentinel (construction time only: this allocates).

Parameters:
Return type:

ResampleTables

resample_channels(ctx, tables, pairs, interp)[source]

Carry per-point channels of a resample’s source loop to its output points.

tables come from a table-writing resample (tables= of resample_constant_spacing() / resample_uniform()). For every (src, dst) pair, with a = src[i0], b = src[i1] and f = frac at each output point, dst receives, by policy:

  • "linear" (float32, vec2f, vec3f): a + f * (b - a);

  • "nearest" (any): a if f <= 0.5 else b (a tie goes to the segment start);

  • "previous" (any): a (zero-order hold);

  • "angle" (float32 radians in [-pi, pi)): wrap(a + f * wrap(b - a)), the short way round, output in [-pi, pi).

Padding slots (i0 < 0) get NaN (floats, every vector component) or -1 (int32). linear and angle blend across the closing segment of the loop. One launch per pair (len(tables.i0) threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • tables (ResampleTables) – the ResampleTables (n output points, n_src source points).

  • pairs (Sequence[tuple[wp.array, wp.array]]) – (src, dst) arrays: src [n_src], dst [n] of the same dtype (float32, int32, vec2f or vec3f); dst must not alias src.

  • interp (Sequence[str]) – one policy per pair.

Raises:

  • ValueError – a wrong type, shape or device, pairs and interp of different lengths, an unknown policy, or dst aliasing src (host-side, before any launch).

  • TypeError – a dst dtype other than src’s, or a policy the dtype does not admit.

Return type:

None

Polygon geometry

self_intersections(ctx, src, stride, out, *, count)[source]

Number of proper self-crossings of each closed polygon.

For env e the polygon is its first N = min(count[e], stride) points (closed: the last edge returns to point 0). out[e] counts the pairs of edges i < j with circular distance > 1 (adjacent edges never count) that cross properly: each segment’s endpoints must straddle the other’s line by more than 1e-3 * len^2 of that segment (a scale-relative collinearity tolerance that rejects float32 sign flips on near-collinear edges). Coordinates are read through nan0() (NaN reads as 0). N <= 3 gives 0. Deterministic (integer atomics).

Two launches (zero out; E*stride threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f polygons; E is src.shape[0] // stride.

  • stride (int) – points per env in src (>= 1).

  • out (wp.array) – [E] int32 crossing counts.

  • count (wp.array) – [E] int32 points per env (keyword-only, required; clamped to stride).

Raises:

ValueError – a wrong type, shape, dtype or device (host-side, before any launch).

Return type:

None

turning_number(ctx, src, stride, out, *, count)[source]

Signed total turning (radians) of each closed polygon.

Sum over the N = min(count[e], stride) vertices of the exterior angle between consecutive edges, each wrapped into (-pi, pi]: about +2*pi for a simple counter-clockwise loop, -2*pi clockwise. A zero-length edge has angle 0. N <= 0 gives 0; a NaN point makes the env’s result NaN. One launch (E threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f polygons; E is src.shape[0] // stride.

  • stride (int) – points per env in src (>= 1).

  • out (wp.array) – [E] float32 output.

  • count (wp.array) – [E] int32 points per env (keyword-only, required; clamped to stride).

Raises:

ValueError – a wrong type, shape, dtype or device.

Return type:

None

thickness(ctx, src, stride, band, out, *, count)[source]

Thickness of each closed polygon: min(1 / max Menger curvature, 0.5 * min |p_i - p_j|) over vertex pairs with circular index distance > band[e].

band is the neighbour half-window excluded from the distance term, in points. The built-in validity gate uses band = max(1, round(2 * half_width / l0)) for a track of half-width half_width whose segments have mean length l0: pairs closer than a track width along the loop are skipped.

Degenerate cases, with N = min(count[e], stride): the curvature maximum is floored at 1e-12 and the distance term starts at 1e30, so a polygon with no curvature and no vertex pair beyond the band (e.g. N < 3 with band[e] >= 1) returns 1e12. NaN points are not guarded: a row containing one gives an unspecified, generally finite, result (the running min/max mostly skip NaN distances and curvatures instead of propagating them). One launch (E threads, O(N^2) work per env); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f polygons; E is src.shape[0] // stride.

  • stride (int) – points per env in src (>= 1).

  • band (wp.array) – [E] int32 excluded neighbour half-window per env.

  • out (wp.array) – [E] float32 output.

  • count (wp.array) – [E] int32 points per env (keyword-only, required; clamped to stride).

Raises:

ValueError – a wrong type, shape, dtype or device.

Return type:

None

normalize_extent(ctx, src, stride, target_extent, *, count=None, size=None)[source]

Centre each env’s bounding box at the origin and scale it uniformly so its longest side equals target_extent (in place).

The extent is floored at 1e-8 before dividing. With count=None all stride points of each env are used (rows are expected to be finite; a NaN coordinate stays NaN). With count, only the first clamp(count[e], 0, stride) points are used and the remaining slots are written NaN. One launch (E threads); no allocation, no host sync.

size is a per-env size multiplier of the target: env e’s longest side becomes target_extent * size[e], each value clamped at use to track_gen.ops.constants.SIZE_BOUNDS (by comparisons: a NaN passes). The size multiplies the finite scalar target before the division (no per-point arithmetic changes), so a size-1 array gives the size=None bytes exactly, on CPU and CUDA; size=None runs the unchanged kernels.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f points, modified in place; E is src.shape[0] // stride.

  • stride (int) – points per env (>= 1).

  • target_extent (float) – longest bbox side after scaling (finite), e.g. ops.constants.TARGET_EXTENT * config.scale.

  • count (wp.array | None) – optional [E] int32 real points per env (keyword-only).

  • size (wp.array | None) – optional [E] float32 per-env size multiplier of the target (keyword-only), on the op’s device; not combinable with count.

Raises:

ValueError – a wrong type, shape, dtype or device, src.shape[0] not a multiple of stride, or count and size both given.

Return type:

None

angle_sort(ctx, src, stride, keys, out, *, count, out_stride=None)[source]

Sort each env’s first points by angle around their centroid.

With m = clamp(count[e], 0, min(stride, out_stride)), env e’s output row is its first m source points ordered ascending by atan2(x - cx, y - cy) (x FIRST: the built-in generators’ ordering), where (cx, cy) is the float64 centroid of those m points; the sort is stable. Slots [m, out_stride) are written NaN, so count[e] <= 0 gives an all-NaN row. One launch (E threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f points; E is src.shape[0] // stride.

  • stride (int) – points per env in src (>= 1).

  • keys (wp.array) – [>= E*out_stride] float32 scratch.

  • out (wp.array) – [E*out_stride] vec2f output.

  • count (wp.array) – [E] int32 points to sort per env (keyword-only, required).

  • out_stride (int | None) – points per env in out (default stride).

Raises:

ValueError – a wrong type, shape, dtype or device.

Return type:

None

Random sampling

Each env’s stream is wp.rand_init(seeds[e] * salt + offset) in int32 wrap-around arithmetic, the keying the built-in generators use. The (salt, offset) pairs the built-ins already use are listed in the generator guide.

grid_cells(min_point_distance)[source]

Grid for sample_grid_points() from the built-in min_point_distance knob (host only).

Returns (num_cells, cell_size) = (int(1 / (2*d)), 2*d): the built-in generators’ grid of num_cells x num_cells cells of side 2*d. The name follows the config knob; it is not a distance guarantee (see sample_grid_points()).

Raises:

ValueError – min_point_distance is not a finite float in (0, 0.5] (the range TrackGenConfig accepts; above 0.5 the grid would be empty).

Parameters:

min_point_distance (float)

Return type:

tuple[int, float]

sample_counts(ctx, seeds, salt, offset, lo, hi, out)[source]

One uniformly drawn integer in [lo, hi] per env.

out[e] = min(lo + int(randf * (hi - lo + 1)), hi) with one randf from wp.rand_init(seeds[e] * salt + offset) (int32 wrap-around). Deterministic per (seed, salt, offset). One launch (E threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • seeds (wp.array) – [E] int32 per-env seeds; E is seeds.shape[0].

  • salt (int) – int32 stream multiplier.

  • offset (int) – int32 stream offset.

  • lo (int) – smallest value (int32).

  • hi (int) – largest value (int32, >= lo).

  • out (wp.array) – [E] int32 output.

Raises:

ValueError – a wrong type, shape, dtype or device, lo > hi or a value out of int32 range (host-side, before any launch).

Return type:

None

sample_grid_points(ctx, seeds, salt, offset, num, num_cells, cell_size, scale, used, out, *, size=None)[source]

num jittered grid points per env, on distinct cells where possible.

Per env, from wp.rand_init(seeds[e] * salt + offset) (int32 wrap-around): for each point c < num draw a cell min(int(randf * num_cells**2), num_cells**2 - 1), redrawing (up to 8 times) while it equals the cell of an earlier point of the env; then draw nx, ny = randf - 0.5 and write ((cell % num_cells) * cell_size + nx, (cell // num_cells) * cell_size + ny) * scale. used is filled with -1 first and receives the chosen cells.

Guarantees: distinct cells are best effort (a point keeps a duplicate cell after 8 redraws), and the jitter is an absolute [-0.5, 0.5) (not scaled by cell_size), so there is no minimum distance between points. Each coordinate lies in [-0.5, (num_cells - 1) * cell_size + 0.5) * scale. Use grid_cells() for the built-in grid. Two launches (fill, E threads); no allocation, no host sync.

size is a per-env size multiplier applied LAST: env e’s points are ((x * cell_size + nx) * scale) * size[e], each value clamped at use to track_gen.ops.constants.SIZE_BOUNDS (by comparisons: a NaN passes). size=None runs the unchanged kernels (same launches, same bytes); a size-1 array gives the size=None bytes exactly, on CPU and CUDA: the product feeds no add / sub, so it creates no FMA contraction, and x * 1.0f is exact. A plugin applying its own size should keep that rule (multiply where no add or subtract consumes the product).

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • seeds (wp.array) – [E] int32 per-env seeds; E is seeds.shape[0].

  • salt (int) – int32 stream multiplier.

  • offset (int) – int32 stream offset.

  • num (int) – points per env (>= 1).

  • num_cells (int) – grid cells per side (>= 1).

  • cell_size (float) – cell side length.

  • scale – float applied to every env, or [E] float32 per-env scale (bit-identical to the scalar form for a uniform array).

  • used (wp.array) – [>= E*num] int32 scratch (chosen cell per point).

  • out (wp.array) – [E*num] vec2f output.

  • size (wp.array | None) – optional [E] float32 per-env size multiplier (keyword-only), on the op’s device.

Raises:

ValueError – a wrong type, shape, dtype or device, or a bad scalar (host-side, before any launch).

Return type:

None

Curves

closed_catmull_rom(ctx, src, stride, per_seg, out, *, count=None, endpoint=False)[source]

Sample the closed uniform Catmull-Rom spline (tension 0.5) through each env’s controls.

out[e, i*per_seg + s] = catmull_rom(p[i-1], p[i], p[i+1], p[i+2]; u) with indices mod K and u = s / (per_seg - endpoint) (see the module docstring). count=None: K = stride for every env. With count: K = min(count[e], stride); segments i >= K (all of them when K < 2) are written NaN, which the NaN-aware arc_resample() drops. One launch (E*stride*per_seg threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f controls; E is src.shape[0] // stride (must divide exactly).

  • stride (int) – controls per env (>= 1).

  • per_seg (int) – samples per segment (>= 1; >= 2 with endpoint=True).

  • out (wp.array) – [E*stride*per_seg] vec2f output.

  • count (wp.array | None) – optional [E] int32 controls per env (keyword-only).

  • endpoint (bool) – include u = 1 in each segment.

Raises:

ValueError – a wrong type, shape, dtype or device, or a bad scalar.

Return type:

None

closed_polyline(ctx, src, stride, per_seg, out, *, count=None, endpoint=False, active=None)[source]

Sample the closed straight polyline through each env’s controls.

out[e, i*per_seg + s] = (1-u)*p[i] + u*p[i+1] (index mod K) with u = s / (per_seg - endpoint). count=None: K = stride; with count, K = min(count[e], stride) and segments i >= K (all when K < 2) are written NaN. Same layout as closed_catmull_rom(). With active, only envs with active[e] > 0 are written and every other row is left untouched: the drop-in “straight” fallback for the envs whose smooth curve self-crossed. One launch (E*stride*per_seg threads); no allocation, no host sync.

Parameters:

  • ctx (LaunchContext) – the LaunchContext to launch on.

  • src (wp.array) – [E*stride] vec2f controls; E is src.shape[0] // stride.

  • stride (int) – controls per env (>= 1).

  • per_seg (int) – samples per segment (>= 1; >= 2 with endpoint=True).

  • out (wp.array) – [E*stride*per_seg] vec2f output.

  • count (wp.array | None) – optional [E] int32 controls per env (keyword-only).

  • endpoint (bool) – include u = 1 in each segment.

  • active (wp.array | None) – optional [E] int32 env mask (keyword-only).

Raises:

ValueError – a wrong type, shape, dtype or device, or a bad scalar.

Return type:

None

Fills, flags and validation

fill(ctx, arr, value)[source]

Set every element of the 1-D array arr to value on ctx’s stream.

One launch of arr.shape[0] threads; no allocation, no host sync.

Parameters:

  • ctx – the LaunchContext to launch on.

  • arr (array) – [n] (1-D) wp.array of int32, uint32, float32, vec2f, vec3f or quatf on ctx.device.

  • value – a scalar; for vec2f/vec3f a vector of that type, a sequence of its components, or one scalar broadcast to every component; for quatf a wp.quatf or a sequence of its four components (no broadcast).

Raises:

  • TypeError – unsupported dtype.

  • ValueError – arr is not a wp.array, not 1-D or not on ctx.device, or a scalar value for a quatf array.

Return type:

None

any_set(ctx, mask, flag_out)[source]

flag_out[0] = 1 if any mask[e] != 0, else 0 (on ctx’s stream).

Two launches, both capturable and allocation-free: zero flag_out, then one thread per element raises it with an integer atomic_max (order-independent, so the result is deterministic).

Parameters:

  • ctx – the LaunchContext to launch on.

  • mask (array) – 1-D int32 array (e.g. an [E] per-env mask).

  • flag_out (array) – [1] int32 array on the same device.

Raises:

  • TypeError – mask or flag_out is not an int32 wp.array.

  • ValueError – a wrong shape, or the arrays and ctx on different devices (host-side, before any launch).

Return type:

None

all_set(ctx, mask, flag_out)[source]

flag_out[0] = 1 if every mask[e] != 0 (an empty mask: 1), else 0 (on ctx’s stream).

Two launches, both capturable and allocation-free: set flag_out to 1, then one thread per element lowers it with an integer atomic_min where mask[e] == 0 (order-independent, so the result is deterministic).

Parameters:

  • ctx – the LaunchContext to launch on.

  • mask (array) – 1-D int32 array (e.g. an [E] per-env mask).

  • flag_out (array) – [1] int32 array on the same device.

Raises:

  • TypeError – mask or flag_out is not an int32 wp.array.

  • ValueError – a wrong shape, or the arrays and ctx on different devices (host-side, before any launch).

Return type:

None

check_array(name, arr, shape, dtype, device)[source]

Validate a wp.array’s type, shape, dtype, and device.

Raises ValueError naming name on the first mismatch. Shared body behind the family’s _validate_inputs / _validate_position / _check_arr wrappers; shape is the expected shape tuple. device is canonicalised before the comparison, so "cuda" matches an array on cuda:0 when that is the current CUDA device.

Parameters:

  • name (str) – argument name used in the error message.

  • arr – the value to check.

  • shape (tuple) – the exact expected shape, e.g. (E * N,).

  • dtype – the expected Warp dtype (compared by identity, e.g. wp.vec2f).

  • device – the expected device (a wp.Device or device string, e.g. "cuda", "cuda:0" or "cpu").

Raises:

ValueError – arr is not a wp.array, or its shape, dtype or device differs from the expected one.

Return type:

None

Hashing

hash32(x: uint32) → uint32[source]

@wp.func: the MurmurHash3 fmix32 finaliser, a uint32 bijection with full avalanche (wrap-around arithmetic). Call it inside a kernel; it is also exported as track_gen.ops.funcs.hash32.

hash32_host(x)[source]

Host twin of hash32() (x is reduced to uint32 first).

Parameters:

x (int)

Return type:

int

Device functions: track_gen.ops.funcs

@wp.func building blocks for plugin kernels.

Each is a wp.Function: call it from inside your own @wp.kernel (or another @wp.func), e.g. from track_gen.ops import funcs then funcs.catmull_rom(...) in the kernel body. They are the device functions the built-in generators use: the same operations in the same order. On CPU a plugin kernel that composes them reproduces the built-in arithmetic bit for bit. On CUDA (the module option fuse_fp, default True; wp.set_module_options({"fuse_fp": ...})) the compiler may contract a multiply and an add into one FMA, and which pairs it contracts depends on the code around the inlined call, so a different surrounding kernel can round differently in the last bit.

The three crossing predicates differ on purpose: segments_cross() is the strict sign test used by gate validity; the polygon crossing counter (track_gen.ops.self_intersections()) adds a scale-relative collinearity tolerance on top of ccw().

ccw(ox: float, oy: float, px: float, py: float, qx: float, qy: float) → float[source]

Orientation of (o, p, q): (qy - oy)*(px - ox) - (py - oy)*(qx - ox), twice the signed triangle area; positive when o -> p -> q turns counter-clockwise.

cross2(a: vec2f, b: vec2f) → float[source]

2-D cross product a.x*b.y - a.y*b.x.

nan0(x: float) → float[source]

x if finite, else 0.0 (a no-op on finite inputs).

safe_normalize2(v: vec2f) → vec2f[source]

v / max(|v|, 1e-8): a unit vector with a length floor. An all-NaN vector (a pruned corner) stays all-NaN. A NaN length is floored to 1e-8 like a short one, so a vector with one NaN component keeps the other component, scaled by 1e8.

catmull_rom(p0: vec2f, p1: vec2f, p2: vec2f, p3: vec2f, u: float) → vec2f[source]

Uniform Catmull-Rom (tension 0.5) segment from p1 to p2 at u in [0, 1]: u = 0 gives p1, u = 1 gives p2. The function behind track_gen.ops.closed_catmull_rom().

segments_cross(a: vec2f, b: vec2f, c: vec2f, d: vec2f) → int[source]

1 if segments ab and cd cross properly, else 0. A strict sign test: touching endpoints and collinear overlap are not crossings. It is the gate validity predicate, not the one track_gen.ops.self_intersections() uses.

hash32(x: uint32) → uint32[source]

The same function as track_gen.ops.hash32().

These are the operations the built-in kernels perform, in the same order. A plugin kernel built from them matches the built-ins bit for bit on CPU. On CUDA, Warp compiles with FMA contraction on by default (the module option fuse_fp, default True; wp.set_module_options({"fuse_fp": ...})), and which multiply-add pairs become FMAs depends on the code around the inlined call, so a differently structured kernel can differ in the last bit.

Constants

Shared numeric conventions of the built-in generators.

TARGET_EXTENT is the longest bounding-box side (per unit config.scale) that the polar, voronoi and checkpoint generators (and their gate generators) centre and scale their output to. bezier and hull land near it by construction of their sampling grid (a longest side of about 1.0 to 1.8 at scale 1, not centred), and repulsive is larger (about 2.2). The downstream defaults (half_width, spacing, N_max, relaxation bands) are tuned for coordinates of that size, so a plugin generator that wants those defaults to behave the same should end with ops.normalize_extent(ctx, out_centerline, N, ops.constants.TARGET_EXTENT * config.scale).

SIZE_BOUNDS is the declared range of the per-env size multiplier every built-in generator except external (and normalize.extent) declares, and the range ops.sample_grid_points(size=) / ops.normalize_extent(size=) clamp each env’s value to (by comparisons: a NaN passes). It is the convention value a plugin generator imports for its own size field:

import dataclasses

from track_gen import plugin
from track_gen.ops.constants import SIZE_BOUNDS

@dataclasses.dataclass(frozen=True, kw_only=True)
class MyParams(plugin.Params):
    size: float = plugin.dynamic(1.0, lo=SIZE_BOUNDS[0], hi=SIZE_BOUNDS[1])

A legality range: the useful range is set by n_max (truncation), num_points, the absolute half-width and the utilities’ capacities (the sizing rule in the tuning guide).

TARGET_EXTENT = 1.44

Longest bbox side of a built-in centerline at config.scale == 1 (the bezier baseline’s typical extent).

SIZE_BOUNDS = (0.1, 10.0)

the declared bounds of every built-in size field and the clamp of the size= arrays of ops.sample_grid_points / ops.normalize_extent.

Type:

(lo, hi) of the per-env size multiplier