Repository navigation
Add inlining control to @wp.func #1849
Description
Description
When calling a function from a kernel, in some cases there are benefits from inlining but in other cases it can be more optimal to not inline (due to e.g. register usage/occupancy and instruction cache behaviour in large kernels). Often this needs manual control to get the best setting. The addition of explicit inlining control to Warp (via "inline=true" and "inline=false" func options) will allow users to specify the best option for their specific case. There is precedent for this via @wp.kernel, which already exposes e.g. launch_bounds.
Context
Cases have been observed when using warp for rendering and simulation associated with RL training, where modifying the codegen to explicitly disable inlining has given significant performance benefits. (Given the symmetry, it also makes sense to enable forceinline as part of the change.)
Activity
[-]Add inlining control (noinline / forceinline) to @wp.func[/-][+]Add inlining control to @wp.func[/+]shi-eric commentedon Aug 27, 2026
This is addressed by f647a18.
The change adds the requested manual inlining control to
@wp.func. The hint is preserved across generated forward, adjoint, custom-gradient, and replay functions and is included in the module cache key. The user guide documents the performance tradeoffs and recommends measuring both options for the target workload.Public API behavior
@wp.func(inline=False)keeps the function out of line on CPU and CUDA.@wp.func(inline=True)requires inlining on CPU and CUDA.inline, or passingNone, leaves the choice to the backend compiler.boolorNoneraiseTypeError.Test coverage
warp/tests/test_func_inline.pycover option validation, CPU and CUDA code generation, runtime execution, generated adjoints and custom hooks, native-snippet compatibility, backend macro definitions, and PTX-level confirmation that both forced and disabled inlining are honored.TestModuleHashing.test_inline_hint_hashedinwarp/tests/test_module_hashing.pyverifies that each inline mode produces a distinct module hash.Behavioral verification
Separate verification compiled and ran functions using
inline=None,inline=False, andinline=Trueon CPU and CUDA. The generated attributes and runtime results matched the requested behavior.Extend wp.func to support an inline option [NVIDIAGH-1849]
Extend wp.func to support an inline option [NVIDIAGH-1849]