v1.17.0

[shi-eric]

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

Important

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

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 PTX for the compute capability used by the external application.
artifacts = wp.compile_aot_module(
    square_kernel.module,
    arch=80,  # Match the external application's target compute capability.
    module_dir=Path(__file__).parent / "generated",
    use_ptx=True,
    strip_hash=True,  # Keep the exported kernel name free of hash suffixes.
)
print(artifacts[0].suffix)  # .ptx

This example stops after generating PTX. It does not include the external application that loads and launches the PTX.

External native value types do not gain arithmetic or differentiation automatically. Register those functions separately.

Platform and toolchain

• Warp can now be built from source to run CPU kernels natively on Windows ARM64 (Support Windows ARM64 CPU builds #1755). We plan to publish Windows ARM64 wheels and add CUDA support in future releases.
• Warp's embedded Clang/LLVM libraries now use version 22.1.8 on every platform. Source builds download the prebuilt Clang/LLVM SDK from GitHub Releases and verify each archive against a pinned checksum.

Additional improvements

• Python transform construction is faster, especially for wp.transform() with no arguments and same-type copies. Unsupported Python arguments now raise TypeError instead of silently returning an all-zero transform, and scalar construction inside kernels fills all seven components (wp.transform() silently returns a zero value on unsupported arguments #1742, wp.transform() fill and copy construction fail at kernel scope #1814).
• Custom-strided wp.empty() and wp.zeros() allocations now account for gaps correctly and reject negative or dimensionally invalid strides (wp.empty() and wp.zeros() can allocate insufficient storage for custom-strided arrays #1703).
• Generated callables used with wp.map() and wp.utils.create_warp_function() now have stable identities across processes. This allows persistent kernel-cache reuse when separate processes discover the callables in different orders (Salted hash() in create_warp_function makes kernel hashing nondeterministic, defeating and eventually poisoning the kernel cache #1696).
• New guides cover reducing compilation and startup time, using Warp tapes safely inside PyTorch autograd functions, and understanding tile shared-memory requirements.

New example

example_fdtd_3d.py is a three-dimensional finite-difference time-domain simulation on a Yee grid. It models a Luneburg lens that collimates waves from a point source and supports both interactive Matplotlib visualization and headless execution (#1772).

Announcements

CUDA 13 PyPI wheel timing

• The tentative CUDA 13.4 PyPI wheel transition did not occur in Warp 1.17. CUDA Toolkit 13.4 is not yet generally available. We plan to try again for Warp 1.18.
• A future switch to CUDA 13 wheels will require an R580-series or newer NVIDIA driver and a Turing-class GPU (compute capability 7.5) or newer. pip does not check the installed driver or GPU architecture and may install a wheel whose CUDA backend cannot run on the system.
• CUDA 12 support will continue after the PyPI transition. CUDA 12 builds will remain available from GitHub Releases, and users can continue to build Warp from source.

This updates the tentative plan announced in Warp v1.16.0. See NVIDIA's CUDA minor-version compatibility table for driver requirements.

Removals and deprecations

• Implicit conversion of Python and Warp numeric scalars to vector, matrix, quaternion, and transform kernel parameters and struct fields has been removed. Construct the expected composite value explicitly, such as wp.mat22(123) (Remove implicit scalar promotion to composite types #1721).
• The same implicit conversion from NumPy numeric scalars and from Python, NumPy, and Warp Boolean values is now deprecated. Use an explicit composite constructor, such as wp.mat22(np.float32(123)). This conversion will be removed in a future feature release under Warp's standard deprecation timeline (Remove implicit scalar promotion to composite types #1721).
• wp.Texture.copy_from_array() and wp.Texture.copy_to_array() have been removed. Use wp.Texture.copy_from() and wp.Texture.copy_to() (Remove deprecated Texture array-copy aliases #1722).
• wp.spatial_jacobian() and wp.spatial_mass() have been removed. Although previously listed as built-ins, neither function was callable from kernels or Python (spatial_jacobian() and spatial_mass() cannot be called from kernels or from Python #1768).

Migration examples:

- wp.launch(consume_matrix, dim=1, inputs=[123])
+ wp.launch(consume_matrix, dim=1, inputs=[wp.mat22(123)])

- wp.launch(consume_matrix, dim=1, inputs=[np.float32(123)])
+ wp.launch(consume_matrix, dim=1, inputs=[wp.mat22(np.float32(123))])

Acknowledgments

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

• @arko92 for fixing the CPU JIT race that could make kernels unavailable after parallel module loading (Concurrent CPU module loads can fail kernel lookup #1705).
• @mehdiataei for adding restart support to CG and CR and correcting iteration counts for graph-captured solver cycles (Fix iteration counts for graph-captured multi-iteration solver cycles #1707, Add restart support to CG and CR #1708).
• @pei-tian for fixing storage underallocation in custom-strided wp.empty() and wp.zeros() arrays (wp.empty() and wp.zeros() can allocate insufficient storage for custom-strided arrays #1703).
• @thomasbbrunner for reporting the closure-generated AOT kernel-name collision and contributing initial implementation work toward @wp.kernel(name=...) (Allow passing a key/name to the @wp.kernel #1561, Add name argument to @wp.kernel (GH-1561) #1570).

---

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

---

This discussion was created from the release v1.17.0.