Configuration#

Warp has settings at the global, module, and kernel level that can be used to fine-tune the compilation and verbosity of Warp programs. In cases in which a setting can be changed at multiple levels (e.g.: enable_backward), the setting at the more-specific scope takes precedence.

Global Settings#

Settings can be modified by direct assignment before or after calling wp.init(), though some settings only take effect if set prior to initialization.

For example, the location of the user kernel cache can be changed with:

import os

import warp as wp

example_dir = os.path.dirname(os.path.realpath(__file__))

# set default cache directory before wp.init()
wp.config.kernel_cache_dir = os.path.join(example_dir, "tmp", "warpcache1")

wp.init()

See warp.config for a complete list of global settings.

Module Settings#

Module-level settings to control runtime compilation and code generation may be changed by passing a dictionary of option pairs to wp.set_module_options().

For example, compilation of backward passes for the kernel in an entire module can be disabled with:

wp.set_module_options({"enable_backward": False})

The options for a module can also be queried using wp.get_module_options().

Field	Type	Default Value	Description
`mode`	String	`None`	A module-level override of the `warp.config.mode` setting. `None` defers to the global setting at compile time.
`optimization_level`	Integer	`None`	A module-level override of the `warp.config.optimization_level` setting. `None` defers to the global setting at compile time.
`max_unroll`	Integer	Global setting	A module-level override of the `warp.config.max_unroll` setting.
`enable_backward`	Boolean	Global setting	A module-level override of the `warp.config.enable_backward` setting.
`fast_math`	Boolean	`False`	If `True`, CUDA kernels will be compiled with the `--use_fast_math` compiler option, which enables some fast math operations that are faster but less accurate.
`fuse_fp`	Boolean	`True`	If `True`, allow compilers to emit fused floating point operations such as fused-multiply-add. This may improve numerical accuracy and is generally recommended. Setting to `False` can help ensuring that functionally equivalent kernels will produce identical results unaffected by the presence or absence of fused operations.
`lineinfo`	Boolean	Global setting	A module-level override of the `warp.config.lineinfo` setting.
`compile_time_trace`	Boolean	Global setting	A module-level override of the `warp.config.compile_time_trace` setting.
`cuda_output`	String	`None`	A module-level override of the `warp.config.cuda_output` setting.
`block_dim`	Integer	256	The number of CUDA threads per block that kernels in the module will be compiled for.
`strip_hash`	Boolean	`False`	If `True`, avoids using a content-based hash to identify the module and its functions.
`extra_build_options`	Object	`None`	Experimental extra build inputs for CPU and CUDA modules. Set to an `warp.ModuleBuildOptions` instance.
`enable_mathdx_gemm`	Boolean	`None`	A module-level override of the `warp.config.enable_mathdx_gemm` setting. `None` defers to the global setting at compile time.
`enable_mathdx_solver`	Boolean	`None`	A module-level override of the `warp.config.enable_mathdx_solver` setting. `None` defers to the global setting at compile time.

Kernel Settings#

Kernel-level settings can be passed as arguments to the @wp.kernel decorator.

Field	Type	Default Value	Description
`name`	String	`None`	Sets the kernel key used for registration and native code generation. If `None`, Warp derives the key from the Python callable passed to `wp.kernel`. A custom name must be a valid C++ identifier. When `strip_hash=True`, Warp uses the key without a hash suffix as the base of the generated native entry-point names.
`enable_backward`	Boolean	`None`	If `False`, the backward pass will not be generated for this kernel. If `None`, inherits from the module/global setting.
`module`	Module \| `"unique"` \| str	`None`	Controls which module the kernel belongs to. If `"unique"`, the kernel is assigned to a new module named after the kernel (with a hash suffix). If a plain string is provided, the kernel is registered in the module with that name. If `None`, the module is inferred from the function’s module.
`launch_bounds`	int \| tuple	`None`	CUDA `__launch_bounds__` attribute for the kernel. Can be an int (`maxThreadsPerBlock`) or a tuple of 1–2 ints `(maxThreadsPerBlock, minBlocksPerMultiprocessor)`. Only applies to CUDA kernels. The `block_dim` parameter in `warp.launch()` must not exceed the `maxThreadsPerBlock` value specified here.
`cuda_max_registers`	int	`None`	CUDA `__maxnreg__` attribute specifying the maximum number of registers allocated per thread. Must be positive and cannot be combined with `launch_bounds`. Ignored on CPU and when Warp was built with CUDA Toolkit earlier than 12.4 or when `wp.config.llvm_cuda` is `True`.
`enable_cuda_smem_spilling`	Boolean	`None`	If `True`, allow the CUDA Toolkit used to build Warp, when version 13.0 or later, to spill registers into shared memory. Silently ignored for entry points that use dynamic shared memory, on CPU, with older CUDA Toolkits, in unsupported device-debug compilation, and when `wp.config.llvm_cuda` is `True`.
`module_options`	dict	`None`	A dict of module-level compilation options to apply to the kernel’s module. Requires `module="unique"`; raises `ValueError` otherwise. Keys are validated against the module’s known options (see Module Settings above). For shared modules, use `wp.set_module_options()` instead.
`entry_point_abi`	String	`None`	Selects the experimental entry-point ABI. `"warp"` uses the regular launch-compatible CPU and CUDA ABI. `"external_constant_params"` is CUDA-only and binds one struct argument from constant memory symbol `params`; it requires `enable_backward=False` and cannot be launched with `warp.launch()`.

@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(cuda_max_registers=64)
def register_limited_kernel(a: wp.array[float]):
    # CUDA __maxnreg__(64) will be set when supported
    tid = wp.tid()
    a[tid] = a[tid] * 2.0


@wp.kernel(enable_cuda_smem_spilling=True, launch_bounds=256)
def smem_spilling_kernel(a: wp.array[float]):
    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

CUDA shared-memory register spilling uses otherwise available shared memory to reduce local-memory spill traffic. Warp evaluates forward and backward entry points independently and enables the optimization only when the corresponding entry point requires no dynamic shared memory, including scratch space required by tile operations and their callees. Explicit launch_bounds are recommended to keep the compiler’s shared-memory estimate aligned with the intended block size. See NVIDIA’s shared-memory register spilling guidance for performance considerations and CUDA limitations.

CUDA Thread Block Clusters#

CUDA Thread Block Clusters group adjacent CTAs into clusters that the hardware co-schedules on a single GPU Processing Cluster (GPC), unlocking features such as distributed shared memory. Clusters are supported on devices with compute capability 9.0 (Hopper) and above.

The cluster size is declared per-kernel via the cluster_dim decorator argument — a positive int up to 16 (the default 1 means no clustering):

@wp.kernel(cluster_dim=2)
def my_kernel(a: wp.array[float]):
    i = wp.tid()
    a[i] = a[i] * 2.0

Warp launches CUDA kernels with a 1D hardware grid, so cluster_dim=N emits CUDA __cluster_dims__(N, 1, 1); multidimensional cluster shapes are not supported. Values 2..8 are portable and work on every cluster-capable device; values 9..16 are non-portable and depend on the GPU’s GPC layout (some integrated parts such as NVIDIA Thor and DGX Spark GB10 cap at 8). Query warp.get_cuda_max_cluster_dim() before selecting a value above 8 if you target a range of devices (it returns 1 on devices that cannot form clusters).

On a genuine sub-cluster device (compute capability < 9.0) and on CPU, cluster_dim is silently ignored and the kernel runs unclustered, so portable code can set it freely. Requesting a cluster while compiling for a target below sm90 on cluster-capable hardware (for example with warp.config.ptx_target_arch set below 90) is an error, because the cluster attribute would be dropped from the generated kernel.

Declaring cluster_dim only forms the cluster; using it — distributed shared memory, cluster barriers, and cluster rank queries — requires native CUDA code. See Thread Block Clusters and Distributed Shared Memory in the C++/CUDA workflows guide for a worked distributed-shared-memory example.

Function Settings#

Function-level settings can be passed as arguments to the @wp.func decorator.

Setting	Type	Default Value	Description
`name`	String	`None`	Sets the function key used for registration and native code generation. If `None`, Warp derives the key from the Python callable.
`module`	Module \| `"unique"` \| str	`None`	Controls which module the function belongs to, following the same rules as the equivalent kernel setting.
`inline`	Boolean \| `None`	`None`	Whether the function is inlined into its call sites. `None` leaves the choice to the backend compiler. `True` requires inlining (CUDA `__forceinline__`), overriding the compiler’s own heuristic. `False` keeps the function out of line (CUDA `__noinline__`).

@wp.func(inline=False)
def expensive_helper(x: wp.vec3) -> wp.vec3:
    # kept out of line in every kernel that calls it
    return wp.normalize(x) * wp.length(x)


@wp.func(inline=True)
def cheap_helper(x: float) -> float:
    # inlined even where the backend compiler would not choose to
    return x * 2.0

By default, the backend compiler decides whether to inline a function at each call site. Set inline=False to prevent inlining or inline=True to require it.

Inlining duplicates the function body at each call site. This can increase register pressure and instruction-cache use. Keeping a function out of line avoids that duplication but adds function-call overhead. The best choice depends on the workload, so measure both options.

The hint covers the generated adjoint as well as the forward function, and is lowered per backend, so the same function remains valid for CPU and CUDA.