Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Support fp8 w8a8 for pt backend #2959

Open
wants to merge 5 commits into
base: main
Choose a base branch
from
Open

Support fp8 w8a8 for pt backend #2959

wants to merge 5 commits into from

Conversation

RunningLeon
Copy link
Collaborator

Motivation

Support fp8 w8a8 for pt backend

Modification

Support fp8 w8a8 for pt backend

BC-breaking (Optional)

None

Use cases (Optional)

fp8 quant

lmdeploy lite smooth_quant \
  meta-llama/Meta-Llama-3-8B-Instruct \
  --work-dir Meta-Llama-3-8B-Instruct-fp8-w8a8 \
  --quant-dtype fp8

chat

lmdeploy  chat Meta-Llama-3-8B-Instruct-fp8-w8a8 --backend pytorch

Checklist

  1. Pre-commit or other linting tools are used to fix the potential lint issues.
  2. The modification is covered by complete unit tests. If not, please add more unit tests to ensure the correctness.
  3. If the modification has a dependency on downstream projects of a newer version, this PR should be tested with all supported versions of downstream projects.
  4. The documentation has been modified accordingly, like docstring or example tutorials.

AllentDan
AllentDan reviewed Dec 27, 2024
View reviewed changes
docs/zh_cn/quantization/w8a8.md Outdated Show resolved Hide resolved
lmdeploy/pytorch/kernels/cuda/w8a8_triton_kernels.py Outdated Show resolved Hide resolved
lmdeploy/pytorch/kernels/cuda/w8a8_triton_kernels.py
accumulator += tl.dot(a, b)
a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_K, other=None)
b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_K, other=None)
accumulator = tl.dot(a, b, accumulator, out_dtype=ACCUMULATOR_DTYPE)
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Does it help improve our other kernels that use tl.dot too? @grimoire

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Most gemm kernels should be benefited from this.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@AllentDan AllentDan AllentDan left review comments

@grimoire grimoire grimoire approved these changes

At least 2 approving reviews are required to merge this pull request.

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
3 participants
@RunningLeon @grimoire @AllentDan