-
Notifications
You must be signed in to change notification settings - Fork 34
[AMDGPU] Feat: per-kernel LLVM function attributes via @qd.kernel(fn_attrs=...) #774
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: main
Are you sure you want to change the base?
Changes from all commits
0f44102
9cb65ab
3880053
c094a4a
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -5,6 +5,7 @@ | |
| from functools import update_wrapper, wraps | ||
| from typing import Any, Callable, TypeVar, cast, overload | ||
|
|
||
| from quadrants._lib import core as _qd_core | ||
| from quadrants.lang import impl | ||
| from quadrants.lang.exception import ( | ||
| QuadrantsCompilationError, | ||
|
|
@@ -153,13 +154,34 @@ def _inside_class(level_of_class_stackframe: int) -> bool: | |
| return False | ||
|
|
||
|
|
||
| def _validate_fn_attrs(fn_attrs: dict[str, dict[str, str]], func_name: str) -> None: | ||
| registry = _qd_core.get_fn_attrs_registry() | ||
| for backend, attrs in fn_attrs.items(): | ||
| if backend not in registry: | ||
| raise QuadrantsSyntaxError( | ||
| f"@qd.kernel(fn_attrs=...) on {func_name!r}: unknown backend " | ||
| f"{backend!r}. Registered backends: {sorted(registry)}." | ||
| ) | ||
| allowed = registry[backend] | ||
| for key in attrs: | ||
| if key not in allowed: | ||
| raise QuadrantsSyntaxError( | ||
| f"@qd.kernel(fn_attrs=...) on {func_name!r}: unknown " | ||
| f"{backend} attribute {key!r}. Registered attributes: " | ||
| f"{sorted(allowed)}." | ||
| ) | ||
|
|
||
|
|
||
| def _kernel_impl( | ||
| _func: Callable, | ||
| level_of_class_stackframe: int, | ||
| verbose: bool = False, | ||
| graph: bool = False, | ||
| checkpoints: bool = False, | ||
| fn_attrs: dict[str, dict[str, str]] | None = None, | ||
| ) -> QuadrantsCallable: | ||
| if fn_attrs: | ||
| _validate_fn_attrs(fn_attrs, _func.__name__) | ||
| # Can decorators determine if a function is being defined inside a class? | ||
| # https://stackoverflow.com/a/8793684/12003165 | ||
| is_classkernel = _inside_class(level_of_class_stackframe + 1) | ||
|
|
@@ -171,6 +193,8 @@ def _kernel_impl( | |
| primal.use_graph = graph | ||
| primal.use_checkpoints = checkpoints | ||
| adjoint.use_checkpoints = checkpoints | ||
| primal.fn_attrs = fn_attrs | ||
| adjoint.fn_attrs = fn_attrs | ||
| # Having |primal| contains |grad| makes the tape work. | ||
| primal.grad = adjoint | ||
|
|
||
|
|
@@ -213,7 +237,12 @@ def wrapped_classkernel(*args, **kwargs): | |
| # TODO: This callable should be Callable[[F], F]. | ||
| # See comments below. | ||
| def kernel( | ||
| _fn: None = None, *, pure: bool = False, graph: bool = False, checkpoints: bool = False | ||
| _fn: None = None, | ||
| *, | ||
| pure: bool = False, | ||
| graph: bool = False, | ||
| checkpoints: bool = False, | ||
| fn_attrs: dict[str, dict[str, str]] | None = None, | ||
| ) -> Callable[[Any], Any]: ... | ||
|
|
||
|
|
||
|
|
@@ -224,7 +253,14 @@ def kernel( | |
| # However, by making it return Any, we can make the pure parameter | ||
| # change now, without breaking pyright. | ||
| @overload | ||
| def kernel(_fn: Any, *, pure: bool = False, graph: bool = False, checkpoints: bool = False) -> Any: ... | ||
| def kernel( | ||
| _fn: Any, | ||
| *, | ||
| pure: bool = False, | ||
| graph: bool = False, | ||
| checkpoints: bool = False, | ||
| fn_attrs: dict[str, dict[str, str]] | None = None, | ||
| ) -> Any: ... | ||
|
|
||
|
|
||
| def kernel( | ||
|
|
@@ -234,6 +270,7 @@ def kernel( | |
| fastcache: bool = False, | ||
| graph: bool = False, | ||
| checkpoints: bool = False, | ||
| fn_attrs: dict[str, dict[str, str]] | None = None, | ||
|
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more.
This exposes Useful? React with 👍 / 👎.
Contributor
Author
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Fixed in commit 9cb65ab. Added a 'Per-kernel LLVM function attributes' section to |
||
| ): | ||
| """ | ||
| Marks a function as a Quadrants kernel. | ||
|
|
@@ -278,7 +315,13 @@ def decorator(fn: F, has_kernel_params: bool = True) -> F: | |
| f"@qd.kernel({fn.__name__!r}, checkpoints=True) requires graph=True; " | ||
| "the checkpoint resume model is only meaningful for graph kernels." | ||
| ) | ||
| wrapped = _kernel_impl(fn, level_of_class_stackframe=level, graph=graph, checkpoints=checkpoints) | ||
| wrapped = _kernel_impl( | ||
| fn, | ||
| level_of_class_stackframe=level, | ||
| graph=graph, | ||
| checkpoints=checkpoints, | ||
| fn_attrs=fn_attrs, | ||
| ) | ||
| wrapped.is_pure = pure is not None and pure or fastcache | ||
| if pure is not None: | ||
| warnings_helper.warn_once( | ||
|
|
||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,33 @@ | ||
| #pragma once | ||
|
|
||
| #include <string> | ||
| #include <unordered_map> | ||
| #include <unordered_set> | ||
|
|
||
| namespace quadrants::lang { | ||
|
|
||
| // Central registry of LLVM function attributes that users may set per-kernel | ||
| // via `@qd.kernel(fn_attrs={...})`. Keyed by backend short name (matching | ||
| // the namespacing used in the Python dict, e.g. "amdgpu"). Adding a new | ||
| // attribute means: (1) register it here, (2) ensure the consuming backend | ||
| // honors it (for AMDGPU: codegen_llvm.cpp applies it via addFnAttr; if the | ||
| // backend has a hardcoded default for the same key, gate that default with | ||
| // `!F.hasFnAttribute(key)` in the corresponding jit_*.cpp). | ||
| inline const std::unordered_map<std::string, std::unordered_set<std::string>> & | ||
| get_fn_attrs_registry() { | ||
| static const std::unordered_map<std::string, | ||
| std::unordered_set<std::string>> | ||
| kRegistry = { | ||
| {"amdgpu", | ||
| { | ||
| "amdgpu-max-num-workgroups", | ||
| "amdgpu-agpr-alloc", | ||
| "amdgpu-waves-per-eu", | ||
| "amdgpu-flat-work-group-size", | ||
| "amdgpu-sched-strategy", | ||
| }}, | ||
| }; | ||
| return kRegistry; | ||
| } | ||
|
|
||
| } // namespace quadrants::lang |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Platform-specific configuration conflicts with the ideal that Quadrants code will run unmodified on all platforms.
The standard approach is to set appropriate defaults - or use appropriate heuristics - within each platform-specific pipeline.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thanks — I agree with the principle, and I think this PR is consistent with it. The sensible AMDGPU defaults already live in the backend pipeline (
jit_amdgpu.cpp, each gated byhasFnAttribute);fn_attrsis just an optional override on top for when a user knows better. Portable code never needs it.Portability is also preserved: attributes are only applied under
arch == Arch::amdgpu(codegen_llvm.cpp:3244), so@qd.kernel(fn_attrs={"amdgpu": {...}})still runs unmodified on CPU/CUDA/Metal — the block is simply ignored there. It's namespaced by backend for exactly that reason, much like CUDA__launch_bounds__or Tritonnum_warps: an optional hint that only tunes the backend that understands it.To land the right scope, which is your concern:
@qd.kernel(...)?If (1), I'll narrow this PR to just the backend-default infrastructure and drop the public
fn_attrs=for now. If (2), I'll move it behind a clearly-marked advanced API. Just say which unblocks the merge.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
My concern is anything platform-specific in the public API.
What I want to avoid is code that runs on one platform, and not on other; or on all platforms except one.
This includes optimizations, ideally. Code should ideally run optimally on each platform without the user needing platform-specific knowledge.