hebiao064
/ per_token_quant_fp8.cu
Created
March 11, 2025 06:34
[v2] Per Token FP8 Quant Kernel handling weird hidden_dim % 8 != 0 and do cast load when hidden_dim % 8 = 0
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
| #include <ATen/cuda/CUDAContext.h> | |
| #include <cmath> | |
| #include <cub/block/block_reduce.cuh> | |
| #include <flashinfer/vec_dtypes.cuh> | |
| #include "utils.h" | |
| template <typename T> | |
| __global__ void per_token_quant_fp8_kernel( |
hebiao064
/ per_token_quant_fp8.cu
Last active
March 11, 2025 06:33
[v1] Per Token FP8 Quant Kernel handling weird hidden_dim % 8 != 0 and do cast load when hidden_dim % 8 = 0
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
| #include <ATen/cuda/CUDAContext.h> | |
| #include <cmath> | |
| #include <cub/block/block_reduce.cuh> | |
| #include <flashinfer/vec_dtypes.cuh> | |
| #include "utils.h" | |
| template <typename T> | |
| __global__ void per_token_quant_fp8_kernel( |