
    3jQ                     
   % 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
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 5      r%\" S 5      r&\" S 5      r'\" S 5      r(\RR                  S 5       r*S r+S r,S r-S r.S r/S r0S  r1S! r2\" S" 5      r3\4\5S#'   \" S$ 5      r6\4\5S%'   \" S& 5      r7\4\5S''   \" S( 5      r8\4\5S)'   \=(       a    \(       + r9\4\5S*'   \" S+ 5      r:\4\5S,'   S- r;S. r<S/ r=\" S0 5      r>\4\5S1'   \" S2 5      r?\4\5S3'   \" S4 5      r@\4\5S5'   \" S6 5      rA\4\5S7'   S8 rB\" S9 5      rC\4\5S:'   S; rDS< rES= rFS> rGS? rH\" S@ 5      rI\4\5SA'   \" SB 5      rJ\4\5SC'   \" SD 5      rK\4\5SE'   \" SF 5      rL\4\5SG'   \" SH 5      rM\4\5SI'   \(       a    SSKNrO\OR                   R                  5       rQOSJrQSJqUSK rV\RR                  SL 5       rW\RR                  SfSM j5       rX\RR                  SN 5       rYSfSOSP.SQ jjrZSR r[SS r\ST r]SU r^SV r_SW r`SX ra\`" 5       rb\a" 5       rcSY\R                  R                  S4SZ jrfSY\R                  \R                  R                  S4S[ jrhS\ riS] rjS^ rkS_ rlS` rmSa rnSb roScrp\R                  " \R                  R                  (       + =(       a)    \R                  R                  SL =(       a
    \]" 5       \p:  SdR                  " \p6 5      rv\(       d(  \R                   R#                  5       (       a  \w" Se5      egg! \R\S\T4 a    SJrQSJr GNlf = f)gz>This file is allowed to initialize CUDA context when imported.    N)LazyVal
TEST_NUMBATEST_WITH_ROCM	TEST_CUDA
IS_WINDOWSIS_MACOSTEST_XPU   zcuda:0c                      [         $ N)r        ]/home/wildlama/miniconda3/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                  (       a<  [        S [         R                  R                  R	                  S5      S S  5       5      $ S$ )Nc              3   8   #    U  H  n[        U5      v   M     g 7fr   int).0vs     r   	<genexpr><lambda>.<locals>.<genexpr>   s     %W6Vc!ff6V   .r
   r   r   )r   r   hiptuplesplitr   r   r   r   r      sE    [`[h[h[l[l%Wemm6G6G6M6Mc6RSUTU6V%W W xrx 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      r3   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N   r   r/   r   r   r   r   r      r3   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N)r8   r-   r/   r   r   r   r   r       r3   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ )N   r   r/   r   r   r   r   r   !   r3   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ Nr<   	   r/   r   r   r   r   r   "   r3   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:  $ N)r@   r   r/   r   r   r   r   r   #   r3   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   %   rF   r   c                     [         R                  R                  5       =(       a    [         R                  R                  S L=(       a    [         R                  R	                  5       S:H  =(       a)    [        [         R                  R                  S S 5      S:  =(       dQ    [         R                  R	                  5       S:H  =(       a)    [        [         R                  R                  S S 5      S:  $ )N)   r   r
      )rE      )r   r0   r1   r   r2   r!   r   r   r   r   r   '   s    %**113 i8J8JRV8V iJJ446'Agc%--J\J\]_^_J`FaegFg hJJ446'Afc%--J\J\]_^_J`FadfFfir   c                      [         R                  R                  5       =(       a.    [         R                  R                  5       S;   =(       d    [        $ )N))r8   r
   )r<   r8   )r   r0   r1   r2   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   +   ,    %**113d

8X8X8Z^d8ddr   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:H  $ rB   r/   r   r   r   r   r   ,   rP   r   c                      [         R                  R                  5       =(       a!    [         R                  R                  5       S:H  $ rD   r/   r   r   r   r   r   -   s,    5::224f9Y9Y9[_f9ffr   c                      [         R                  R                  5       =(       a$    [         R                  R                  5       S   S:H  $ )Nr   rH   r/   r   r   r   r   r   .   s2    5::224d9Y9Y9[\]9^bd9ddr   c              #   f  #    [         R                  R                  R                  5       n[         R                  R                  R                  U 5         S v   [         R                  R                  R                  U5        g ! [         R                  R                  R                  U5        f = f7fr   )r   r   r0   preferred_blas_library)backendprev_backends     r   blas_library_contextrX   0   sn     >>&&==?L	NN..w7A22<@22<@s   AB1B *B1+B..B1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      $ )NFr0   /PYTORCH_DEBUG_FLASH_ATTENTION_GCN_ARCH_OVERRIDEc              3   ,   >#    U  H	  oT;   v   M     g 7fr   r   )r"   archeffective_archs     r   r$   +evaluate_gfx_arch_within.<locals>.<genexpr>@   s     <)$~%)s   )	r   r0   r1   get_device_propertiesgcnArchNameosenvirongetany)	arch_listgcn_arch_namer]   s     @r   evaluate_gfx_arch_withinrg   9   s\    ::""$$JJ44V<HHMZZ^^$UWdeN <)<<<r   c                      [        SS/5      $ )Ngfx942gfx950rg   r   r   r   CDNA3OrLaterrl   B       #Xx$899r   c                      [        SS/5      $ )Ngfx90ari   rk   r   r   r   CDNA2OrLaterrp   E   rm   r   c                      [         (       a:  / SQn [        R                  R                  SS5      S:w  a  U / SQ-  n [	        U 5      $ [
        (       a  [        (       + =(       a    [        $ [        (       a  ggN)ro   ri   gfx1100gfx1201rj   'TORCH_ROCM_AOTRITON_ENABLE_EXPERIMENTAL0)gfx1101gfx1102gfx1150gfx1151gfx1200TF)	r   ra   rb   rc   rg   r   r   SM80OrLaterr	   re   s    r   *evaluate_platform_supports_flash_attentionr~   H   sW    ~H	::>>CSISPPPI'	22y~-+-xr   c                  j    [         (       a(  [        R                  R                  R	                  5       $ g)NF)r   r   r   r0   is_ck_sdpa_availabler   r   r   "evaluate_platform_supports_ck_sdpar   T   s#    ~~~""7799r   c                      [         (       a:  / SQn [        R                  R                  SS5      S:w  a  U / SQ-  n [	        U 5      $ [
        (       a  g[        (       a  ggrr   )r   ra   rb   rc   rg   r   r	   r}   s    r   .evaluate_platform_supports_efficient_attentionr   Z   sM    ~H	::>>CSISPPPI'	22yxr   c                  R    [         (       + =(       a    [        =(       a	    [        S:  $ )Ni_ )r   r|   TEST_CUDNN_VERSIONr   r   r   *evaluate_platform_supports_cudnn_attentionr   f   s    QKQ5G55PQr   c                     [         (       a  g[        5       S:  d  g[        R                  R                  R                  [        R                  R                  R                  5      n U c  g[        U R                  S5      S   5      S:  $ )NFrH   r<   r'   r   i:  	r   _get_torch_cuda_versionr   utilscollect_envget_nvidia_driver_versionrunr!   r+   driver_versions    r   (evaluate_platform_supports_green_contextr   i   k    z"$/[[,,FFu{{G^G^GbGbcN~##C(+,33r   c                      [        5       $ r   )r~   r   r   r   r   r   s       :d:fr   !PLATFORM_SUPPORTS_FLASH_ATTENTIONc                      [        5       $ r   )r   r   r   r   r   r   t   s    <j<lr   #PLATFORM_SUPPORTS_MEM_EFF_ATTENTIONc                      [        5       $ r   )r   r   r   r   r   r   u   r   r   !PLATFORM_SUPPORTS_CUDNN_ATTENTIONc                  B    [         =(       d    [        =(       d    [        $ r   )r   r   r   r   r   r   r   r   w   s    :[ ;V2S;V2U;Vr   !PLATFORM_SUPPORTS_FUSED_ATTENTIONPLATFORM_SUPPORTS_FUSED_SDPAc                      [        5       $ r   )r   r   r   r   r   r   }       2T2Vr   PLATFORM_SUPPORTS_CK_SDPAc                      [         R                  R                  (       a  [        $ [         R                  R                  (       a  g[
        (       a  gg)NTF)r   r   r0   r|   r)   r	   r   r   r   evaluate_platform_supports_bf16r      s0    }}				r   c                      [         R                  R                  (       a  [        $ [         R                  R                  (       a	  [
        S:  $ g)Nr;   F)r   r   r0   r|   r)   ROCM_VERSIONr   r   r   'evaluate_platform_supports_bf16_atomicsr      s0    }}			v%%r   c                  T    [         R                  R                  (       a	  [        S:  $ g)Nr;   T)r   r   r)   r   r   r   r   'evaluate_platform_supports_half_atomicsr      s    }}v%%r   c                      [        5       $ r   )r   r   r   r   r   r      s    /N/Pr   PLATFORM_SUPPORTS_BF16c                      [        5       $ r   )r   r   r   r   r   r          7^7`r   PLATFORM_SUPPORTS_BF16_ATOMICSc                      [        5       $ r   )r   r   r   r   r   r      r   r   PLATFORM_SUPPORTS_HALF_ATOMICSc                      [        5       $ r   )r   r   r   r   r   r      s    8`8br   PLATFORM_SUPPORTS_GREEN_CONTEXTc                     [         (       a  g[        5       S:  d  g[        R                  R                  R                  [        R                  R                  R                  5      n U c  g[        U R                  S5      S   5      S:  $ )NF)rK   rL   r'   r   iN  r   r   s    r   +evaluate_platform_supports_workqueue_configr      r   r   c                      [        5       $ r   )r   r   r   r   r   r          ;f;hr   "PLATFORM_SUPPORTS_WORKQUEUE_CONFIGc                     [         R                  R                  5       (       a  [         R                  R                  (       as  S/n [
        S:  a  U R                  S/5        [
        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  $ [         R                  R                  5       (       a  gg)
Ngfx94)r5   r.   gfx120)r5   r-   gfx95r   TFr?   )r   r0   r1   r   r)   r   extendappendr_   r`   SM90OrLaterr2   xpuarchsr\   s     r   evaluate_platform_supports_fp8r      s    zz  ==IEv%hZ(v%W%5::;;A>JJJ  N%**"B"B"D"NNyyr   c                  x   [         R                  R                  5       (       a  [         R                  R                  (       a`  S[         R
                  R                  5       ;  a  gSS/n U  H2  nU[         R                  R                  S5      R                  ;   d  M2    g   g[        =(       a    [        (       + $ g)NUSE_MSLKFri   rj   r   T)r   r0   r1   r   r)   
__config__showr_   r`   r   SM100OrLaterr   s     r   +evaluate_platform_supports_fp8_grouped_gemmr      s    zz  ==!1!1!6!6!88x(E5::;;A>JJJ 
  3|#33r   c                     [         R                  R                  5       (       a]  [         R                  R                  (       a8  [
        S:  a,  S[         R                  R                  S5      R                  ;   $  g[        $ g)Nr7   rj   r   F)	r   r0   r1   r   r)   r   r_   r`   r   r   r   r   "evaluate_platform_supports_mx_gemmr      s_    zz  ==v%5::#C#CA#F#R#RRR &   r   c                      [         R                  R                  5       (       aO  [         R                  R                  (       d0  S[         R
                  R                  5       ;   n U =(       a    [        $ g)Nr   F)r   r0   r1   r   r)   r   r   IS_SM100)built_with_mslks    r   -evaluate_platform_supports_mxfp8_grouped_gemmr      sH    zz  ):):$(8(8(=(=(??+8+r   c                     [         R                  R                  5       (       a  [         R                  R                  (       a,  S[         R                  R                  S5      R                  ;   $ [        =(       d!    [         R                  R                  5       S:H  =(       aZ    [         R                  R                  R                  5       =(       a+    [         R                  R                  R                  5       S:  $ g)Nrj   r   r?   iZ  F)r   r0   r1   r   r)   r_   r`   r   r2   r   
cusparseltr   r   r   %evaluate_platform_supports_fp8_sparser      s    zz  ==uzz??BNNNN L

 @ @ Bf L ?NN--::<?NN--5573>
 r   c                      [        5       $ r   )r   r   r   r   r   r      r   r   PLATFORM_SUPPORTS_MX_GEMMc                      [        5       $ r   )r   r   r   r   r   r      s    .L.Nr   PLATFORM_SUPPORTS_FP8c                      [        5       $ r   )r   r   r   r   r   r      s    5Z5\r   PLATFORM_SUPPORTS_FP8_SPARSEc                      [        5       $ r   )r   r   r   r   r   r      r   r   "PLATFORM_SUPPORTS_FP8_GROUPED_GEMMc                      [        5       $ r   )r   r   r   r   r   r      s    =j=lr   $PLATFORM_SUPPORTS_MXFP8_GROUPED_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_rngrL   zcuda:r   T)r   AssertionError__cuda_ctx_rng_initializedranger   r0   device_countrandn)is    r   initialize_cuda_context_rngr      sT    9^__%%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   r0   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              #   6  #    [         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        U[         R                  R                  R                  l        X0l        g ! , (       d  f       N>= f! U[         R                  R                  R                  l        X0l        f = f7f)NTr   )r   r   r0   r   r   	precisionr   r   )selftf32_precisionr   old_precisions       r   tf32_onr     s     !NN//66AANNM'04""-'^^!!''TXei'j k 1F""-&	 kj 1F""-&s5   ;DAC% CC% $0D
C"C% %1DD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   T)only_ifc                0   ^ ^^^ S mU 4S jmU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_disabledM  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_enabledQ  s    T>*O +**s   
-c                    >^ ^ [         R                  " T 5      R                  n[        UR	                  5       5      m[
        R                  " T 5      UU UUU4S j5       nU$ )Nc                    >^ TR                  [        TU SS95        [        R                  R	                  5       =(       a    T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 )
NFstrictr   r0   dtyper   c                     > T " S0 TD6$ Nr   r   fkwargss   r   r   Ctf32_on_and_off.<locals>.wrapper.<locals>.wrapped.<locals>.<lambda>b  s    1;v;r   c                     > T " S0 TD6$ r  r   r  s   r   r   r
  c  s    !+f+r   r   )	updatezipr   r0   is_tf32_supportedr   typefloat32	complex64)argsr	  cond	arg_namesr  r   r   r   s    ` r   wrapped1tf32_on_and_off.<locals>.wrapper.<locals>.wrappedY  s    MM#ie<=:://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   wrapper tf32_on_and_off.<locals>.wrapperU  sP    ""1%00&++-(				 	 
	 r   r   )r   r   r  r   r   s   `` @@r   tf32_on_and_offr   L  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>.wrappedp  s    Zd%f% ZZs   
,)r  r  )r  r  s   ` r   with_tf32_offr$  o  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Magmar(   zMagma 
r   c              3   8   #    U  H  n[        U5      v   M     g 7fr   r    r"   xs     r   r$   %_get_magma_version.<locals>.<genexpr>|  s     8!7AQ!7r&   r'   )r   r   r   findlenr+   r*   )positionversion_strs     r   _get_magma_versionr0  w  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   r    r)  s     r   r$   *_get_torch_cuda_version.<locals>.<genexpr>       9!8AQ!8r&   r'   )r   r   r0   strr*   r+   )cuda_versions    r   r   r   ~  sE    }}!u}}))*L9!3!3C!8999r   c                      [         (       a  [        R                  R                  c  g[	        [        R                  R                  5      n U R                  SSS9S   n [        S U R                  S5       5       5      $ )Nr(   -rL   maxsplitr   c              3   8   #    U  H  n[        U5      v   M     g 7fr   r    r)  s     r   r$   *_get_torch_rocm_version.<locals>.<genexpr>  r4  r&   r'   r   r   r   r)   r5  r+   r*   )rocm_versions    r   _get_torch_rocm_versionr?    sc    >U]]..6u}}(()L%%cA%6q9L9!3!3C!8999r   c                      [         (       d  g  [        R                  R                  5       n U b  U S:X  a  g U S-  nU S-  S-  nU S-  nXU4$ ! [        [
        4 a     g f = f)Nr   i'  d   )r   r   _C_cuda_getHipblasltVersionAttributeErrorRuntimeError)version_intmajorminorpatchs       r   _get_torch_hipblaslt_versionrJ    sx    > hh88:+"2u$u$,c!e$$L) s   'A A AAc                      [         (       + $ r   )r   r   r   r   !_check_cusparse_generic_availablerL    s    r   c                  2   [         (       d  g[        R                  R                  (       d  g[	        [        R                  R                  5      n U R                  SSS9S   n [        S U R                  S5       5       5      nUS L =(       d    US:  (       + $ )	NFr8  rL   r9  r   c              3   8   #    U  H  n[        U5      v   M     g 7fr   r    r)  s     r   r$   5_check_hipsparse_generic_available.<locals>.<genexpr>  s     G/F!s1vv/Fr&   r'   )r-   rL   r=  )r>  rocm_version_tuples     r   "_check_hipsparse_generic_availablerQ    s}    >==u}}(()L%%cA%6q9LG|/A/A#/FGG"d*I.@6.IJJr   r0   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       SS9 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   Tr  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_optimizersrb    s/    ((%%ehhooa&;UXX__QPQ=RSVV^dVeK((%%ehhooa&;UXX__QPQ=RSVV^dVeK	..0+2H2H2JSWXDAGGAJ Y 
 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   rZ  r[  )r   r   rT  MSELossrW  rb  )r   r  rZ  r[  dataloss_fn	skip_iters          r   _create_scaling_caserh    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   xfailIfSM89ro        w4BH$<$<T$BBr   c                 H    [         (       d  U $ [        R                  " U 5      $ r   )IS_SM90rk  rl  rm  s    r   xfailIfSM90rs    rp  r   c                 d    [         (       a$  [        5       S:  a  [        R                  " U 5      $ U $ )zUxfail on SM89 only for CUDA < 13. On CUDA 13+, test should pass on all architectures.)rK   r   )rj  r   rk  rl  rm  s    r   xfailIfSM89PreCUDA13ru    s(    w*,w6''--Kr   c                 H    [         (       d  U $ [        R                  " U 5      $ r   )r   rk  rl  rm  s    r   xfailIfSM100OrLaterrw        #|4G)A)A$)GGr   c                 H    [         (       d  U $ [        R                  " U 5      $ r   )SM120OrLaterrk  rl  rm  s    r   xfailIfSM120OrLaterr{    rx  r   c                 H    [         (       d  U $ [        R                  " U 5      $ r   )IS_SM12Xrk  rl  rm  s    r   xfailIfSM12Xr~    s    x4CX%=%=d%CCr   c                 ^    [         (       d  [        (       d  U $ [        R                  " U 5      $ r   )r   	IS_JETSONrk  rl  rm  s    r   xfailIfDistributedNotSupportedr    s      II4RH4L4LT4RRr   r   z2Requires CUDA {}.{} to match Tritons ptxas versionz(CUDA should not be initialized on import)gh㈵>)x__doc__r  r   
torch.cuda$torch.testing._internal.common_utilsr   r   r   r   r   r   r	   r  
contextlibra   rk  r0   is_initialized"CUDA_ALREADY_INITIALIZED_ON_IMPORTr   TEST_MULTIGPUr   r   r   r   r   SM53OrLaterSM60OrLaterSM70OrLaterSM75OrLaterr|   SM89OrLaterr   r   rz  rN   r  rj  rr  r   r}  contextmanagerrX   rg   rl   rp   r~   r   r   r   r   r   bool__annotations__r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   
numba.cudanumbar1   TEST_NUMBA_CUDAImportErrorrE  OSErrorr   r   r   r   r   r   r$  r0  r   r?  rJ  rL  rQ  TEST_CUSPARSE_GENERICTEST_HIPSPARSE_GENERICoptimSGDrb  floatrh  ro  rs  ru  rw  r{  r~  r  TRITON_PTXAS_VERSIONskipIfr   r   r)   formatrequires_triton_ptxas_compatr   r   r   r   <module>r     s   F         	  &+ZZ%>%>%@ " <ejj5571<(1ell8$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
 i j }~	
d
e
d
efgdeA A=::

R4 +22f*g !4 g,34l,m #T m*12f*g !4 g*1 3V +W !4 W &/%E~3E d E")*V"W 4 W  ''PQ  Q'./`'a  a'./`'a  a(/0b(c  c4 ,33h+i "D i&
 #**V"W 4 W%&NO t O%,-\%] d ]+23h+i "D i-45l-m $d m**113
 O # * F F 
' 
' F FXD F9:: 	K :; ;=  .4EKKOOfj >$ !'ekk%++//lp #CCHHDS  '5==3D3D/D 0d38==3D3D3L0d3J3LOc3c/c/j/j  mA  0B C  *zz  ""GHH # *A w/ 
s   4O1 1PP