[CUDA] FpA IntB Gemm Kernel Test #25109
Open
+857
−152
Conversation
This file contains hidden or 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
| @@ -39,7 +40,9 @@ | |||
| #include "contrib_ops/cuda/llm/cutlass_heuristic.h" | |||
| #include "contrib_ops/cuda/llm/cutlass_type_conversion.h" | |||
| #include "contrib_ops/cuda/llm/fpA_intB_gemm/fpA_intB_gemm.h" | |||
| #ifndef EXCLUDE_SM_90 | |||
There was a problem hiding this comment.
Why is sm90 particularly slow during compilation and not newer ones?
There was a problem hiding this comment.
sm90 uses a new set of files here.
Blackwell GPU will fall back to use sm80 kernel.
Since CI pipeline uses sm75 or sm86, so no need to compile sm90. This skip those files and might speed up build.
nenad1002
reviewed
Jun 19, 2025
| @@ -374,14 +375,18 @@ void CutlassFpAIntBGemmRunner<ActivationType, WeightType, QuantOp, ScaleZeroType | |||
| dispatch_gemm_to_cutlass<ActivationType, WeightType, ScaleZeroType, BiasType, OutputType, cutlass::arch::Sm89, | |||
| QuantOp, EpilogueTag>(A, B, weight_scales, weight_zero_points, biases, alpha, C, m, n, k, group_size, | |||
| workspace_ptr, workspace_bytes, gemm_config, stream, occupancy); | |||
| #ifndef EXCLUDE_SM_90 | |||
There was a problem hiding this comment.
Is this needed since you already do a check without macro?
There was a problem hiding this comment.
This can avoid compiling sm90_dispatch_gemm_to_cutlass in CI pipeline, or when your GPU is not H100/H200.
| @@ -67,6 +58,18 @@ void kernel_launcher(int arch, Params& params, cudaStream_t s) { | |||
|
|
|||
| EXEC(KernelType::BF16Int8Groupwise, BF16DetailsA, Int8DetailsW, ColumnMajorInterleavedForHopper, true); | |||
| EXEC(KernelType::BF16Int4Groupwise, BF16DetailsA, Int4DetailsW, ColumnMajorInterleavedForHopper, true); | |||
| #endif | |||
| } else { | |||
| // if (arch >= 89) | |||
There was a problem hiding this comment.
This is reserved for a new op that supports alpha.
nenad1002
reviewed
Jun 19, 2025
| #define PRETTY_FUNCTION __PRETTY_FUNCTION__ | ||
| #endif | ||
|
|
||
| #define ORT_LLM_VERBOSE 0 // Set to 1 for verbose, 2 for max verbosity |
There was a problem hiding this comment.
#ifndef ORT_LLM_VERBOSE
#define ORT_LLM_VERBOSE 0
#endif
So we can externally modify it?
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.
Enhance MatMulNBits CUDA kernel testing:
(1) Add a kernel testing for different cuda kernels used in MatMulNBits.
(2) Refactoring the gemm profiler to use cuda allocator
(2) Add verbose logging macros.
(3) Adjustments to speed up compiling when sm90 is excluded from build.
Example kernel test output:
