Skip to content

Remove custom fp4/fp8 classes - #3212

Merged
zcbenz merged 1 commit into
ml-explore:mainfrom
zcbenz:remove-cuda-fp4-h
Mar 10, 2026
Merged

zcbenz merged 1 commit into
ml-explore:mainfrom
zcbenz:remove-cuda-fp4-h

Conversation

@zcbenz

@zcbenz zcbenz commented Mar 6, 2026

Copy link
Copy Markdown
Member

We have some custom fp4/fp8 classes in cuda/quantized/cuda_fp4.h for building with not so recent CUDA toolkits (like 12.6), in which case there is no builtin equivalents.

This PR replaces them with the equivalents in CUTLASS, which generates same code with CUDA's builtin __nv_fp8_xxs/__nv_fp8_xxx and provides decent fallbacks for old CUDA toolkits.

The main benefit is we can just rely on cutlass::NumericArrayConverter to do vectorized number conversions, instead of dealing with various CUDA types.

@nastya236

Copy link
Copy Markdown
Collaborator

Thank you, looks like a great alternative to custom header. Do you mind if I put myself as a reviewer?

@zcbenz

zcbenz commented Mar 6, 2026

Copy link
Copy Markdown
Member Author

It is always a good thing having more people reviewing this 🙏

@nastya236
nastya236 self-requested a review March 6, 2026 19:02

@nastya236 nastya236 left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Looks awesome, thank you for adding this! Much cleaner than with custom classes.

@zcbenz
zcbenz merged commit 8d022bc into ml-explore:main Mar 10, 2026
29 of 32 checks passed
@zcbenz
zcbenz deleted the remove-cuda-fp4-h branch March 10, 2026 07:08
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants