Grok-1 optimization #164
Closed
Grok-1 optimization #164
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…iling code and fix the gemm shape could not be dumped correctly in multiple-gpu
…fused_moe accuracy check test file 4) sync optimization from branch MLPerf-4.1
* First version * Revert error. While there, add missing finalize. * Use the correct defaults for ROCm. Increase sampling area to capture crossover. * Scope end_sync as well. * Guard only volatile keyword for ifndef USE_ROCM * Document crossover
* remove scoping * while there fix a typo * while there remove unused variable
…the different TP setting
There was a problem hiding this comment.
Would you mind renaming this file to not use a '+' symbol?
maleksan85
reviewed
Sep 3, 2024
| @@ -328,7 +332,7 @@ if (VLLM_PUNICA_GPU_ARCHES) | |||
| DESTINATION vllm | |||
| LANGUAGE ${VLLM_GPU_LANG} | |||
| SOURCES ${VLLM_PUNICA_EXT_SRC} | |||
| COMPILE_FLAGS ${VLLM_PUNICA_GPU_FLAGS} | |||
| COMPILE_FLAGS ${VLLM_PUNICA_GPU_eLAGS} | |||
rasmith
reviewed
Sep 3, 2024
| #triton_col_output = fused_moe_col_major(a, w1, w2, score, topk, renormalize=False) | ||
| #print(f"triton_col_output: {triton_col_output}") | ||
|
|
||
| assert torch.allclose(triton_output, torch_output, atol=1e-2, rtol=0) |
| @@ -293,6 +293,8 @@ __device__ void paged_attention_kernel( | |||
| // This includes a reduction across the threads in the same thread group. | |||
| float qk = scale * Qk_dot<scalar_t, THREAD_GROUP_SIZE>::dot( | |||
| q_vecs[thread_group_offset], k_vecs); | |||
| float max_attn_val = 30.0; //hardcoded for grok | |||
There was a problem hiding this comment.
Does this affect any other model? Can this be moved to a #define or constexpr or something?
| @@ -998,4 +1000,4 @@ void paged_attention_v2( | |||
| #undef WARP_SIZE | |||
| #undef MAX | |||
| #undef MIN | |||
| #undef DIVIDE_ROUND_UP | |||
There was a problem hiding this comment.
Does this affect anything else? Still need everything else to work.
| @@ -328,7 +332,7 @@ if (VLLM_PUNICA_GPU_ARCHES) | |||
| DESTINATION vllm | |||
| LANGUAGE ${VLLM_GPU_LANG} | |||
| SOURCES ${VLLM_PUNICA_EXT_SRC} | |||
| COMPILE_FLAGS ${VLLM_PUNICA_GPU_FLAGS} | |||
| COMPILE_FLAGS ${VLLM_PUNICA_GPU_eLAGS} | |||
There was a problem hiding this comment.
This looks to be a typo, probably does not build.
|
Thanks for the feedback, we will correct in our new PR "#181" based on the vllm 0.6.0. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR adds GROK-FP8 support in vLLM.