.. _Configuration: 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: Global Settings --------------- Settings can be modified by direct assignment before or after calling :func:`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: .. code-block:: python 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 :doc:`../api_reference/warp_config` for a complete list of global settings. See :doc:`execution_and_performance/reducing_compilation_and_startup_time` for guidance on settings that affect compilation and startup. .. _module-settings: Module Settings --------------- Module-level settings to control runtime compilation and code generation may be changed by passing a dictionary of option pairs to :func:`wp.set_module_options() `. For example, compilation of backward passes for the kernel in an entire module can be disabled with: .. code:: python wp.set_module_options({"enable_backward": False}) The options for a module can also be queried using :func:`wp.get_module_options() `. +--------------------------------------+---------+-------------+--------------------------------------------------------------------------+ | Field | Type |Default Value| Description | +======================================+=========+=============+==========================================================================+ |``mode`` | String | ``None`` | A module-level override of the :attr:`warp.config.mode` setting. | | | | | ``None`` defers to the global setting at compile time. | +--------------------------------------+---------+-------------+--------------------------------------------------------------------------+ |``optimization_level`` | Integer | ``None`` | A module-level override of the :attr:`warp.config.optimization_level` | | | | | setting. ``None`` defers to the global setting at compile time. | +--------------------------------------+---------+-------------+--------------------------------------------------------------------------+ |``max_unroll`` | Integer | Global | A module-level override of the :attr:`warp.config.max_unroll` setting. | | | | setting | | +--------------------------------------+---------+-------------+--------------------------------------------------------------------------+ |``enable_backward`` | Boolean | Global | A module-level override of the :attr:`warp.config.enable_backward` | | | | setting | 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 | A module-level override of the :attr:`warp.config.lineinfo` setting. | | | | setting | | +--------------------------------------+---------+-------------+--------------------------------------------------------------------------+ |``compile_time_trace`` | Boolean | Global | A module-level override of the :attr:`warp.config.compile_time_trace` | | | | setting | setting. | +--------------------------------------+---------+-------------+--------------------------------------------------------------------------+ |``cuda_output`` | String | ``None`` | A module-level override of the :attr:`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 | | | | | :class:`warp.ModuleBuildOptions` instance. | +--------------------------------------+---------+-------------+--------------------------------------------------------------------------+ |``enable_mathdx_gemm`` | Boolean | ``None`` | A module-level override of the :attr:`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 :attr:`warp.config.enable_mathdx_solver` | | | | | setting. ``None`` defers to the global setting at compile time. | +--------------------------------------+---------+-------------+--------------------------------------------------------------------------+ .. _kernel-settings: Kernel Settings --------------- Kernel-level settings can be passed as arguments to the :func:`@wp.kernel ` decorator. .. list-table:: :header-rows: 1 :widths: 20 20 10 50 * - 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 :func:`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 :func:`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 :func:`warp.launch`. .. code-block:: 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(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. .. _kernel-cluster-dim: 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): .. code-block:: python @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 :func:`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 :ref:`thread-block-clusters` in the C++/CUDA workflows guide for a worked distributed-shared-memory example.