cuda: add sqrt_softplus in topk-moe for dsv4#25896
Conversation
github-actions
Bot
added
testing
ORippler
left a comment
ORippler
left a comment
There was a problem hiding this comment.
Left some non-blocking comments
| __device__ void sqrt_softplus_warp_inplace(float (&vals)[experts_per_thread], const int limit, const int lane) { | ||
| #pragma unroll | ||
| for (int i = 0; i < experts_per_thread; i++) { | ||
| const int idx = lane + i * WARP_SIZE; |
There was a problem hiding this comment.
| const int idx = lane + i * WARP_SIZE; | |
| const int idx = lane + i * ggml_cuda_get_physical_warp_size(); |
This should resolve to 32 on CUDA and 64 on HIP afaik
There was a problem hiding this comment.
Edit: Looking at the kernel-context, this is something that could/should be done in a separate follow-up PR for the whole op
| } | ||
| } | ||
|
|
||
| template <int experts_per_thread, bool use_limit> |
There was a problem hiding this comment.
| template <int experts_per_thread, bool use_limit> | |
| template <int experts_per_thread> |
It seems like the function is instantiated with use_limit=false always, so we may remove this template parameter (other activations seem to be sharing this potentially -> may be worth follow-up PR)
| bool sigmoid{}; | ||
| bool sqrt_softplus{}; | ||
| bool softmax{}; |
There was a problem hiding this comment.
May consider adding an enum in future, as sigmoid/sqrt_softplus and softmax seem to be mutually exclusive
|
@ggml-org/maintainers can I get another approval please? |
3 participants
Overview
Add sqrt_softplus to the list of supported topk-moe functions.
Additional information
Requirements