[Improvement] Re-structure sonic MoE kernel interface to use a Handle object - #47316
[Improvement] Re-structure sonic MoE kernel interface to use a Handle object#47316remi-or wants to merge 8 commits into
Conversation
|
The docs for this PR live here. All of your documentation changes will be reflected on that endpoint. The docs are available until 30 days after the last update. |
CI recapDashboard: View test results in Grafana |
| if missing: | ||
| raise ImportError( | ||
| f"sonic-moe kernel is missing required symbols: {', '.join(missing)}. " | ||
| "Make sure you have the `kernels` package and `nvidia-cutlass-dsl` installed." |
There was a problem hiding this comment.
can you please also add the specific versions that worked for you of tvm ffi and cutedsl, i didn't expect it to start failing so fast with newer versions, when we tell users to install them they'll just do "pip install nvidia-cutlass-dsl" which would then result in version mismatch between the vendored quack and the installed cutedsl
cc @vasqu
There was a problem hiding this comment.
💯 we just had discussions around that for the past few days so let's guard it on our side properly
There was a problem hiding this comment.
Here you go:
nvidia-cutlass-dsl 4.5.2
apache-tvm-ffi 0.1.9
| # check if cutlass-dsl is installed | ||
| from cutlass.utils.hardware_info import HardwareInfo | ||
|
|
||
| # sonic-moe JIT-builds CuteDSL kernels; bail early if the driver can't load their device image. | ||
| # HardwareInfo uses the CUDA driver API directly and needs a current context, which torch only | ||
| # creates lazily — without this, the probe fails with CUDA_ERROR_INVALID_CONTEXT. | ||
| torch.cuda.synchronize() | ||
| try: | ||
| HardwareInfo().get_max_active_clusters(1) | ||
| except Exception as e: # cutlass wraps the CUDA driver error in a bare RuntimeError | ||
| raise ImportError( | ||
| f"Image error: sonic-moe's CuteDSL kernels cannot load a device image on this GPU/driver " | ||
| f"({type(e).__name__}: {e}). This usually means cutlass-dsl's bundled CUDA toolchain is " | ||
| f"newer than the driver (e.g. cu13 libs on a CUDA 12.x driver)." | ||
| ) |
There was a problem hiding this comment.
sorry but i am once again totally against syncs and tensor creation in the loader (also does this not break cuda graph capture ?) when a more direct alternative exist, in this case checking cutedsl exists and explicitly verifying its version. here the user sees import error no package named cutlass.
There was a problem hiding this comment.
Yeah the synchronize should be guarded against, mb. As for tensor creation, sure I don't mind avoiding them, I will look into the version check.
| def sonicmoe_experts_forward( | ||
| self, | ||
| module: torch.nn.Module, | ||
| hidden_states: torch.Tensor, | ||
| top_k_index: torch.Tensor, | ||
| top_k_weights: torch.Tensor, | ||
| ) -> torch.Tensor: |
There was a problem hiding this comment.
the scope of the handle is a bit fuzzy, it doesn't handle the full loading (lets _load_sonicmoe_kernel do part of it) and here it implements something very implementation specific.. With FinegrainedFP8 for example which implements multiple implementations, it will be a very big handle. I think the implementation part should stay outside of it.
There was a problem hiding this comment.
Yea I think the goal here was to only handle the import / cache structure; in that case we really should be strict and dont mix functionality
There was a problem hiding this comment.
Makes a lot of sense, will change it up
IlyasMoutawwakil
left a comment
There was a problem hiding this comment.
love the idea of the handle and standardizing the interaction with the kernel package (loading, checking for availability, completeness and usability), i just think it can be scoped better and also it should generalize to other kernel integrations. for example we can make it so that it has a class attr of required packages and min/max versions etc, a bit like we do in exporters' validate_environment. the list of required attributes as well can be programatically checked. ofc some custom checks will be necessary for deepgemm and its cuda toolkit for example.
vasqu
left a comment
There was a problem hiding this comment.
I feel like the pain points are valid but the direction is not ideal
I think we already have existing patterns we could apply to solve all these problems, notably
- the lru cache (with the loading structure of deepgemm for example)
- the availability check can very well be moved to import utils (and it makes things a lot clearer, not relying on try catch)
| def sonicmoe_experts_forward( | ||
| self, | ||
| module: torch.nn.Module, | ||
| hidden_states: torch.Tensor, | ||
| top_k_index: torch.Tensor, | ||
| top_k_weights: torch.Tensor, | ||
| ) -> torch.Tensor: |
There was a problem hiding this comment.
Yea I think the goal here was to only handle the import / cache structure; in that case we really should be strict and dont mix functionality
|
|
||
|
|
||
| # Singleton object: this should be the only SonicMoeHandle object instantiated in the codebase | ||
| SONIC_MOE_HANDLE = SonicMoeHandle() |
There was a problem hiding this comment.
Tbh I'm not a fan of singleton objects (but that's more of a personal opinion) - it feels like OOP for the sake of OOP 😅
I'm still a bit confused about the actual benefits and why we may not adopt a generalized pattern similar to flash attention
transformers/src/transformers/modeling_flash_attention_utils.py
Lines 124 to 272 in a689565
I think the biggest point is is_sonicmoe_available and the avoidance of the lru cache
is_sonicmoe_availablecould be properly moved to import utils and we make proper versioning guards, e.g. is kernels available and the respective cute version.- Re lru cache, I feel like this is valid but we can restructure to an FA like structure as well (I think deepgemm does something similar)
There was a problem hiding this comment.
Haha my personal would be a singleton object over globals, but then again it comes back to the same results (pretty much) and FA already uses globals so might as well do the same. Will reformat into the same pattern and we can re-evaluate from there.
|
Lots of feedback, thanks! This was more of a rough PoC so lots to take in, and I did not plan to spend too much time on this, it was more because tests were crashing on main that I took the opportunity to look into it. Will update when I have the bandwidth! |
|
No worries, we also nudged the kernels team so maybe we can simplify a bit more in the future. It's a bit awkward to handle these kernels 😢 |
In #47126 and offline we talked about how we want to take care of loading kernels like
DeepGEMM.Notably, we can put everything related to loading and calling the kernel inside a singleton
Handleobject. This object would only keep track of a few state variable and provide a single entry point to access the kernel. This has the following advantages:@lru_cachewhich can be opaque and cause memory leaks@lru_cachedoes not cache exceptions, so the cache never worked in a machine / env that did not supportsonicmoeresetmethodsonicmoe_is_availableproperty, that accesses the same underlying cached kernel as the function call itself (quite convenient, as is shown in the reworked_test_eager_matches_batched_and_grouped_inference)The downside of the Handle API are that transient failures are also cached (hence the
resetmethod, but it is not called automatically) and it's beefier than@lru_cache. IMO this is an ok tradeoff, but I would love some feedback on this.Since I was fixing a failure in
sonicmoe, I thought this was a good occasion to try out the Handle API. The fail was an image failure, which is now guarded against.I also added a test for sonicmoe + compile and it passes on B200. Also tested the sonicmoe regular tests on a qwen and they pass.