Warp v1.18.0

Warp v1.18 introduces the warp.geometry module, with Python and kernel APIs for extracting meshes from implicit fields and sampled rigid motion, and improving 2D triangulations directly on the device. Experimental CPU block execution improves tile portability by supporting the same explicit block_dim on CPU and CUDA. The release also supports automatic differentiation through vector and matrix component updates in arrays and avoids unnecessary GPU-to-CPU transfers of sparse-matrix block counts.

New features

Sparse marching cubes

You can now extract an isosurface from a custom sparse volume or other geometry representation without converting it to a dense grid. warp.geometry.sparse_marching_cubes() accepts a callable that samples a scalar field at batches of query points, so your sampler can read directly from the existing data structure. Warp builds a Lipschitz octree to prune cells away from the requested surface and runs marching cubes on the remaining cells (#1803).

nx, ny, and nz specify grid-node counts, as in dense marching cubes. The field callable accepts wp.vec3 query points in batches and returns wp.float32 values on the same device. The output contains wp.vec3 vertices and flat triangle indices, with three indices per face. This example evaluates a sphere field in a Warp kernel on the GPU and compares its evaluation count with sampling every node of a 257³ grid:

import warp as wp
import warp.geometry

@wp.kernel
def sample_sphere(points: wp.array[wp.vec3], values: wp.array[float]):
    i = wp.tid()
    values[i] = wp.length(points[i]) - 0.5

def field(points):
    values = wp.empty(points.shape, dtype=wp.float32, device=points.device)
    wp.launch(sample_sphere, dim=points.shape[0], inputs=[points],
              outputs=[values], device=points.device)
    return values

n = 257
vertices, indices, stats = warp.geometry.sparse_marching_cubes(
    field, n, n, n,
    lower=(-1.0, -1.0, -1.0), upper=(1.0, 1.0, 1.0),
    return_stats=True,
)
print(f"Mesh: {vertices.shape[0]:,} vertices, {indices.shape[0] // 3:,} triangles")
print(f"Sparse field evaluations: {stats['field_evaluations']:,}")
print(f"Dense grid samples: {n**3:,}")

Output:

Mesh: 77,094 vertices, 154,184 triangles
Sparse field evaluations: 400,259
Dense grid samples: 16,974,593

Here, sparse extraction uses about 2.4% of the field evaluations required by dense sampling.

You can differentiate extracted vertex positions with respect to sampled field values. Allocate the field output with requires_grad=True and extract inside wp.Tape(). Cell selection and mesh connectivity remain fixed during differentiation.

The default lipschitz_bound=1.0 works for a true signed distance function. Fields that change faster with distance need a larger bound. Setting it too low can discard cells containing the surface.

Parts of the surface outside lower and upper are omitted, so the mesh may be open at those bounds. See the sparse marching cubes example for a mesh-query field and a comparison with dense extraction.

If you already have a list of grid cells and field values at their eight corners, use warp.geometry.sparse_marching_cubes_from_cells(). It extracts the mesh from those cells without building an octree.

Delaunay edge flipping

warp.geometry.delaunay_edge_flip() updates a 2D triangle mesh's connectivity on the device after its vertices move. It flips non-Delaunay interior edges in place. warp.geometry.tri_tri_adjacency() identifies each edge's neighboring triangle and the matching edge index within that neighbor (#1797).

import warp as wp
import warp.geometry

positions = wp.array([[-3.0, 0.0], [3.0, 0.0], [0.0, 1.0], [0.0, -1.0]],
                     dtype=wp.vec2)
triangles = wp.array([[0, 1, 2], [1, 0, 3]], dtype=wp.int32)
flips = warp.geometry.delaunay_edge_flip(positions, triangles)
neighbors, neighbor_edges = warp.geometry.tri_tri_adjacency(triangles, vertex_count=4)
print(int(flips.numpy()[0]))  # 1
faces = triangles.numpy().tolist()
print(sorted(set(faces[0]) & set(faces[1])))  # [2, 3]

The shared diagonal changes from (0, 1) to (2, 3).

Edge flipping requires a manifold mesh whose triangle vertices are consistently ordered counterclockwise. It stops when no more flips are needed or max_passes is reached. If you supply reference positions, the algorithm skips flips that would create degenerate triangles in the reference mesh. Both APIs support CUDA graph capture. Supply vertex_count to the adjacency helper during capture, and read the flip count after replay.

The elastic shape optimization example uses device-side edge flipping while optimizing a deforming mesh.

Swept-volume meshing

warp.geometry.swept_volume_mesh() builds a single mesh of the space occupied by a rigid assembly across sampled poses. It samples a signed distance field on a dense grid. Each grid point queries every mesh at every supplied pose, so finer grids and more pose samples increase computation. Use the result for motion visualization or as an input to robotics clearance checks (#1824).

import numpy as np
import warp as wp
import warp.geometry

points = wp.array([[0, 0, 0], [1, 0, 0], [0, 1, 0], [0, 0, 1]], dtype=wp.vec3)
faces = wp.array([0, 2, 1, 0, 3, 2, 0, 1, 3, 1, 2, 3], dtype=wp.int32)
mesh = wp.Mesh(points, faces, support_winding_number=True)
poses = np.zeros((1, 8, 7), dtype=np.float32)
poses[..., 6] = 1.0  # Identity rotations, quaternion order xyzw.
poses[0, :, 0] = np.linspace(0.0, 1.0, 8)
vertices, indices = warp.geometry.swept_volume_mesh(
    [mesh], poses, voxel_size=0.1, threshold=0.09
)
v = vertices.numpy()
print(np.round(v.min(axis=0), 2))  # [-0.09 -0.09 -0.09]
print(np.round(v.max(axis=0), 2))  # [2.09 1.09 1.09]

The sampled tetrahedron spans (0, 0, 0) to (2, 1, 1); the positive threshold expands the mesh beyond those bounds.

Transforms have shape (mesh_count, sample_count, 7), with translation followed by an xyzw quaternion. The default sign mode requires each input mesh to be built with support_winding_number=True. The new mesh.support_winding_number property lets you check that requirement.

The envelope covers the supplied poses. It does not conservatively bound motion between samples, so fast or thin objects need sufficiently dense pose sampling. To enclose the sampled poses despite grid discretization, use a positive threshold greater than half the grid-cell diagonal. The example uses 0.09 with a cubic cell size of 0.1.

See the swept-volume example for a procedural robot arm, animated USD input, and alternative sign modes.

Dense isosurface extraction

Dense isosurface extraction now has a common warp.geometry.IsoSurfaceBase interface, implemented by warp.geometry.IsoSurfaceMarchingCubes. Call extract() on the class to return vertices and triangle indices without creating an extractor instance (#1614):

import numpy as np
import warp as wp
import warp.geometry

x, y, z = np.meshgrid(*([np.linspace(-1.0, 1.0, 5)] * 3), indexing="ij")
field = wp.array(np.sqrt(x*x + y*y + z*z) - 0.6, dtype=wp.float32)
vertices, indices = warp.geometry.IsoSurfaceMarchingCubes.extract(
    field, lower=(-1.0, -1.0, -1.0), upper=(1.0, 1.0, 1.0)
)
print(vertices.shape[0], indices.shape[0] // 3)  # 30 56

The input must be a three-dimensional wp.float32 array with at least two nodes on every axis. You can differentiate the extracted vertex positions with respect to field values, with mesh connectivity held fixed. The isosurface example has moved into warp/examples/geometry/.

Tile programming

Cooperative CPU block execution

Tile kernels that require multiple cooperating logical threads within a block can now produce matching results on CPU and CUDA using the same explicit block_dim. Set wp.config.enable_cpu_blocks = True to enable CPU blocks with 2 to 1024 logical kernel threads. Each logical thread has its own index within the block. Threads in the same block can synchronize and share tile data (#1638).

import warp as wp

wp.config.enable_cpu_blocks = True

@wp.kernel
def sum_thread_indices(out: wp.array[float]):
    block_index, thread_index = wp.tid()
    values = wp.tile(float(thread_index))
    wp.tile_store(out, wp.tile_sum(values))

out = wp.empty(1, dtype=wp.float32, device="cpu")
wp.launch_tiled(sum_thread_indices, dim=1, outputs=[out], block_dim=4, device="cpu")
print(out.numpy())  # [6.]

CPU execution defaults to one logical kernel thread per block. Cooperative CPU blocks schedule their logical threads on a single operating-system thread and can be substantially slower. They do not add CPU parallelism or vectorize execution across those logical threads. AddressSanitizer builds reject enabled CPU launches with block_dim > 1. See [CPU tile execution](https://github.com/NVIDIA/warp/blob/v1.18.0/docs/user_guide/programming_model/tiles.rst#...

Warp v1.17.0

Warp v1.17 expands geometry queries with sphere and capsule searches over BVHs, exact sphere queries against mesh triangles, and direct access to a mesh's BVH. Tiles now support matrix-row indexing, CG and CR solvers can restart periodically, and new controls let you tune and inspect CUDA kernel resource use. The release also includes experimental native build hooks for external C++ and CUDA integrations, along with native CPU support when building Warp from source on Windows ARM64.

If you are upgrading, note that implicit conversion of Python and Warp numeric scalars to composite types has been removed; see Removals and deprecations for migration guidance.

New features

Sphere and capsule spatial queries

A BVH can now be queried with a sphere or capsule instead of first converting the search region to an AABB. wp.bvh_query_sphere() finds item bounds that overlap a sphere using an exact sphere-AABB test. wp.bvh_query_capsule() searches for item bounds that overlap the volume swept by moving a sphere along a line segment. Capsule queries are conservative: they do not miss bounds within the requested radius, but they can return extra candidates near AABB corners (#1741).

A capsule query takes a start point and a direction rather than two endpoints. To query the segment from p0 to p1, pass p0 as the start and p1 - p0 as the direction to wp.bvh_query_capsule(). By default, traversal continues indefinitely along that direction. Pass max_dist=1.0 to wp.bvh_query_next() to limit the query to the full segment, including both endpoints. If p0 == p1, use wp.bvh_query_sphere() instead.

wp.mesh_get_bvh() exposes a mesh's internal BVH to the general wp.bvh_query_*() APIs. This makes BVH-only operations such as wp.bvh_query_capsule() available for meshes, with returned bound indices corresponding to triangle faces. Use wp.mesh_query_sphere() when you need exact triangle-sphere intersections rather than broad-phase candidates. wp.MeshQuery is now the common base type for AABB and sphere mesh queries, with wp.mesh_query_next() as the canonical iterator for both. The existing wp.MeshQueryAABB type and wp.mesh_query_aabb_next() alias remain available.

Tune and inspect CUDA kernel resource use

CUDA kernels now have controls for register allocation and shared-memory spilling, and their resource use can be inspected before launch. cuda_max_registers requires Warp to have been built with CUDA Toolkit 12.4 or newer and cannot be combined with launch_bounds. The Linux and Windows warp-lang wheels on PyPI are built with CUDA Toolkit 12.9, so they meet this requirement. enable_cuda_smem_spilling takes effect only when Warp itself was built with CUDA Toolkit 13.0 or newer. Warp 1.17's PyPI wheels use CUDA Toolkit 12.9 and ignore this option. CUDA 13.0 wheels for Linux x86-64, Linux ARM64, and Windows x86-64 are available from GitHub Releases, or you can build Warp from source with CUDA Toolkit 13.0 or newer. CUDA 13 builds require an R580-series or newer NVIDIA driver and a Turing-class GPU (compute capability 7.5) or newer. See CUDA 13 PyPI wheel timing for the planned PyPI transition. Shared-memory spilling is also ignored when a kernel uses dynamic shared memory (#1671).

wp.get_cuda_kernel_properties() compiles the requested kernel variant if necessary, without launching it, and reports its per-thread register count and local-memory use (#1805).

import warp as wp


@wp.kernel(cuda_max_registers=64, enable_backward=False)
def update(values: wp.array[float]):
    i = wp.tid()
    values[i] = wp.sin(values[i]) + wp.cos(values[i])


properties = wp.get_cuda_kernel_properties(
    update,
    device="cuda:0",
    block_dim=128,
)
print(sorted(properties))  # ['local_memory_size', 'register_count']

Resource counts depend on the GPU, toolchain, compiler options, and block size. Use them to investigate occupancy and spilling, then profile the kernel before changing its configuration.

Tile programming

Index matrix rows inside tiles

Matrix-valued tiles now support chained row indexing such as tile[i, j, k][row]. Reads, writes, negative row indices, and adjoints work for tiles with one through four logical dimensions (#1028).

import numpy as np
import warp as wp

TILE_SIZE = 8


@wp.kernel
def extract_last_row(matrices: wp.array[wp.mat33], rows: wp.array[wp.vec3]):
    i = wp.tid()

    # Load eight 3x3 matrices into a one-dimensional tile.
    matrix_tile = wp.tile_load(matrices, shape=(TILE_SIZE,))

    # Select matrix i from the tile, then use -1 to select its last row.
    rows[i] = matrix_tile[i][-1]


# Repeat [[1, 2, 3], [4, 5, 6], [7, 8, 9]] eight times.
data = np.tile(
    np.arange(1.0, 10.0, dtype=np.float32).reshape(3, 3),
    (TILE_SIZE, 1, 1),
)
matrices = wp.array(data, dtype=wp.mat33, device="cuda:0")
rows = wp.zeros(TILE_SIZE, dtype=wp.vec3, device="cuda:0")
wp.launch(
    extract_last_row,
    dim=TILE_SIZE,
    inputs=[matrices],
    outputs=[rows],
    block_dim=TILE_SIZE,
    device="cuda:0",
)
print(rows.numpy()[0])  # Last row of the first matrix: [7. 8. 9.]

Iterative solvers

Limit finite-precision drift with CG and CR restarts

Over a long CG or CR solve, the recursively updated residual can drift away from the true residual b - A x. Set restart=N on warp.optim.linear.cg() or warp.optim.linear.cr() to recompute it and reset the search direction every N iterations. This matters most for float32 CUDA workloads, including batched and matrix-free solves. The restart path also works with CUDA graph capture (#1708).

A restart requires one extra matrix-vector product. Warp checks convergence and invokes callbacks only at cycle boundaries, so the solve can run up to restart - 1 iterations past maxiter. Omit restart to keep the existing recursive behavior.

Batched CUDA solves also get more accurate dot products as subproblem size grows. When the largest subproblem is known, pass max_batch_length to LinearOperator or aslinearoperator() to avoid unnecessary reduction work (#1700). Reusable CG, CR, BiCGSTAB, and GMRES states returned with run=False now allocate their temporary device memory once during construction instead of on each solve.

Compilation and tooling

Name generated kernels explicitly

Kernel factories can now give closure-generated kernels distinct, predictable names for registration and ahead-of-time compilation. Set name on @wp.kernel to assign the registration key and the base of the generated native entry-point name (#1561).

import warp as wp


def make_scaler(factor: float, kernel_name: str):
    @wp.kernel(name=kernel_name)
    def scale(values: wp.array[float]):
        i = wp.tid()
        values[i] *= factor

    return scale


double = make_scaler(2.0, "scale_by_two")
triple = make_scaler(3.0, "scale_by_three")
print(double.key, triple.key)  # scale_by_two scale_by_three

Names must be valid C++ identifiers. With strip_hash=True, Warp uses the custom key without a hash suffix as the base of generated entry-point names.

Experimental native build hooks

Add-on packages can now attach native C++ or CUDA headers to a Warp module and make header-defined types and functions available to Warp kernels. wp.ModuleBuildOptions supplies include directories, preambles, and dependency files. wp.build_experimental.add_native_type() and add_builtin() register the matching ABI types and functions, while wp.compile_aot_module() returns artifact paths for an external build or runtime system (#1575). We have not yet validated this as a complete production workflow.

In this example, addon_math.h comes from the add-on package, not Warp. The example assumes this package layout:

my_addon/
├── build_addon.py
└── include/
    └── addon_math.h

addon_math.h defines the native function that the add-on exposes to Warp kernels. Warp includes its native headers first, so the add-on header can use CUDA_CALLABLE:

include/addon_math.h:

#pragma once

namespace addon {
CUDA_CALLABLE inline float square(float value)
{
    return value * value;
}
}  // namespace addon

build_addon.py locates the header relative to its own file, registers addon::square as wp.addon_square, and tells Warp where to find the header when compiling the module:

from pathlib import Path

import warp as wp


# Resolve the header shipped with this add-on package.
header = (Path(__file__).parent / "include/addon_math.h").resolve()

# Map the C++ function to the name and signature used in Warp kernels.
wp.build_experimental.add_builtin(
    "addon_square",
    {"value": wp.float32},
    wp.float32,
    native_name="addon::square",
)


# add_builtin() registers no adjoint, so generate only the forward kernel.
@wp.kernel(enable_backward=False)
def square_kernel(values: wp.array[float], output: wp.array[float]):
    i = wp.tid()
    output[i] = wp.addon_square(values[i])


build_options = wp.ModuleBuildOptions(
    extra_cuda_include_dirs=[header.parent],  # Let #include find addon_math.h.
    extra_cuda_preamble='#include "addon_math.h"',  # Include it in generated CUDA source.
    extra_build_dependencies=[header],  # Recompile when the header changes.
)

# Apply the add-on's build inputs before compiling this kernel's module.
wp.set_module_options({"extra_build_options": build_options}, module=square_kernel.module)

# Generate ...

Warp v1.16.0

Warp v1.16 adds in-place rebuilding for fixed-capacity NanoVDB volumes, including during CUDA graph capture. This lets fluid simulations update sparse-grid topologies without allocating new volumes. The release also adds grouped HashGrid queries for multi-environment workloads, NumPy-style tile slicing, CPU support for JAX FFI wrappers, experimental support for replaying more operations in CPU graphs and saved .wrp graphs, and CUDA profiler range controls.

New features

Rebuild sparse NanoVDB volumes in place

Previously, a volume's sparse topology was fixed at allocation time. Simulations whose active grid changed from one step to the next, such as the affine particle-in-cell (APIC) fluid example, had to allocate a new volume. That allocation required host synchronization and could not be replayed inside a CUDA graph. Warp 1.16 adds fixed-capacity, in-place rebuilding so simulations can reuse the same volume and dependent FEM topology buffers across graph replays (#1606).

The excerpt below follows that example's allocation, rebuild, and replay flow. It omits particle initialization, capacity estimation, FEM space construction, and the APIC transfer and solve.

import warp as wp
import warp.fem as fem

# API-shape excerpt. Simulation setup and solver work are omitted.
# particle_q, voxel_size, grid_capacity, and frame_count come from that setup.

grid_status = wp.zeros(1, dtype=wp.uint32, device=particle_q.device)
volume = wp.Volume.allocate_by_voxels(
    voxel_points=particle_q,
    voxel_size=voxel_size,
    device=particle_q.device,
    rebuildable=True,
    **grid_capacity,
    status=grid_status,
)
grid = fem.Nanogrid(volume, rebuildable=True)

# ... build linear_basis_space and strain_space once from grid ...


def simulate():
    # Rebuild from the particle positions updated by the previous step.
    grid.rebuild(particle_q, status=grid_status)
    linear_basis_space.topology.rebuild()
    strain_space.topology.rebuild()
    # ... transfer particles to the grid, solve, and advect particles ...


with wp.ScopedCapture(particle_q.device) as capture:
    simulate()

for _ in range(frame_count):
    wp.capture_launch(capture.graph)

# Host-side status and topology queries stay outside capture.
status = int(grid_status.numpy()[0])
stats = volume.get_active_stats()
print(f"rebuild status=0x{status:x}, active voxels={stats.voxel_count}")

Capacity does not grow automatically. Pass a one-element wp.uint32 status array to the allocation method and to rebuild(), then check for wp.Volume.REBUILD_SUCCESS, a REBUILD_*_CAPACITY_EXCEEDED flag, or REBUILD_INVALID_INPUT. A tile-allocated volume must be rebuilt from tile positions, while a voxel-allocated volume must be rebuilt from voxel positions. The point_mask parameter can exclude input points, and Warp deduplicates repeated positions.

On CUDA, rebuildable allocation and rebuilding can run inside graph capture when memory-pool allocation is enabled. Exact, non-rebuildable allocation and host-side queries such as get_active_stats() must remain outside capture.

See example_apic_fluid.py for the complete capacity estimation, topology construction, and captured simulation loop.

API changes:

Multi-environment spatial queries

Grouped HashGrid construction and queries

Independent simulation environments can now share one spatial grid without visiting candidates from neighboring worlds. Assign each point an int32 group in HashGrid.build(), then pass the desired group to wp.hash_grid_query() (#1579).

import warp as wp


@wp.kernel
def count_neighbors(
    grid: wp.uint64,
    points: wp.array[wp.vec3],
    groups: wp.array[wp.int32],
    counts: wp.array[wp.int32],
):
    i = wp.tid()
    count = int(0)
    # Restrict the query to candidates in the current point's environment.
    for j in wp.hash_grid_query(grid, points[i], 0.5, groups[i]):
        if wp.length(points[j] - points[i]) <= 0.5:
            count += 1
    counts[i] = count


points = wp.array([(0.0, 0.0, 0.0), (0.4, 0.0, 0.0)] * 2, dtype=wp.vec3)
# The repeated two-point scene represents environments 0 and 1.
groups = wp.array([0, 0, 1, 1], dtype=wp.int32)
counts = wp.zeros(4, dtype=wp.int32)
grid = wp.HashGrid(8, 8, 8)
grid.build(points, radius=0.5, groups=groups)
wp.launch(count_neighbors, dim=4, inputs=[wp.uint64(grid.id), points, groups], outputs=[counts])
print(f"same-group neighbor counts: {counts.numpy()}")  # same-group neighbor counts: [2 2 2 2]

Group IDs may be any int32 values and can change between rebuilds, including during graph replay. If the query omits the group argument, it visits all candidates. Before capturing a grouped CUDA rebuild without a warm-up build, call grid.reserve(num_points, with_groups=True). Hash-grid queries return cell candidates, so kernels should still test the actual distance.

Tile programming

NumPy-style tile slicing

Tile kernels now use NumPy-style indexing to select, reverse, stride, gather, and assign subregions. An integer index collapses a dimension and may be negative. Use wp.tile_slice_indexed() for one-axis gathers such as tile[indices, :] (#1176).

import numpy as np
import warp as wp


@wp.kernel
def slice_tile(src: wp.array2d[float], dst: wp.array2d[float]):
    tile = wp.tile_load(src, shape=(4, 4))
    # Reverse the rows and keep columns 0 and 2.
    wp.tile_store(dst, tile[::-1, ::2])


src = wp.array(np.arange(16, dtype=np.float32).reshape(4, 4), device="cuda:0")
dst = wp.zeros((4, 2), dtype=float, device="cuda:0")
# Launch one cooperative 32-thread block for the single tile.
wp.launch_tiled(slice_tile, dim=[1], inputs=[src], outputs=[dst], block_dim=32, device="cuda:0")
print(f"sliced tile (reversed rows, even columns):\n{dst.numpy()}")

Output:

sliced tile (reversed rows, even columns):
[[12. 14.]
 [ 8. 10.]
 [ 4.  6.]
 [ 0.  2.]]

The source tile in a slice assignment must have a matching shape. Tile views do not support compound assignment, and slices do not support scalar broadcasting. Indexed gathers select one axis at a time and require full : slices on the other axes.

The out-of-place wp.tile_lower_solve() now supports gradients for vector and matrix right-hand sides. This makes lower-triangular solves available in wp.Tape backward passes (#1378).

JAX integration

Run Warp FFI callbacks on CPU

JAX can now run wp.jax_kernel() and wp.jax_callable() wrappers on CPU as well as CUDA. JAX selects the device, so a CPU jax.jit program invokes the Warp callback without routing buffers through a GPU (#1661).

import jax
import jax.numpy as jnp
import numpy as np
import warp as wp


@wp.kernel
def triple(x: wp.array[float], out: wp.array[float]):
    i = wp.tid()
    out[i] = 3.0 * x[i]


run = wp.jax_kernel(triple)
with jax.default_device(jax.devices("cpu")[0]):
    # JAX places both the arrays and the Warp FFI callback on CPU.
    (result,) = jax.jit(run)(jnp.arange(4, dtype=jnp.float32))
print(f"Warp FFI result on CPU: {np.asarray(result)}")  # Warp FFI result on CPU: [0. 3. 6. 9.]

JAX 0.5.0 or newer is required. JaxCallableGraphMode.NONE and JaxCallableGraphMode.JAX work on CPU. The Warp-managed graph modes remain CUDA-only. example_jax_kernel.py shows more call shapes.

Set CUDA block sizes for JAX tile kernels

JAX FFI wrappers created with wp.jax_kernel() now accept block_dim, letting a CUDA tile kernel choose how many threads cooperate on each tile. The value is fixed when the wrapper is constructed and is also used for generated adjoint launches (#1436).

import jax
import jax.numpy as jnp
import numpy as np
import warp as wp

ROW_COUNT = 4
TILE_SIZE = 256
TILE_THREADS = 64


@wp.kernel
def row_sum(values: wp.array2d[float], output: wp.array[float]):
    row = wp.tid()
    # Threads in one CUDA block cooperate to reduce each row.
    tile = wp.tile_load(values[row], shape=TILE_SIZE)
    wp.tile_store(output, wp.tile_sum(tile), offset=row)


# JAX FFI uses wp.launch(), so append the block width to the logical row count.
jax_row_sum = wp.jax_kernel(
    row_sum,
    ...

Prebuilt LLVM/Clang SDKs that Warp's CPU-JIT compiler (warp-clang) links against. This is not a Warp release, and nothing here is required: if you install Warp from PyPI you can ignore this page, and if you build from source you can compile LLVM yourself with build_lib.py --build-llvm, which takes about an hour. These archives are the shortcut.

Warp builds its own rather than using the upstream clang+llvm binaries because the SDK ends up inside a PyPI wheel, and the upstream builds are far too large for that. These are cut down to clang and static libraries with no external dependencies, targeting the host backend and NVPTX, compiled for size.

Built from the LLVM 22.1.8 source release by the Build LLVM SDK workflow. The Linux archives are produced inside manylinux containers, which puts their glibc floor at 2.28 on x86_64 and 2.34 on aarch64. windows-x86_64 is built with the v142 toolset so that it remains linkable from VS2019, while windows-arm64 uses v143 and is still experimental. Per-platform build details are in llvm-source-and-build-info.json.

sha256sum -c SHA256SUMS --ignore-missing verifies the downloads. Once extracted, the tree works directly as build_lib.py --llvm-path or WARP_LLVM_PATH. See tools/llvm/README.md if you want to build one yourself.

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).

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

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_snippe...

Warp v1.14.0

Warp v1.14 expands serialized CPU capture support: captured graphs can now include backward launches, tiled kernels, richer launch arguments, and .wrp files with arrays nested inside structs or wp.indexedarray arguments. This release also adds multi-environment FEM support for batched simulations, reusable and batched linear solvers, pluggable logging, portable tile FFT and solver fallbacks, stable JAX integration APIs, and relaxed CPU/GPU array access for Heterogeneous Memory Management (HMM) and Address Translation Services (ATS) systems.

New features

API Capture expands to cover more workflows

Building on the initial API Capture serialization support in Warp v1.13, Warp v1.14 primarily broadens the set of CPU graph patterns that can be saved and replayed. CPU captures can now include forward execution and reverse-mode passes from wp.Tape().backward(), wp.launch_tiled() kernels, and scalar parameters of any size (#1431). The shared .wrp serialization format also now supports @wp.struct arguments that contain arrays and wp.indexedarray arguments that carry data, gradient, and index buffers.

Saved APIC graphs can still be consumed from standalone C++ through the C API declared in warp/native/apic.h. Native replay behavior is unchanged apart from the explicit pointer spelling for APIC handles.

Key capabilities:

• Reverse-mode replay on CPU: adjoint launches emitted by wp.Tape().backward() are recorded into the CPU APIC stream and replayed from live captures or loaded graphs.
• Richer launch arguments: APIC now relocates array data pointers, gradient pointers, index pointers for wp.indexedarray arguments, and handles inside serialized launch value blobs.
• Tiled kernels: CPU captures can replay kernels that use tiles, including reductions and scans that previously fell outside the captured operation set.

Known limitations:

• wp.utils.array_scan() is still not recorded into CPU APIC and raises NotImplementedError in CPU capture.
• Nonzero array.fill_() operations on CPU are not recorded.
• APIC serializes wp.Mesh handles, but wp.Volume and wp.Bvh handles are not yet supported.
• Loading CPU .wrp graphs requires the warp-clang backend and the companion _modules/ directory with the recorded CPU kernel objects.

Multi-environment warp.fem

warp.fem can now represent many independent simulation environments inside one geometry and one solve setup (#1407). Colocated Grid2D and Grid3D geometries expose an env_count, sparse Nanogrid and AdaptiveNanogrid geometries can pack per-environment voxels into one NanoVDB volume, and unstructured meshes can carry per-cell environment metadata through cell_env and env_count.

This feature changes two positional call signatures. See the FEM migration diff if your code passes requires_grad, device, or temporary_store positionally.

import warp as wp
import warp.fem as fem

geo = fem.Grid3D(res=(8, 8, 8), bounds_lo=wp.vec3(0.0), bounds_hi=wp.vec3(1.0), env_count=4)
pressure_space = fem.make_polynomial_space(geo, degree=0, discontinuous=True)
partition = fem.make_space_partition(space_topology=pressure_space.topology, environment_first=True)

# For scalar pressure spaces, these node offsets can be passed directly to
# warp.optim.linear.LinearOperator(batch_offsets=...).
pressure_batch_offsets = partition.env_offsets

Environment-aware lookup keeps colocated environments from interacting accidentally. When a geometry has more than one environment, pass an environment index to fem.lookup() and pass env_indices to fem.PicQuadrature when particles are binned from world-space positions. The new warp/examples/fem/example_apic_fluid_multi_env.py example uses these APIs to run colocated APIC fluid environments with environment-aware particle quadrature and batched pressure solves.

Key capabilities:

• Colocated environments: grid environments can overlap in world coordinates while remaining topologically independent.
• Sparse packed environments: Nanogrid.from_environment_voxels() and AdaptiveNanogrid.from_environment_voxels() build packed sparse grids with per-cell cell_env metadata and hidden offsets for packed grids.
• Mesh environments: FEM mesh constructors can use cell_env and env_count so grouped BVH lookup only traverses the requested environment.
• Batched solves: make_space_partition(..., environment_first=True) exposes env_offsets that line up with the new linear-solver batch_offsets support for scalar spaces.
• Known limitations: environment_first=True does not support halo nodes for partitions that do not cover a whole geometry. Mesh environment indices are lookup and partition metadata, so callers must still provide disconnected mesh topology for independent mesh environments.

Reusable and batched linear solvers

The iterative solvers in warp.optim.linear can now preallocate their temporary buffers and reuse them across compatible solves (#1391). Passing run=False to cg(), cr(), bicgstab(), or gmres() returns a solver object that can be called repeatedly with replacement operands that have the same shape, dtype, device, and batch layout.

import numpy as np
import warp as wp
from warp.optim.linear import aslinearoperator, cg

diag = wp.array([2.0, 2.0, 5.0, 5.0], dtype=float, device="cpu")
b = wp.array([2.0, 4.0, 10.0, 15.0], dtype=float, device="cpu")
x = wp.zeros_like(b)
offsets = wp.array(np.array([0, 2, 4], dtype=np.int32), dtype=int, device="cpu")

A = aslinearoperator(diag, batch_offsets=offsets)
state = cg(A, b, x, maxiter=100, run=False)
state()  # solve the original system

b2 = wp.array([4.0, 2.0, 15.0, 10.0], dtype=float, device="cpu")
x2 = wp.zeros_like(b2)
state(b=b2, x=x2)  # reuse temporary buffers for a compatible system

batch_offsets is independent of solver-state reuse. It partitions a LinearOperator into scalar degree-of-freedom intervals that are solved as independent subproblems in one solver launch sequence, with convergence checked per batch and reported through the worst residual. For one-shot solves, call the solver directly as before.

Pluggable logging

Warp's Python-side diagnostics now flow through a configurable logger (#1315, #1434). Install a custom logger with wp.set_logger(), scope it with wp.ScopedLogger, and control verbosity through wp.config.log_level or wp.ScopedLogLevel.

wp.config.log_level accepts the standard Warp log-level constants:

• wp.LOG_DEBUG: most verbose, including code-generation details and module loads.
• wp.LOG_INFO: the default, including the init banner and compile timings.
• wp.LOG_WARNING: warnings and errors only.
• wp.LOG_ERROR: errors only.

import warp as wp

wp.config.log_level = wp.LOG_WARNING

class ListLogger:
    def __init__(self):
        self.records = []
    def debug(self, message):
        self.records.append(("debug", message))
    def info(self, message):
        self.records.append(("info", message))
    def warning(self, message, category=None, stacklevel=1):
        self.records.append(("warning", message))
    def error(self, message):
        self.records.append(("error", message))

logger = ListLogger()
with wp.ScopedLogger(logger), wp.ScopedLogLevel(wp.LOG_INFO):
    wp.get_logger().info("captured by application logger")
print(logger.records)  # [('info', 'captured by application logger')]

The default Warp logger now routes warnings emitted by Warp's Python code through Python's warnings.warn(), so application warning filters can suppress Warp deprecation warnings again. wp.config.verbose and wp.config.quiet still work during the deprecation window, but they now emit one-time DeprecationWarnings and map to wp.config.log_level.

Relaxed CPU/GPU array access

Warp no longer enforces the old rule that every Warp array passed to a kernel must be allocated on the launch device. wp.config.launch_array_access_mode now defaults to wp.config.LaunchArrayAccessMode.RELAXED, so CUDA kernels can receive CPU arrays on hardware where CUDA reports pageable CPU memory as GPU-accessible (#1461). The main targets are Linux Heterogeneous Memory Management (HMM) systems and NVIDIA Address Translation Services (ATS) platforms such as GH200, GB200, DGX Spark / GB10, and Jetson Thor, where ordinary CPU allocations can be GPU-accessible without an explicit .to(device) copy. Use wp.can_access() for concrete arrays and Device.can_access() for coarse device checks when code needs to choose a direct launch or an explicit copy at runtime.

import warp as wp

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

device = wp.get_device("cuda:0")
cpu_data = wp.empty(1024, dtype=float, device="cpu")

if wp.can_access(device, cpu_data):
    wp.launch(increment, dim=cpu_data.size, inputs=[cpu_data], device=device)
else:
    gpu_data = cpu_data.to(device)
    wp.launch(increment, dim=gpu_data.size, inputs=[gpu_data], device=device)

Choose the validation mode based on how much pre-launch checking you want:

• wp.config.LaunchArrayAccessMode.RELAXED is the default...

Warp v1.13.0

Warp v1.13 introduces experimental graph capture serialization with CPU replay, letting captured simulations roundtrip through a portable .wrp file and load from standalone C++ on either GPU or CPU. It also adds an experimental cuBQL BVH backend for wp.Mesh that accelerates ray-heavy mesh queries, the wp.bfloat16 scalar type, a pluggable CUDA allocator interface with built-in RAPIDS Memory Manager (RMM) integration, scoped memory tracking with C++-layer call-site attribution, and a batch of new tile primitives (tile_dot, tile_axpy, tile_stack, scatter helpers).

New features

Graph capture serialization and CPU replay

Warp v1.13 introduces a portable serialized-graph format. Operations recorded during wp.capture_begin(apic=True) / wp.capture_end() can be saved to a .wrp file with wp.capture_save() and replayed from either Python or standalone C++ via wp.capture_load(), enabling cross-process and cross-language graph reuse (#1349). CPU graph capture is also new in this release: the same wp.Graph object now replays on CPU through wp.capture_launch(), and the underlying APIC operation log is what gets serialized. A new wp.handle (a uint64 alias) carries wp.Mesh handles across save and load so kernels can keep referencing meshes after deserialization.

import warp as wp

with wp.ScopedDevice("cpu"):
    a = wp.zeros(64, dtype=float)
    b = wp.zeros(64, dtype=float)

    wp.capture_begin(apic=True)
    wp.copy(b, a)
    graph = wp.capture_end()

    wp.capture_save(graph, "demo", inputs={"a": a}, outputs={"b": b})

# Later (in the same process or a fresh one): replay from disk.
with wp.ScopedDevice("cpu"):
    loaded = wp.capture_load("demo")
    loaded.set_param("a", wp.array([1.0] * 64, dtype=float))
    wp.capture_launch(loaded)

Loading and replaying from standalone C++ (CPU device shown). The full example also walks the _modules/ directory, loads each .o via wp_load_obj, resolves kernel symbols, and registers them with wp_apic_register_loaded_cpu_kernel before the first replay. The snippet below elides that boilerplate:

#include "apic.h"
#include "warp.h"

wp_init(nullptr);
APICGraph graph = wp_apic_load_graph(nullptr, "demo.wrp", 1);  // 1 = CPU device

// (Walk demo_modules/, load each .o, and register kernels. See linked example.)

wp_apic_set_param(graph, "a", a_buffer, a_size);
wp_apic_cpu_replay_graph(graph);              // For CUDA: cudaGraphLaunch(wp_apic_get_cuda_graph_exec(graph), stream)
wp_apic_get_param(graph, "b", b_buffer, b_size);

wp_apic_destroy_graph(graph);

See warp/examples/cpp/02_apic_visualization (CUDA replay) and warp/examples/cpp/03_apic_visualization_cpu (CPU replay) for end-to-end demos with OpenGL visualization.

What gets written:

demo.wrp                  # operation byte stream + region snapshots + metadata
demo_modules/
    <hash>.cubin / .meta  # one per CUDA kernel module, arch-pinned
    <hash>.o              # one per CPU kernel module (CPU capture)

Key capabilities:

• wp.capture_save(graph, path, inputs=..., outputs=...) registers named bindings so the consumer side can swap in fresh inputs and read outputs by name without touching the graph topology.
• wp.capture_load() + wp.capture_launch() support replay on both CPU and CUDA. Loaded graphs expose set_param, get_param, and get_param_ptr for each registered binding, plus params and is_loaded properties on wp.Graph.
• wp.handle scalar type and wp.Mesh remap let kernels accept mesh handles whose underlying objects are reconstructed on load. APIC walks @wp.struct fields recursively to find handle pointers and remap them.

Stability and known gaps:

API Capture is experimental, and we plan to keep adding capabilities and closing gaps over future releases (tracker: #1388). For now, regenerate .wrp artifacts when upgrading Warp. The current operation set, handle types, and platform constraints are documented in the Graphs section of the user guide.

cuBQL BVH backend for wp.Mesh

wp.Mesh now accepts bvh_constructor="cubql" to build its acceleration structure with cuBQL, an Apache 2.0-licensed header-only CUDA library for fast BVH construction and traversal (#1286). For ray-heavy workloads on dense static meshes, where the existing SAH builder's exhaustive construction dominates setup time and where ray traversal sits on the simulation hot path, cuBQL typically delivers faster ray queries alongside consistently lower build times than the SAH, median, and LBVH builders. As one specific data point, a Warp-based renderer benchmark on an RTX 4090 (Franka Emika Panda visual mesh, 8192 parallel worlds) saw simulation time drop from 1.41 s to 0.98 s after switching the constructor with no other changes. Speedups depend heavily on mesh size, query mix, and how much of the frame the mesh queries occupy, so benchmark on your own scene before relying on a particular win.

The "cubql" backend currently only routes wp.mesh_query_ray() through cuBQL's traversal kernels. Extending it to point queries, AABB queries, grouped queries, and winding-number support is future work. Today, passing groups=... or support_winding_number=True to a cuBQL wp.Mesh raises a RuntimeError at construction. Calling wp.mesh_query_point_* or wp.mesh_query_aabb_* against a cuBQL mesh silently returns no results. Stick with the default SAH/median/LBVH builders for kernels that mix query types or aren't ray-bound.

import warp as wp

mesh = wp.Mesh(
    points=points,           # wp.array of wp.vec3
    indices=tri_indices,     # wp.array of wp.int32, shape (num_tris * 3,)
    bvh_constructor="cubql",
)

Pluggable CUDA allocator and RMM integration

CUDA device-memory allocations can now be routed through any object implementing the wp.Allocator protocol via wp.set_cuda_allocator(), wp.set_device_allocator(), or scoped with wp.ScopedAllocator (#781). The built-in wp.utils.AllocatorRmm delegates to RAPIDS Memory Manager so Warp can share a memory pool with PyTorch, CuPy, or any other RMM-aware framework, eliminating duplicate caching on GPUs that train and simulate in the same process.

import rmm
import warp as wp

rmm.reinitialize(pool_allocator=True, initial_pool_size=2**30)
wp.set_cuda_allocator(wp.utils.AllocatorRmm())

a = wp.zeros(1024, dtype=wp.float32, device="cuda:0")  # served from the RMM pool

wp.utils.AllocatorRmm requires the rmm package on Linux (pip install rmm-cu12).

Memory tracking

wp.ScopedMemoryTracker and the wp.config.track_memory flag enable allocation tracking with call-site attribution and per-category reports across GPU, host, and pinned-host memory (#1269). Tracking is implemented in the C++ native layer by intercepting all wp_alloc_* / wp_free_* calls, so internal allocations from BVH, hash-grid, mesh, volume, and sparse subsystems show up alongside Python-originated arrays, labeled with their subsystem (e.g. (native:bvh), (native:hashgrid), (native:mesh)).

import warp as wp


@wp.kernel
def fill(x: wp.array[float]):
    i = wp.tid()
    x[i] = float(i)


with wp.ScopedMemoryTracker("training step"):
    a = wp.zeros(1_000_000, dtype=wp.float32, device="cuda:0")
    b = wp.zeros(2_000_000, dtype=wp.float32, device="cuda:0")
    wp.launch(fill, dim=a.size, inputs=[a])
    wp.launch(fill, dim=b.size, inputs=[b])

Output:

Allocation Tracking Report
  Total allocations:
    cuda:0           2 (11.44 MB)
  Peak usage:
    cuda:0           11.44 MB
  Live allocations:
    (none)

Nest trackers to form hierarchical scopes (e.g. "simulation/collision"), and pass report_func=... to redirect the report to a logger or test-assert callback instead of stdout. For unscoped tracking across an entire process, set wp.config.track_memory = True before wp.init() and call wp.print_memory_report() at any time.

wp.bfloat16 scalar data type

wp.bfloat16 joins wp.float16, wp.float32, and wp.float64 as a first-class Warp scalar type (#1332). It supports array allocation, kernel execution, autodiff, DLPack, PyTorch (wp.from_torch(t) for torch.bfloat16 tensors), JAX, and optional NumPy interop via the ml_dtypes package. wp.tile_matmul() and atomic operations also accept bfloat16 tiles, so transformer-style mixed-precision kernels can stay in Warp end to end.

import warp as wp


@wp.kernel
def axpy(a: wp.array[wp.bfloat16], b: wp.array[wp.bfloat16], out: wp.array[wp.bfloat16]):
    i = wp.tid()
    out[i] = a[i] * wp.bfloat16(2.0) + b[i]


a = wp.array([1.0, 2.0, 3.0], dtype=wp.bfloat16)
b = wp.array([0.5, 0.5, 0.5], dtype=wp.bfloat16)
out = wp.zeros_like(a)
wp.launch(axpy, dim=a.size, inputs=[a, b], outputs=[out])
print(out.numpy())  # [2.5 4.5 6.5]

Warp v1.12.1

Warp v1.12.1 is a bugfix release following v1.12.0. For a complete list of changes, see the changelog.

Highlights

• Tile Correctness: Fixed several tile kernel issues, including kernel dispatch using incorrect block_dim across devices (#1254), wp.tile_matmul() and wp.tile_fft() ignoring module-level enable_backward (#1320), and @wp.func with tile parameters failing to compile with shared-memory tiles (#1313). Tile parameters in @wp.func are now passed by reference for both register and shared storage, matching Python's semantics for mutable objects.

• Silent Correctness Bugs: Fixed compile-time constants silently losing precision when passed to 64-bit scalar constructors like wp.float64() (#485), wp.HashGrid neighbor queries missing results for negative coordinates (#1256), and augmented assignments with subscript or attribute targets (e.g., s.field += expr, arr[i] *= expr) double-evaluating the target expression (#1233).

• Type System and Tooling: Fixed struct field assignments converting Warp scalar types to plain Python types (#1288), wp.array[dtype] not being recognized by mypy (#1278), and array annotation repr() displaying raw internal class paths (#1341).

• Kit Extensions Removed: The Omniverse Kit extensions have been removed from this repository (#1296).

• New Examples: Added a differentiable 2-D Navier-Stokes optimization example and three warp.fem examples for Taylor-Green vortex, Kelvin-Helmholtz instability, and shallow water equations.

Announcements

Upcoming removals

The following deprecations will be finalized in Warp 1.13.0:

• Python 3.9 support will be removed. Python 3.10 becomes the minimum supported version. If you are still on Python 3.9, please plan to upgrade.
• wp.isfinite(), wp.isnan(), and wp.isinf() will no longer accept integer types. These functions only produce meaningful results for floating-point inputs and have been emitting deprecation warnings for integer arguments since 1.12.0.
• Private API deprecation warnings will be removed. Code that imports from warp._src directly (e.g., from warp._src.math import ...) will stop working without warning. Please migrate to the public warp.* API.

Acknowledgments

Thanks to @thomasbbrunner for fixing boolean vector element assignment (#1302).

Warp v1.12.0

Warp v1.12 adds experimental hardware-accelerated texture sampling on CUDA GPUs, extends tile programming with element-wise arithmetic operators and differentiable FFT, and broadens JAX interoperability with jax.vmap support. This release also introduces subscript-style type hints for better IDE integration, new quaternion and approximate-math builtins, B-spline shape functions in warp.fem, and a collection of utility and diagnostics APIs.

New features

Hardware-accelerated textures

Experimental. This API may change without a formal deprecation cycle.

Warp v1.12 introduces wp.Texture1D, wp.Texture2D, and wp.Texture3D classes that leverage CUDA texture memory for hardware-accelerated interpolation directly inside Warp kernels. On GPU, texture reads are routed through dedicated texture units that perform filtered lookups in a single instruction, making them ideal for rendering, volume sampling, signed-distance-field queries, and simulation lookup tables. On CPU, a software fallback provides identical semantics so the same kernel code runs on both devices.

import warp as wp
import numpy as np

wp.init()

# 64x64 single-channel height map
data = np.random.rand(64, 64).astype(np.float32)

# Create a 2D texture with bilinear filtering
tex = wp.Texture2D(data, filter_mode=wp.Texture.FILTER_LINEAR)

@wp.kernel
def sample_texture(tex: wp.Texture2D, coords: wp.array[wp.vec2f], out: wp.array[float]):
    i = wp.tid()
    # Coordinates are in [0, 1]; bilinear interpolation is automatic
    out[i] = wp.texture_sample(tex, coords[i], dtype=float)

coords = wp.array(np.random.rand(1024, 2).astype(np.float32), dtype=wp.vec2f)
result = wp.zeros(1024, dtype=float)
wp.launch(sample_texture, dim=1024, inputs=[tex, coords, result])

print(f"Sampled {result.shape[0]} points, range: [{result.numpy().min():.4f}, {result.numpy().max():.4f}]")
# Example output: Sampled 1024 points, range: [0.0069, 0.9793]

Key capabilities:

• 1D / 2D / 3D texture classes (wp.Texture1D, wp.Texture2D, wp.Texture3D) with matching wp.texture_sample() overloads that accept scalar, vec2f, or vec3f coordinates.
• Filter modes: FILTER_POINT for nearest-neighbor sampling and FILTER_LINEAR for bilinear (2D) or trilinear (3D) interpolation.
• Address modes: ADDRESS_WRAP, ADDRESS_CLAMP, ADDRESS_MIRROR, and ADDRESS_BORDER control how out-of-range texture coordinates are handled, configurable per axis.
• Array interop: Texture objects provide copy_from_array() and copy_to_array() methods to transfer data between wp.array objects and texture memory. A cuda_surface property exposes the CUDA surface handle for advanced interop.
• Broad dtype support: Textures accept integer and floating-point data types with 1, 2, or 4 channels. Integer types are automatically normalized to floating-point values on read.

Subscript-style type hints

When annotating kernel parameters with call-syntax forms like wp.array(dtype=float), static type checkers such as Pyright and Pylance flag these as errors because the expressions look like constructor calls rather than type annotations. Warp v1.12 adds subscript-style alternatives that are recognized as valid generic aliases (#1216):

# Before (flagged as error by Pyright/Pylance):
@wp.kernel
def my_kernel(a: wp.array(dtype=float), b: wp.array2d(dtype=wp.vec3)):
    ...

# After (clean subscript syntax):
@wp.kernel
def my_kernel(a: wp.array[float], b: wp.array2d[wp.vec3]):
    ...

The subscript syntax is supported for all array dimensionalities (wp.array[dtype] through wp.array4d[dtype]) as well as wp.tile[dtype] for tile-typed arguments.

Warp's static type checking compatibility is being improved incrementally, and you may encounter other Pyright/Pylance diagnostics that are not yet resolved. If you run into type checking issues, please report them as sub-issues of #549.

Diagnostics utility

The new wp.print_diagnostics() function displays a comprehensive snapshot of the Warp build and runtime environment (software versions, CUDA information, build flags, and available devices) in a single call (#1221). Two companion helpers, wp.get_cuda_toolkit_version() and wp.get_cuda_driver_version(), return the CUDA toolkit and driver versions as integer tuples (#1172). Together these are useful for debugging environment issues, capturing context in CI logs, and providing system information when filing bug reports.

Quaternion and spatial helpers

Warp v1.12 adds quaternion and spatial transformation helpers: wp.quat_from_euler(), wp.quat_to_euler(), wp.transform_twist(), and wp.transform_wrench() (#1237). The Euler conversion functions accept axis indices (0 = X, 1 = Y, 2 = Z) so you can specify arbitrary rotation-order conventions such as ZYX or XYZ, making them suitable for robotics and animation pipelines:

euler = wp.vec3(0.0, wp.PI / 4.0, 0.0)
q = wp.quat_from_euler(euler, 2, 1, 0)  # ZYX convention
print(q)  # [0.0, 0.3826834559440613, 0.0, 0.9238795042037964]

Approximate math intrinsics

wp.div_approx() and wp.inverse_approx() expose GPU hardware fast-math instructions (div.approx.f32 and rcp.approx.ftz.f64) for approximate floating-point division and reciprocal, offering higher throughput at reduced precision (#1199). Only floating-point types are supported. On CPU, both functions fall back to exact arithmetic so the same kernel code runs correctly on either device.

Marching cubes lookup tables

The internal marching cubes lookup tables are now exposed as public class attributes on wp.MarchingCubes: CUBE_CORNER_OFFSETS, EDGE_TO_CORNERS, CASE_TO_TRI_RANGE, and TRI_LOCAL_INDICES (#1151). These tables enable custom marching cubes implementations for advanced use cases such as sparse volume extraction or procedural mesh generation without having to duplicate the standard lookup data.

Graph coloring API

wp.utils.graph_coloring_assign(), wp.utils.graph_coloring_balance(), and wp.graph_coloring_get_groups() are now part of the public API (#1145). These graph coloring utilities were originally introduced in warp.sim in v1.5.0 for use with VBDIntegrator and were removed along with the warp.sim module in v1.10.0. They are now re-introduced as standalone functions in wp.utils, independent of any physics module. They partition a graph into independent color groups, which is useful for parallel constraint solving, conflict-free mesh updates, and other tasks that require concurrent writes to non-adjacent elements.

Tile programming enhancements

Tile arithmetic operators

Tiles now support native Python * and / operators for element-wise multiplication and division, including broadcast between tiles and scalar constants (#1006, #1009). The supported forms are tile * tile, tile * constant, constant * tile for multiplication, and tile / tile, tile / constant, constant / tile for division. All combinations are differentiable and work with scalar, vector, and matrix element types.

import warp as wp

TILE_SIZE = wp.constant(64)

@wp.kernel
def scale_and_normalize(
    a: wp.array[float],
    b: wp.array[float],
    out: wp.array[float],
):
    i = wp.tid()
    ta = wp.tile_load(a, shape=TILE_SIZE, offset=i * TILE_SIZE)
    tb = wp.tile_load(b, shape=TILE_SIZE, offset=i * TILE_SIZE)

    product = ta * tb          # element-wise multiply
    scaled = product * 0.5     # broadcast scalar multiply
    result = scaled / tb       # element-wise divide

    wp.tile_store(out, result, offset=i * TILE_SIZE)

N = 256
a = wp.ones(N, dtype=float)
b = wp.full(N, value=2.0, dtype=float)
out = wp.zeros(N, dtype=float)
wp.launch_tiled(scale_and_normalize, dim=[N // 64], inputs=[a, b, out], block_dim=64)

print(out.numpy()[:8])  # [0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5]

wp.tile_from_thread()

wp.tile_from_thread() broadcasts a scalar or vector value held by a single thread to a shared tile visible to all threads in the block (#1178). This is useful when one thread computes a value (e.g., a reduction result or a loop-invariant parameter) that the entire block needs to use in subsequent tile operations. The function accepts a thread_idx argument to specify which thread's value is broadcast, and supports both "shared" and "register" storage modes.

Differentiable FFT

wp.tile_fft() and wp.tile_ifft() now support reverse-mode automatic differentiation when recorded on a wp.Tape() (#1138). Warp automatically provides the correct gradient implementations for both transforms, so gradients propagate seamlessly through frequency-domain operations. This enables end-to-end gradient computation through pipelines that mix spatial and spectral steps, which is useful for differentiable signal processing, spectral methods, and PDE solvers.

MathDx GEMM toggle

Setting wp.config.enable_mathdx_gemm = False (or passing "enable_mathdx_gemm": False as a module option) disables cuBLASDx for wp.tile_matmul(), falling back to an optimized scalar GEMM implementation (#1228). This avoids the slow link-time optimization (LTO) step required by libmathdx during development iteration, while keeping libmathdx available for operations that have no scalar fallback, such as Cholesky factorization and FFT. The scalar fallback may be slower than cuBLASDx depending on tile sizes, data types, and block_dim, so this option is primarily intended for faster compile–edit–run cycles during development rather than production use.

Accelerated tile load/store

Shared-memory tile loads and stores via wp.tile_load() / wp.tile_store() have been accelerated for non-power-of-two tile sizes (#1239). The improvement is most pronounced when source arrays fit within the GPU L2 cache, reducing...

Warp v1.11.1

Warp v1.11.1 is a bugfix release following v1.11.0. For a complete list of changes, see the changelog.

Highlights

This is primarily a bugfix release with no major new features. Key fixes include:

• Tile Operations: Fixed wp.tile_matmul() sometimes producing NaN results when using the c = wp.tile_matmul(a, b) form due to reading uninitialized output memory (#1180). Also fixed tile multiplication with scalar constants when one operand is a vector or matrix type (#1175), and enabled scalar, vector, and matrix arguments in wp.tile_map() (#1136).

• Code Generation: Fixed wp.static() incorrectly resolving loop variables to same-named global Python variables when used for static loop unrolling in kernels (e.g., wp.static(i) inside for i in range(n) would use a global i if one existed, instead of the loop iteration value) (#1139). Also fixed a segfault in conditional expressions (ternary if/else) when one branch accesses an array element and the other branch is taken.

• CUDA Graphs: Fixed CUDA graphs with multiple temporary allocations using more memory than necessary. Previously, memory freed during graph capture wasn't properly sequenced for reuse by later allocations, causing memory to accumulate (e.g., three sequential 1GB allocations would consume 3GB instead of reusing the same 1GB).

• Developer Experience: Fixed @wp.func decorated functions showing generic _Wrapped types in Pyright/Pylance instead of their actual signatures on Python 3.10+. Also fixed multiple issues with IDE autocomplete stubs that caused type checker errors in mypy and Pyright, including incorrect @overload usage, shadowed bool type references, and missing Literal[] syntax for integer type parameters.

• Documentation: Added missing docstrings across the API, including built-in functions, type constructors, constants, and warp.fem symbols (#1159). Added example_particle_repulsion.py demonstrating how to use wp.grad() (#1137).

Announcements

Upcoming removals

The following feature is deprecated and will be removed in v1.12:

• Constructing matrices from column vectors via wp.matrix(): The ability to construct matrices by passing column vectors to wp.matrix() has been deprecated at both Python and kernel scopes. Use wp.matrix_from_cols() instead.

• Constructing matrices from vectors via wp.matrix(): The ability to construct matrices by passing vectors to wp.matrix() has been deprecated at both Python and kernel scopes. The replacement depends on the scope because the behavior was inconsistent:

  # Kernel scope (vectors were interpreted as columns)
  # Deprecated (will be removed in v1.12):
  m = wp.mat33(col0, col1, col2)
  # Use instead:
  m = wp.matrix_from_cols(col0, col1, col2)
  
  # Python scope (vectors were interpreted as rows)
  # Deprecated (will be removed in v1.12):
  m = wp.mat22(wp.vec2(1, 2), wp.vec2(3, 4))
  # Use instead:
  m = wp.matrix_from_rows(wp.vec2(1, 2), wp.vec2(3, 4))

Acknowledgments

We thank the following contributors:

• @Adityakk9031 for fixing the JAX FFI deadlock when using cached graphs across multiple GPUs (#1181).
• @Cucchi01 for fixing bsr_get_diag() not zeroing the output buffer when provided (#1170).
• @liblaf for fixing the inverted verbose flag in wp.capture_debug_dot_print() (#1202).