[Refactor] Clean up JIT kernel utilites

[DarkSharpness]

Motivation

Currently, the code structure of our JIT kernel is terrible. Also, there're many diffusion PRs that may rely on some common utilities of our JIT kernel #14717, #16404

Modifications

1. Major clean up JIT code structure, make them more organized and prepare for future AMD support
2. Rewrite and simplify per_tensor_quant cc @BBuf
3. Provide a simple norm template that supports hidden_size=64, 128, 256 * n. cc @jianyingzhu @attack204 (see rmsnorm.cuh and qknorm.cuh for a quick example)

Future works:

1. More documents to be added. Our C++ code now lacks docs explaining their design
2. More utilities to be added (e.g. add operator overload for our AlignedVector type)

Accuracy Tests

Benchmarking and Profiling

Checklist

• Format your code according to the Format code with pre-commit.
• Add unit tests according to the Run and add unit tests.
• Update documentation according to Write documentations.
• Provide accuracy and speed benchmark results according to Test the accuracy and Benchmark the speed.
• Follow the SGLang code style guidance.

Review Process

1. Ping Merge Oncalls to start the PR flow. See the PR Merge Process.
2. Get approvals from CODEOWNERS and other reviewers.
3. Trigger CI tests with comments or contact authorized users to do so.
   • /tag-run-ci-label, /rerun-failed-ci, /tag-and-rerun-ci
4. After green CI and required approvals, ask Merge Oncalls to merge.

[gemini-code-assist]

Summary of Changes

Hello @DarkSharpness, I'm Gemini Code Assist¹! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!

This pull request significantly refactors the JIT kernel utilities to improve code organization, reusability, and prepare for future hardware support. By introducing a modular set of CUDA primitives and a flexible norm template, it streamlines the development of new kernels and simplifies existing ones, such as per-tensor quantization and various normalization layers. The changes aim to enhance the maintainability and extensibility of the JIT kernel infrastructure.

Highlights

• JIT Code Structure Refactoring: The internal JIT kernel code structure has undergone a major cleanup, reorganizing utilities into more modular and reusable components (e.g., atomic, cta, math, tile, type, norm specific implementations). This prepares the codebase for future extensions, such as AMD support.
• Simplified Per-Tensor Quantization: The per_tensor_quant implementation for FP8 has been rewritten and simplified, leveraging the new internal utility modules for improved clarity and maintainability. This includes updates to its benchmark and test files.
• New Norm Template and RMSNorm Implementation: A generic norm template has been introduced (sgl_kernel/impl/norm.cuh) that supports various hidden sizes (e.g., 64, 128, 256, and multiples of 256 up to 8192). This template is utilized by a newly added RMSNorm kernel (rmsnorm.cuh) and has been integrated into the existing QKNorm kernel (qknorm.cuh), along with new benchmarks and tests for RMSNorm.
• Standardized CUDA Type Aliases and Utilities: Common CUDA types (__half, __nv_bfloat16, __nv_fp8_e4m3, etc.) are now aliased to fp16_t, bf16_t, fp8_e4m3_t, etc., respectively, within sgl_kernel/utils.cuh. New dtype_trait and cast utilities simplify type handling and conversions across the kernel.

🧠 New Feature in Public Preview: You can now enable Memory to help Gemini Code Assist learn from your team's feedback. This makes future code reviews more consistent and personalized to your project's style. Click here to enable Memory in your admin console.

Using Gemini Code Assist

The full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips.

Invoking Gemini

You can request assistance from Gemini at any point by creating a comment using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands on the current page.

Feature	Command	Description
Code Review	/gemini review	Performs a code review for the current pull request in its current state.
Pull Request Summary	/gemini summary	Provides a summary of the current pull request in its current state.
Comment	@gemini-code-assist	Responds in comments when explicitly tagged, both in pull request comments and review comments.
Help	/gemini help	Displays a list of available commands.

Customization

To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a .gemini/ folder in the base of the repository. Detailed instructions can be found here.

Limitations & Feedback

Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with 👍 and 👎 on @gemini-code-assist comments. If you're interested in giving your feedback about your experience with Gemini Code Assist for Github and other Google products, sign up here.

You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the Gemini Code Assist IDE Extension.

Footnotes

1. Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configure Gemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution. ↩

[gemini-code-assist]

Code Review

This pull request is a significant and well-executed refactoring of the JIT kernel utilities. It introduces several new abstractions (tile, type, math, atomic, cta) that greatly improve code organization, readability, and maintainability. The existing per_tensor_quant and qknorm kernels have been cleanly rewritten to leverage these new components, and a new rmsnorm kernel has been added along with corresponding benchmarks and tests. The overall quality of the changes is high. I've identified a couple of minor issues: a potential for silent failure in the new rmsnorm kernel and a naming inconsistency in a new test file. Addressing these will further solidify this excellent contribution.

[DarkSharpness]

/tag-and-rerun-ci

[BBuf]

Can you post per_tensor_quant and qk_norm kernel's benchmark with main and pr?

[BBuf]

Is the RMSNorm added here necessary? @HydraQYH is about to finish a high-performance RMSNorm kernel, preparing to put it in the JIT kernel, there might be some conflicts. Code is here

[DarkSharpness]

It is just a demo of our cta norm template, and we don't plan to land it in deployment since there's no much perf gain against flashinfer baseline.

[DarkSharpness]

Feel free to update the norm template if it is not fastest.

[BBuf]

OK.

[DarkSharpness]

perf of per_tensor_quant on H200 cc @BBuf
This implemenation can be further optimized on B200 by leveraging 256bit load/store instructions

After

   element_count   SGL Kernel
0         1024.0     9.986603
1         2048.0    19.749636
2         4096.0    39.195474
3         8192.0    75.126296
4        16384.0   143.118794
5        32768.0   279.161743
6        65536.0   696.613938
7       131072.0  1386.467563
8       262144.0  2768.661393

Before:

0         1024.0     9.491280
1         2048.0    18.408333
2         4096.0    37.905891
3         8192.0    79.710076
4        16384.0   151.951233
5        32768.0   292.911992
6        65536.0   717.469187
7       131072.0  1423.264952
8       262144.0  2836.233974

Slightly worse performance for smaller batch size, and slightly better for larger batch size. I guess this could be further optimized.

[DarkSharpness]

the only ci break will be fixed by #16912

[BBuf]

/tag-and-rerun-ci

[DarkSharpness]

/rerun-failed-ci