-
Notifications
You must be signed in to change notification settings - Fork 36
[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.
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.
Agreed as the default, and it holds for the common case. For peak perf I don't think it's attainable, though — no production GPU stack manages it. CUDA and HIP ship launch_bounds precisely so the user can state an occupancy/register tradeoff the compiler can't infer; Triton has num_warps/num_stages/@autotune; Halide/TVM make the schedule explicit. The vendors ship these knobs because the "sufficiently smart compiler" doesn't exist for this.
On the "AMD-only" objection: I'd argue that's expected, not a defect. Tuning is vendor-owned — AMD owns the AMD backend and its lowering and can't be expected to implement the NVIDIA/Intel/Apple equivalents; each vendor adds their own. That's how the ecosystem is already structured (torch.backends.cuda/.cudnn/.mps, LLVM's namespaced amdgpu-/nvvm. attributes).
And crucially, the current design already answers the "runs on one platform, breaks on another" worry:
It's namespaced by backend — fn_attrs={"amdgpu": {...}}, applied only under arch == Arch::amdgpu (codegen_llvm.cpp:3244). @qd.kernel(fn_attrs={"amdgpu": {...}}) runs unmodified on CPU/CUDA/Metal; the block is just ignored. Portable code is unaffected; only the backend that understands the hint tunes on it — same model as launch_bounds.
It's registry-validated, not an open escape hatch — _validate_fn_attrs rejects unknown backends and unknown attributes against get_fn_attrs_registry(), so the surface is curated and enumerable rather than "any LLVM attribute goes."
So it's an optional, explicitly-scoped, validated per-backend hint — opt-in, ignored elsewhere, and owned by the vendor whose backend it tunes. That seems to satisfy the underlying concern (no silent cross-platform breakage, no unbounded surface) while accepting that vendor-specific tuning is a legitimate and necessary thing to expose.
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.
Interesting. Currently, in genesis-world, we are not - I think - doing such tuning, for either CPU or CUDA. Perhaps we should. Or perhaps we are doing so implicitly.
I would much prefer to ground design and needs on real-world needs, ideally.
Do you see anywhere in genesis-world where we would be able to improve our performance, concretely, on our rigid benchmarks https://github.com/Genesis-Embodied-AI/genesis-world/blob/main/tests/benchmarks/test_rigid.py , for CUDA and AMD, by doing such tuning? And crucially, do you see anywhere where such tuning would use different numbers for each of CUDA and AMD?
This would give a concrete example to inspect and discuss, rather than discussing in the abstract.
(You could also provide a non genesis-world example, but, if it's about as much work for you to use genesis-world vs some other example, genesis-world will be a lot easier for me personally to understand and gain intuition and insight from).
Uh oh!
There was an error while loading. Please reload this page.
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.
Benchmark:
amdgpu-waves-per-euvia this PR's per-kernelfn_attrsI benchmarked the knob this PR exposes on a Genesis workload — the
anymal_randombenchmark in
tests/benchmarks/test_rigid.py, 30,000 parallel envs, MI308X GPU.I applied
amdgpu-waves-per-eu="2,4"to the rigid-body solver kernels via@qd.kernel(fn_attrs={"amdgpu": {"amdgpu-waves-per-eu": "2,4"}})and compared against thesame kernels with no attribute (compiler default occupancy):
anymal_random(30k envs, GPU)This node runs the CG constraint solver (the parametrization passes
solver=None). Eachround runs default and tuned back-to-back and takes the ratio within the round, averaged
with a geometric mean (30 s warmup + 30 s record, 5 rounds).
What was modified: the feature under test is this PR's
fn_attrs=API in Quadrants. Toexercise it, the only change on the Genesis side is passing
fn_attrs=...on therigid-solver kernel decorators — the
anymal_randomworkload and theruntime_fpsmetricare the stock upstream
test_rigid.py. The 5-round interleaved averaging is an externalwrapper around repeated stock runs (to cancel GPU clock drift); the longer warmup/record
windows are an optional tweak and not required to see the effect.
amdgpu-waves-per-eudirectly tells the register allocator the target occupancy(resident wavefronts per EU), trading latency-hiding against registers-per-wave. This is
a more direct handle than
__launch_bounds__(which HIP and CUDA already expose):__launch_bounds__constrains occupancy only indirectly, as a side effect of a block-sizecap and an optional min-blocks hint, whereas
amdgpu-waves-per-eusets the occupancytarget itself. That LLVM attribute isn't otherwise reachable from the JIT kernel path —
this PR is what makes it settable per kernel. The optimal point is kernel- and
workload-specific, so per-kernel granularity is the right level rather than a single
global flag, and the results show it has first-order performance impact on a real, public
workload — so exposing it is clearly worthwhile.
Reproduction
Requirements: Quadrants built from this PR, an AMD gfx942 GPU (MI3xx), and a Genesis build.
test_rigid.pyis run unmodified — the only source edit is adding the attribute to the rigid-solver kernel decorators (that's the feature under test).Step 1. Add
fn_attrs={"amdgpu": {"amdgpu-waves-per-eu": "2,4"}}to these@qd.kernel(...)decorators —kernel_step_1andkernel_step_2ingenesis/engine/solvers/rigid/rigid_solver.py, andfunc_solve_initand_kernel_solve_monolithingenesis/engine/solvers/rigid/constraint/solver.py. For example,@qd.kernel(fastcache=True)becomes@qd.kernel(fastcache=True, fn_attrs={"amdgpu": {"amdgpu-waves-per-eu": "2,4"}}).Step 2. Run the stock benchmark node once with the edit (tuned) and once without it (baseline):
Step 3. Compare
runtime_fps. A single run of each already shows the gain; averaging a few back-to-back rounds just tightens it against GPU clock noise.