Enable preshuffled mixed dtype Cutlass Gemm #3722
Closed
+452
−0
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
|
This pull request was exported from Phabricator. Differential Revision: D69955197 |
✅ Deploy Preview for pytorch-fbgemm-docs ready!
To edit notification comments on pull requests, go to your Netlify site configuration. |
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Feb 21, 2025
Summary: WIP to enable new optimized preshuffled fp8xint4 gemm. Differential Revision: D69955197
jwfromm
force-pushed
the
export-D69955197
branch
from
254c644 to
24a824b
Compare
|
This pull request was exported from Phabricator. Differential Revision: D69955197 |
jwfromm
force-pushed
the
export-D69955197
branch
from
24a824b to
dc741e7
Compare
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Feb 21, 2025
Summary: WIP to enable new optimized preshuffled fp8xint4 gemm. While the example compiles and runs, it runs into a variety of problems. The outputs are either completely incorrect, contain NaNs, or the kernel hits an Illegal Memory Access. I'm not yet sure why. Differential Revision: D69955197
|
This pull request was exported from Phabricator. Differential Revision: D69955197 |
|
@IwakuraRein Despite this compiling and running, I'm getting incorrect outputs and very poor performance (even slower than the legacy f8i4 without packing or shuffling). Can you take a look and see if I'm doing something obviously wrong? Ignore files besides f8i4_shuffled.cu and mixed_dtype_utils.cu as the others just fix cutlass v3.8 compatibility. |
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Feb 22, 2025
Summary: WIP to enable new optimized preshuffled fp8xint4 gemm. While the example compiles and runs, it runs into a variety of problems. The outputs are either completely incorrect, contain NaNs, or the kernel hits an Illegal Memory Access. I'm not yet sure why. Differential Revision: D69955197
jwfromm
force-pushed
the
export-D69955197
branch
from
dc741e7 to
70477ce
Compare
|
This pull request was exported from Phabricator. Differential Revision: D69955197 |
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Feb 22, 2025
Summary: WIP to enable new optimized preshuffled fp8xint4 gemm. While the example compiles and runs, it runs into a variety of problems. The outputs are either completely incorrect, contain NaNs, or the kernel hits an Illegal Memory Access. I'm not yet sure why. Differential Revision: D69955197
jwfromm
force-pushed
the
export-D69955197
branch
from
70477ce to
bbca782
Compare
|
This pull request was exported from Phabricator. Differential Revision: D69955197 |
|
@jwfromm Are there negative values in the scale factors? This might be the reason for the accuracy drop after enabling lookup table, and can be easily fixed by applying this change to |
|
@IwakuraRein The scales are all positive and I'm running with the latest cutlass head commit (as of yesterday). The link you posted doesnt seem to include any changes to sm90_mma_tma_gmma_rs_warpspecialized_mixed_input.hpp, did you mean to paste a differint one? |
|
@jwfromm Sorry I mean the changes in |
|
fbgemm_gpu/experimental/gen_ai/bench/quantize_ops.py:1145: - scales = scales.view(x.shape[0], -1)
+ scales = scales.view(x.shape[0], -1).t().contiguous()fbgemm_gpu/experimental/gen_ai/src/quantize/cutlass_extensions/mixed_dtype_utils.cu:59: - StrideB stride_B;
+ StrideB stride_B = cutlass::make_cute_packed_stride(StrideB{}, shape_B);These should fix the bugs. |
jwfromm
force-pushed
the
export-D69955197
branch
from
bbca782 to
15a5738
Compare
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 5, 2025
Summary: Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Differential Revision: D69955197
|
This pull request was exported from Phabricator. Differential Revision: D69955197 |
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 5, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Differential Revision: D69955197
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 5, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 5, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jwfromm
pushed a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 5, 2025
Summary: Pull Request resolved: pytorch#3722 WIP to enable new optimized preshuffled fp8xint4 gemm. Differential Revision: D69955197
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 6, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jiawenliu64
pushed a commit
to jiawenliu64/FBGEMM
that referenced
this pull request
Mar 6, 2025
Summary: Pull Request resolved: pytorch#3722 WIP to enable new optimized preshuffled fp8xint4 gemm. Differential Revision: D69955197
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 6, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jwfromm
force-pushed
the
export-D69955197
branch
from
15a5738 to
3e2ddfa
Compare
|
This pull request was exported from Phabricator. Differential Revision: D69955197 |
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 6, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 6, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jwfromm
force-pushed
the
export-D69955197
branch
from
3e2ddfa to
8407869
Compare
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 7, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
|
This pull request was exported from Phabricator. Differential Revision: D69955197 |
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 7, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 7, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jwfromm
added a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 7, 2025
Summary: X-link: facebookresearch/FBGEMM#846 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197
jwfromm
pushed a commit
to jwfromm/FBGEMM
that referenced
this pull request
Mar 7, 2025
Summary: Pull Request resolved: pytorch#3722 WIP to enable new optimized preshuffled fp8xint4 gemm. Differential Revision: D69955197
Summary: Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Differential Revision: D69955197
jwfromm
force-pushed
the
export-D69955197
branch
from
8407869 to
b225a2a
Compare
|
This pull request has been merged in 27724d9. |
q10
pushed a commit
to q10/FBGEMM
that referenced
this pull request
Apr 10, 2025
Summary: Pull Request resolved: facebookresearch/FBGEMM#846 X-link: pytorch#3722 Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete. Reviewed By: jiawenliu64 Differential Revision: D69955197 fbshipit-source-id: 151b5dd96728b82b4bf5a4c3967310d055d56094
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.
Enable new preshuffled FP8 x I4 kernels. These are the most performant mixed dtype kernels to date and dramatically outperform prior approaches including those in FBGEMM, marlin, and Machete.
Differential Revision: D69955197