o
    TÃiÂ  ã                   @   s„   d dl Z d dlZd dlmZ ejdejdejdejdejdejf
dd„ƒZejdejdejdejdejdejf
d	d
„ƒZdd„ Z	dS )é    NÚBLOCK_SIZE_MÚBLOCK_SIZE_NÚBLOCK_SIZE_KÚGROUP_SIZE_MÚquantization_group_sizec           +      C   sÐ  t jdd}t  ||¡}t  ||¡}|| }|| }|| }t|| |ƒ}||| |  }|| | }|| t  d|¡ | }|| t  d|¡ | }t  d|¡}| |d d …d f | |d d d …f |   }||d d …d f |	 |d d d …f |
   }|d d …d f |	|  || |  } t j||ft jd}!tdt  ||¡ƒD ]y}"t j||d d d …f ||"|  k dd}#t j||d d …d f |k dd}$t  ||  |"| |	 |  ¡}%|$d@  	t j
¡d> }&|&|$d@  	t j
¡d	> B }&|&d
  	t j
¡}&|&j	t jdd 	t j¡|%  	t j¡}&||| 7 }|||	 7 }|!t  |#|&¡7 }!q¢|! 	t j¡}'|| t  d|¡ }(|| t  d|¡ })|||(d d …d f   ||)d d d …f   }*t j|*|'|(d d …d f |k |)d d d …f |k @ d d S )Nr   ©Úaxis©Údtypeç        ©ÚmaskÚotheré€   é   é   é   i <  T©Úbitcast©r   )ÚtlÚ
program_idÚcdivÚminÚarangeÚzerosÚfloat32ÚrangeÚloadÚtoÚuint16Úbfloat16ÚdotÚstore)+Úinp_ptrÚ
weight_ptrÚout_ptrÚ	scale_ptrÚMÚNÚKÚ	stride_amÚ	stride_akÚ	stride_bkÚ	stride_bnÚ	stride_cmÚ	stride_cnr   r   r   r   r   ÚpidÚ	num_pid_mÚ	num_pid_nÚnum_pid_in_groupÚgroup_idÚfirst_pid_mÚgroup_size_mÚpid_mÚpid_nÚoffs_amÚoffs_bnÚoffs_kÚinp_dataÚweight_dataÚweight_ptrs_offsetÚaccumulatorÚkÚinpÚweightÚscaleÚwÚoutÚoffs_cmÚoffs_cnÚout_data© rJ   ú^/home/ubuntu/.local/lib/python3.10/site-packages/deepspeed/ops/fp_quantizer/fp8_gemm_triton.pyÚmatmul_kernel_fp8_bf16   sD   ,,
ÿ( $,8rL   c           +      C   sú  t jdd}t  ||¡}t  ||¡}|| }|| }|| }t|| |ƒ}||| |  }|| | }|| t  d|¡ | }|| t  d|¡ | }t  d|¡}| |d d …d f | |d d d …f |   }||d d …d f |	 |d d d …f |
   }|d d …d f |	|  || |  } t j||d d …d f |k dd}!t  ||  ¡}"t j||ft jd}#tdt  ||¡ƒD ]w}$t j||d d d …f ||$|  k dd}%|!d@ d> |!d@ d	> B  	t j
¡}&|&d
  	t j
¡}&|&j	t jdd|"  	t j¡}&||| 7 }|||	 7 }t j||d d …d f ||$d |  k dd}!t  || |$d | |	 |   ¡}"|#t  |%|&¡7 }#q¹|# 	t j¡}'|| t  d|¡ }(|| t  d|¡ })|||(d d …d f   ||)d d d …f   }*t j|*|'|(d d …d f |k |)d d d …f |k @ d d S )Nr   r   r   r   r	   r   r   r   é   i    Tr   é   r   )r   r   r   r   r   r   r   r   r   r   r    Úfloat16r"   r#   )+r$   r%   r&   r'   r(   r)   r*   r+   r,   r-   r.   r/   r0   r   r   r   r   r   r1   r2   r3   r4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   rC   rD   r@   rA   rB   rE   rF   rG   rH   rI   rJ   rJ   rK   Úmatmul_kernel_fp8_fp16B   sJ   ,,
ÿ ( ,ÿ,8rP   c                    s<  | j d |j d ksJ d| j › d|j › dƒ‚| j \‰ }|j \}‰tjˆ ˆf| j| jd}ˆ dkr3dnˆ dkr9dnˆ d	kr?d	nd
}d	}td	|ƒ}d}	d}
d}ˆ dkr_d}d
}td
|ƒ}d}
d}‡ ‡fdd„}| jtjkrntnt}|| | |||ˆ ˆ||  	d¡|  	d¡| 	d¡| 	d¡| 	d¡| 	d¡|||||	|
|d |S )NrN   r   z Incompatible dimensions (input: z
, weight: ú))Údevicer
   é   é    é@   r   r   r   é   é   c                    s"   t  ˆ | d ¡t  ˆ| d ¡ fS )Nr   r   )Útritonr   )ÚMETA©r(   r)   rJ   rK   Ú<lambda>   s   " z#matmul_fp8_triton.<locals>.<lambda>)r   r   r   r   r   Ú
num_stagesÚ	num_warps)
ÚshapeÚtorchÚemptyrR   r
   Úmaxr!   rL   rP   Ústride)rB   rC   rD   r   r*   rF   r   r   r   r   r\   r]   ÚgridÚkernelrJ   rZ   rK   Úmatmul_fp8_tritonv   sT   ÿ

(

íre   )
r_   rX   Útriton.languageÚlanguager   ÚjitÚ	constexprrL   rP   re   rJ   rJ   rJ   rK   Ú<module>   s8   ÿÿþþý/ÿÿþþý3