
    pi?                        U d Z ddlZddlZddlZddlmZmZmZmZm	Z	m
Z
 ddlZddlZddlZddlZej                  j!                         Zexr ej                  j%                         dk\  Zer ej(                  d      ndZer
 ed       Zn	 ed       Z ed       Z ed	       Z ed
       Z ed       Z ed       Z ed       Z ed       Z ed       Z ed       Z ed       Z  ed       Z! ed       Z" ed       Z# ed       Z$ ed       Z% ed       Z&d Z'd Z(d Z)d Z*d Z+d Z, ed       Z-e.e/d<    ed        Z0e.e/d!<    ed"       Z1e.e/d#<    ed$       Z2e.e/d%<   exr e Z3e.e/d&<    ed'       Z4e.e/d(<   d) Z5d* Z6d+ Z7d, Z8 ed-       Z9e.e/d.<    ed/       Z:e.e/d0<    ed1       Z;e.e/d2<    ed3       Z9e.e/d.<    ed4       Z<e.e/d5<   er 	 ddl=Z>e>j                  j                         Z@nd6Z@d6aBd7 ZCej                  d8        ZEej                  dId9       ZFej                  d:        ZGdJd;ZHd< ZId= ZJd> ZKd? ZLd@ ZMdA ZN eM       ZO eN       ZPdBej                  j                  dfdCZSdBej                  ej                  j                  dfdDZUdE ZVdF ZWdG ZXdH ZYesej                  j!                         rJ yy# eA$ r d6Z@d6ZY w xY w)Kz>This file is allowed to initialize CUDA context when imported.    N)LazyVal
TEST_NUMBATEST_WITH_ROCM	TEST_CUDA
IS_WINDOWSIS_MACOS   zcuda:0c                      t         S N)r        e/opt/services/ai/voice_agent/venv/lib/python3.12/site-packages/torch/testing/_internal/common_cuda.py<lambda>r      s     r   c                      t         xrB t        j                  j                  j	                  t        j
                  dt                    S )N      ?device)r   torchbackendscudnnis_acceptabletensorCUDA_DEVICEr   r   r   r   r      s0    !wu~~/C/C/Q/QRWR^R^_ajuRv/w r   c                  b    t         r(t        j                  j                  j	                         S dS )Nr   )
TEST_CUDNNr   r   r   versionr   r   r   r   r      s     zU^^%9%9%A%A%C WX r   c                      t         j                  j                  r<t        d t         j                  j                  j	                  d      d d D              S dS )Nc              3   2   K   | ]  }t        |        y wr   int).0vs     r   	<genexpr>z<lambda>.<locals>.<genexpr>   s     %Wc!f%W   .r	   r   r   )r   r   hiptuplesplitr   r   r   r   r      sC    [`[h[h[l[l%Wemm6G6G6M6Mc6RSUTU6V%W W rx r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N)      r   cudais_availableget_device_capabilityr   r   r   r   r      +    ejj557hEJJ<\<\<^bh<h r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N)   r   r-   r   r   r   r   r      r1   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N   r   r-   r   r   r   r   r      r1   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N)r6   r+   r-   r   r   r   r   r       r1   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N)   r   r-   r   r   r   r   r   !   r1   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S Nr9   	   r-   r   r   r   r   r   "   r1   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S N)r=   r   r-   r   r   r   r   r   #   r1   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S N)
   r   r-   r   r   r   r   r   $   +    uzz668jUZZ=]=]=_cj=j r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N)   r   r-   r   r   r   r   r   %   rC   r   c                      t         j                  j                         xrJ t         j                  j                         d   dk(  xr$ t         j                  j                         d   dkD  S )Nr   rB      r-   r   r   r   r   r   '   sW    %**113 @

8X8X8Z[\8]ac8c @jj668;a? r   c                      t         j                  j                         xr( t         j                  j                         dv xs t        S )N))r6   r	   )r9   r6   )r   r.   r/   r0   IS_THORr   r   r   r   r   )   s1    EJJ335}5::;[;[;]aq;q;|u| r   c                      t         j                  j                         xr! t         j                  j                         dk(  S r;   r-   r   r   r   r   r   *   +    %**113d

8X8X8Z^d8d r   c                      t         j                  j                         xr! t         j                  j                         dk(  S r?   r-   r   r   r   r   r   +   rK   r   c                      t         j                  j                         xr! t         j                  j                         dk(  S rA   r-   r   r   r   r   r   ,   s+    5::224f9Y9Y9[_f9f r   c                     t         j                  j                         syt         j                  j                  d      j                  }t
        j                  j                  d|      t        fd| D              S )NFr.   /PYTORCH_DEBUG_FLASH_ATTENTION_GCN_ARCH_OVERRIDEc              3   &   K   | ]  }|v  
 y wr   r   )r!   archeffective_archs     r   r#   z+evaluate_gfx_arch_within.<locals>.<genexpr>5   s     <$t~%<s   )	r   r.   r/   get_device_propertiesgcnArchNameosenvirongetany)	arch_listgcn_arch_namerR   s     @r   evaluate_gfx_arch_withinr[   .   sY    ::""$JJ44V<HHMZZ^^$UWdeN <)<<<r   c                      t        g d      S )N)gfx940gfx941gfx942gfx950r[   r   r   r   CDNA3OrLaterrb   7   s    #$LMMr   c                      t        ddg      S )Ngfx90ar_   ra   r   r   r   CDNA2OrLaterre   :   s    #Xx$899r   c                      t         r9g d} t        j                  j                  dd      dk7  r| g dz  } t	        |       S t
        rt         xr t        S y)Nrd   r_   gfx1100gfx1201r`   'TORCH_ROCM_AOTRITON_ENABLE_EXPERIMENTAL0gfx1101gfx1150gfx1151gfx1200F)r   rU   rV   rW   r[   r   r   SM80OrLaterrY   s    r   *evaluate_platform_supports_flash_attentionrs   =   sK    H	::>>CSISPEEI'	22~-+-r   c                      t         r9g d} t        j                  j                  dd      dk7  r| g dz  } t	        |       S t
        ryy)Nrg   rj   rk   rl   TF)r   rU   rV   rW   r[   r   rr   s    r   .evaluate_platform_supports_efficient_attentionru   G   sA    H	::>>CSISPEEI'	22r   c                  6    t          xr t        xr	 t        dk\  S )Ni_ )r   rq   TEST_CUDNN_VERSIONr   r   r   *evaluate_platform_supports_cudnn_attentionrx   Q   s    QKQ5G55PQr   c                      t               S r   )rs   r   r   r   r   r   T   
    :d:f r   !PLATFORM_SUPPORTS_FLASH_ATTENTIONc                      t               S r   )ru   r   r   r   r   r   U   s
    <j<l r   #PLATFORM_SUPPORTS_MEM_EFF_ATTENTIONc                      t               S r   )rx   r   r   r   r   r   V   rz   r   !PLATFORM_SUPPORTS_CUDNN_ATTENTIONc                  .    t         xs t        xs t        S r   )r{   r   r}   r   r   r   r   r   X   s    :[ ;V2S;V2U r   !PLATFORM_SUPPORTS_FUSED_ATTENTIONPLATFORM_SUPPORTS_FUSED_SDPAc                      t         xr t        S r   )r   rq   r   r   r   r   r   ^   s    y/H[ r   PLATFORM_SUPPORTS_BF16c                     t         j                  j                         rt         j                  j                  rmdg} t
        dk\  r| j                  dg       t
        dk\  r| j                  d       | D ]/  }|t         j                  j                  d      j                  v s/ y y	t        xs! t         j                  j                         dk(  S y	)
Ngfx94)r3   r,   gfx120)r3   r+   gfx95r   Tr<   F)r   r.   r/   r   r'   ROCM_VERSIONextendappendrS   rT   SM90OrLaterr0   archsrQ   s     r   evaluate_platform_supports_fp8r   `   s    zz ==IEv%hZ(v%W%  5::;;A>JJJ 
  N%**"B"B"D"NNr   c                  D   t         j                  j                         rt         j                  j                  rYdt         j
                  j                         vrydg} | D ]/  }|t         j                  j                  d      j                  v s/ y yt        xr t         S y)NUSE_FBGEMM_GENAIFr_   r   T)r   r.   r/   r   r'   
__config__showrS   rT   r   SM100OrLaterr   s     r   +evaluate_platform_supports_fp8_grouped_gemmr   o   s    zz ==!)9)9)>)>)@@JE  5::;;A>JJJ 
  3|#33r   c                      t         j                  j                         rUt         j                  j                  r5t
        dk\  r+dt         j                  j                  d      j                  v S yt        S y)Nr5   r`   r   F)	r   r.   r/   r   r'   r   rS   rT   r   r   r   r   "evaluate_platform_supports_mx_gemmr   |   sW    zz ==v%5::#C#CA#F#R#RRR   r   c                      t         j                  j                         rDt         j                  j                  s*dt         j
                  j                         v } | xr t        S y)Nr   F)r   r.   r/   r   r'   r   r   IS_SM100)built_with_fbgemm_genais    r   -evaluate_platform_supports_mxfp8_grouped_gemmr      sF    zz ):):"48H8H8M8M8O"O&383r   c                      t               S r   )r   r   r   r   r   r      s
    2T2V r   PLATFORM_SUPPORTS_MX_GEMMc                      t               S r   )r   r   r   r   r   r      s
    .L.N r   PLATFORM_SUPPORTS_FP8c                      t               S r   )r   r   r   r   r   r      s
    ;f;h r   "PLATFORM_SUPPORTS_FP8_GROUPED_GEMMc                      t         xr t        S r   )r   r   r   r   r   r   r      s    )2L r   c                      t               S r   )r   r   r   r   r   r      s
    =j=l r   $PLATFORM_SUPPORTS_MXFP8_GROUPED_GEMMFc                      t         sJ d       t        sIt        t        j                  j                               D ]  } t        j                  dd|          day y )Nz?CUDA must be available when calling initialize_cuda_context_rngrG   zcuda:r   T)r   __cuda_ctx_rng_initializedranger   r.   device_countrandn)is    r   initialize_cuda_context_rngr      sS    WWW9%uzz..01 	/AKKE!+.	/%)"	 &r   c               #     K   t         j                  j                  j                  j                  } 	 dt         j                  j                  j                  _        t         j                  j
                  j                  d d d d      5  d  d d d        | t         j                  j                  j                  _        y # 1 sw Y   3xY w# | t         j                  j                  j                  _        w xY ww)NFenabled	benchmarkdeterministic
allow_tf32r   r   r.   matmulr   r   flagsold_allow_tf32_matmuls    r   tf32_offr      s     !NN//66AAF05""-^^!!''TXej'k 		 1F""-	 	 1F""-5   /C9AC B?C *C9?CC +C66C9c              #     K   t         j                  j                  r3t        j                  j                  dd       }dt        j                  d<   t         j                  j                  j                  j                  }| j                  }	 dt         j                  j                  j                  _	        || _
        t         j                  j                  j                  d d d d      5  d  d d d        t         j                  j                  r'|t        j                  d<   nt        j                  d= |t         j                  j                  j                  _	        || _
        y # 1 sw Y   {xY w# t         j                  j                  r'|t        j                  d<   nt        j                  d= |t         j                  j                  j                  _	        || _
        w xY ww)NHIPBLASLT_ALLOW_TF321Tr   )r   r   r'   rU   rV   rW   r   r.   r   r   	precisionr   r   )selftf32_precisionhip_allow_tf32r   old_precisions        r   tf32_onr      s[    }}(>E-0

)*!NN//66AANNM'04""-'^^!!''TXei'j 		 ==)5C

12JJ560E""-&	 	 ==)5C

12JJ560E""-&s9   BG)AE3 (E'-E3 5A2G)'E0,E3 3A3G&&G)c               #     K   t         j                  j                  j                  j                  } 	 dt         j                  j                  j                  _        t         j                  j
                  j                  dddd      5  d ddd       | t         j                  j                  j                  _        y# 1 sw Y   3xY w# | t         j                  j                  j                  _        w xY ww)z
    Context manager to temporarily enable TF32 for CUDA operations.
    Restores the previous TF32 state after exiting the context.
    TNr   r   r   s    r   tf32_enabledr      s      "NN//66AAF04""-^^!!''D ( 
 	 	
 1F""-	 	
 1F""-r   c                 ,     d  fdfd}|S )Nc                 P    t               5   |        d d d        y # 1 sw Y   y xY wr   r   )r   function_calls     r   with_tf32_disabledz+tf32_on_and_off.<locals>.with_tf32_disabled   s!    Z 	O	 	 	s   %c                 V    t        |       5   |        d d d        y # 1 sw Y   y xY wr   )r   )r   r   r   s     r   with_tf32_enabledz*tf32_on_and_off.<locals>.with_tf32_enabled   s'    T>* 	O	 	 	s   (c                      t        j                         j                  }t        |j	                               t        j                          fd       }|S )Nc                     j                  t        |              t        j                  j	                         xr }dv r)|xr% t        j
                  d         j                  dk(  }dv r)|xr% d   t        j                  t        j                  hv }|r! d   fd        d   fd       y  di  y )Nr   r.   dtyper   c                        di S Nr   r   fkwargss   r   r   zCtf32_on_and_off.<locals>.wrapper.<locals>.wrapped.<locals>.<lambda>  s    1;v; r   c                        di S r   r   r   s   r   r   zCtf32_on_and_off.<locals>.wrapper.<locals>.wrapped.<locals>.<lambda>  s    !+f+ r   r   )	updatezipr   r.   is_tf32_supportedr   typefloat32	complex64)argsr   cond	arg_namesr   only_ifr   r   s    ` r   wrappedz1tf32_on_and_off.<locals>.wrapper.<locals>.wrapped  s    MM#i./:://1=gD6!OfX.>!?!D!D!N& UU]]EOO4T!T"6&>3FG!&.2EFFr   )inspect	signature
parametersr(   keys	functoolswraps)r   paramsr   r   r   r   r   s   `  @r   wrapperz tf32_on_and_off.<locals>.wrapper  sK    ""1%00&++-(				 
	 r   r   )r   r   r   r   r   s   `` @@r   tf32_on_and_offr      s    & Nr   c                 B     t        j                          fd       }|S )Nc                  T    t               5   | i |cd d d        S # 1 sw Y   y xY wr   r   )r   r   r   s     r   r   zwith_tf32_off.<locals>.wrapped  s*    Z 	&d%f%	& 	& 	&s   ')r   r   )r   r   s   ` r   with_tf32_offr     s%    __Q& & Nr   c                  ^   dt         j                  j                         vryt         j                  j                         j                  d      } t         j                  j                         | t	        d      z   d  j                  d      d   }t        d |j                  d      D              S )NMagmar&   zMagma 
r   c              3   2   K   | ]  }t        |        y wr   r   r!   xs     r   r#   z%_get_magma_version.<locals>.<genexpr>)  s     8AQ8r$   r%   )r   r   r   findlenr)   r(   )positionversion_strs     r   _get_magma_versionr   $  s    e&&++--$$&++H5H""'')(S]*B*CDJJ4PQRSK8!2!23!7888r   c                      t         j                  j                  yt        t         j                  j                        } t	        d | j                  d      D              S )Nr&   c              3   2   K   | ]  }t        |        y wr   r   r   s     r   r#   z*_get_torch_cuda_version.<locals>.<genexpr>/       9AQ9r$   r%   )r   r   r.   strr(   r)   )cuda_versions    r   _get_torch_cuda_versionr   +  sE    }}!u}}))*L9!3!3C!8999r   c                      t         rt        j                  j                  yt	        t        j                  j                        } | j                  dd      d   } t        d | j                  d      D              S )Nr&   -rG   maxsplitr   c              3   2   K   | ]  }t        |        y wr   r   r   s     r   r#   z*_get_torch_rocm_version.<locals>.<genexpr>6  r   r$   r%   r   r   r   r'   r   r)   r(   )rocm_versions    r   _get_torch_rocm_versionr  1  sa    U]]..6u}}(()L%%cA%6q9L9!3!3C!8999r   c                      t          S r   )r   r   r   r   !_check_cusparse_generic_availabler  8  s    r   c                     t         syt        j                  j                  syt	        t        j                  j                        } | j                  dd      d   } t        d | j                  d      D              }|d u xs |dk   S )	NFr   rG   r   r   c              3   2   K   | ]  }t        |        y wr   r   r   s     r   r#   z5_check_hipsparse_generic_available.<locals>.<genexpr>C  s     G!s1vGr$   r%   )r+   rG   r  )r  rocm_version_tuples     r   "_check_hipsparse_generic_availabler	  ;  s{    ==u}}(()L%%cA%6q9LG|/A/A#/FGG"d*I.@6.IJJr   r.   c                    t         j                  j                  t         j                  j                  dd      t         j                  j                  dd            j	                  |       }t         j                  j                  t         j                  j                  dd      t         j                  j                  dd            j	                  |       }t        j
                         5  t        |j                         |j                               D ]  \  }}|j                  |        	 d d d        ddi}||j                  |        ||j                         fi |} ||j                         fi |}	||||	fS # 1 sw Y   TxY w)Nr9   r   lrr   )
r   nn
SequentialLineartono_gradr   r   copy_r   )
r   optimizer_ctoroptimizer_kwargsmod_controlmod_scalingcsr   opt_controlopt_scalings
             r   !_create_scaling_models_optimizersr  K  s>    ((%%ehhooa&;UXX__QPQ=RSVV^dVeK((%%ehhooa&;UXX__QPQ=RSVV^dVeK	 ..0+2H2H2JK 	DAqGGAJ	 C[F#&' !7!7!9DVDK !7!7!9DVDK[+== s   -AFFc           	         t        j                  d||       t        j                  d||       ft        j                  d||       t        j                  d||       ft        j                  d||       t        j                  d||       ft        j                  d||       t        j                  d||       fg}t         j                  j                         j	                  |       }d}t        | ||      |||fz   S )N)r9   r9   )r   r   r	   )r   r  r  )r   r   r  MSELossr  r  )r   r   r  r  dataloss_fn	skip_iters          r   _create_scaling_caser   ]  s    [[uV<ekk&X]fl>mn[[uV<ekk&X]fl>mn[[uV<ekk&X]fl>mn[[uV<ekk&X]fl>mnpD
 hh ##F+GI,nGW	w	"# #r   c                 <    t         s| S t        j                  |       S r   )IS_SM89unittestexpectedFailurefuncs    r   xfailIfSM89r'  l  s    4BH$<$<T$BBr   c                 <    t         s| S t        j                  |       S r   )r   r#  r$  r%  s    r   xfailIfSM100OrLaterr)  o      #4G)A)A$)GGr   c                 <    t         s| S t        j                  |       S r   )SM120OrLaterr#  r$  r%  s    r   xfailIfSM120OrLaterr-  r  r*  r   c                 H    t         st        s| S t        j                  |       S r   )r   	IS_JETSONr#  r$  r%  s    r   xfailIfDistributedNotSupportedr0  u  s     I4RH4L4LT4RRr   )h㈵>)r1  T)Z__doc__r   r   
torch.cuda$torch.testing._internal.common_utilsr   r   r   r   r   r   r   
contextlibrU   r#  r.   is_initialized"CUDA_ALREADY_INITIALIZED_ON_IMPORTr   TEST_MULTIGPUr   r   r   rw   r   SM53OrLaterSM60OrLaterSM70OrLaterSM75OrLaterrq   SM89OrLaterr   r   r,  rI   r/  r"  IS_SM90r   r[   rb   re   rs   ru   rx   r{   bool__annotations__r}   r   r   r   r   r   r   r   r   r   r   r   r   
numba.cudanumbar/   TEST_NUMBA_CUDA	Exceptionr   r   contextmanagerr   r   r   r   r   r   r   r  r  r	  TEST_CUSPARSE_GENERICTEST_HIPSPARSE_GENERICoptimSGDr  floatr   r'  r)  r-  r0  r   r   r   <module>rK     s   F    u u   	  &+ZZ%>%>%@ " <ejj5571<(1lell8$t*+JwxJXY xyhihihihihihihijkjk
 @ A}~	
d
e
d
efg=N:R +22f*g !4 g,34l,m #T m*12f*g !4 g*1 3V +W !4 W &/%E~3E d E&'HI  I #**V"W 4 W%&NO t O+23h+i "D i")*L"M 4 M-45l-m $d m**113
 O # * F F ' '* F FXF9::	K :; ;=  .4EKKOOfj >$ !'ekk%++//lp #CHHS *zz((**** *I  
s   K 	KK