o Zh;@sddlmZmZmZddlmZmZmZmZer3ddl Z ddl m Z ddl Z ddl mZddl mZer.grid)r ) is_contiguousshapetorchZ empty_like float8_e4m3fn new_emptysizerr)rr!rrr*rr)r act_quant0s 2 r1 BLOCK_SIZE_M BLOCK_SIZE_N BLOCK_SIZE_K GROUP_SIZE_Mc6Cstjdd}t||}t||}||}||}||}t|||}|||}|||} ||td||}!| |td||}"td|}#||!dddf| |#dddf| }$||#dddf| |"dddf| }%||!|}&|"|}'||'|}(tj||ftjd})tdt||D]h}*tj|$|#dddf||*|kdd}+tj|%|#dddf||*|kdd},|*|}-|-| }.t|&|.|}/t|(|.|}0|)t |+|,|/dddf|0dddf7})|$|| 7}$|%|| 7}%q|j j tj kr|) tj }1n|j j tjkr%|) tj}1n|) tj}1||td|}2| |td|}3|||2dddf||3dddf}4|2dddf|k|3dddf|k@}5tj|4|1|5ddS)zTriton-accelerated function used to perform linear operations (dot product) on input tensors `A` and `B` with block-wise quantization, and store the result in output tensor `C`. rr Nr$g)maskother)r6)rrr&minrzerosrrangerdotrrZbfloat16rZfloat16r)6ABCAsBsMNKZgroup_nZgroup_kZ stride_amZ stride_akZ stride_bkZ stride_bnZ stride_cmZ stride_cnZ stride_As_mZ stride_As_kZ stride_Bs_kZ stride_Bs_nr2r3r4r5rZ num_pid_mZ num_pid_nZnum_pid_in_groupZgroup_idZ first_pid_mZ group_size_mZpid_mZpid_nZoffs_amZoffs_bnZoffs_kZa_ptrsZb_ptrsZAs_ptrsZoffs_bsnZBs_ptrsZ accumulatorkabZk_startZoffs_ksZa_sZb_scZoffs_cmZoffs_cnZc_ptrsZc_maskrrr_w8a8_block_fp8_matmul>sL %     ,,  ((0  ,(rHr<r=r?r@ output_dtypecst|dksJ|d|d}}|jd|jdksJ|jdd|jddkr/|s1Jt|jd||jdksAJ||jd|jdkrX|rX|jdksZJ|j\}t||jdkslJt|||jdksyJ|jddf} |j| |d} d} | krt} t | d} |} || dksJ|} fd d }t |||| |||||| d | d| d| d| d | d| d | d| d| d| | | d d | S)aThis function performs matrix multiplication with block-wise quantization. It takes two input tensors `A` and `B` with scales `As` and `Bs`. The output is returned in the specified `output_dtype`. Args: A: The input tensor, e.g., activation. B: The input tensor, e.g., weight. As: The per-token-group quantization scale for `A`. Bs: The per-block quantization scale for `B`. block_size: The block size for per-block quantization. It should be 2-dim, e.g., [128, 128]. output_dytpe: The dtype of the returned tensor. Returns: torch.Tensor: The result of matmul. rrr#Nr$r cs"t|dt|dfS)Nr2r3)r%r&)ZMETArArBrrr*s"z*w8a8_block_fp8_matmul_triton..grid)r2r3r4r5) lenr,r+r%r&r'ndimr/Znext_power_of_2rrHZstride)r<r=r?r@r!rIZblock_nZblock_krCZC_shaper>r2r4r3r*rrLrw8a8_block_fp8_matmul_tritons^(     rQinput_qweight_q input_scale weight_scalec Cs|jdkr|jn d|jd|jdf\}}}|jd} |d|} ||jdd} | |d} ||d} tj||| ftj|jd}t| D]k}||d}||d}t| D]X}||d}||d}| dd||f}|||||f}| dd||df}|||f}tj|| tj dtj|jd||d|}|dd||f|7<qZqH|||| }| |S)a Performs blocked matrix multiplication with FP8 quantized matrices. Args: input_q: Quantized input tensor with 1x128 block quantization weight_q: Quantized weight tensor with 128x128 block quantization input_scale: Scaling factors for input blocks weight_scale: Scaling factors for weight blocks block_size: Tuple of (M, N) for weight block dimensions output_dtype: Desired output dtype rJrr#rdeviceN)Zscale_aZscale_bZ out_dtype) rPr,viewr-r9rrXr:Z _scaled_mmtZtensorr)rRrSrTrUr!rIZ batch_sizeZseq_lenZ hidden_dim out_featuresZinput_reshapedZinput_scale_reshapedZnum_weight_blocks_mZnum_weight_blocks_noutputiZm_startZm_endjZn_startZn_endZ input_blockZ weight_blockZcurr_input_scaleZcurr_weight_scaleZ block_resultrrrw8a8_block_fp8_matmul_compiles>,             r_c sbeZdZejZ     ddedededee eefffdd Z d ej d ej fd d Z Z S) FP8LinearFNdynamic in_featuresr[biasr!c st||||_||_tjtj||tj |d|_ |j dkrJ||dd|d}||dd|d} ttj|| tj |d|_ n|dd||_||_|rdtt|j|_dS|dddS)NrWrJrweight_scale_invrc)super__init__rbr[r-nn Parameteremptyr`rweight element_sizerrdZregister_parameterr!activation_schemerc) selfrbr[rcrr!rXrlZscale_out_featuresZscale_in_features __class__rrrf)s   zFP8Linear.__init__inputr"c Cs|jdkrt||j|jStrtjj nd}t t|tj }| |j  t ||jd\}}t||j||j|j|jd}Wdn1sKwY||jdur^||j}|j|jdS)NrJcuda)rIr$)rjrkFZlinearrcrr-Z acceleratorZcurrent_acceleratortypegetattrrqrXr1r!rQrdrZ synchronizer)rmrpZ device_typeZtorch_accelerator_moduleZqinputscaler\rrrforwardKs&  zFP8Linear.forward)FNNNra)__name__ __module__ __qualname__r-r.rintboolrrrfTensorrv __classcell__rrrnrr`&s""r`Fc s|durg}|D]p\}}||t|tjr_||pgvr_d|tfdd|p-gDs_t#t|j |j |j du|j j |j j|j|jd|j|<d}Wdn1sZwYtt|dkrut||||||d\}}|d q ||fS) z%Replace Linear layers with FP8Linear.N.c3s|]}|vVqdS)Nr).0keyZcurrent_key_name_strrr usz+_replace_with_fp8_linear..)rbr[rcrXrrlr!Tr)has_been_replacedr#)Znamed_childrenappend isinstancergLinearjoinanyr r`rbr[rcrjrXrrlZweight_block_sizeZ_modulesrOlistchildren_replace_with_fp8_linearpop) modeltp_planmodules_to_not_convertZcurrent_key_namequantization_configrnamemodule_rrrrds<       rcCs\|durdgn|}|jdur||jtt|}t||j||d\}}|s,td|S)z:Helper function to replace model layers with FP8 versions.NZlm_head)rrrzYou are loading your model using fp8 but no linear modules were found in your model. Please double check your model architecture.)rextendrsetrZ_tp_planloggerwarning)rrrrrrrreplace_with_fp8_linears    r)r )NNNNF)NN)'typingrrrutilsrrrr r-Ztorch.nnrgr%Ztriton.languagelanguagerr rrZ accelerater Z get_loggerrwrZjitZ constexprrr|rzr1rHrrrQcompiler_rr`rrrrrrs     &  Z Q A@ 0