introduce kernel for converting e4m3fn kv_cache to e4m3fnuz#4113
Closed
bradleyhd wants to merge 1 commit into
Closed
introduce kernel for converting e4m3fn kv_cache to e4m3fnuz#4113bradleyhd wants to merge 1 commit into
bradleyhd wants to merge 1 commit into
Conversation
✅ Deploy Preview for pytorch-fbgemm-docs ready!
To edit notification comments on pull requests, go to your Netlify site configuration. |
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D74502334 |
bradleyhd
force-pushed
the
export-D74502334
branch
from
29a8228 to
887eadb
Compare
bradleyhd
added a commit
to bradleyhd/FBGEMM
that referenced
this pull request
May 12, 2025
…4113) Summary: X-link: facebookresearch/FBGEMM#1197 As titled. NV uses e4m3fn for fp8 KV cache, and AMD uses e4m3fnuz. If we want to do HH serving with NV prefill and AMD decode, we need a way to convert between the two. In practice this can be done with two operations: overwriting -0 (which are valid in e4m3fn but not in uz) with +0, and multiplying the scale qparam by 2, since the exponent bias differs by 1 between the two formats. The scale qparam is packed as the low bits of a __half2 (fp16x2). This is a first step of introducing the kernel only. A follow up step will be to wire it into DisaggDecodeCommunicator and ensure the qparam tensor shape matches, etc. **NB: this is my first CUDA kernel ever, so please review carefully** Differential Revision: D74502334
bradleyhd
added a commit
to bradleyhd/FBGEMM
that referenced
this pull request
May 12, 2025
…4113) Summary: X-link: facebookresearch/FBGEMM#1197 As titled. NV uses e4m3fn for fp8 KV cache, and AMD uses e4m3fnuz. If we want to do HH serving with NV prefill and AMD decode, we need a way to convert between the two. In practice this can be done with two operations: overwriting -0 (which are valid in e4m3fn but not in uz) with +0, and multiplying the scale qparam by 2, since the exponent bias differs by 1 between the two formats. The scale qparam is packed as the low bits of a __half2 (fp16x2). Differential Revision: D74502334
bradleyhd
force-pushed
the
export-D74502334
branch
from
887eadb to
03f1e79
Compare
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D74502334 |
2 similar comments
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D74502334 |
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D74502334 |
bradleyhd
added a commit
to bradleyhd/FBGEMM
that referenced
this pull request
May 12, 2025
…4113) Summary: Pull Request resolved: pytorch#4113 X-link: https://github.com/facebookresearch/FBGEMM/pull/1197 As titled. NV uses e4m3fn for fp8 KV cache, and AMD uses e4m3fnuz. If we want to do HH serving with NV prefill and AMD decode, we need a way to convert between the two. In practice this can be done with two operations: overwriting -0 (which are valid in e4m3fn but not in uz) with +0, and multiplying the scale qparam by 2, since the exponent bias differs by 1 between the two formats. The scale qparam is packed as the low bits of a __half2 (fp16x2). Differential Revision: D74502334
bradleyhd
force-pushed
the
export-D74502334
branch
2 times, most recently
from
4e780a4 to
9931f52
Compare
bradleyhd
added a commit
to bradleyhd/FBGEMM
that referenced
this pull request
May 13, 2025
…4113) Summary: X-link: facebookresearch/FBGEMM#1197 As titled. NV uses e4m3fn for fp8 KV cache, and AMD uses e4m3fnuz. If we want to do HH serving with NV prefill and AMD decode, we need a way to convert between the two. In practice this can be done with two operations: overwriting -0 (which are valid in e4m3fn but not in uz) with +0, and multiplying the scale qparam by 2, since the exponent bias differs by 1 between the two formats. The scale qparam is packed as the low bits of a __half2 (fp16x2). Differential Revision: D74502334
bradleyhd
force-pushed
the
export-D74502334
branch
from
9931f52 to
99075a7
Compare
bradleyhd
added a commit
to bradleyhd/FBGEMM
that referenced
this pull request
May 13, 2025
…4113) Summary: X-link: facebookresearch/FBGEMM#1197 As titled. NV uses e4m3fn for fp8 KV cache, and AMD uses e4m3fnuz. If we want to do HH serving with NV prefill and AMD decode, we need a way to convert between the two. In practice this can be done with two operations: overwriting -0 (which are valid in e4m3fn but not in uz) with +0, and multiplying the scale qparam by 2, since the exponent bias differs by 1 between the two formats. The scale qparam is packed as the low bits of a __half2 (fp16x2). Differential Revision: D74502334
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D74502334 |
1 similar comment
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D74502334 |
…4113) Summary: Pull Request resolved: pytorch#4113 X-link: https://github.com/facebookresearch/FBGEMM/pull/1197 As titled. NV uses e4m3fn for fp8 KV cache, and AMD uses e4m3fnuz. If we want to do HH serving with NV prefill and AMD decode, we need a way to convert between the two. In practice this can be done with two operations: overwriting -0 (which are valid in e4m3fn but not in uz) with +0, and multiplying the scale qparam by 2, since the exponent bias differs by 1 between the two formats. The scale qparam is packed as the low bits of a __half2 (fp16x2). Differential Revision: D74502334
bradleyhd
force-pushed
the
export-D74502334
branch
from
99075a7 to
612052b
Compare
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D74502334 |
Contributor
|
This pull request has been merged in 40259b2. |
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
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.
Summary:
X-link: https://github.com/facebookresearch/FBGEMM/pull/1197
As titled. NV uses e4m3fn for fp8 KV cache, and AMD uses e4m3fnuz. If we want to do HH serving with NV prefill and AMD decode, we need a way to convert between the two.
In practice this can be done with two operations: overwriting -0 (which are valid in e4m3fn but not in uz) with +0, and multiplying the scale qparam by 2, since the exponent bias differs by 1 between the two formats. The scale qparam is packed as the low bits of a __half2 (fp16x2).
Differential Revision: D74502334