o iS3 @sddlZddlmZddlZdadadadadada ddZ ddej dej d ej d ej d e d ej f d dZ dejfdej d e dejd eej ej ffddZdejfdej d e d eej ej ffddZ ddej dej d e d ej fddZdS)N)TupleFcs^trdSdaddlm}|sdadSdaddlddlmddlmfdddD}j|gd d j d j d j d j dj dj dj f fdd t j j ddd d(dt jdt jdt jdt jdtdt jf fdd }|jd(dd }|aj d!j ffd"d# }|aj d!j ffd$d% }|aj d!j ffd&d' }|adS))NTr) has_tritonFConfigc s4g|]}dD]}dD] }||d|ddq qqS)) @)) BLOCK_SIZE_M BLOCK_SIZE_N) num_stages num_warps).0block_mblock_nrrrY/home/ubuntu/.local/lib/python3.10/site-packages/torchao/kernel/blockwise_quantization.py /s z%_lazy_init_triton..)rrr)NKM_BUCKET BLOCK_SIZE_K)configskeyrrrr rrc s<jdd} jdd} || }| | d| |}| | d| |}d| }||dddf||dddf}||dddf||dddf}|||}||| |}j| | fjd}t|D]\}j||dddf||| kdd}j||dddf||| kdd}|}|}||||dddf|dddf7}|| 7}|| 7}|d7}|d7}qt||j j }| | d| }| | d| }||dddf||dddf}|dddf|k|dddf|k@}j |||ddS)Nraxisdtypeg)maskotherr$) program_idcdivarangezerosfloat32rangeloaddottor# element_tystore)a_ptrb_ptrc_ptra_s_ptrb_s_ptrMrrrr rrpid_mpid_nkoffs_moffs_noffs_ka_ptrsb_ptrsa_s_ptrsb_s_ptrs accumulatoriaba_sb_scc_ptrsr$tlrrblockwise_fp8_gemm_kernel:s6    ((  ((  0 ((z4_lazy_init_triton..blockwise_fp8_gemm_kernelzao::blockwise_fp8_gemmr) mutates_argsrrDrFrErG block_sizereturnc s|sJ|s J|sJ|sJ|d}|||dtt}|jg|ddRdtji}fdd}|||||||||d |S)Nrr#cs |d|dfS)Nr rr()METAr7rtritonrr{zC_lazy_init_triton.._blockwise_fp8_gemm_op..)r) is_contiguoussizenumelmathceillog2 new_emptytorchbfloat16) rDrFrErGrNrrrHgrid)rLrT)r7rr_blockwise_fp8_gemm_opjs       (z1_lazy_init_triton.._blockwise_fp8_gemm_opcSs6|d}|jg|dd|Rdtji}|S)NrrPr#)rXr]r^r_)rDrFrErGrNrrHrrr_s (z_lazy_init_triton.._ BLOCK_SIZEc sjdd}||d|}||j}|d}||}||jj} ||| |||dS)a! Quantizes the input tensor `x_ptr` and stores the result in `y_ptr` and the scaling factor in `s_ptr`. Args: x_ptr (triton.Pointer): Pointer to the input tensor. y_ptr (triton.Pointer): Pointer to the output tensor where quantized values will be stored. s_ptr (triton.Pointer): Pointer to the output tensor where scaling factors will be stored. BLOCK_SIZE (tl.constexpr): The size of the block to be processed by each program instance. Returns: None rr|@N) r'r)r-r/r+maxabsr#r0r1) x_ptry_ptrs_ptrrcpidoffsxsyrJrr$_fp8_blockwise_act_quant_kernel_impls z?_lazy_init_triton.._fp8_blockwise_act_quant_kernel_implcs jdd}jdd}||}||d|} ||d|} | dddf|| dddf} | dddf|k| dddf|k@} j|| | dj} | d}| |}||jj }j || || d |||||dS)aj Quantizes the input tensor `x_ptr` and stores the result in `y_ptr` and the scaling factors in `s_ptr`. Args: x_ptr (tl.pointer): Pointer to the input tensor. y_ptr (tl.pointer): Pointer to the output tensor where quantized values will be stored. s_ptr (tl.pointer): Pointer to the output tensor where scaling factors will be stored. M (int): Number of rows in the weight matrix. N (int): Number of columns in the weight matrix. BLOCK_SIZE (tl.constexpr): The size of the block to be processed by each program instance. rrr!Nr&rd) r'r(r)r-r/r+rerfr#r0r1)rgrhrir7rrcr8r9nr;r<rkr$rlrmrnrJrr'_fp8_blockwise_weight_quant_kernel_impls   $(zB_lazy_init_triton.._fp8_blockwise_weight_quant_kernel_implcsjdd}jdd}||}||d|} ||d|} | dddf|| dddf} | dddf|k| dddf|k@} j|| | dj} ||||}| |}j|| || ddS)a% Dequantizes weights using the provided scaling factors and stores the result. Args: x_ptr (tl.pointer): Pointer to the quantized weights. s_ptr (tl.pointer): Pointer to the scaling factors. y_ptr (tl.pointer): Pointer to the output buffer for dequantized weights. M (int): Number of rows in the weight matrix. N (int): Number of columns in the weight matrix. BLOCK_SIZE (tl.constexpr): Size of the block for tiling. Returns: None rrr!Nr&)r'r(r)r-r/r+r1)rgrirhr7rrcr8r9rpr;r<rkr$rlrmrnrJrr)_fp8_blockwise_weight_dequant_kernel_impls   $(zD_lazy_init_triton.._fp8_blockwise_weight_dequant_kernel_implr)_triton_initializedtorch.utils._tritonr_triton_availablerTtriton.languagelanguagerautotunejit constexprr^library custom_opTensorint register_fake_blockwise_fp8_gemm_impl_fp8_blockwise_act_quant_kernel"_fp8_blockwise_weight_quant_kernel$_fp8_blockwise_weight_dequant_kernel)rfp8_gemm_configsrarbrorqrrr)rrLrKrTr_lazy_init_tritons         , rrrDrFrErGrNrOcCs"tts tdt|||||S)Nunsupported without triton)rrvAssertionErrorr)rDrFrErGrNrrrblockwise_fp8_gemmsrrlr#cstts tdddlsJdd|dks&Jd|d|tjtjfvs2Jdtj |d }j gddd|Rd tj i}fd d }t ||||d ||fS)a Quantizes the input tensor `x` using block-wise quantization with block size being BLOCK_SIZEx1. Args: x (torch.Tensor): The input tensor to be quantized. Must be contiguous and its last dimension size must be divisible by `block_size`. block_size (int, optional): The size of the blocks to be used for quantization. Default is 128. dtype (torch.dtype, optional): The dtype to use for the quantized tensor. Default is `torch.float8_e4m3fn`. Returns: Tuple[torch.Tensor, torch.Tensor]: A tuple containing: - The quantized tensor with dtype `dtype`. - A tensor of scaling factors with dtype `torch.float32`. rrNInput tensor must be contiguousrPz@Last dimension size must be divisible by block_size (block_size=)6dtype must be torch.float8_e4m3fn or torch.float8_e5m2r"r#cs|dfSNrc)r(rYmetarTrlrrrUsz)fp8_blockwise_act_quant..rc) rrvrrTrWrXr^ float8_e4m3fn float8_e5m2 empty_liker]r+rrlrNr#rnrmr`rrrfp8_blockwise_act_quants& 2rcstts tdddl|sJd|dksJd|d|dkr1|d|dks9Jd|d |tjtj fvsEJd |\tj ||d }|j ||tj d }fd d }t |||||d||fS)aa Quantizes the given weight tensor using block-wise quantization with block size being BLOCK_SIZExBLOCK_SIZE. Args: x (torch.Tensor): The weight tensor to be quantized. block_size (int, optional): The block size to use for quantization. Defaults to 128. dtype (torch.dtype, optional): The dtype to use for the quantized tensor. Defaults to `torch.float8_e4m3fn`. Returns: Tuple[torch.Tensor, torch.Tensor]: A tuple containing: - The quantized weight tensor with dtype `dtype`. - A tensor of scaling factors with dtype `torch.float32`. rrNrz#Input tensor must have 2 dimensionsr!zABoth dimensions of x must be divisible by block_size (block_size=rrr"c |d|dfSrrQrrSrrrU<rVz,fp8_blockwise_weight_quant..r)rrvrrTrWdimrXr^rrrr]r+rrrrSrfp8_blockwise_weight_quants*&  rrmcstts tdddl|r|sJd|dkr%|dks)Jd|\tj|t d}fdd }t |||||d |S) a Dequantizes the given weight tensor using the provided scale tensor. Args: x (torch.Tensor): The quantized weight tensor of shape (M, N). s (torch.Tensor): The scale tensor of shape (M, N). block_size (int, optional): The block size to use for dequantization. Defaults to 128. Returns: torch.Tensor: The dequantized weight tensor of the same shape as `x`. Raises: AssertionError: If `x` or `s` are not contiguous or if their dimensions are not 2. rrNz Input tensors must be contiguousrz$Input tensors must have 2 dimensionsr"crrrQrrSrrrU_rVz.fp8_blockwise_weight_dequant..r) rrvrrTrWrrXr^rget_default_dtyper)rlrmrNrnr`rrSrfp8_blockwise_weight_dequantDs  rrs)rZtypingrr^rtrvrrrrrr~rrrr#rrrrrrrsj X  ' +