Fix NVRTC include path option storage

[fallintoplace]

Description

wp_cuda_compile_program() built the main NVRTC include option in a fixed-size stack buffer. The guard checked only include_dir, but the buffer also had to hold the --include-path= prefix and the null terminator. That allowed paths that passed the check to overflow the stack buffer while constructing the option.

This changes the main include option to use the same dynamic string storage pattern used for the other generated NVRTC options, while reserving capacity up front so the c_str() pointers passed to NVRTC stay stable until compilation.

Changes

• Store the primary --include-path= option in std::vector<std::string> instead of a fixed-size char buffer.
• Reserve storage for all generated option strings before taking c_str() pointers.
• Remove the incorrect length check and the strcpy()/strcat() construction.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.
• CHANGELOG.md is updated for any user-facing changes under the Unreleased section.

Validation summary

I verified the focused source diff and formatting locally:

• git diff --check
• uvx pre-commit run --files warp/native/warp.cu
• uv run build_lib.py --quick

The quick build succeeded on macOS arm64 in CPU-only mode. This machine does not have a CUDA toolchain available, so I could not run an NVRTC/CUDA compile-path regression test locally.

Bug fix

The issue can be triggered when Warp is installed or checked out under a very long path and a CUDA kernel compile uses that path as warp_home/native for the include directory. Without this change, the fixed stack buffer can overflow while prepending --include-path=.

import warp as wp


@wp.kernel
def noop():
    pass


wp.init()
wp.launch(noop, dim=1, inputs=[], device="cuda:0")

New feature / enhancement

Not applicable.

Summary by CodeRabbit

Release Notes

• Refactor
  • Improved CUDA compilation memory management by transitioning include path handling from fixed-size buffers to dynamic allocation. This enhancement increases system reliability, removes previous capacity constraints, and provides better flexibility for projects requiring extensive include path configurations.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yml

Review profile: CHILL

Plan: Enterprise

Run ID: 975bb3f7-a648-451a-b74f-cbef7907dddd

📥 Commits

Reviewing files that changed from the base of the PR and between 52f756b and e1d59bc.

📒 Files selected for processing (1)

• warp/native/warp.cu

---

📝 Walkthrough

Walkthrough

In wp_cuda_compile_program (warp/native/warp.cu), the NVRTC include-path option is no longer built into a fixed-size stack buffer. Instead, stored_options is created earlier, the --include-path= string is pushed into it, and opts references stored_options.back().c_str(). The buffer-size guard and the later redundant stored_options declaration are removed.

Changes

NVRTC include-path option storage refactor

Layer / File(s)	Summary
Include-path option: stack buffer → dynamic vector warp/native/warp.cu	Removes the include_dir length guard and stack include_opt buffer. Initializes stored_options earlier with reserved capacity, pushes the --include-path= string into it, updates opts to use stored_options.back().c_str(), and deletes the now-redundant later stored_options declaration.

Estimated code review effort

🎯 1 (Trivial) | ⏱️ ~3 minutes

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title 'Fix NVRTC include path option storage' directly and accurately summarizes the main change: replacing fixed-size buffer storage with dynamic string vector for NVRTC include path options.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

[greptile-apps]

Greptile Summary

This PR fixes a stack buffer overflow in wp_cuda_compile_program() by replacing a fixed-size char include_opt[max_path] buffer with a std::vector<std::string> for dynamic storage of the primary --include-path= NVRTC option, matching the pattern already used for other generated options.

• Removes the incorrect length guard (which checked only include_dir length, ignoring the --include-path= prefix and null terminator) along with the unsafe strcpy/strcat construction.
• Moves stored_options declaration before option assembly and calls reserve() upfront with the correct total capacity (num_cuda_include_dirs + 3) so that subsequent push_back calls never trigger reallocation and all c_str() pointers pushed into opts remain stable through nvrtcCompileProgram.

Confidence Score: 5/5

The change is a well-scoped buffer-overflow fix with correct reserve math and proper pointer-stability management; safe to merge.

The reserve capacity of max(num_cuda_include_dirs, 0) + 3 precisely accounts for all conditional stored_options push_backs (1 for the main include path, 1 for an optional pch-dir, 1 for an optional compile-time-trace, plus the loop), so no reallocation can occur after c_str() pointers are stored in opts. The removal of the incorrect max_path guard and the unsafe strcpy/strcat is clean. No unrelated code is touched.

No files require special attention.

Important Files Changed

Filename	Overview
warp/native/warp.cu	Replaces fixed-size stack buffer and unsafe strcpy/strcat with std::vectorstd::string for the primary --include-path= NVRTC option; reserve() math is correct and pointer stability is properly maintained.

_{Reviews (1): Last reviewed commit: "Fix NVRTC include path option storage" | Re-trigger Greptile}

[shi-eric]

Thanks for the PR. Before we consider this further, can you clarify what real circumstance led you to this change?

Was this based on an actual Warp failure you encountered, or was it found by static analysis / code inspection? If it was an observed failure, please share the platform, install path layout, CUDA environment, reproduction steps, and validation showing that this change fixes it.

For non-trivial changes, especially in CUDA runtime internals, we want contributions to be driven by demonstrated user issues, production failures, or clearly validated feature work. We generally do not want to establish a pattern of accepting speculative hardening PRs for hypothetical edge cases, even when the patch itself is small.

If this is only a theoretical issue involving an unusually long Warp install path, we are unlikely to prioritize it. If there is a concrete real-world case behind it, please add those details so we can evaluate the impact.