Add name argument to @wp.kernel (GH-1561)#1570
Conversation
Signed-off-by: Thomas B. Brunner <thomas.brunner@helsing.ai>
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (3)
✅ Files skipped from review due to trivial changes (1)
📝 WalkthroughWalkthroughAdds a Changes
Estimated code review effort🎯 2 (Simple) | ⏱️ ~12 minutes Possibly related issues
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
Greptile SummaryThis PR adds an optional
Confidence Score: 5/5Safe to merge — the change is additive, validation is in place, and the overload-lookup fallback is well-tested including the closure pattern. The new No files require special attention. Important Files Changed
Flowchart%%{init: {'theme': 'neutral'}}%%
flowchart TD
A["@wp.kernel decorator"] --> B{name provided?}
B -- No --> C["key = qualified name of func"]
B -- Yes --> D{valid C identifier?}
D -- No --> E["raise ValueError"]
D -- Yes --> F["key = name"]
C --> G["Kernel.__init__"]
F --> G
G --> H["store self.qualname = qualified name"]
H --> I["store self.key = key"]
I --> J["module.kernels[key] = self"]
K["@wp.overload as decorator"] --> L["module._find_kernel(fn)"]
L --> M["compute fn qualname"]
M --> N{kernel keyed by qualname exists?}
N -- Yes --> O["return kernel"]
N -- No --> P["scan all kernels by .qualname attr"]
P --> Q["return last matching kernel"]
%%{init: {'theme': 'base', 'themeVariables': {"darkMode": true, "background": "#0d1117", "primaryColor": "#21262d", "primaryTextColor": "#e6edf3", "primaryBorderColor": "#8b949e", "lineColor": "#8b949e", "textColor": "#e6edf3", "edgeLabelBackground": "#161b22", "actorBkg": "#21262d", "actorBorder": "#8b949e", "actorTextColor": "#e6edf3", "actorLineColor": "#8b949e", "signalColor": "#8b949e", "signalTextColor": "#e6edf3", "noteBkgColor": "#373320", "noteBorderColor": "#d4a72c", "noteTextColor": "#f0e6c0", "labelBoxBkgColor": "#21262d", "labelBoxBorderColor": "#8b949e", "labelTextColor": "#e6edf3", "loopTextColor": "#e6edf3", "activationBkgColor": "#30363d", "activationBorderColor": "#8b949e"}}}%%
flowchart TD
A["@wp.kernel decorator"] --> B{name provided?}
B -- No --> C["key = qualified name of func"]
B -- Yes --> D{valid C identifier?}
D -- No --> E["raise ValueError"]
D -- Yes --> F["key = name"]
C --> G["Kernel.__init__"]
F --> G
G --> H["store self.qualname = qualified name"]
H --> I["store self.key = key"]
I --> J["module.kernels[key] = self"]
K["@wp.overload as decorator"] --> L["module._find_kernel(fn)"]
L --> M["compute fn qualname"]
M --> N{kernel keyed by qualname exists?}
N -- Yes --> O["return kernel"]
N -- No --> P["scan all kernels by .qualname attr"]
P --> Q["return last matching kernel"]
Reviews (2): Last reviewed commit: "Handle renamed kernels and wp.overload" | Re-trigger Greptile |
![greptile-apps[bot]](https://avatars.githubusercontent.com/in/867647?s=60&v=4){kind=small}
| def is_valid_c_identifier(name: str) -> bool: | ||
| return _C_IDENTIFIER_RE.match(name) is not None |
There was a problem hiding this comment.
is_valid_c_identifier only checks the structural regex pattern — it accepts any syntactically valid identifier, including C reserved words (int, float, void, if, for, while, return, struct, etc.) and CUDA built-ins (threadIdx, blockDim, etc.). A user who calls @wp.kernel(name="int") will pass this check but still get a cryptic compiler error, which is the exact failure mode this check was added to prevent.
Note: If this suggestion doesn't match your team's coding style, reply to this and let me know. I'll remember it for next time!
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
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 `@warp/_src/codegen.py`:
- Around line 5135-5139: The is_valid_c_identifier function currently only
validates the syntactic shape of identifiers using regex pattern matching, but
fails to reject reserved C/C++ keywords like "for", "class", "int", and "return"
which are syntactically valid identifiers but cannot be used as function
identifiers in generated C/CUDA code. Add a check within is_valid_c_identifier
to verify that the name is not a C/C++ reserved keyword by comparing it against
a comprehensive set of reserved keywords, and return False if the name matches
any keyword before returning the regex match result.
🪄 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: 63571386-5138-4cd6-adc4-1cd1135347c2
📒 Files selected for processing (5)
CHANGELOG.mdwarp/_src/codegen.pywarp/_src/context.pywarp/tests/test_codegen.pywarp/tests/test_context.py
| _C_IDENTIFIER_RE = re.compile(r"^[A-Za-z_][A-Za-z0-9_]*$") | ||
|
|
||
|
|
||
| def is_valid_c_identifier(name: str) -> bool: | ||
| return _C_IDENTIFIER_RE.match(name) is not None |
There was a problem hiding this comment.
🧩 Analysis chain
🌐 Web query:
Are reserved C/C++ keywords (for example for, while, class, template) legal identifiers for function names?
💡 Result:
No, reserved C/C++ keywords (such as for, while, class, and template) are not legal identifiers for function names [1][2][3]. In C and C++, keywords are reserved by the language specification to provide specific syntactic and semantic meaning to the compiler [1][4]. Because these terms are essential to the language's core structures, they cannot be redefined, overloaded, or used as names for identifiers such as functions, variables, classes, or templates [1][2][5]. The only exception to this rule is the use of these tokens within an attribute-token context (introduced in C++11), where they are not treated as reserved keywords [1][4]. Attempting to use a keyword as a function name will result in a compilation error [3]. Furthermore, beyond explicit keywords, certain other identifiers are reserved for use by the implementation (such as those containing a double underscore or beginning with an underscore followed by an uppercase letter), and using these as function names can lead to undefined behavior [1][2][6].
Citations:
- 1: https://en.cppreference.com/cpp/keywords
- 2: https://cppreference.com/cpp/language/name
- 3: https://www.tpointtech.com/difference-between-tokens-identifiers-and-keywords-in-cpp
- 4: https://eel.is/c++draft/lex.key
- 5: https://www.programiz.com/cpp-programming/keywords-identifiers
- 6: https://timsong-cpp.github.io/cppwp/n4868/lex.name
🏁 Script executed:
# Check the current code at the specified location
head -5140 warp/_src/codegen.py | tail -15Repository: NVIDIA/warp
Length of output: 572
🏁 Script executed:
# Search for any existing reserved keyword handling in codegen.py
rg -n "keyword|reserved|RESERVED" warp/_src/codegen.py | head -20Repository: NVIDIA/warp
Length of output: 883
🏁 Script executed:
# Check if this validator is used for kernel names or function validation
rg -n "is_valid_c_identifier\|_C_IDENTIFIER_RE" warp/_src/codegen.pyRepository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# The is_valid_c_identifier might not be in use yet. Let's check if it's called anywhere
rg "is_valid_c_identifier" warp/_src/Repository: NVIDIA/warp
Length of output: 209
🏁 Script executed:
# Check if there's any validation of kernel names in the kernel decorator or related code
rg -n "def kernel\|`@kernel`" warp/_src/codegen.py | head -10Repository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# Search for where kernel names are processed or validated
rg -n "name=" warp/_src/codegen.py | grep -i "kernel\|func" | head -15Repository: NVIDIA/warp
Length of output: 952
🏁 Script executed:
# Check where is_valid_c_identifier is called in context.py and what it validates
rg -B 5 -A 5 "is_valid_c_identifier" warp/_src/context.pyRepository: NVIDIA/warp
Length of output: 492
🏁 Script executed:
# Also check the kernel decorator implementation to understand the kernel name validation flow
fd "kernel.py\|__init__.py" warp/_src/ --type f | head -5Repository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# Check if there's any existing keyword filtering or handling anywhere in warp codebase
rg -i "keyword|reserved" warp/_src/ --type py | grep -v "\.pyc" | head -20Repository: NVIDIA/warp
Length of output: 2137
🏁 Script executed:
# Look at the make_full_qualified_name function to understand how kernel names are generated
rg -B 3 -A 8 "make_full_qualified_name" warp/_src/codegen.pyRepository: NVIDIA/warp
Length of output: 473
🏁 Script executed:
# Check C/C++ codegen targets to understand which keywords Warp needs to avoid
rg -n "cuda\|c_\+\+\|CPP\|GLSL" warp/_src/codegen.py | grep -i "target\|backend\|language" | head -10Repository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# Check if Warp targets CUDA or other languages beyond C/C++
rg -n "class.*CodeGen\|def.*codegen" warp/_src/codegen.py | head -10Repository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# Verify if CUDA has additional reserved keywords beyond C++
web_search
CUDA C/C++ reserved keywords additional to C++ standardRepository: NVIDIA/warp
Length of output: 149
🏁 Script executed:
# Check if there are existing tests for identifier validation
fd "test.*codegen\|test.*kernel" warp/tests/ --type fRepository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# Find test files related to codegen or kernel
find warp/tests -name "*test*codegen*" -o -name "*test*kernel*" | head -10Repository: NVIDIA/warp
Length of output: 294
🏁 Script executed:
# Check if there are existing tests that might validate identifier handling
rg "is_valid_c_identifier\|_C_IDENTIFIER_RE" warp/tests/Repository: NVIDIA/warp
Length of output: 37
🌐 Web query:
CUDA C reserved keywords beyond C standard library
💡 Result:
In CUDA C/C++ (as implemented by nvcc), the main “reserved keywords beyond the C standard library” are the CUDA language extension keywords/specifiers—i.e., identifiers starting with double underscores that nvcc recognizes and gives special meaning. Key CUDA-reserved specifiers (language extensions) documented by NVIDIA include: - host / device / global (execution space specifiers for where code runs, or to declare kernels) [1] - restrict (pointer restrict qualifier) [2][1] - launch_bounds (function qualifier for launch bounds on global functions) [2][1] - grid_constant (qualifier for parameters) [2][1] - noinline, forceinline, inline_hint (inlining control qualifiers) [2][1] - cluster_dims (cluster dimensions qualifier) [2][1] - device, constant, shared, managed (memory-space/storage specifiers for variables) [1] Additionally, NVIDIA’s CUDA programming guide documents that CUDA source files use these extensions and that any source file containing CUDA extensions must be compiled with nvcc. [3]
Citations:
- 1: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html.md
- 2: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html
- 3: https://docs.nvidia.com/cuda/cuda-programming-guide/
🏁 Script executed:
# Check test_codegen.py for any identifier validation tests
rg -n "is_valid_c_identifier\|C_IDENTIFIER" warp/tests/test_codegen.pyRepository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# Look at test_codegen.py to understand what's already tested
head -100 warp/tests/test_codegen.pyRepository: NVIDIA/warp
Length of output: 1991
🏁 Script executed:
# Check if there are any actual kernel name validations in tests
rg -B 3 -A 3 "name.*=.*`@wp`\.kernel\|`@wp`\.kernel.*name" warp/tests/ | head -30Repository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# Let's check the context where the validator is used to understand the full picture
rg -B 10 -A 5 "is_valid_c_identifier" warp/_src/context.pyRepository: NVIDIA/warp
Length of output: 582
🏁 Script executed:
# Verify there are no existing keyword checks elsewhere in the kernel path
rg -n "keyword\|reserved.*keyword" warp/_src/context.pyRepository: NVIDIA/warp
Length of output: 37
🏁 Script executed:
# Check what happens when a keyword is used as kernel name currently
rg -n "ValueError.*valid.*C.*identifier" warp/_src/Repository: NVIDIA/warp
Length of output: 192
Reject reserved C/C++ keywords in the identifier validator
Line 5138 currently accepts keyword tokens like for/class because it only checks identifier shape. Those are invalid function identifiers in generated C/CUDA code and will fail at compile time, defeating the purpose of early validation.
Proposed fix
_C_IDENTIFIER_RE = re.compile(r"^[A-Za-z_][A-Za-z0-9_]*$")
+_C_RESERVED_KEYWORDS = {
+ # include the full C/C++ keyword set used by Warp codegen targets
+ "auto", "break", "case", "char", "const", "continue", "default", "do",
+ "double", "else", "enum", "extern", "float", "for", "goto", "if", "int",
+ "long", "register", "return", "short", "signed", "sizeof", "static",
+ "struct", "switch", "typedef", "union", "unsigned", "void", "volatile",
+ "while", "class", "namespace", "template", "typename",
+}
def is_valid_c_identifier(name: str) -> bool:
- return _C_IDENTIFIER_RE.match(name) is not None
+ return _C_IDENTIFIER_RE.fullmatch(name) is not None and name not in _C_RESERVED_KEYWORDS🤖 Prompt for 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.
In `@warp/_src/codegen.py` around lines 5135 - 5139, The is_valid_c_identifier
function currently only validates the syntactic shape of identifiers using regex
pattern matching, but fails to reject reserved C/C++ keywords like "for",
"class", "int", and "return" which are syntactically valid identifiers but
cannot be used as function identifiers in generated C/CUDA code. Add a check
within is_valid_c_identifier to verify that the name is not a C/C++ reserved
keyword by comparing it against a comprehensive set of reserved keywords, and
return False if the name matches any keyword before returning the regex match
result.
| else: | ||
| if not warp._src.codegen.is_valid_c_identifier(name): | ||
| raise ValueError(f"@wp.kernel for '{f.__name__}': name '{name}' is not a valid C identifier.") | ||
| key = name |
There was a problem hiding this comment.
Please fix decorator-form @wp.overload with renamed generic kernels.
This branch registers a kernel only under the explicit name= key, but the documented overload-stub path still resolves the kernel from the stub's Python qualified name via Module._find_kernel(). That means this now fails:
from typing import Any
import warp as wp
@wp.kernel(name="custom_scale")
def scale(x: wp.array[Any], s: Any):
i = wp.tid()
x[i] = s * x[i]
@wp.overload
def scale(x: wp.array[wp.float32], s: wp.float32):
...Observed on this branch:
RuntimeError: Failed to find a kernel named 'scale' in module __main__
Direct overload registration still works:
wp.overload(scale, [wp.array[wp.float32], wp.float32])so the regression is specifically in the decorator-form lookup. I think the fix should preserve lookup by the original Python qualified name, or otherwise teach _find_kernel() / @wp.overload how to find kernels whose key was overridden by @wp.kernel(name=...).
Please also add a regression test covering @wp.kernel(name="...") plus decorator-form @wp.overload.
There was a problem hiding this comment.
Indeed, thanks for spotting this issue.
From looking at the code, there are a couple ways to address this. I took the liberty of implementing one of them.
For completeness, these are the options considered:
- Key self.kernels by kernel.qualname instead of kernel.key. This makes the _find_kernel() implementation pretty simple, but messes with the AOT and causes collisions deeper in the code. Not really a valid solution.
- Store the kernel.qualname next to kernel.key. If _find_kernel fails to find a kernel from the func's qualname, assume it was renamed. Then iterate over the kernels in the module, and compare the qualnames. This is the one implemented.
- Keep another dict self.kernels_by_qualname next to self.kernels in the Module class. _find_kernel() uses this other dict. Drawback is that we need to keep these two data structures in sync.
Open to iterating over this.
| - Add a `name` argument to `@wp.kernel` to override the kernel's key, mirroring `@wp.func`. | ||
| ([GH-1561](https://github.com/NVIDIA/warp/issues/1561)). |
There was a problem hiding this comment.
Nitpick: Please revise the CHANGELOG.md entry to avoid saying this mirrors @wp.func.
@wp.func(name=...) is useful precedent for the API shape, but the changelog should describe the kernel behavior directly rather than compare it to a different decorator.
Suggested wording:
Add a `name` argument to `@wp.kernel` to override the kernel's generated key.There was a problem hiding this comment.
Agreed, done!
|
Added a fix for the Now, it can happen that when using the closure pattern with wp.overload, we get multiple possible matches in _find_kernel(). For instance, in code below, the overload would find two kernels with the qualname def make(suffix):
@wp.kernel(name=f"kernel_{suffix}")
def scale(x: Any):
return
@wp.overload # <-- this overload will match the first registered `scale` kernel
def scale(x: float): ...
return scale
first = make("a")
second = make("b")I've added some logic to address this, namely, picking the last match, which should correspond to the correct target kernel. Admittedly, this check is fragile. It assumes that the Alternatively, we can raise an error if multiple matches are found. In this case, users would need to pass a custom name that is different for each closure case. Note: I've added the fix on a separate commit for easier review. I'd squash everything before merging. |
2 participants
Description
Addresses #1561.
@wp.kernelderives a kernel's key frommake_full_qualified_name(func), which in turn uses the function's__qualname__. There's no way to override it.The
@wp.funcdecorator already exposes anameargument for the same purpose.Changes
Here I'm adding support for a
nameargument for thewp.kerneldecorator, enabling overriding the kernel's name.I chose
nameto be consistent with thewp.funcAPI. However, internally this is calledkey.A potential pitfall is if the user passes a name that is not a valid C name. I've added an explicit check for this, otherwise users are faced with a somewhat cryptic compiler error.
New feature / enhancement
This is especially useful when generating kernels with the closure pattern. In this case, different kernels have the same key and only differ by their hashes:
With these changes, users can now specify a name for the kernel:
Checklist
Unreleasedsection.Validation summary
Ran all tests and linter.
Summary by CodeRabbit
@wp.kerneldecorator now accepts an optionalnameargument to customize the kernel’s identifiernameoverrides the auto-generated key@wp.overloadstubs