
    Th9                     t   % S r SSKrSSKrSSKrSSKJrJrJrJrJ	r	J
r
  SSKrSSKrSSKrSSKr\R                  R!                  5       r\=(       a    \R                  R%                  5       S:  r\(       a  \R(                  " S5      OSr\(       a
  \" S 5      rO	\" S 5      r\" S 5      r\" S	 5      r\" S
 5      r\" S 5      r\" S 5      r\" S 5      r\" S 5      r\" S 5      r\" S 5      r\" S 5      r \" S 5      r!\" S 5      r"\" S 5      r#S r$S r%S r&S r'S r(S r)\" S 5      r*\+\,S'   \" S 5      r-\+\,S'   \" S 5      r.\+\,S '   \" S! 5      r/\+\,S"'   \=(       a    \(       + r0\+\,S#'   \" S$ 5      r1\+\,S%'   S& r2\" S' 5      r3\+\,S('   \" S) 5      r4\+\,S*'   \(       a    SSK5r6\6R                  Ro                  5       r8OS+r8S+q:S, r;\Rx                  S- 5       r=\Rx                  S>S. j5       r>\Rx                  S/ 5       r?S>S0 jr@S1 rAS2 rBS3 rCS4 rDS5 rES6 rF\E" 5       rG\F" 5       rHS7\R                  R                  S4S8 jrKS7\R                  \R                  R                  S4S9 jrMS: rNS; rOS< rPS= rQ\(       d"  \R                  R!                  5       (       a   egg! \9 a    S+r8S+r Nf = f)?z>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                      [         $ N)r        ]/var/www/fran/franai/venv/lib/python3.13/site-packages/torch/testing/_internal/common_cuda.py<lambda>r      s    r   c                      [         =(       a@    [        R                  R                  R	                  [        R
                  " S[        S95      $ )N      ?device)r   torchbackendscudnnis_acceptabletensorCUDA_DEVICEr   r   r   r   r      s1    !wu~~/C/C/Q/QRWR^R^_ajuRv/w!wr   c                  l    [         (       a(  [        R                  R                  R	                  5       $ S$ )Nr   )
TEST_CUDNNr   r   r   versionr   r   r   r   r      s$    zzU^^%9%9%A%A%C%XWX%Xr   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)      r   cudais_availableget_device_capabilityr   r   r   r   r      ,    ejj557hEJJ<\<\<^bh<hhr   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)   r   r    r   r   r   r   r      r$   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)   r   r    r   r   r   r   r      r$   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)r(   r   r    r   r   r   r   r      r$   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)   r   r    r   r   r   r   r       r$   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ Nr+   	   r    r   r   r   r   r   !   r$   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)r/   r   r    r   r   r   r   r   "   r$   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)
   r   r    r   r   r   r   r   #   ,    uzz668jUZZ=]=]=_cj=jjr   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)   r   r    r   r   r   r   r   $   r3   r   c                      [         R                  R                  5       =(       aO    [         R                  R                  5       S   S:H  =(       a$    [         R                  R                  5       S   S:  $ )Nr   r2      r    r   r   r   r   r   &   sX    %**113 @

8X8X8Z[\8]ac8c @jj668;a?@r   c                      [         R                  R                  5       =(       a.    [         R                  R                  5       S;   =(       d    [        $ )N))r(   r	   )r+   r(   )r   r!   r"   r#   IS_THORr   r   r   r   r   (   s2    EJJ335}5::;[;[;]aq;q;|u|}r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:H  $ r-   r    r   r   r   r   r   )   s,    %**113d

8X8X8Z^d8ddr   c                   ^ [         R                  R                  5       (       d  g[         R                  R                  S5      R                  n[
        R                  R                  SU5      m[        U4S jU  5       5      $ )NFr!   /PYTORCH_DEBUG_FLASH_ATTENTION_GCN_ARCH_OVERRIDEc              3   ,   >#    U  H	  oT;   v   M     g 7fr   r   ).0archeffective_archs     r   	<genexpr>+evaluate_gfx_arch_within.<locals>.<genexpr>2   s     <)$~%)s   )	r   r!   r"   get_device_propertiesgcnArchNameosenvirongetany)	arch_listgcn_arch_namer@   s     @r   evaluate_gfx_arch_withinrK   +   s\    ::""$$JJ44V<HHMZZ^^$UWdeN <)<<<r   c                      [        / SQ5      $ )N)gfx940gfx941gfx942gfx950rK   r   r   r   CDNA3OrLaterrR   4   s    #$LMMr   c                      [        SS/5      $ )Ngfx90arO   rQ   r   r   r   CDNA2OrLaterrU   7   s    #Xx$899r   c                      [         (       a:  / SQn [        R                  R                  SS5      S:w  a  U / SQ-  n [	        U 5      $ [
        (       a  [        (       + =(       a    [        $ g)NrT   rO   gfx1100gfx1201rP   'TORCH_ROCM_AOTRITON_ENABLE_EXPERIMENTAL0gfx1101gfx1150gfx1151gfx1200F)r   rE   rF   rG   rK   r   r   SM80OrLaterrI   s    r   *evaluate_platform_supports_flash_attentionrc   :   sO    ~H	::>>CSISPEEI'	22y~-+-r   c                      [         (       a:  / SQn [        R                  R                  SS5      S:w  a  U / SQ-  n [	        U 5      $ [
        (       a  gg)NrW   rZ   r[   r\   TF)r   rE   rF   rG   rK   r   rb   s    r   .evaluate_platform_supports_efficient_attentionre   D   sE    ~H	::>>CSISPEEI'	22yr   c                  R    [         (       + =(       a    [        =(       a	    [        S:  $ )Ni_ )r   ra   TEST_CUDNN_VERSIONr   r   r   *evaluate_platform_supports_cudnn_attentionrh   N   s    QKQ5G55PQr   c                      [        5       $ r   )rc   r   r   r   r   r   Q       :d:fr   !PLATFORM_SUPPORTS_FLASH_ATTENTIONc                      [        5       $ r   )re   r   r   r   r   r   R   s    <j<lr   #PLATFORM_SUPPORTS_MEM_EFF_ATTENTIONc                      [        5       $ r   )rh   r   r   r   r   r   S   rj   r   !PLATFORM_SUPPORTS_CUDNN_ATTENTIONc                  B    [         =(       d    [        =(       d    [        $ r   )rk   ro   rm   r   r   r   r   r   U   s    :[ ;V2S;V2U;Vr   !PLATFORM_SUPPORTS_FUSED_ATTENTIONPLATFORM_SUPPORTS_FUSED_SDPAc                  (    [         =(       a    [        $ r   )r   ra   r   r   r   r   r   [   s    y/H[/Hr   PLATFORM_SUPPORTS_BF16c                  2   [         R                  R                  5       (       a  [         R                  R                  (       a  [        S [         R                  R                  R                  S5      S S  5       5      n S/nU S:  a  UR                  S/5        U S:  a  UR                  S5        U H2  nU[         R                  R                  S	5      R                  ;   d  M2    g
   g[        =(       d!    [         R                  R                  5       S:H  $ g)Nc              3   8   #    U  H  n[        U5      v   M     g 7fr   int)r>   vs     r   rA   1evaluate_platform_supports_fp8.<locals>.<genexpr>`   s      R1QAQ1Q   .r	   gfx94)r&   r   gfx120)r&   r   gfx95r   Tr.   F)r   r!   r"   r   hiptuplesplitextendappendrC   rD   SM90OrLaterr#   )ROCM_VERSIONarchsr?   s      r   evaluate_platform_supports_fp8r   ]   s    zz  ==  R1B1B1H1H1Mbq1Q RRLIEv%hZ(v%W%5::;;A>JJJ 
  N%**"B"B"D"NNr   c                      [        5       $ r   )r   r   r   r   r   r   m   s    .L.Nr   PLATFORM_SUPPORTS_FP8c                  (    [         =(       a    [        $ r   )r   SM100OrLaterr   r   r   r   r   o   s    )2L2Lr   PLATFORM_SUPPORTS_MX_GEMMFc                      [         (       d   S5       e[        (       dI  [        [        R                  R                  5       5       H  n [        R                  " SSU  3S9  M     Sqg g )Nz?CUDA must be available when calling initialize_cuda_context_rngr7   zcuda:r   T)r   __cuda_ctx_rng_initializedranger   r!   device_countrandn)is    r   initialize_cuda_context_rngr      sS    9WWW9%%uzz..01AKKE!+. 2%)"	 &r   c               #     #    [         R                  R                  R                  R                  n  S[         R                  R                  R                  l        [         R                  R
                  R                  S S S SS9   S v   S S S 5        U [         R                  R                  R                  l        g ! , (       d  f       N8= f! U [         R                  R                  R                  l        f = f7f)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 l 1F""- lk 1F""-5   /C;AC B<
C *C;<
C
C +C88C;c              #     #    [         R                  R                  (       a3  [        R                  R                  SS 5      nS[        R                  S'   [         R                  R                  R                  R                  nU R                  n S[         R                  R                  R                  l	        Xl
        [         R                  R                  R                  S S S SS9   S v   S S S 5        [         R                  R                  (       a(  Wb  U[        R                  S'   O[        R                  S	 U[         R                  R                  R                  l	        X@l
        g ! , (       d  f       N= f! [         R                  R                  (       a(  Wb  U[        R                  S'   O[        R                  S	 U[         R                  R                  R                  l	        X@l
        f = f7f)NHIPBLASLT_ALLOW_TF321Tr   )r   r   r   rE   rF   rG   r   r!   r   r   	precisionr   r   )selftf32_precisionhip_allow_tf32r   old_precisions        r   tf32_onr      sO    }}(>E-0

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

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

12JJ560E""-&s9   BG9AE> )E-.E> 6A7G9-
E;7E> >A8G66G9c               #     #    [         R                  R                  R                  R                  n  S[         R                  R                  R                  l        [         R                  R
                  R                  SSSSS9   Sv   SSS5        U [         R                  R                  R                  l        g! , (       d  f       N8= f! U [         R                  R                  R                  l        f = f7f)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                 ,   ^ ^^ S mU 4S jmUU4S jnU$ )Nc                 Z    [        5          U" 5         S S S 5        g ! , (       d  f       g = fr   r   )r   function_calls     r   with_tf32_disabled+tf32_on_and_off.<locals>.with_tf32_disabled   s    ZO ZZs   
*c                 `   > [        U T5         U" 5         S S S 5        g ! , (       d  f       g = fr   )r   )r   r   r   s     r   with_tf32_enabled*tf32_on_and_off.<locals>.with_tf32_enabled   s    T>*O +**s   
-c                    >^ ^ [         R                  " T 5      R                  n[        UR	                  5       5      m[
        R                  " T 5      UU UU4S j5       nU$ )Nc                    >^ TR                  [        TU 5      5        [        R                  R	                  5       nST;   a/  U=(       a&    [        R
                  " TS   5      R                  S:H  nST;   a/  U=(       a&    TS   [        R                  [        R                  1;   nU(       a#  T" TS   UU4S j5        T" TS   UU4S j5        g T" S0 TD6  g )Nr   r!   dtyper   c                     > T " S0 TD6$ Nr   r   fkwargss   r   r   Ctf32_on_and_off.<locals>.wrapper.<locals>.wrapped.<locals>.<lambda>   s    1;v;r   c                     > T " S0 TD6$ r   r   r   s   r   r   r      s    !+f+r   r   )	updatezipr   r!   is_tf32_supportedr   typefloat32	complex64)argsr   cond	arg_namesr   r   r   s    ` r   wrapped1tf32_on_and_off.<locals>.wrapper.<locals>.wrapped   s    MM#i./:://1D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   s   `  @r   wrapper tf32_on_and_off.<locals>.wrapper   sK    ""1%00&++-(				 
	 r   r   )r   r   r   r   s   ` @@r   tf32_on_and_offr      s    & Nr   c                 F   ^  [         R                  " T 5      U 4S j5       nU$ )Nc                  ^   > [        5          T" U 0 UD6sS S S 5        $ ! , (       d  f       g = fr   r   )r   r   r   s     r   r   with_tf32_off.<locals>.wrapped   s    Zd%f% ZZs   
,)r   r   )r   r   s   ` r   with_tf32_offr      s%    __Q& & Nr   c                  b   S[         R                  R                  5       ;  a  g[         R                  R                  5       R                  S5      n [         R                  R                  5       U [	        S5      -   S  R                  S5      S   n[        S UR                  S5       5       5      $ )NMagmar   r   zMagma 
r   c              3   8   #    U  H  n[        U5      v   M     g 7fr   rw   r>   xs     r   rA   %_get_magma_version.<locals>.<genexpr>
  s     8!7AQ!7r{   r|   )r   
__config__show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                      [         R                  R                  c  g[        [         R                  R                  5      n [	        S U R                  S5       5       5      $ )Nr   c              3   8   #    U  H  n[        U5      v   M     g 7fr   rw   r   s     r   rA   *_get_torch_cuda_version.<locals>.<genexpr>       9!8AQ!8r{   r|   )r   r   r!   strr   r   )cuda_versions    r   _get_torch_cuda_versionr     sE    }}!u}}))*L9!3!3C!8999r   c                      [         (       a  [        R                  R                  c  g[	        [        R                  R                  5      n U R                  S5      S   n [        S U R                  S5       5       5      $ )Nr   -r   c              3   8   #    U  H  n[        U5      v   M     g 7fr   rw   r   s     r   rA   *_get_torch_rocm_version.<locals>.<genexpr>  r   r{   r|   r   r   r   r   r   r   r   )rocm_versions    r   _get_torch_rocm_versionr     s^    >U]]..6u}}(()L%%c*1-L9!3!3C!8999r   c                      [         (       + $ r   )r   r   r   r   !_check_cusparse_generic_availabler     s    r   c                  4   [         (       d  g[        R                  R                  (       d  g[	        [        R                  R                  5      n U R                  S5      S   n [        S U R                  S5       5       5      nUS L =(       d    US:  (       + $ )NFr   r   c              3   8   #    U  H  n[        U5      v   M     g 7fr   rw   r   s     r   rA   5_check_hipsparse_generic_available.<locals>.<genexpr>$  s     G/F!s1vv/Fr{   r|   )r   r7   r   )r   rocm_version_tuples     r   "_check_hipsparse_generic_availabler     sx    >==u}}(()L%%c*1-LG|/A/A#/FGG"d*I.@6.IJJr   r!   c                    [         R                  R                  [         R                  R                  SS5      [         R                  R                  SS5      5      R	                  U S9n[         R                  R                  [         R                  R                  SS5      [         R                  R                  SS5      5      R	                  U S9n[         R
                  " 5          [        UR                  5       UR                  5       5       H  u  pVUR                  U5        M     S S S 5        SS0nUb  UR                  U5        U" UR                  5       40 UD6nU" UR                  5       40 UD6n	X4X4$ ! , (       d  f       NX= f)Nr+   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  ,  s,    ((%%ehhooa&;UXX__QPQ=RSVV^dVeK((%%ehhooa&;UXX__QPQ=RSVV^dVeK	..0+2H2H2JKDAGGAJ L 
 C[F#&' !7!7!9DVDK !7!7!9DVDK[== 
s   (AE==
Fc           
         [         R                  " SXS9[         R                  " SXS94[         R                  " SXS9[         R                  " SXS94[         R                  " SXS9[         R                  " SXS94[         R                  " SXS9[         R                  " SXS94/n[         R                  R                  5       R	                  U 5      nSn[        XUS9XEU4-   $ )N)r+   r+   )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    [[u<ekk&X]>mn[[u<ekk&X]>mn[[u<ekk&X]>mn[[u<ekk&X]>mnpD
 hh ##F+GI,GW		"# #r   c                 H    [         (       d  U $ [        R                  " U 5      $ r   )IS_SM89unittestexpectedFailurefuncs    r   xfailIfSM89r  M  s    w4BH$<$<T$BBr   c                 H    [         (       d  U $ [        R                  " U 5      $ r   )r   r  r  r  s    r   xfailIfSM100OrLaterr  P      #|4G)A)A$)GGr   c                 H    [         (       d  U $ [        R                  " U 5      $ r   )SM120OrLaterr  r  r  s    r   xfailIfSM120OrLaterr   S  r  r   c                 ^    [         (       d  [        (       d  U $ [        R                  " U 5      $ r   )r   	IS_JETSONr  r  r  s    r   xfailIfDistributedNotSupportedr#  V  s      II4RH4L4LT4RRr   )gh㈵>)R__doc__r   r   
torch.cuda$torch.testing._internal.common_utilsr   r   r   r   r   r   r   
contextlibrE   r  r!   is_initialized"CUDA_ALREADY_INITIALIZED_ON_IMPORTr   TEST_MULTIGPUr   r   r   rg   SM53OrLaterSM60OrLaterSM70OrLaterSM75OrLaterra   SM89OrLaterr   r   r  r9   r"  r  rK   rR   rU   rc   re   rh   rk   bool__annotations__rm   ro   rq   rr   rt   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   r#  r   r   r   <module>r<     s#   F    u u   	  &+ZZ%>%>%@ " <ejj5571<(1ell8$t*+JwxJXY hihihihihihihijkjk
 @ A}~	
d
e=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  &&NO t O")*L"M 4 M**113
 O # * F F ' '* F FXH9::	K :; ;=  .4EKKOOfj >$ !'ekk%++//lp #CHHS *zz((***** *K  
s   .J* *
J76J7