
    ^jV                        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
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' ed       Z( ed       Z)ejT                  d        Z+d Z,d Z-d Z.d Z/d Z0d  Z1d! Z2d" Z3 ed#       Z4e5e6d$<    ed%       Z7e5e6d&<    ed'       Z8e5e6d(<    ed)       Z9e5e6d*<   exr e Z:e5e6d+<    ed,       Z;e5e6d-<   d. Z<d/ Z=d0 Z> ed1       Z?e5e6d2<    ed3       Z@e5e6d4<    ed5       ZAe5e6d6<    ed7       ZBe5e6d8<   d9 ZC ed:       ZDe5e6d;<   d< ZEd= ZFd> ZGd? ZHd@ ZI edA       ZJe5e6dB<    edC       ZKe5e6dD<    edE       ZLe5e6dF<    edG       ZMe5e6dH<    edI       ZNe5e6dJ<   er 	 ddlOZPePj                   j                         ZRndKZRdKaVdL ZWejT                  dM        ZXejT                  didN       ZYejT                  dO        ZZdidPdQdRZ[dS Z\dT Z]dU Z^dV Z_dW Z`dX ZadY Zb ea       Zc eb       ZddZej                  j                  dfd[ZgdZej                  ej                  j                  dfd\Zid] Zjd^ Zkd_ Zld` Zmda Zndb Zodc Zpdd Zqde ZrdfZs ej                  ej                  j                   xr$ ej                  j                  du xr
  e^       esk   dgj                  es       Zyes#ej                   j#                         r ezdh      yy# eSeTeUf$ r dKZRdKZY Tw xY w)jz>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                      t         S N)r        n/var/www/ramen.bs-engineer-server.com/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      r2   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N   r   r.   r   r   r   r   r      r2   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N)r7   r,   r.   r   r   r   r   r       r2   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S )N   r   r.   r   r   r   r   r   !   r2   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S Nr;   	   r.   r   r   r   r   r   "   r2   r   c                      t         j                  j                         xr! t         j                  j                         dk\  S N)r?   r   r.   r   r   r   r   r   #   r2   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   %   rE   r   c                     t         j                  j                         xr t         j                  j                  d uxr t         j                  j	                         dk(  xr) t        t         j                  j                  d d       dk\  xsL t         j                  j	                         dk(  xr) t        t         j                  j                  d d       dk  S )N)   r   r
      )rD      )r   r/   r0   r   r1   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 r   c                      t         j                  j                         xr( t         j                  j                         dv xs t        S )N))r7   r
   )r;   r7   )r   r/   r0   r1   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 rA   r.   r   r   r   r   r   ,   rO   r   c                      t         j                  j                         xr! t         j                  j                         dk(  S rC   r.   r   r   r   r   r   -   s+    5::224f9Y9Y9[_f9f r   c                      t         j                  j                         xr$ t         j                  j                         d   dk(  S )Nr   rD   r.   r   r   r   r   r   .   1    5::224d9Y9Y9[\]9^bd9d r   c                      t         j                  j                         xr$ t         j                  j                         d   dk(  S )Nr   rG   r.   r   r   r   r   r   /   rS   r   c              #   f  K   t         j                  j                  j                         }t         j                  j                  j                  |        	 d  t         j                  j                  j                  |       y # t         j                  j                  j                  |       w xY wwr   )r   r   r/   preferred_blas_library)backendprev_backends     r   blas_library_contextrY   1   so     >>&&==?L	NN..w7A22<@22<@s   AB1B *B1+B..B1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>A   s     <$t~%<s   )	r   r/   r0   get_device_propertiesgcnArchNameosenvirongetany)	arch_listgcn_arch_namer^   s     @r   evaluate_gfx_arch_withinrg   :   sY    ::""$JJ44V<HHMZZ^^$UWdeN <)<<<r   c                      t        ddg      S )Ngfx942gfx950rg   r   r   r   CDNA3OrLaterrl   C   s    #Xx$899r   c                      t        g d      S )N)gfx90ari   rj   rk   r   r   r   CDNA2OrLaterro   F   s    #$BCCr   c                      t         r9g d} t        j                  j                  dd      dk7  r| g dz  } t	        |       S t
        rt         xr t        S t        ryyN)rn   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}   I   sQ    H	::>>CSISPPPI'	22~-+-r   c                  `    t         r(t        j                  j                  j	                         S y)NF)r   r   r   r/   is_ck_sdpa_availabler   r   r   "evaluate_platform_supports_ck_sdpar   U   s!    ~~""7799r   c                      t         r9g d} t        j                  j                  dd      dk7  r| g dz  } t	        |       S t
        ryt        ryyrq   )r   ra   rb   rc   rg   r   r	   r|   s    r   .evaluate_platform_supports_efficient_attentionr   [   sG    H	::>>CSISPPPI'	22r   c                  6    t          xr t        xr	 t        dk\  S )Ni_ )r   r{   TEST_CUDNN_VERSIONr   r   r   *evaluate_platform_supports_cudnn_attentionr   g   s    QKQ5G55PQr   c                     t         ryt               dk\  syt        j                  j                  j                  t        j                  j                  j                        } | yt        | j                  d      d         dk\  S )NFrG   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   j   i    "$/[[,,FFu{{G^G^GbGbcN~##C(+,33r   c                      t               S r   )r}   r   r   r   r   r   t   
    :d:f r   !PLATFORM_SUPPORTS_FLASH_ATTENTIONc                      t               S r   )r   r   r   r   r   r   u   s
    <j<l r   #PLATFORM_SUPPORTS_MEM_EFF_ATTENTIONc                      t               S r   )r   r   r   r   r   r   v   r   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               S r   )r   r   r   r   r   r   ~   
    2T2V r   PLATFORM_SUPPORTS_CK_SDPAc                      t         j                  j                  rt        S t         j                  j                  ryt
        ryy)NTF)r   r   r/   r{   r(   r	   r   r   r   evaluate_platform_supports_bf16r      s.    }}				r   c                      t         j                  j                  rt        S t         j                  j                  r	t
        dk\  S y)Nr:   F)r   r   r/   r{   r(   ROCM_VERSIONr   r   r   'evaluate_platform_supports_bf16_atomicsr      s0    }}			v%%r   c                  J    t         j                  j                  r	t        dk\  S y)Nr:   T)r   r   r(   r   r   r   r   'evaluate_platform_supports_half_atomicsr      s    }}v%%r   c                      t               S r   )r   r   r   r   r   r      s
    /N/P r   PLATFORM_SUPPORTS_BF16c                      t               S r   )r   r   r   r   r   r      
    7^7` r   PLATFORM_SUPPORTS_BF16_ATOMICSc                      t               S r   )r   r   r   r   r   r      r   r   PLATFORM_SUPPORTS_HALF_ATOMICSc                      t               S r   )r   r   r   r   r   r      s
    8`8b r   PLATFORM_SUPPORTS_GREEN_CONTEXTc                     t         ryt               dk\  syt        j                  j                  j                  t        j                  j                  j                        } | yt        | j                  d      d         dk\  S )NF)rJ   rK   r&   r   iN  r   r   s    r   +evaluate_platform_supports_workqueue_configr      r   r   c                      t               S r   )r   r   r   r   r   r      
    ;f;h r   "PLATFORM_SUPPORTS_WORKQUEUE_CONFIGc                     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 t         j                  j                         ryy)
Ngfx94)r4   r-   gfx120)r4   r,   gfx95r   TFr>   )r   r/   r0   r   r(   r   extendappendr_   r`   SM90OrLaterr1   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                  F   t         j                  j                         rt         j                  j                  rZdt         j
                  j                         vryddg} | D ]/  }|t         j                  j                  d      j                  v s/ y yt        xr t         S y)NUSE_MSLKFri   rj   r   T)r   r/   r0   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                      t         j                  j                         rUt         j                  j                  r5t
        dk\  r+dt         j                  j                  d      j                  v S yt        S y)Nr6   rj   r   F)	r   r/   r0   r   r(   r   r_   r`   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/   r0   r   r(   r   r   IS_SM100)built_with_mslks    r   -evaluate_platform_supports_mxfp8_grouped_gemmr      sE    zz ):):$(8(8(=(=(??+8+r   c                     t         j                  j                         rt         j                  j                  r+dt         j                  j                  d      j                  v S t        xs! t         j                  j                         dk(  xrU t         j                  j                  j                         xr+ t         j                  j                  j                         dk\  S y)Nrj   r   r>   iZ  F)r   r/   r0   r   r(   r_   r`   r   r1   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                      t               S r   )r   r   r   r   r   r      r   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
    5Z5\ r   PLATFORM_SUPPORTS_FP8_SPARSEc                      t               S r   )r   r   r   r   r   r      r   r   "PLATFORM_SUPPORTS_FP8_GROUPED_GEMMc                      t               S r   )r   r   r   r   r   r      s
    =j=l r   $PLATFORM_SUPPORTS_MXFP8_GROUPED_GEMMFc                      t         st        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_rngrK   zcuda:r   T)r   AssertionError__cuda_ctx_rng_initializedranger   r/   device_countrandn)is    r   initialize_cuda_context_rngr      sT    ^__%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              #   8  K   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                  j                  _        || _        y # 1 sw Y   :xY w# |t         j                  j                  j                  _        || _        w xY ww)NTr   )r   r   r/   r   r   	precisionr   r   )selftf32_precisionr   old_precisions       r   tf32_onr     s     !NN//66AANNM'04""-'^^!!''TXei'j 		 1F""-&		 	 1F""-&s5   ;DAC% C C% (1DC"C% %2DD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   T)only_ifc                ,     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_disabledN  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_enabledR  s'    T>* 	O	 	 	s   (c                      t        j                         j                  }t        |j	                               t        j                          fd       }|S )Nc                     j                  t        | d             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 )
NFstrictr   r/   dtyper   c                        di S Nr   r   fkwargss   r   r   zCtf32_on_and_off.<locals>.wrapper.<locals>.wrapped.<locals>.<lambda>c  s    1;v; r   c                        di S r  r   r  s   r   r   zCtf32_on_and_off.<locals>.wrapper.<locals>.wrapped.<locals>.<lambda>d  s    !+f+ r   r   )	updatezipr   r/   is_tf32_supportedr   typefloat32	complex64)argsr  cond	arg_namesr  r   r   r   s    ` r   wrappedz1tf32_on_and_off.<locals>.wrapper.<locals>.wrappedZ  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   wrapperz tf32_on_and_off.<locals>.wrapperV  sK    ""1%00&++-(				 
	 r   r   )r   r   r  r   r   s   `` @@r   tf32_on_and_offr  M  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>.wrappedq  s*    Z 	&d%f%	& 	& 	&s   'r  r  )r  r  s   ` r   with_tf32_offr  p  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)  x  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   r   r     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'   -rK   )maxsplitr   c              3   2   K   | ]  }t        |        y wr   r    r#  s     r   r$   z*_get_torch_rocm_version.<locals>.<genexpr>  r,  r%   r&   )r   r   r   r(   r-  r*   r)   )rocm_versions    r   _get_torch_rocm_versionr4    sa    U]]..6u}}(()L%%cA%6q9L9!3!3C!8999r   c                      t         sy 	 t        j                  j                         } | | dk(  ry | dz  }| dz  dz  }| dz  }|||fS # t        t
        f$ r Y y w xY w)Nr   i'  d   )r   r   _C_cuda_getHipblasltVersionAttributeErrorRuntimeError)version_intmajorminorpatchs       r   _get_torch_hipblaslt_versionr?    sx     hh88:+"2u$u$,c!ue$$L) s   %A A AAc                      t          S r   )r   r   r   r   !_check_cusparse_generic_availablerA    s    r   c                  J    t         xr t        j                  j                  d uS r   )r   r   r   r(   r   r   r   "_check_hipsparse_generic_availablerC    s    ;emm//t;;r   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      D ]  \  }}|j                  |        	 d d d        ddi}||j                  |        ||j                         fi |} ||j                         fi |}	||||	fS # 1 sw Y   TxY w)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_optimizersrT    sA    ((%%ehhooa&;UXX__QPQ=RSVV^dVeK((%%ehhooa&;UXX__QPQ=RSVV^dVeK	 ..0+2H2H2JSWX 	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)r;   r;   )r  r   r
   )r   rL  rM  )r   r   rF  MSELossrI  rT  )r   r  rL  rM  dataloss_fn	skip_iters          r   _create_scaling_caserZ    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   xfailIfSM89ra        4BH$<$<T$BBr   c                 <    t         s| S t        j                  |       S r   )IS_SM90r]  r^  r_  s    r   xfailIfSM90re    rb  r   c                 V    t         r"t               dk  rt        j                  |       S | S )zUxfail on SM89 only for CUDA < 13. On CUDA 13+, test should pass on all architectures.)rJ   r   )r\  r   r]  r^  r_  s    r   xfailIfSM89PreCUDA13rg    s&    *,w6''--Kr   c                 L    t         r| S t        s| S t        j                  |       S r   )r   r   r]  r^  r_  s    r   xfailIfSM100OrLaterri    s$    
 #4G)A)A$)GGr   c                 L    t         r| S t        s| S t        j                  |       S r   )r   SM120OrLaterr]  r^  r_  s    r   xfailIfSM120OrLaterrl    s"    #4G)A)A$)GGr   c                 <    t         s| S t        j                  |       S r   )IS_SM12Xr]  r^  r_  s    r   xfailIfSM12Xro    s    4CX%=%=d%CCr   c                 H    t         st        s| S t        j                  |       S r   )r   	IS_JETSONr]  r^  r_  s    r   xfailIfDistributedNotSupportedrr    s     I4RH4L4LT4RRr   c                 B     t        j                          fd       }|S )a  Wraps a test to xfail when running on a CUDA device.

    Works for both device-parameterized tests (device_type == "cuda") and plain
    TestCase CUDA tests (device_type is None). Only passes through for tests
    explicitly running on a non-CUDA device type (e.g. "cpu", "xpu").
    c                    dd l }t        | dd       }||dk7  r | g|i |S 	  | g|i | t        dj                   d      # t        $ r+}|j                  dj                   d|        Y d }~Hd }~ww xY w)Nr   device_typer/   zExpected failure for z: zTest z$ was expected to fail but succeeded.)pytestgetattr	Exceptionxfail__name__r   )slfr  r  rv  ru  etest_fns         r   r  z/_xfail_cuda_on_windows_wrapper.<locals>.wrapper  s    c=$7"{f'<30000	JC)$)&) uW%5%5$66Z[\\  	JLL01A1A0B"QCHII	Js   A 	B !A;;B r  )r}  r  s   ` r   _xfail_cuda_on_windows_wrapperr~    s)     __W
] 
] Nr   c                 4    t         rt        rt        |       S | S )zMark a CUDA test as expected failure on Windows with SM >= 8.9.

    Works for both device-parameterized tests and plain TestCase CUDA tests.
    CPU/XPU variants of device-parameterized tests are unaffected.
    )r   SM89OrLaterr~  )r}  s    r   xfailCUDAIfSM89OrLaterOnWindowsr    s     7A[)'2]V]]r   r   z2Requires CUDA {}.{} to match Tritons ptxas versionz(CUDA should not be initialized on import)gh㈵>){__doc__r  r   
torch.cuda$torch.testing._internal.common_utilsr   r   r   r   r   r   r	   r  
contextlibra   r]  r/   is_initialized"CUDA_ALREADY_INITIALIZED_ON_IMPORTr   TEST_MULTIGPUr   r   r   r   r   SM53OrLaterSM60OrLaterSM70OrLaterSM75OrLaterr{   r  r   r   rk  rM   rq  r\  rd  r   IS_SM10Xrn  contextmanagerrY   rg   rl   ro   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numbar0   TEST_NUMBA_CUDAImportErrorr:  OSErrorr   r   r   r   r   r  r  r)  r   r4  r?  rA  rC  TEST_CUSPARSE_GENERICTEST_HIPSPARSE_GENERICoptimSGDrT  floatrZ  ra  re  rg  ri  rl  ro  rr  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<(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
 i j }~	
d
e
d
efgdedeA A=:D

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:: < :; ;=  .4EKKOOfj >$ !'ekk%++//lp #CCHH
DS,^  .x5==3D3D/D 0d38==3D3D3L0d3J3LOc3c/j/c/j/j  mA  0B C  *zz  "GHH # * w/ 
s    O OO