NVIDIA Warp

https://nvidia.github.io/warp/stable/
Admin
NVIDIA Warp is a Python framework for GPU-accelerated simulation, robotics, and machine learning,...
Tokens:421,412
Snippets:3,970
Trust Score:9.7
Update:1 month ago
Show doc for...
Tokens:
Raw
### Install Warp with examples

Source: https://nvidia.github.io/warp/stable/user_guide/faq.html

Installs the Warp package along with additional dependencies required for running examples or USD rendering.

```bash
warp-lang[examples]
```

--------------------------------

### Run Warp examples

Source: https://nvidia.github.io/warp/stable/_sources/index.rst.txt

Command-line syntax to execute specific examples after installing optional dependencies.

```sh
python -m warp.examples.<example_subdir>.<example>
```

--------------------------------

### Benchmark output example

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/interoperability/pytorch.rst.txt

Example output showing performance differences between conversion methods.

```text
5095 ms  from_torch(...)
2113 ms  from_torch(..., return_ctype=True)
2950 ms  direct from torch
```

--------------------------------

### Install clang-format locally

Source: https://nvidia.github.io/warp/stable/project/contribution_guide.html

Install the specific version of clang-format required for IDE integration.

```bash
# Ubuntu/Debian - Install from apt.llvm.org
# See https://apt.llvm.org/ for repository setup instructions
# Check .pre-commit-config.yaml for the current version
sudo apt-get install clang-format-21
```

--------------------------------

### Example ScopedTimer output with synchronization

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/execution_and_performance/profiling.rst.txt

Example output format when synchronize=True is enabled.

```console
Demo took 4.91 ms
```

--------------------------------

### Warp Initialization Output

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/basics.rst.txt

Example of the startup information printed when Warp initializes.

```bat
Warp 1.2.0 initialized:
CUDA Toolkit 12.5, Driver 12.5
Devices:
    "cpu"      : "x86_64"
    "cuda:0"   : "NVIDIA GeForce RTX 3090" (24 GiB, sm_86, mempool enabled)
    "cuda:1"   : "NVIDIA GeForce RTX 3090" (24 GiB, sm_86, mempool enabled)
CUDA peer access:
    Supported fully (all-directional)
Kernel cache:
    /home/nvidia/.cache/warp/1.2.0
```

--------------------------------

### Install Warp via pip

Source: https://nvidia.github.io/warp/stable/_sources/index.rst.txt

Standard installation command for the Warp framework from PyPI.

```sh
$ pip install warp-lang
```

--------------------------------

### Install clang-format on Ubuntu/Debian

Source: https://nvidia.github.io/warp/stable/_sources/project/contribution_guide.rst.txt

Installs the clang-format tool locally for IDE integration, matching the version required by the project.

```bash
sudo apt-get install clang-format-21
```

--------------------------------

### Install Specific Platform Wheels

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/installation.rst.txt

Direct installation commands for specific platform binaries from GitHub releases.

```sh
pip install https://github.com/NVIDIA/warp/releases/download/v1.16.0/warp_lang-1.16.0+cu13-py3-none-manylinux_2_34_aarch64.whl
```

```sh
pip install https://github.com/NVIDIA/warp/releases/download/v1.16.0/warp_lang-1.16.0+cu13-py3-none-manylinux_2_28_x86_64.whl
```

```sh
pip install https://github.com/NVIDIA/warp/releases/download/v1.16.0/warp_lang-1.16.0+cu13-py3-none-win_amd64.whl
```

--------------------------------

### Launch a 2D kernel

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/basics.rst.txt

Example of launching a kernel with a 2D grid of threads.

```python
wp.launch(kernel=compute_image, dim=(1024, 1024), inputs=[img], device="cuda")
```

--------------------------------

### Output of mesh_get_point example

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.mesh_get_point.html

The expected output when running the provided mesh_get_point example.

```text
[0. 0. 0.]
```

--------------------------------

### Install Nightly Builds

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/installation.rst.txt

Commands to install the latest development version from the NVIDIA package index.

```sh
$ pip install -U --pre warp-lang --extra-index-url=https://pypi.nvidia.com/
```

```text
export PIP_EXTRA_INDEX_URL="https://pypi.nvidia.com"
```

--------------------------------

### Expected Output of tile_stack Example

Source: https://nvidia.github.io/warp/stable/_sources/language_reference/_generated/warp.tile_stack.rst.txt

The console output resulting from the execution of the provided kernel example.

```text
[6, 7, 8, 9]
```

--------------------------------

### CUDA Driver Warning Example

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/installation.rst.txt

Example of the warning message displayed when the installed CUDA driver version is insufficient.

```text
Warp UserWarning:
   Insufficient CUDA driver version.
   The minimum required CUDA driver version is 12.0, but the installed CUDA driver version is 11.8.
   Visit https://nvidia.github.io/warp/stable/user_guide/installation.html for guidance.
```

--------------------------------

### Running Test Suite via CLI

Source: https://nvidia.github.io/warp/stable/_sources/project/contribution_guide.rst.txt

Example output when executing a test module configured for multiple devices.

```bash
uv run test_amazing_code.py
Warp 1.11.1 initialized:
CUDA Toolkit 12.9, Driver 13.1
Devices:
    "cpu"      : "x86_64"
    "cuda:0"   : "NVIDIA RTX 6000 Ada Generation" (48 GiB, sm_89, mempool enabled)
    "cuda:1"   : "NVIDIA RTX 6000 Ada Generation" (48 GiB, sm_89, mempool enabled)
CUDA peer access:
    Supported fully (all-directional)
Kernel cache:
    /home/nvidia/.cache/warp/1.11.1
test_amazing_code_test_one_cpu (__main__.TestAmazingCode) ... ok
test_amazing_code_test_one_cuda_0 (__main__.TestAmazingCode) ... ok
test_amazing_code_test_one_cuda_1 (__main__.TestAmazingCode) ... ok

----------------------------------------------------------------------
Ran 3 tests in 0.001s

OK
```

--------------------------------

### Launch Tiled Kernel

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/tiles.rst.txt

Example of launching a reduction kernel using tiles, with an optional block_dim parameter to tune performance.

```python
data = wp.array(rng.random((4096, 4096), dtype=np.float64), dtype=wp.float64)
result = wp.zeros((1,), dtype=wp.float64)

wp.launch(reduce_array_tile, (4096, 4096), inputs=[data], outputs=[result])
```

```python
wp.launch(reduce_array_tile, (4096, 4096), inputs=[data], outputs=[result], block_dim=128)
```

--------------------------------

### Docker Output Example

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/installation.rst.txt

Expected output after initializing Warp in a container.

```text
==========
== CUDA ==
==========

CUDA Version 13.0.0

Container image Copyright (c) 2016-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.

This container image and its contents are governed by the NVIDIA Deep Learning Container License.
By pulling and using the container, you accept the terms and conditions of this license:
https://developer.nvidia.com/ngc/nvidia-deep-learning-container-license

A copy of this license is made available in this container at /NGC-DL-CONTAINER-LICENSE for your convenience.

Warp 1.10.0.dev0 initialized:
CUDA Toolkit 13.0, Driver 13.0
Devices:
    "cpu"      : "x86_64"
    "cuda:0"   : "NVIDIA L40S" (47 GiB, sm_89, mempool enabled)
Kernel cache:
  /root/.cache/warp/1.10.0.dev0
```

--------------------------------

### Install Scoped Logger

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/debugging.rst.txt

Temporarily install a logger that restores the previous configuration upon context exit.

```python
with wp.ScopedLogger(my_capture_logger):
    wp.launch(my_kernel, ...)  # diagnostics flow to my_capture_logger
```

--------------------------------

### Warp Initialization Output

Source: https://nvidia.github.io/warp/stable/user_guide/debugging.html

Example output from wp.init() showing the kernel cache directory location.

```text
Warp 0.8.1 initialized:
    CUDA Toolkit: 11.8, Driver: 11.8
    Devices:
    "cpu"    | AMD64 Family 25 Model 33 Stepping 0, AuthenticAMD
    "cuda:0" | NVIDIA GeForce RTX 3090 (sm_86)
    "cuda:1" | NVIDIA GeForce RTX 2080 Ti (sm_75)
    Kernel cache: C:\Users\LukasW\AppData\Local\NVIDIA Corporation\warp\Cache\0.8.1
```

--------------------------------

### Inclusive Scan Example

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.tile_scan_inclusive.html

Demonstrates performing an inclusive scan on a tile created with tile_arange.

```python
@wp.kernel
def scan_example():
    t = wp.tile_arange(1, 5, dtype=int)
    s = wp.tile_scan_inclusive(t)
    print(s)

wp.launch_tiled(scan_example, dim=[1], inputs=[], block_dim=16)
```

--------------------------------

### Example output

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.tile_max.html

The expected output format for the tile maximum operation.

```text
[127] = tile(shape=(1), storage=register)
```

--------------------------------

### Tile argmax output example

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.tile_argmax.html

Shows the expected output format for a tile containing the argmax index.

```text
[63] = tile(shape=(1), storage=register)
```

--------------------------------

### Inclusive Min Scan Example

Source: https://nvidia.github.io/warp/stable/_sources/language_reference/_generated/warp.tile_scan_min_inclusive.rst.txt

Demonstrates loading an array into a tile and performing an inclusive min scan operation within a kernel.

```python
@wp.kernel
def scan_example(input: wp.array[int]):
    t = wp.tile_load(input, shape=(4,))
    s = wp.tile_scan_min_inclusive(t)
    print(s)

input = wp.array([3, 1, 4, 2], dtype=int)
wp.launch_tiled(scan_example, dim=[1], inputs=[input], block_dim=16)
```

```text
[3, 1, 1, 1] = tile(shape=(4), storage=register)
```

--------------------------------

### Build Warp from source

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/compatibility.rst.txt

Commands to build the library and install it in editable mode.

```console
$ python build_lib.py
$ pip install -e .
```

--------------------------------

### Configure Nsight Compute for API-based Capture

Source: https://nvidia.github.io/warp/stable/user_guide/execution_and_performance/profiling.html

Disable profiling at startup using --profile-from-start off to ensure collection only occurs when the profiler start API is invoked.

```bash
ncu --profile-from-start off -o my_app python my_app.py
```

--------------------------------

### CPU Execution Output

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/tiles.rst.txt

Example output when launching a kernel on the CPU where block_dim is 1.

```text
[ 0  1  2  3  4  5  6  7  8  9 10 11]
```

--------------------------------

### Warp Initialization Output

Source: https://nvidia.github.io/warp/stable/user_guide/execution_and_performance/memory_management.html

Example of the console output during Warp initialization showing peer access support status.

```text
Warp 0.15.1 initialized:
   CUDA Toolkit 11.5, Driver 12.2
   Devices:
     "cpu"      : "x86_64"
     "cuda:0"   : "NVIDIA L40" (48 GiB, sm_89, mempool enabled)
     "cuda:1"   : "NVIDIA L40" (48 GiB, sm_89, mempool enabled)
     "cuda:2"   : "NVIDIA L40" (48 GiB, sm_89, mempool enabled)
     "cuda:3"   : "NVIDIA L40" (48 GiB, sm_89, mempool enabled)
   CUDA peer access:
     Supported fully (all-directional)
```

--------------------------------

### Synchronizing Streams for Kernel Execution

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/execution_and_performance/concurrency.rst.txt

Provides the initial setup for synchronizing streams before kernel execution.

```python
a = wp.zeros(n, dtype=float, device="cuda:0")

s = wp.Stream("cuda:0")
```

--------------------------------

### Observe module reloading output

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/generics.rst.txt

Example output showing the module load time and execution for the intermixed kernel definitions.

```text
Module __main__ load on device 'cuda:0' took 155.73 ms
17
Module __main__ load on device 'cuda:0' took 164.83 ms
42
```

--------------------------------

### Verify Warp Installation

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/installation.rst.txt

Run the diagnostic utility to ensure native libraries are loaded correctly.

```console
$ uv run python -c "import warp; warp.print_diagnostics()"
```

--------------------------------

### Load a Volume from a NanoVDB File

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/runtime.rst.txt

Example of creating a Warp Volume object from a file on disk.

```python
# open NanoVDB file on disk
file = open("mygrid.nvdb", "rb")

# create Volume object
volume = wp.Volume.load_from_nvdb(file, device="cpu")
```

--------------------------------

### Naive Branching Implementation

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/code_generation.rst.txt

Initial setup for a kernel using runtime branching, which may lead to performance issues.

```python
import warp as wp
import numpy as np
```

--------------------------------

### Warp in-place operation output

Source: https://nvidia.github.io/warp/stable/user_guide/differentiability.html

The expected output for the in-place operation example.

```text
[ 5. 11. 19.]
```

--------------------------------

### Warp gradient output

Source: https://nvidia.github.io/warp/stable/user_guide/differentiability.html

The expected output for the Warp gradient calculation example.

```text
[5. 7. 9.]
```

--------------------------------

### Checking Initialization Status

Source: https://nvidia.github.io/warp/stable/user_guide/execution_and_performance/memory_management.html

Example output from Warp initialization showing mempool status for devices.

```text
Warp 0.15.1 initialized:
CUDA Toolkit 11.5, Driver 12.2
Devices:
    "cpu"      : "x86_64"
    "cuda:0"   : "NVIDIA GeForce RTX 4090" (24 GiB, sm_89, mempool enabled)
    "cuda:1"   : "NVIDIA GeForce RTX 3090" (24 GiB, sm_86, mempool enabled)
```

--------------------------------

### Test Warp in Docker

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/installation.rst.txt

Verify the Warp installation within a container by running the initialization command.

```sh
docker run --rm --gpus all warp-github-clone:example python3 -c "import warp as wp; wp.init()"
```

--------------------------------

### Build Documentation

Source: https://nvidia.github.io/warp/stable/_sources/project/contribution_guide.rst.txt

Builds the project documentation from the root directory.

```bash
uv run --extra docs build_docs.py
```

--------------------------------

### Define Warp kernels for autograd

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/interoperability/pytorch.rst.txt

Example setup for forward and backward kernels used within a custom PyTorch autograd function.

```python
import warp as wp
import torch

device = 'cuda'
N = 300_000_000

@wp.kernel(enable_backward=False)
def forward_kernel(
    a: wp.array[float],
    b: wp.array[float],
    output: wp.array[float]
):
    i = wp.tid()
    x = a[i]
    y = b[i]
    output[i] = x*x + y*y


@wp.kernel(enable_backward=False)
def backward_kernel(
    grad_output: wp.array[float],
    a: wp.array[float],
    b: wp.array[float],
    grad_a: wp.array[float],
    grad_b: wp.array[float]
):
    i = wp.tid()
    x = a[i]
    y = b[i]
    adj_z = grad_output[i]

    grad_a[i] = 2.0 * x * adj_z
    grad_b[i] = 2.0 * y * adj_z


class WarpFunction(torch.autograd.Function):
```

--------------------------------

### Querying and using launch configuration for a SAXPY kernel

Source: https://nvidia.github.io/warp/stable/api_reference/_generated/warp.get_suggested_block_size.html

Demonstrates how to retrieve a suggested block size for a kernel and apply it during a launch.

```python
@wp.kernel
def saxpy(alpha: float, x: wp.array[float], y: wp.array[float]):
    i = wp.tid()
    y[i] = alpha * x[i] + y[i]


n = 1000000
x = wp.ones(n, dtype=float)
y = wp.zeros(n, dtype=float)

block_size, min_grid_size = wp.get_suggested_block_size(saxpy)
wp.launch(saxpy, dim=n, inputs=[2.0, x, y], block_dim=block_size)
```

--------------------------------

### Example output

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.mesh_query_ray_count_intersections.html

The expected output from the provided intersection counting example.

```text
crossings: 2
```

--------------------------------

### Inclusive Max Scan Example

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.tile_scan_max_inclusive.html

Demonstrates loading an array into a tile and performing an inclusive max scan operation within a kernel.

```python
@wp.kernel
def scan_example(input: wp.array[int]):
    t = wp.tile_load(input, shape=(4,))
    s = wp.tile_scan_max_inclusive(t)
    print(s)

input = wp.array([3, 1, 4, 2], dtype=int)
wp.launch_tiled(scan_example, dim=[1], inputs=[input], block_dim=16)
```

--------------------------------

### Example ScopedTimer output

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/execution_and_performance/profiling.rst.txt

Example output format for a ScopedTimer measurement.

```console
Demo took 0.52 ms
```

--------------------------------

### Initialize and Build a Grouped HashGrid

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/runtime.rst.txt

Demonstrates the initialization and construction of a HashGrid with group support.

```python
grid = wp.HashGrid(dim_x=8, dim_y=8, dim_z=8, device="cuda:0")
grid.build(points=points, radius=0.5, groups=groups)
```

--------------------------------

### Module Compilation Output

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/code_generation.rst.txt

Example output showing module compilation times for interleaved kernel definitions.

```text
Module __main__ 6cd1d53 load on device 'cuda:0' took 168.19 ms  (compiled)
foo
Module __main__ c7c0e9a load on device 'cuda:0' took 160.35 ms  (compiled)
bar
```

--------------------------------

### Build and test Docker image

Source: https://nvidia.github.io/warp/stable/user_guide/installation.html

Commands to build a Docker image from a Dockerfile and verify the Warp installation.

```bash
docker build -t warp-github-clone:example .
```

```bash
docker run --rm --gpus all warp-github-clone:example python3 -c "import warp as wp; wp.init()"
```

--------------------------------

### Print Warp diagnostics

Source: https://nvidia.github.io/warp/stable/user_guide/faq.html

Displays the installed Warp build information and detected hardware devices.

```python
wp.print_diagnostics()
```

--------------------------------

### Configure CMake with custom LLVM path

Source: https://nvidia.github.io/warp/stable/user_guide/installation.html

Uses an existing LLVM installation for building warp-clang.

```bash
$ cmake -S . -B _build/cmake -G Ninja -DWARP_LLVM_PATH=/opt/llvm
$ cmake --build _build/cmake --parallel
```

--------------------------------

### Optimizing PyTorch and Warp Data Conversion

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/interoperability/pytorch.rst.txt

Examples showing how to avoid redundant wp.from_torch calls inside loops to improve performance.

```python
x_t = torch.arange(n, dtype=torch.float32, device=device)
y_t = torch.ones(n, dtype=torch.float32, device=device)

for i in range(10):
    x_w = wp.from_torch(x_t)
    y_w = wp.from_torch(y_t)
    wp.launch(saxpy, dim=n, inputs=[x_w, y_w, 1.0], device=device)
```

```python
x_t = torch.arange(n, dtype=torch.float32, device=device)
y_t = torch.ones(n, dtype=torch.float32, device=device)

x_w = wp.from_torch(x_t)
y_w = wp.from_torch(y_t)

for i in range(10):
    wp.launch(saxpy, dim=n, inputs=[x_w, y_w, 1.0], device=device)
```

--------------------------------

### Atomic Tile Addition Example

Source: https://nvidia.github.io/warp/stable/_sources/language_reference/_generated/warp.tile_atomic_add_indexed.rst.txt

Demonstrates a blocked, row-wise reduction using tile_atomic_add_indexed within a kernel.

```python
TILE_M = wp.constant(2)
TILE_N = wp.constant(2)

@wp.kernel
def tile_atomic_add_indexed(x: wp.array2d[float], y: wp.array2d[float]):
    i, j = wp.tid()

    t = wp.tile_load(x, shape=(TILE_M, TILE_N), offset=(i*TILE_M, j*TILE_N), storage="register")

    zeros = wp.tile_zeros(TILE_M, dtype=int, storage="shared")

    wp.tile_atomic_add_indexed(y, indices=zeros, t=t, offset=(i, j*TILE_N), axis=0)

M = TILE_M * 2
N = TILE_N * 2

arr = np.arange(M * N, dtype=float).reshape(M, N)

x = wp.array(arr, dtype=float, requires_grad=True, device=device)
y = wp.zeros((2, N), dtype=float, requires_grad=True, device=device)

wp.launch_tiled(tile_atomic_add_indexed, dim=[2,2], inputs=[x], outputs=[y], block_dim=32, device=device)

print(x.numpy())
print(y.numpy())
```

```text
[[ 0.  1.  2.  3.]
    [ 4.  5.  6.  7.]
    [ 8.  9. 10. 11.]
    [12. 13. 14. 15.]]

[[ 4.  6.  8. 10.]
    [20. 22. 24. 26.]]
```

--------------------------------

### Configure CMake with custom CUDA path

Source: https://nvidia.github.io/warp/stable/user_guide/installation.html

Specifies a custom installation path for the CUDA Toolkit.

```bash
$ cmake -S . -B _build/cmake -G Ninja -DWARP_CUDA_PATH=/usr/local/cuda
$ cmake --build _build/cmake --parallel
```

--------------------------------

### Create and Sample Warp Textures

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/runtime.rst.txt

Shows how to initialize 1D, 2D, and 3D textures from arrays and sample them inside a kernel.

```python
import warp as wp
import numpy as np

# Create a 256-element RGBA 1D texture
data1d = np.random.rand(256, 4).astype(np.float32)
tex1d = wp.Texture1D(data1d, filter_mode=wp.TextureFilterMode.LINEAR, device="cuda:0")

# Create a 256x256 RGBA 2D texture
data = np.random.rand(256, 256, 4).astype(np.float32)
tex2d = wp.Texture2D(data, filter_mode=wp.TextureFilterMode.LINEAR, device="cuda:0")

# Create a 64x64x64 single-channel 3D texture
data3d = np.random.rand(64, 64, 64).astype(np.float32)
tex3d = wp.Texture3D(data3d, filter_mode=wp.TextureFilterMode.LINEAR, device="cuda:0")

# Explicitly allocate a texture using Warp dtypes
tex_alloc = wp.Texture2D(width=256, height=256, num_channels=4, dtype=wp.float32, device="cuda:0")
```

```python
@wp.kernel
def sample_texture(
    tex: wp.Texture2D,
    uvs: wp.array[wp.vec2f],
    output: wp.array[float],
):
    tid = wp.tid()
    uv = uvs[tid]
    output[tid] = wp.texture_sample(tex, uv, dtype=float)
```

--------------------------------

### Inclusive Max Scan Example

Source: https://nvidia.github.io/warp/stable/_sources/language_reference/_generated/warp.tile_scan_max_inclusive.rst.txt

Demonstrates performing an inclusive max scan on a tile within a Warp kernel.

```python
@wp.kernel
def scan_example(input: wp.array[int]):
    t = wp.tile_load(input, shape=(4,))
    s = wp.tile_scan_max_inclusive(t)
    print(s)

input = wp.array([3, 1, 4, 2], dtype=int)
wp.launch_tiled(scan_example, dim=[1], inputs=[input], block_dim=16)
```

```text
[3, 3, 4, 4] = tile(shape=(4), storage=register)
```

--------------------------------

### Build and Run Workflow

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/runtime.rst.txt

Standard commands for capturing a graph, building the C++ viewer, and executing the application.

```bash
cd warp/examples/cpp/02_apic_visualization     # or 03_apic_visualization_cpu

# 1. Capture the graph (requires Warp + Python)
uv run capture_wave.py

# 2. Build the C++ viewer
cmake -B build -DCMAKE_BUILD_TYPE=Release
cmake --build build --config Release

# 3. Run
./build/02_apic_visualization                  # or 03_apic_visualization_cpu
```

--------------------------------

### Assertion Failure Output

Source: https://nvidia.github.io/warp/stable/user_guide/debugging.html

Example of the console output generated when an assertion fails during kernel execution.

```text
default_program:49: void expect_ones_133f9859_cuda_kernel_forward(wp::launch_bounds_t, wp::array_t<int>): block: [0,0,0], thread: [0,0,0] Assertion `("assert a[i] == 1, \"Array element must be 1\"",var_3)` failed.
```

--------------------------------

### Install pre-commit hooks

Source: https://nvidia.github.io/warp/stable/_sources/project/contribution_guide.rst.txt

Configures the local git repository to run linting and formatting checks automatically during commits.

```bash
uvx pre-commit install
```

--------------------------------

### Module Reloading Log Output

Source: https://nvidia.github.io/warp/stable/user_guide/programming_model/generics.html

Example log output showing the overhead caused by implicit kernel instantiation.

```text
Module __main__ load on device 'cuda:0' took 170.37 ms
Module __main__ load on device 'cuda:0' took 171.43 ms
Module __main__ load on device 'cuda:0' took 179.49 ms
```

--------------------------------

### Install NVTX Python package

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/execution_and_performance/profiling.rst.txt

Command to install the NVTX dependency required for NVTX integration.

```bash
pip install nvtx
```

--------------------------------

### Compare launch_tiled and launch

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/tiles.rst.txt

Demonstrates that wp.launch_tiled and wp.launch are equivalent when block_dim is explicitly provided.

```python
# These are equivalent:
wp.launch_tiled(kernel, dim=[M, N],     inputs=[...], block_dim=64)
wp.launch(      kernel, dim=[M, N, 64], inputs=[...], block_dim=64)
```

--------------------------------

### setup_tiled_rendering()

Source: https://nvidia.github.io/warp/stable/api_reference/_generated/warp.render.OpenGLRenderer.html

Configures the renderer for tiled rendering, allowing the render buffer to be split into multiple tiles with independent view and projection matrices.

```APIDOC
### setup_tiled_rendering(instances, rescale_window=False, tile_width=None, tile_height=None, tile_ncols=None, tile_nrows=None, tile_positions=None, tile_sizes=None, projection_matrices=None, view_matrices=None)

Set up tiled rendering where the render buffer is split into multiple tiles that can visualize different shape instances of the scene.

#### Parameters
- **instances** (list[list[int]]) - Required - A list of lists of shape instance ids.
- **rescale_window** (bool) - Optional - If True, the window will be resized to fit the tiles.
- **tile_width** (int | None) - Optional - The width of each tile in pixels.
- **tile_height** (int | None) - Optional - The height of each tile in pixels.
- **tile_ncols** (int | None) - Optional - Number of tiles horizontally.
- **tile_nrows** (int | None) - Optional - Number of tiles vertically.
- **tile_positions** (list[tuple[int]] | None) - Optional - List of (x, y) tuples for tile positions.
- **tile_sizes** (list[tuple[int]] | None) - Optional - List of (width, height) tuples for tile sizes.
- **projection_matrices** (list | None) - Optional - List of projection matrices for each tile.
- **view_matrices** (list | None) - Optional - List of view matrices for each tile.
```

--------------------------------

### Generated SSA Code Example

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/code_generation.rst.txt

Example of the static-single-assignment (SSA) form generated by Warp for a simple function.

```c++
static CUDA_CALLABLE wp::float32 my_func_0(
    wp::float32 var_a,
    wp::float32 var_b)
{
    //---------
    // primal vars
    wp::float32 var_0;
    wp::float32 var_1;
    //---------
    // forward
    // def my_func(a: float, b: float):                                                       <L 6>
    // c = wp.sin(b) * a                                                                      <L 7>
    var_0 = wp::sin(var_b);
    var_1 = wp::mul(var_0, var_a);
    // return c                                                                               <L 8>
    return var_1;
}
```

--------------------------------

### Demonstrating Deferred Gradient Allocation Synchronization Issues

Source: https://nvidia.github.io/warp/stable/user_guide/interoperability/pytorch.html

This example illustrates how calling wp.from_torch() on tensors with requires_grad=True triggers immediate gradient allocation, causing expensive device-wide synchronization.

```python
import warp as wp
import torch

device = 'cuda'
N = 300_000_000

@wp.kernel(enable_backward=False)
def forward_kernel(
    a: wp.array[float],
    b: wp.array[float],
    output: wp.array[float]
):
    i = wp.tid()
    x = a[i]
    y = b[i]
    output[i] = x*x + y*y


@wp.kernel(enable_backward=False)
def backward_kernel(
    grad_output: wp.array[float],
    a: wp.array[float],
    b: wp.array[float],
    grad_a: wp.array[float],
    grad_b: wp.array[float]
):
    i = wp.tid()
    x = a[i]
    y = b[i]
    adj_z = grad_output[i]

    grad_a[i] = 2.0 * x * adj_z
    grad_b[i] = 2.0 * y * adj_z


class WarpFunction(torch.autograd.Function):

    @staticmethod
    def forward(ctx, a, b):
        ctx.save_for_backward(a, b)

        device = wp.device_from_torch(a.device)

        output = torch.empty(N, device=a.device)
        wp.launch(
            kernel=forward_kernel,
            dim=(N),
            device=device,
            inputs=[
                wp.from_torch(a),      # ⚠️ Triggers gradient allocation
                wp.from_torch(b),      # ⚠️ Triggers gradient allocation
                wp.from_torch(output),
            ],
        )

        return output

    @staticmethod
    def backward(ctx, grad_output):
        a, b = ctx.saved_tensors

        device = wp.device_from_torch(a.device)

        grad_a = torch.empty_like(a)
        grad_b = torch.empty_like(b)

        wp.launch(
            kernel=backward_kernel,
            dim=(N),
            device=device,
            inputs=[
                wp.from_torch(grad_output.contiguous()),
                wp.from_torch(a),
                wp.from_torch(b),
                wp.from_torch(grad_a),
                wp.from_torch(grad_b),
            ],
        )

        return grad_a, grad_b


a = torch.randn(N, device=device, dtype=torch.float32, requires_grad=True)
b = torch.randn(N, device=device, dtype=torch.float32, requires_grad=True)

torch.cuda.synchronize()

for i in range(TRIALS):
    with wp.ScopedTimer(f"TRIAL {i}", use_nvtx=True, synchronize=True):
        with wp.ScopedTimer("Create Tensors", use_nvtx=True, synchronize=True):
            a_torch = a.clone().detach().requires_grad_(True)
            b_torch = b.clone().detach().requires_grad_(True)
        with wp.ScopedTimer("Forward", use_nvtx=True, synchronize=True):
            output_warp = WarpFunction.apply(a_torch, b_torch)
        with wp.ScopedTimer("Loss", use_nvtx=True, synchronize=True):
            loss_warp = output_warp.sum()
        with wp.ScopedTimer("Backward", use_nvtx=True, synchronize=True):
            loss_warp.backward()
```

--------------------------------

### Configure Kernel Settings

Source: https://nvidia.github.io/warp/stable/user_guide/configuration.html

Examples of using kernel decorator arguments to control backward pass generation, module isolation, launch bounds, and module-specific compilation options.

```python
@wp.kernel(enable_backward=False)
def scale_2(
    x: wp.array[float],
    y: wp.array[float],
):
    y[0] = x[0] ** 2.0


@wp.kernel(module="unique")
def isolated_kernel(a: wp.array[float], b: wp.array[float]):
    # This kernel will be registered in a new unique module created
    # just for this kernel and its dependent functions and structs
    tid = wp.tid()
    b[tid] = a[tid] + 1.0


@wp.kernel(launch_bounds=(256, 1))
def bounded_kernel(a: wp.array[float]):
    # CUDA __launch_bounds__ will be set to (256, 1)
    tid = wp.tid()
    a[tid] = a[tid] * 2.0


@wp.kernel(module_options={"fast_math": True}, module="unique")
def fast_kernel(a: wp.array[float], b: wp.array[float]):
    # fast_math is applied to this kernel's unique module
    tid = wp.tid()
    b[tid] = a[tid] + 1.0
```

--------------------------------

### Capture and Replay CPU Kernels

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/runtime.rst.txt

Demonstrates using wp.ScopedCapture to record kernel launches on the CPU and replaying them with wp.capture_launch to minimize dispatch overhead.

```python
import warp as wp

@wp.kernel(module="unique")
def my_kernel(a: wp.array[float], b: wp.array[float]):
    i = wp.tid()
    b[i] = a[i] + 1.0

N = 1024
steps = 10

a = wp.zeros(N, dtype=float, device="cpu")
b = wp.zeros(N, dtype=float, device="cpu")

with wp.ScopedCapture(device="cpu") as capture:
    for _ in range(steps):
        wp.launch(my_kernel, dim=N, inputs=[a, b], device="cpu")

# Replay the entire batch with one native call into Warp
wp.capture_launch(capture.graph)
```

--------------------------------

### Configure CMake with custom libmathdx path

Source: https://nvidia.github.io/warp/stable/user_guide/installation.html

Uses an existing libmathdx installation instead of fetching it via Packman.

```bash
$ cmake -S . -B _build/cmake -G Ninja -DWARP_LIBMATHDX_PATH=/path/to/libmathdx
$ cmake --build _build/cmake --parallel
```

--------------------------------

### Install RMM for Warp

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/execution_and_performance/memory_management.rst.txt

Install the RMM library on Linux systems to enable high-performance pooled CUDA allocators.

```bash
pip install rmm-cu12
```

--------------------------------

### Create a linear sequence tile in a kernel

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.tile.html

Demonstrates constructing a tile from thread indices and launching the kernel.

```python
@wp.kernel
def compute():
    i = wp.tid()
    t = wp.tile(i*2)
    print(t)

wp.launch(compute, dim=16, inputs=[], block_dim=16)
```

--------------------------------

### Launch a 2D grid with tiled threads

Source: https://nvidia.github.io/warp/stable/api_reference/_generated/warp.launch_tiled.html

Demonstrates using launch_tiled to assign a specific number of threads per element in a 2D grid.

```python
warp.launch_tiled(kernel, [M, N], inputs=[...], block_dim=64)
```

```python
warp.launch(kernel, [M, N, 64], inputs=[...], block_dim=64)
```

--------------------------------

### Configure Nsight Systems for API-based Capture

Source: https://nvidia.github.io/warp/stable/user_guide/execution_and_performance/profiling.html

Run the application with the --capture-range flag set to cudaProfilerApi to honor the start and stop calls within the code.

```bash
nsys profile --capture-range=cudaProfilerApi python my_app.py
```

--------------------------------

### Install Warp in editable mode

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/installation.rst.txt

Install the built Warp package in editable mode to reflect local source changes.

```console
$ pip install -e .
```

--------------------------------

### Profile simulation with Nsight Systems

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/execution_and_performance/profiling.rst.txt

Run the simulation for a limited number of frames using Nsight Systems to identify performance bottlenecks.

```bash
nsys profile --stats=true python example_sph.py --stage-path None --num-frames 10
```

--------------------------------

### Install Warp via Conda

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/installation.rst.txt

Commands to install Warp using the conda-forge channel, including specific build variants.

```sh
$ conda install conda-forge::warp-lang
```

```sh
# CPU-only (no CUDA dependencies)
$ conda install conda-forge::warp-lang=*=*cpu*

# CUDA 12.9
$ conda install conda-forge::warp-lang=*=*cuda129*
```

--------------------------------

### warp.cuda_profiler_start(device=None)

Source: https://nvidia.github.io/warp/stable/api_reference/_generated/warp.cuda_profiler_start.html

Starts the CUDA profiler for the specified device. This function is used to define the start of a profiling region for tools like Nsight Systems or Nsight Compute.

```APIDOC
## warp.cuda_profiler_start(device=None)

### Description
Begin CUDA profiler data collection for the specified device. This is equivalent to the driver API `cuProfilerStart` and operates on the CUDA context of the provided device.

### Parameters
- **device** (Device | str | None) - Optional - The device whose CUDA context profiling is started for. Defaults to the current device.

### Usage Example
```python
import warp as wp

wp.cuda_profiler_start()
run_workload()
wp.cuda_profiler_stop()
```
```

--------------------------------

### warp.cuda_profiler_start()

Source: https://nvidia.github.io/warp/stable/_sources/api_reference/_generated/warp.cuda_profiler_start.rst.txt

Starts the CUDA profiler for the current process.

```APIDOC
## warp.cuda_profiler_start()

### Description
Starts the CUDA profiler. This function is used to begin profiling operations within the CUDA environment.

### Signature
`warp.cuda_profiler_start()`
```

--------------------------------

### Configure Kernel Modules and Launch Bounds

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/configuration.rst.txt

Examples of using kernel decorators to specify unique modules, set CUDA launch bounds, and enable fast math optimizations.

```python
@wp.kernel(module="unique")
def isolated_kernel(a: wp.array[float], b: wp.array[float]):
    # This kernel will be registered in a new unique module created
    # just for this kernel and its dependent functions and structs
    tid = wp.tid()
    b[tid] = a[tid] + 1.0
```

```python
@wp.kernel(launch_bounds=(256, 1))
def bounded_kernel(a: wp.array[float]):
    # CUDA __launch_bounds__ will be set to (256, 1)
    tid = wp.tid()
    a[tid] = a[tid] * 2.0
```

```python
@wp.kernel(module_options={"fast_math": True}, module="unique")
def fast_kernel(a: wp.array[float], b: wp.array[float]):
    # fast_math is applied to this kernel's unique module
    tid = wp.tid()
    b[tid] = a[tid] + 1.0
```

--------------------------------

### Ray-Mesh Intersection Output

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.mesh_query_ray.html

Expected output from the ray-mesh intersection example.

```text
t = 2.0 normal = [ 0.  0. -1.]
```

--------------------------------

### Measure execution time with CUDA events

Source: https://nvidia.github.io/warp/stable/user_guide/execution_and_performance/profiling.html

Use wp.Event with enable_timing=True to record start and end points for timing arbitrary code sections. Ensure events are recorded on the same device and synchronized before retrieving elapsed time.

```python
with wp.ScopedDevice("cuda:0") as device:

    # ensure the module is loaded
    wp.load_module(device=device)

    # create events with enabled timing
    e1 = wp.Event(enable_timing=True)
    e2 = wp.Event(enable_timing=True)

    n = 10000000

    # start timing...
    wp.record_event(e1)

    a = wp.zeros(n, dtype=float)
    wp.launch(inc, dim=n, inputs=[a])

    # ...end timing
    wp.record_event(e2)

    # get elapsed time between the two events
    elapsed = wp.get_event_elapsed_time(e1, e2)
    print(elapsed)
```

--------------------------------

### Tile AXPY output example

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.tile_axpy.html

The expected output format for the tile operation.

```text
[17 17 17 17] = tile(shape=(4), storage=register)
```

--------------------------------

### Tiled BVH Ray Casting Example

Source: https://nvidia.github.io/warp/stable/language_reference/_generated/warp.bvh_query_ray_tiled.html

Demonstrates initializing a tiled BVH ray query within a kernel and processing results cooperatively across a thread block.

```python
@wp.kernel
def tiled_cast(bvh_id: wp.uint64, lowers: wp.array[wp.vec3], uppers: wp.array[wp.vec3],
               origin: wp.vec3, dir: wp.vec3, centers: wp.array[wp.vec3]):
    query = wp.bvh_query_ray_tiled(bvh_id, origin, dir)
    while wp.tile_query_valid(query):
        result = wp.bvh_query_next_tiled(query)
        item = wp.untile(result)
        if item >= 0:
            centers[item] = 0.5 * (lowers[item] + uppers[item])

lowers = wp.array([[0, 0, 0], [2, 0, 0], [4, 0, 0]], dtype=wp.vec3)
uppers = wp.array([[1, 1, 1], [3, 1, 1], [5, 1, 1]], dtype=wp.vec3)
bvh = wp.Bvh(lowers=lowers, uppers=uppers)

centers = wp.zeros(3, dtype=wp.vec3)
wp.launch_tiled(tiled_cast, dim=[1], inputs=[bvh.id, lowers, uppers, wp.vec3(-1.0, 0.5, 0.5), wp.vec3(1.0, 0.0, 0.0)], outputs=[centers], block_dim=32)
print(centers.numpy().tolist())
```

```text
[[0.5, 0.5, 0.5], [2.5, 0.5, 0.5], [4.5, 0.5, 0.5]]
```

--------------------------------

### Launch Kernels with Device-Specific Arrays

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/execution_and_performance/memory_management.rst.txt

Demonstrates launching kernels on specific devices using arrays allocated on the corresponding device.

```python
cpu_array = wp.zeros(1024, dtype=float, device="cpu")
gpu_array = wp.zeros(1024, dtype=float, device="cuda:0")

wp.launch(kernel, dim=cpu_array.size, inputs=[cpu_array], device="cpu")
wp.launch(kernel, dim=gpu_array.size, inputs=[gpu_array], device="cuda:0")
```

--------------------------------

### Warp Function Definition

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/code_generation.rst.txt

Example of a Warp function that will be translated into CUDA code.

```python
@wp.func
def my_func(a: float, b: float):
    c = wp.sin(b) * a
    return c
```

--------------------------------

### Resolve C++ runtime library errors in Conda

Source: https://nvidia.github.io/warp/stable/user_guide/installation.html

Example of an error message encountered when C++ runtime versions mismatch between build and execution environments.

```text
OSError: <...>/libstdc++.so.6: version `GLIBCXX_3.4.30' not found (required by <...>/warp/warp/bin/warp.so)
```

--------------------------------

### Perform lane-based store with wp.launch

Source: https://nvidia.github.io/warp/stable/_sources/user_guide/programming_model/tiles.rst.txt

Shows how wp.launch executes threads as separate tasks, leading to overlapping full-tile stores on CPU.

```python
TILE_SIZE = wp.constant(8)

@wp.kernel
def lane_store(x: wp.array[float], out: wp.array[float]):
    i = wp.tid()
    t = wp.tile_zeros(shape=(TILE_SIZE,), dtype=float)
    t[i] = x[i] * 2.0
    wp.tile_store(out, t)

wp.launch(lane_store, dim=8, inputs=[x], outputs=[out], block_dim=TILE_SIZE)
```

--------------------------------

### Registering Warp Kernels with Various Configurations

Source: https://nvidia.github.io/warp/stable/api_reference/_generated/warp.kernel_decorator.html

Examples demonstrating the use of the @wp.kernel decorator with different parameters such as disabling backward passes, specifying modules, setting launch bounds, and applying module-level compilation options.

```python
@wp.kernel
def my_kernel(a: wp.array[float], b: wp.array[float]):
    tid = wp.tid()
    b[tid] = a[tid] + 1.0


@wp.kernel(enable_backward=False)
def my_kernel_no_backward(a: wp.array2d[float], x: float):
    # the backward pass will not be generated
    i, j = wp.tid()
    a[i, j] = x


@wp.kernel(module="unique")
def my_kernel_unique_module(a: wp.array[float], b: wp.array[float]):
    # the kernel will be registered in new unique module created just for this
    # kernel and its dependent functions and structs
    tid = wp.tid()
    b[tid] = a[tid] + 1.0


@wp.kernel(launch_bounds=(256, 1))
def my_kernel_with_launch_bounds(a: wp.array[float]):
    # CUDA __launch_bounds__ will be set to (256, 1)
    tid = wp.tid()
    a[tid] = a[tid] * 2.0


@wp.kernel(module_options={"fast_math": True}, module="unique")
def my_kernel_fast(a: wp.array[float], b: wp.array[float]):
    # fast_math is a module-level option, so module="unique" is required
    tid = wp.tid()
    b[tid] = a[tid] + 1.0
```
Tokens:
Raw