
    ^(j                        d dl Z d dlZd dlmZ 	 d dlZd dlmZmZm	Z
mZmZmZmZ dZe j"                  j$                  ej&                  j)                  d       nj e ee ee j"                  j$                        j3                  d                  Zedk  r-ej&                  j)                  d        ej6                  d       ej8                  r#	 d dlZ ej<                  d	ej>                         nej&                  j)                  d        ejF                         jI                         D ]  \  Z%Z& ej<                  de% de&         	 dZ'er		 d dlm(Z) dZ'e's
 G d d      Z)d dl*Z+ G d de
      Z, G d  d!e)      Z( G d" d#e      Z G d$ d%e,      Z- G d& d'e,      Z.e-Z	 ed(e	        ed%e-        ed'e.        ed#e       e'r	 ed!e(       e j^                  d)d*hd%d+e j`                  d)d*hd'd+e jb                  h d,d#d-d.d/Z2e're j^                  d)d*hd!d0d.e2d1<   g d2Z3y# e $ r<Z! ejD                  d
e! d       ej&                  j)                  d       Y dZ![!UdZ![!ww xY w# e $ rGZ! ejD                  de! d       dZ G d d      Z G d d      Z
 G d d      Zd Zd ZY dZ![!ddZ![!ww xY w# e $ r  ej6                  d       Y yw xY w)3    N)args)QuantizedTensorQuantizedLayoutTensorCoreFP8LayoutTensorCoreNVFP4Layoutregister_layout_opregister_layout_classget_layout_classTcuda.)   zPWARNING: You need pytorch with cu130 or higher to use optimized CUDA operations.z7Found triton %s. Enabling comfy-kitchen triton backend.z Failed to import triton, Error: z9, the comfy-kitchen triton backend will not be available.tritonzFound comfy_kitchen backend z: z'Failed to import comfy_kitchen, Error: z,, fp8 and fp4 support will not be available.Fc                       e Zd Zy)r   N__name__
__module____qualname__     9/Users/danicosta/Desktop/Flux2/ComfyUI/comfy/quant_ops.pyr   r   )       r   r   c                       e Zd Zy)_CKFp8LayoutNr   r   r   r   r   r   ,   r   r   r   c                       e Zd Zy)_CKNvfp4LayoutNr   r   r   r   r   r   /   r   r   r   c                      y Nr   )nameclss     r   r	   r	   2   s    r   c                      y r   r   )r   s    r   r
   r
   5   s    r   )TensorCoreMXFP8LayoutzBcomfy_kitchen does not support MXFP8, please update comfy_kitchen.c                       e Zd Zy)_CKMxfp8LayoutNr   r   r   r   r#   r#   A   r   r   r#   c                   "    e Zd ZdZedd       Zy)_TensorCoreFP8LayoutBaseNc                    | j                   t        | j                   d      |j                  }t	        |j
                        }t        |t              r|dk(  rt        j                  |j                               j                  t        j                        t        j                  | j                         j                  z  }|j                  t        j                  t        j                  fvrQt        j                  |j                        }dt        j                   d|z  |j"                  |j                        z  }|0t        j$                  d|j&                  t        j                        }t        |t        j(                        s0t        j*                  ||j&                  t        j                        }|dkD  rr|r"|d|z  j                  |j                        z  }n!|d|z  j                  |j                        z  }t,        j.                  j1                  || j                   |	      }n!t3        j4                  ||| j                         }| j7                  |j/                         ||
      }	||	fS )Nz must define FP8_DTYPErecalculate)dtypeg      ?)minmaxr   devicer(   r   )r(   seedscale
orig_dtype
orig_shape)	FP8_DTYPENotImplementedErrorr   r(   tupleshape
isinstancestrtorchamaxabstofloat32finfor*   bfloat16clampr)   onesr,   Tensortensorcomfyfloatstochastic_roundingckquantize_per_tensor_fp8Params)
r   rB   r/   rE   inplace_opsr0   r1   tensor_infoqdataparamss
             r   quantizez!_TensorCoreFP8LayoutBase.quantizeM   s   == %6L&MNN\\
6<<(
eS!e}&<JJvzz|,//emm/Du{{SVS`S`GaGeGeeE||EMM5>>#BB#kk&,,7u{{C%KkooS^SbSbcc=JJr&--u}}ME%.LLv}}EMMRE"3;**6<<883;"2"26<<"@@KK33F#--Vi3jE..vucmmLE%++-JS]^f}r   Nr   F)r   r   r   r2   classmethodrM   r   r   r   r%   r%   J   s    I r   r%   c                       e Zd Zedd       Zy)r!   Nc                    |j                         dk7  rt        d|j                          d      |j                  }t        |j                        }| j                  |      }||k7  }|dkD  r&t        j                  j                  |||      \  }	}
nt        j                  ||      \  }	}
| j                  |
||      }|	|fS )N   zMXFP8 requires 2D tensor, got Dr   )pad_32xr-   )rT   r.   )dim
ValueErrorr(   r4   r5   get_padded_shaperC   rD   (stochastic_round_quantize_mxfp8_by_blockrF   quantize_mxfp8rH   r   rB   r/   rE   rI   r0   r1   padded_shapeneeds_paddingrK   block_scalerL   s               r   rM   zTensorCoreMXFP8Layout.quantizen   s    ::<1=fjjl^1MNN\\
6<<(
++J7$
2"!&!U!UV\fs  {N!U  "OE;!#!2!26=!QE;!!  

 f}r   rN   r   r   r   rO   rM   r   r   r   r!   r!   m   s     r   r!   c                       e Zd Zedd       Zy)r   Nc                 $   |j                         dk7  rt        d|j                          d      |j                  }t        |j                        }|t        |t              r^|dk(  rYt        j                  |j                               t        j                  j                  t        j                  j                  z  z  }t        |t        j                        st        j                  |      }|j!                  |j"                  t        j$                        }| j'                  |      }||k7  }|dkD  r't(        j*                  j-                  ||||      \  }	}
nt        j.                  |||      \  }	}
| j1                  ||||
	      }|	|fS )
NrR   zNVFP4 requires 2D tensor, got rS   r'   r+   r   )pad_16xr-   )ra   )r/   r0   r1   r]   )rU   rV   r(   r4   r5   r6   r7   r8   r9   r:   rF   float_utilsF8_E4M3_MAXF4_E2M1_MAXrA   rB   r;   r,   r<   rW   rC   rD   (stochastic_round_quantize_nvfp4_by_blockquantize_nvfp4rH   rZ   s               r   rM   zTensorCoreNVFP4Layout.quantize   sR   ::<1=fjjl^1MNN\\
6<<(
=Zs38NJJvzz|,0J0JR^^MgMg0ghE%.LL'EU]]C++J7$
2"!&!U!UV\^cmz  BU!U  "VE;!#!2!265-!XE;!!#	  
 f}r   rN   r^   r   r   r   r   r      s     r   r   c                   $    e Zd Zej                  Zy)TensorCoreFP8E4M3LayoutN)r   r   r   r8   float8_e4m3fnr2   r   r   r   rh   rh      s    ##Ir   rh   c                   $    e Zd Zej                  Zy)TensorCoreFP8E5M2LayoutN)r   r   r   r8   float8_e5m2r2   r   r   r   rk   rk      s    !!Ir   rk   r   weight_scaleinput_scale)	storage_t
parameterscomfy_tensor_layout>   rn   rm   weight_scale_2   )ro   rp   rq   
group_size)ri   rl   nvfp4    mxfp8)r   r   r   rh   rk   r   QUANT_ALGOSr   )4r8   loggingcomfy.cli_argsr   comfy_kitchenrF   comfy_kitchen.tensorr   r   r   r   r   r   r   r	   r
   _CK_AVAILABLEversionr   registrydisabler4   mapintr7   splitcuda_versionwarningenable_triton_backendr   info__version__ImportErroreerrorlist_backendsitemskv_CK_MXFP8_AVAILABLEr!   r#   comfy.floatrC   r%   rh   rk   ri   rl   uint8rx   __all__r   r   r   <module>r      s     0   M}}!
F#Sc%--*<*<&=&C&CC&HIJ%KK'GOOno!!	*GLLRTZTfTfg
 	H%   "((*13A3b<= +*  ^P"    |  FN 2N B$6 $"6 "
 .  +-@ A /1H I /1H I -/D E13HI ((%}58 &&%}58 [[G6	& ((%}56	K	A  	*GMM<QC?xyzKK))	*  GMM;A3>jklM    !0  ^\]^sU   C I+ !H' 2AI+ J: 'I(,1I#I+ #I((I+ +J70<J22J7:KK