v1.15.0

[shi-eric]

Warp v1.15.0

Warp v1.15 adds opt-in deterministic execution for supported atomic patterns, broader replay for CPU and saved CUDA graphs, CUDA managed-memory allocation, and mipmapped textures. It also extends Warp functions with in-place references and function arguments, adds struct support to tile operations, and lets sparse matrices reserve BSR row capacity.

New features

Reproducible atomic execution

Simulation, validation, and regression-test workloads can now trade some performance for reproducible ordering of common atomic updates. wp.DeterministicMode.RUN_TO_RUN produces bit-exact repeated results on the same GPU architecture, while wp.DeterministicMode.GPU_TO_GPU uses a more conservative path intended to preserve results across GPU architectures (#1443).

Note

Deterministic mode changes the execution algorithm and uses temporary storage, so its performance impact depends on the atomic pattern and contention. Highly contended accumulations may run faster because deterministic reduction replaces a serialized atomic hotspot with parallel sort-and-reduce work. Sparse accumulations usually slow down because normal atomics have little contention, while deterministic mode still pays for recording and sorting. Counter-based slot allocation also requires counting, scanning, and replaying the kernel to assign stable slots, so it usually slows down. The more conservative GPU_TO_GPU mode can also be slower than RUN_TO_RUN. Benchmark affected kernels before enabling deterministic mode broadly.

The following kernel compacts even thread IDs into an output array. Since each thread uses the return value from wp.atomic_add() as its slot, deterministic mode assigns the same output ordering across repeated launches.

import warp as wp

wp.config.deterministic = wp.DeterministicMode.RUN_TO_RUN


@wp.kernel
def compact_even(count: wp.array[wp.int32], output: wp.array[wp.int32]):
    tid = wp.tid()
    if tid % 2 == 0:
        slot = wp.atomic_add(count, 0, 1)
        output[slot] = tid


num_threads = 4_096
count = wp.zeros(1, dtype=wp.int32, device="cuda:0")
output = wp.empty(num_threads // 2, dtype=wp.int32, device="cuda:0")

runs = []
for _ in range(5):
    count.zero_()
    wp.launch(compact_even, dim=num_threads, inputs=[count], outputs=[output], device="cuda:0")
    runs.append(output.numpy())

reference = runs[0].tobytes()
if any(result.tobytes() != reference for result in runs[1:]):
    raise RuntimeError("Deterministic launches produced different results")

print(runs[0][:8])  # [ 0  2  4  6  8 10 12 14]

Key capabilities:

• Set the application-wide default with wp.config.deterministic before defining or importing kernel modules. Use wp.set_module_options() to override an existing module, or module="unique" with module_options for intentional per-kernel isolation.
• With deterministic mode enabled, Warp records supported wp.atomic_add(), wp.atomic_sub(), wp.atomic_min(), and wp.atomic_max() updates, sorts them by destination and thread, and combines the updates for each destination with a deterministic reduction instead of applying them in CUDA's scheduler-dependent atomic order. It handles += and -= updates the same way, while counter-based allocation uses count, scan, and replay passes to assign stable output slots.
• Generated and custom-adjoint accumulation patterns are supported. Backward replay of a consumed-return counter needs an explicit replay function and stored slot mapping.
• Dynamic-loop capacity: If a data-dependent loop can revisit the same deterministic atomic site, set wp.config.deterministic_max_records before defining or importing affected modules. The value is the maximum number of records one thread can emit for one atomic target, such as an output array. The default 0 uses Warp's code-generated estimate. For example, use wp.config.deterministic_max_records = 8 when each thread can execute that atomic at most eight times. A bound that is too small raises an overflow error outside CUDA graph capture. Capture and replay cannot report overflows, so size the bound conservatively. Larger bounds allocate more temporary storage and can increase captured-launch work.
• Unsupported paths: wp.atomic_cas(), wp.atomic_exch(), and tile atomics are not rewritten. Deterministic CUDA kernels cannot currently run inside conditional body graphs or APIC serialization. See the deterministic execution guide for the full limits.

Replay more operations in CPU and saved CUDA graphs

Important

This is an experimental feature. The API may change without a formal deprecation cycle.

CPU graphs can now replay fills, scans, sorts, run-length encoding, sparse topology updates, adjoint launches, reusable recorded launches, and conditional graph nodes. Saved CUDA .wrp graphs can now serialize and rebuild supported helper operations, conditional nodes, adjoint launches, and reusable record_cmd=True launches (#1431).

The captured scan below reads the current contents of src each time the graph launches. Changing the input produces a new prefix sum without recapturing the graph.

import numpy as np
import warp as wp

src = wp.array(np.ones(8, dtype=np.int32), device="cpu")
prefix = wp.zeros(8, dtype=wp.int32, device="cpu")

with wp.ScopedCapture(device="cpu") as capture:
    wp.utils.array_scan(src, prefix, inclusive=True)

wp.capture_launch(capture.graph)
print(prefix.numpy())  # [1 2 3 4 5 6 7 8]

src.fill_(2)  # Change the input without recapturing the graph.
wp.capture_launch(capture.graph)
print(prefix.numpy())  # [ 2  4  6  8 10 12 14 16]

Notable replay limitations:

• CPU graphs: Capture is scoped to one device, and referenced arrays must remain alive for the graph's lifetime. Scans support int32, float32, int64, and float64 scalar and vector arrays, but not negative strides. Non-contiguous copies and fills, along with the host-return form of wp.utils.runlength_encode(), are unsupported.
• Saved CUDA graphs: Cross-device and non-contiguous copies, non-contiguous fills, the host-return form of wp.utils.runlength_encode(), and CUDA wp.sparse.bsr_set_from_triplets() are unsupported. Stream events and texture array copies are not recorded, and only wp.Mesh object handles can be serialized.

See the runtime guide sections on CPU Graphs and Saving and Loading Graphs for the complete lists of recorded operations and limitations.

Shared CPU and GPU arrays with managed memory

Applications can now allocate CUDA Unified Memory explicitly. These allocations remain CUDA arrays in Warp, while the CUDA driver manages physical page placement and may migrate pages between host memory and CUDA devices as they are accessed. On supported systems, CPU kernels can access them directly. wp.CudaManagedAllocator() works with the existing allocator scopes, and array.memory_kind identifies the allocation as wp.MemoryKind.CUDA_MANAGED (#1523).

import warp as wp


@wp.kernel
def increment(data: wp.array[float]):
    i = wp.tid()
    data[i] += 1.0


managed = wp.CudaManagedAllocator()
with wp.ScopedAllocator("cuda:0", managed):
    data = wp.array([1.0, 2.0, 3.0], dtype=float, device="cuda:0")

if data.memory_kind is not wp.MemoryKind.CUDA_MANAGED:
    raise RuntimeError("Expected a CUDA managed-memory allocation")

cpu_data = data if wp.can_access("cpu", data) else data.to("cpu")
wp.launch(increment, dim=data.size, inputs=[cpu_data], device="cpu")
print(cpu_data.numpy())  # [2. 3. 4.]

Direct CPU access depends on the system's managed-memory capabilities, so portable code should check wp.can_access(). Allocate managed arrays before CUDA graph capture and reuse them inside the graph; Warp does not currently support allocating managed arrays during capture. Managed arrays cannot be exported through CUDA IPC because CUDA does not provide IPC handles for cudaMallocManaged() allocations.

Language and kernel programming

Modify values in place with wp.ref

User-defined Warp functions previously received scalar, vector, matrix, quaternion, transform, and struct arguments by value. A helper that changed one of these values had to return it and rely on its caller to assign it back. Functions can now declare wp.ref[T] parameters to update caller-owned storage directly. This removes the return-and-reassign step from reusable helpers, especially when they update several values (#1277).

import warp as wp


@wp.func
def add_in_place(value: wp.ref[float], delta: float):
    value += delta


@wp.kernel(enable_backward=False)
def add_ten(values: wp.array[float]):
    i = wp.tid()
    add_in_place(values[i], 10.0)


values = wp.array([1.0, 2.0, 3.0], dtype=float, device="cpu")
wp.launch(add_ten, dim=values.size, inputs=[values], device="cpu")
print(values.numpy())  # [11. 12. 13.]

When calling a function with a wp.ref[T] parameter, pass writable storage such as a local variable, array element, or struct field. Literals, arithmetic results, function return values, and compile-time constants are rejected because there is no original storage to update. wp.address_of(expr) exposes a raw pointer to the same kinds of writable expressions for C++/CUDA snippets; use array.ptr for the base pointer of an entire array. Warp does not automatically differentiate @wp.func helpers that use wp.ref[T], so declare the calling kernel with enable_backward=False. For differentiable code, define the helper with @wp.func_native using a C++/CUDA snippet and provide its derivative in adj_snippet.

Pass Warp functions as arguments

User-defined Warp functions can now accept another Warp function as an argument, so one helper can apply different operations without defining a separate wrapper for each one. Annotate the parameter with wp.Function and pass a user-defined @wp.func or a simple builtin such as wp.sin or wp.sqrt (#1424).

import warp as wp


@wp.func
def square(x: float):
    return x * x


@wp.func
def apply(function: wp.Function, x: float):
    return function(x)


@wp.kernel
def evaluate(output: wp.array[float]):
    output[0] = apply(square, 3.0)
    output[1] = apply(wp.sqrt, 9.0)


output = wp.empty(2, dtype=float, device="cpu")
wp.launch(evaluate, dim=1, outputs=[output], device="cpu")
print(output.numpy())  # [9. 3.]

The target is chosen when Warp generates the kernel, not while it runs. Warp compiles a specialized version of the helper for each distinct target, so using many target combinations can increase compilation work. Defaults and keyword arguments are supported. Arbitrary Python callables and builtins such as wp.printf are not valid targets. A helper with a wp.Function parameter cannot also use @wp.func_grad or @wp.func_replay.

Lean CUDA kernel launches

Performance-sensitive CUDA kernels can now opt out of Warp's grid-stride loop. @wp.kernel(grid_stride=False) uses a lean launch path with one thread per work item, reducing loop-control overhead and potentially lowering register pressure for lightweight or bandwidth-bound kernels (#1270).

Use the "default_grid_stride": False module option or set wp.config.default_grid_stride = False to opt in more broadly. Grid-stride launches remain the default. Lean kernels support launches of any practical size, but they cannot use max_blocks to cap the block count. Passing max_blocks > 0 raises an error.

Other CUDA kernel controls

• @wp.kernel(cluster_dim=...) and wp.get_cuda_max_cluster_dim() expose CUDA Thread Block Clusters to native functions on cluster-capable GPUs (Support Thread Block Cluster group specification #1401).

Array utility updates

• wp.utils.array_scan() now accepts 64-bit scalar and vector types. wp.utils.radix_sort_pairs() accepts 32-bit and 64-bit signed, unsigned, and floating-point keys with 4-byte or 8-byte values (Extend scan and radix sort type support #1538).

Textures and geometry

Mipmap sampling

Important

The texture API is experimental and subject to change without a formal deprecation cycle. Pass texture constructor options by keyword when possible.

Textures can now store a full level-of-detail chain and sample a chosen resolution from Warp kernels. Set num_mip_levels=0 to generate the full chain, or request a fixed number of levels, then pass lod to wp.texture_sample() (#1409).

import numpy as np
import warp as wp


@wp.kernel
def sample_levels(texture: wp.Texture2D, output: wp.array[float]):
    level = wp.tid()
    output[level] = wp.texture_sample(
        texture,
        wp.vec2(0.125, 0.125),
        dtype=float,
        lod=float(level),
    )


pixels = np.arange(16, dtype=np.float32).reshape(4, 4)
texture = wp.Texture2D(data=pixels, num_mip_levels=0, device="cpu")

num_levels = texture.num_mip_levels
output = wp.empty(num_levels, dtype=float, device="cpu")
wp.launch(sample_levels, dim=num_levels, inputs=[texture], outputs=[output], device="cpu")
print(output.numpy())  # [0.  2.5 7.5]

Each output element samples the same normalized coordinate at LOD 0, 1, and 2, so [0.0, 2.5, 7.5] shows how the sampled value changes across the mip chain.

Warp builds lower levels from construction data with a 2x box filter. Mipmapped textures are immutable after construction, cannot wrap an external CUDA array, and cannot use surface access. CUDA mipmapped textures also require normalized coordinates.

Expanded cuBQL construction and queries

Important

This is an experimental feature. The API may change without a formal deprecation cycle.

Warp 1.15 expands the experimental cuBQL BVH backend with direct wp.Bvh construction and broader query support. Select it with constructor="cubql" for wp.Bvh or bvh_constructor="cubql" for wp.Mesh. The resulting structures support AABB, closest-point, furthest-point, and ray queries through Warp's existing query APIs (#1467).

Grouped BVHs and meshes are not supported, and wp.Mesh winding-number queries remain unavailable. A cuBQL BVH also cannot switch to or from another constructor during an in-place rebuild.

Tile and sparse programming

Structs in tile operations

Tiles can now contain Warp struct values, so tile kernels no longer need a separate array for each struct field. wp.tile_map() can process struct elements and broadcast non-tile struct arguments. Struct tiles support field-wise addition and subtraction, reductions, atomic addition, gradients, and use as wp.tile_sort() value payloads (#573).

import warp as wp

TILE_SIZE = wp.constant(4)


@wp.struct
class Particle:
    mass: float
    velocity: wp.vec3


@wp.func
def double_mass(particle: Particle) -> Particle:
    particle.mass *= 2.0
    return particle


@wp.kernel
def update(input_particles: wp.array[Particle], output_particles: wp.array[Particle]):
    tile = wp.tile_load(input_particles, shape=TILE_SIZE)
    wp.tile_store(output_particles, wp.tile_map(double_mass, tile))


particles = []
for mass in (1.0, 2.0, 3.0, 4.0):
    particle = Particle()
    particle.mass = mass
    particles.append(particle)

input_particles = wp.array(particles, dtype=Particle, device="cuda:0")
output_particles = wp.empty_like(input_particles)
wp.launch_tiled(
    update,
    dim=1,
    inputs=[input_particles],
    outputs=[output_particles],
    block_dim=32,
    device="cuda:0",
)
print(output_particles.numpy()["mass"])  # [2. 4. 6. 8.]

wp.tile_sort() treats structs as value payloads rather than keys, and sorting is forward-only for struct values. Struct tiles do not define a canonical wp.tile_ones() value, bitwise in-place operations, or ordering reductions.

Reserved BSR row capacity

Sparse assembly can now reserve a bounded number of blocks per row, update topology without reallocating the whole matrix, and compact candidate entries in place. warp.sparse.bsr_zeros(row_capacity=...) creates padded storage, bsr_compress() sorts and coalesces active entries, and matrix status reports whether an operation exceeded its row capacity (#1537).

import warp as wp
from warp.sparse import BSR_STATUS_SUCCESS, bsr_compress, bsr_set_from_triplets, bsr_zeros

A = bsr_zeros(2, 3, float, device="cpu", row_capacity=2)
bsr_set_from_triplets(
    A,
    rows=wp.array([0, 0, 1, 1], dtype=int, device="cpu"),
    columns=wp.array([0, 2, 1, 1], dtype=int, device="cpu"),
    values=wp.array([2.0, 3.0, 4.0, -1.0], dtype=float, device="cpu"),
    topology="padded",
)
print(A.row_counts.numpy())
print(A.status_sync() == BSR_STATUS_SUCCESS)
bsr_compress(A, inplace=True, topology="compact")
print(A.nnz_sync())

Output:

[2 1]
True
3

The inplace=True compression path supports float32 and float64 scalar matrices and is not differentiable. The default out-of-place path remains differentiable. warp.fem.integrate() and warp.fem.interpolate() can use in-place compression during matrix assembly to reduce peak memory use.

Building and debugging from source

CMake source builds

Warp's source tree now includes ready-to-use CMake configuration for parallel and incremental local builds of the native warp and warp-clang libraries (#1495). This optional developer workflow does not build or replace Warp's Python wheels, which remain the primary way to install Warp. See the CMake build documentation.

AddressSanitizer for CPU kernels

Source builds can instrument native libraries and JIT-compiled CPU kernels with AddressSanitizer (#1387):

uv run build_lib.py --sanitize=address

Run the target kernel in release mode so Warp's debug bounds assertion does not fire before AddressSanitizer. Linux processes must preload the AddressSanitizer runtime before importing Warp. Standard pip installations are unchanged.

Breaking changes

wp.copy() validates offsets, counts, and element sizes

Warp now rejects non-integer or negative offsets, a negative count, and contiguous copies whose source and destination element sizes differ. Warp v1.14 treated a negative count as a request to copy the whole array, which could hide mistakes (#1584).

import warp as wp

source = wp.array([1, 2, 3, 4], dtype=wp.int32, device="cpu")
destination = wp.zeros_like(source)

# Warp v1.14 treated a negative count as "copy all".
# Warp v1.15 rejects it before any pointer arithmetic.
try:
    wp.copy(destination, source, count=-1)
except RuntimeError as error:
    print(error)

# Migration: pass the intended non-negative element count explicitly.
wp.copy(destination, source, count=source.size)
print(destination.numpy())

Output:

Copy count must be non-negative, got -1
[1 2 3 4]

Native library version mismatches stop initialization

Warp now stops initialization with RuntimeError when the Python package version does not match warp.dll, warp-clang.dll, or their platform equivalents. Warp previously warned and continued, which could load an incompatible native ABI. Rebuild or reinstall Warp so the native libraries match the Python package (#1508).

Announcements

Tentative CUDA 13 PyPI wheel builds

• We are tentatively planning to build Warp's PyPI wheels with a CUDA 13.x Toolkit starting with Warp 1.17, currently targeted for September 2026. CUDA Toolkit 13.0 was released in August 2025, more than a year before that target, and we recommend moving to an R580-series or newer NVIDIA driver in advance. If this plan proceeds, R580 will be the minimum driver branch for using Warp's CUDA backend from those wheels. pip does not check the installed driver and can install the wheels on systems with older drivers, but the CUDA backend will not work on those systems. The driver must support the CUDA 13.x family, but it does not need to match the exact CUDA 13.x Toolkit version used to build Warp. See NVIDIA's CUDA minor-version compatibility table.

• CUDA 12 support will continue beyond Warp 1.17. After the PyPI wheels move to CUDA 13.x, users who need CUDA 12 can build Warp from source or install a pre-built CUDA 12 wheel from GitHub Releases.

Removals in this release

• Deprecated warp.fem arguments have been removed. Pass a warp.fem.Quadrature or warp.fem.GeometryDomain through at instead of the old quadrature or domain arguments to warp.fem.interpolate(). Use space_topology for partitions and space_topology or space_partition for restrictions.
• The non-functional warp.render.UsdRenderer.update_body_transforms() method has been removed. It referenced attributes that UsdRenderer never defined and could not update USD transforms. Use the supported shape-rendering methods to author time-varying USD transforms instead.

Upcoming removals

• Legacy wp.MarchingCubes sizing arguments and compatibility attributes are deprecated and scheduled for removal in Warp 1.19. The max_verts and max_tris values are retained only for compatibility and do not limit the dynamically sized output arrays. Likewise, device does not select the extraction device. Extraction always uses the input field's device. id no longer identifies a native resource, while runtime is not used during extraction. Neither has a replacement. Stop passing max_verts, max_tris, or device to the constructor, stop passing max_verts or max_tris to resize(), and remove accesses to these compatibility attributes (Deprecate legacy MarchingCubes arguments and state #1594).
• masked=True in warp.sparse topology-changing operations is deprecated. Pass topology="masked" instead. The old argument will be removed in a future feature release under the standard deprecation timeline (Add 4-array BSR storage with row capacity #1537).

Acknowledgments

We also thank the following contributors from outside the core Warp development team:

• @astefanutti for surfacing and diagnosing the free-threaded Python _warp_fastcall crash, which informed the removal of fastcall support (Fix/free threaded python fastcall #1578).
• @adityasingh2400 for fixing Warp function binding to kernel-local variables (Fix binding a Warp function to a kernel-local (GH-1423) #1476).
• @FabienPean-Virtonomy for adding PEP 604 union syntax to Warp array annotations (Enable Union notation at Python-scope for warp arrays #1549).
• @snico432 for making Clang CUDA compilation failures report immediately (Report Clang CUDA compile errors (GH-1325) #1503).

---

For a complete list of changes, see the full changelog.

---

This discussion was created from the release v1.15.0.