Skip to content

fix: kernel return validation#1411

Open
pumpkin-bit wants to merge 9 commits into
NVIDIA:mainfrom
pumpkin-bit:fix-kernel-return-validation
Open

fix: kernel return validation#1411
pumpkin-bit wants to merge 9 commits into
NVIDIA:mainfrom
pumpkin-bit:fix-kernel-return-validation

Conversation

@pumpkin-bit
Copy link
Copy Markdown

@pumpkin-bit pumpkin-bit commented Apr 30, 2026

Description

Kernels are architecturally prohibited from returning values. Previously, the code generation system allowed return type annotations (e.g., '-> float') for kernels, but the function body did not explicitly return a value. This PR adds a check for rejecting return type hints for kernels to prevent misleading function signatures and avoid user confusion.

This also addresses an existing TODO and includes a corresponding test.

Checklist

  • I am familiar with the Contributing Guidelines.
  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

Test plan

Verified by running the codegen test suite: test_codegen.py

The updated test_error_kernel_return_value explicitly checks and passes the WarpCodegenTypeError assertion when a return type hint is provided without a return value.

Bug fix

import warp as wp

wp.init()

# Prior to this PR, this compiled silently without raising an error
@wp.kernel
def f(x: float) -> float:
    return

# Should raise a WarpCodegenTypeError
wp.launch(f, dim=1, inputs=[3.0], device="cpu")


<!-- This is an auto-generated comment: release notes by coderabbit.ai -->
## Summary by CodeRabbit

* **Bug Fixes**
  * Strengthened validation to reject kernels that declare or imply non-None return values, preventing erroneous launches and related runtime errors.

* **Tests**
  * Re-enabled and added tests that assert kernels with bare returns are rejected while kernels annotated to return None remain accepted.
<!-- end of auto-generated comment: release notes by coderabbit.ai -->

@copy-pr-bot
Copy link
Copy Markdown

copy-pr-bot Bot commented Apr 30, 2026

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
Copy link
Copy Markdown

coderabbitai Bot commented Apr 30, 2026

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review
📝 Walkthrough

Walkthrough

ModuleBuilder now raises WarpCodegenTypeError when a kernel indicates a return value either via kernel.adj.return_var or via a "return" entry in the kernel argument map. Tests were updated to re-enable a case asserting this error and to confirm a -> None return remains accepted.

Changes

Kernel Codegen Validation

Layer / File(s) Summary
Core Validation
warp/_src/context.py
Composite condition added: raise WarpCodegenTypeError when kernel.adj.return_var is not None OR "return" is present in kernel.adj.arg_types.
Tests
warp/tests/test_codegen.py
Re-enabled f4(x: float) -> float test asserting launch raises WarpCodegenTypeError for bare return; added f5(x: float) -> None test confirming return with None is accepted.
Changelog
CHANGELOG.md
Unreleased changelog updated with many Added items and Fixed bullets including the WarpCodegenTypeError behavior; release anchor adjusted.

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~8 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: kernel return validation' clearly describes the main change: adding validation to reject non-None return type hints on kernel functions, which is the primary modification across the codebase.
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

Tip

💬 Introducing Slack Agent: The best way for teams to turn conversations into code.

Slack Agent is built on CodeRabbit's deep understanding of your code, so your team can collaborate across the entire SDLC without losing context.

  • Generate code and open pull requests
  • Plan features and break down work
  • Investigate incidents and troubleshoot customer tickets together
  • Automate recurring tasks and respond to alerts with triggers
  • Summarize progress and report instantly

Built for teams:

  • Shared memory across your entire org—no repeating context
  • Per-thread sandboxes to safely plan and execute work
  • Governance built-in—scoped access, auditability, and budget controls

One agent for your entire SDLC. Right inside Slack.

👉 Get started


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

@greptile-apps
Copy link
Copy Markdown

greptile-apps Bot commented Apr 30, 2026

Greptile Summary

This PR tightens kernel validation so that a non-None return type annotation on a @wp.kernel function is now a compile-time error, not just silently ignored. The core change is a one-line extension in build_kernel that ORs the existing return_var guard with a check on whether "return" is present in arg_types (which codegen.py already filters to exclude None-valued annotations).

  • warp/_src/context.py: Adds or "return" in kernel.adj.arg_types to the existing return_var check so that annotating a kernel -> float without a real return value is caught during build.
  • warp/tests/test_codegen.py: Replaces the long-standing TODO with two active test cases — one confirming the new error for -> float with a bare return, and one confirming -> None remains accepted.
  • CHANGELOG.md: Records the fix under Fixed.

Confidence Score: 5/5

Safe to merge — the change is a targeted one-line guard that closes a loophole in kernel signature validation without touching any code-generation paths.

The single changed line in context.py piggybacks on an already-existing arg_types dictionary that codegen.py populates and filters (-> None is excluded from arg_types at line 1058, so the -> None acceptance case is inherently correct). Both new test cases are well-scoped and directly cover the added branch.

No files require special attention.

Important Files Changed

Filename Overview
warp/_src/context.py Single-line logic extension: adds "return" in kernel.adj.arg_types to the existing return_var guard, correctly catching non-None return type annotations on kernels before code generation.
warp/tests/test_codegen.py Replaces a commented-out TODO block with two concrete test cases: f4 (bare return with -> float annotation, expects error) and f5 (-> None with bare return, expects no error).
CHANGELOG.md Adds a changelog entry under Fixed; entry wording "Fix Raise a..." is slightly awkward but content is correct.

Flowchart

%%{init: {'theme': 'neutral'}}%%
flowchart TD
    A[build_kernel called] --> B[kernel.adj.build]
    B --> C{return_var is not None?}
    C -- Yes --> E[raise WarpCodegenTypeError]
    C -- No --> D{return annotation present and non-None?}
    D -- Yes --> E
    D -- No --> F[kernel built successfully]
Loading

Reviews (8): Last reviewed commit: "Adding a fix to Unreleased" | Re-trigger Greptile

Signed-off-by: falker <arsenija111mot@gmail.com>
Signed-off-by: falker <arsenija111mot@gmail.com>
@pumpkin-bit pumpkin-bit force-pushed the fix-kernel-return-validation branch from 4be775d to 0d95b60 Compare April 30, 2026 09:47
@shi-eric shi-eric requested review from c0d1f1ed and nvlukasz May 1, 2026 15:17
Copy link
Copy Markdown
Contributor

shi-eric commented May 5, 2026

Thanks for the contribution. This addresses the existing TODO in test_error_kernel_return_value: kernels already raise when they actually return a value, but a kernel with a non-None return annotation and only a bare return was still accepted, which leaves a misleading signature that Warp cannot honor.

Could you please make a few cleanup changes before this can move forward?

  • Add a CHANGELOG.md entry under the Unreleased section, since this changes user-visible codegen behavior from silently accepted to reported at compile time.
  • Add a test confirming that -> None remains accepted for kernels.
  • Clean up the commit history by squashing the duplicate/fixup commits and merge commit into one signed-off commit. We need this before merging because we preserve contributor authorship when merging through our internal GitLab workflow, so we cannot squash these commits on your behalf.

Also, for future non-trivial codegen behavior changes, please open an issue or discussion first so we can agree on the desired behavior before implementation.

@c0d1f1ed @nvlukasz, could you weigh in on whether a non-None kernel return annotation should be a hard error, as implemented here, or a warning?

@pumpkin-bit
Copy link
Copy Markdown
Author

Alright, I created a changelog entry and wrote a test regarding the commits, I'm thinking of doing a squash after all the necessary maintainers have shared their feedback, so that all changes are agreed upon if needed also, the single squash commit will be signed as I understand it, intermediate commits don't need signatures since the squash will bring them into one? And next time for similar fixes I will create an issue or a discussion. Thanks.

Copy link
Copy Markdown

@coderabbitai coderabbitai Bot left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@CHANGELOG.md`:
- Around line 201-210: Move the listed bullet items (the duplicate 1.13.0 fixes
starting with "Fix `ValueError: Cell is empty`..." through the
"WarpCodegenTypeError" bullet) out of the released 1.13.0 section and append
them to the end of the Unreleased section, removing the duplicate bullets from
1.13.0; ensure each moved entry uses imperative present tense and retains its GH
reference format (e.g., "([GH-913](...))"), and delete the redundant 1.13.0
copies so Unreleased contains the authoritative entries.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yml

Review profile: CHILL

Plan: Enterprise

Run ID: 0946e968-9275-4e2e-b1fe-a8d75c4b0f52

📥 Commits

Reviewing files that changed from the base of the PR and between 983cdb6 and 7ebfdf1.

📒 Files selected for processing (3)
  • CHANGELOG.md
  • warp/_src/context.py
  • warp/tests/test_codegen.py

Comment thread CHANGELOG.md Outdated
Comment thread CHANGELOG.md Outdated
Comment on lines +2507 to +2510

- Initial publish for alpha testing

[Unreleased]: https://github.com/NVIDIA/warp/compare/v1.13.0...HEAD
[1.13.0]: https://github.com/NVIDIA/warp/releases/tag/v1.13.0
[1.12.1]: https://github.com/NVIDIA/warp/releases/tag/v1.12.1
[Unreleased]: https://github.com/NVIDIA/warp/compare/v1.12.0...HEAD
Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Branch appears to pre-date the 1.13.0 release

The diff removes the entire [1.13.0] - 2026-05-04 section and changes the [Unreleased] comparison link from v1.13.0...HEAD to v1.12.0...HEAD. If this PR is targeting main and 1.13.0 has already been cut, merging as-is will silently drop all 1.13.0 release notes. Please rebase onto the current main tip (or at minimum cherry-pick this commit onto the correct base) so the changelog history is preserved.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants