Ë
    ÌÓ^j™V  ã                   ó²  — 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   € ¤Ò!w¬u¯~©~×/CÑ/C×/QÑ/QÔRW×R^ÑR^Ð_aÔjuÔRvÓ/w€ r   c                  ób   — t         r(t        j                  j                  j	                  «       S dS )Nr   )Ú
TEST_CUDNNr   r   r   Úversionr   r   r   r   r      s    € Åz¤U§^¡^×%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¤Ñ%W´e·m±m×6GÑ6G×6MÑ6MÈcÓ6RÐSUÐTUÐ6VÔ%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      ó+   € œeŸj™j×5Ñ5Ó7Òh¼E¿J¹J×<\Ñ<\Ó<^Ð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   $   ó+   € œuŸz™z×6Ñ6Ó8Òj¼U¿Z¹Z×=]Ñ=]Ó=_Ð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¼¿¹×8JÑ8JÐRVÐ8Vò iÜ—J‘J×4Ñ4Ó6¸'ÑAÒgÄcÌ%Ï-É-×J\ÑJ\Ð]_Ð^_ÐJ`ÓFaÐegÑFgò hÜ—J‘J×4Ñ4Ó6¸'ÑAÒfÄcÌ%Ï-É-×J\ÑJ\Ð]_Ð^_ÐJ`ÓFaÐdfÑFfð 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   € œEŸJ™J×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¼¿
¹
×8XÑ8XÓ8ZÐ^dÑ8d€ 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¼¿¹×9YÑ9YÓ9[Ð_fÑ9f€ 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¼¿¹×9YÑ9YÓ9[Ð\]Ñ9^ÐbdÑ9d€ 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Ü	‡N�N×Ñ×.Ñ.¨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   ø€ Ü�:‰:×"Ñ"Ô$ØÜ—J‘J×4Ñ4°VÓ<×HÑH€MÜ—Z‘Z—^‘^Ð$UÐWdÓe€Nô Ó<°)Ô<Ó<Ð<r   c                  ó   — t        ddg«      S )NÚgfx942Úgfx950©rg   r   r   r   ÚCDNA3OrLaterrl   C   s   € Ü# X¨xÐ$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ˆ	Ü�:‰:�>‰>ÐCÀSÓIÈSÒPØÒPÑPˆIÜ'¨	Ó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ˆ	Ü�:‰:�>‰>ÐCÀSÓIÈSÒPØÒPÑPˆIÜ'¨	Ó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   € ÜÐÒQ¤KÒQÔ5GÈ5Ñ5PÐ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ÑFÄuÇ{Á{×G^ÑG^×GbÑGbÓc€NØÐØÜˆ~×#Ñ# 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   ~   ó
   € Ô2TÓ2V€ 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                  «       r°t         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·   € Ü‡z�z×ÑÔ Ü�=‰=×ÒØ�IˆEÜ˜vÒ%Ø—‘˜h˜ZÔ(Ü˜vÒ%Ø—‘˜WÔ%Øò  �Øœ5Ÿ:™:×;Ñ;¸AÓ>×JÑJÒJÙð ð äÒN¤%§*¡*×"BÑ"BÓ"DÈÑ"NÐNÜ‡y�y×ÑÔØàr   c                  óF  — t         j                  j                  «       rƒt         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Š   € Ü‡z�z×ÑÔ Ü�=‰=×ÒØ¤×!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   € Ü‡z�z×ÑÔ Ü�=‰=×ÒÜ˜vÒ%Ø¤5§:¡:×#CÑ#CÀAÓ#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   € Ü‡z�z×ÑÔ ¬¯©×):Ò):Ø$¬×(8Ñ(8×(=Ñ(=Ó(?Ð?ˆØÒ+¤8Ð+Ør   c                  óÊ  — t         j                  j                  «       rÅt         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    € Ü‡z�z×ÑÔ Ü�=‰=×ÒØœuŸz™z×?Ñ?ÀÓB×NÑNÐNÐNô ÒL¤§
¡
× @Ñ @Ó BÀfÑ Lò ?Ü—N‘N×-Ñ-×:Ñ:Ó<ò?ä—N‘N×-Ñ-×5Ñ5Ó7¸3Ñ>ðð
 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
   € Ô5ZÓ5\€ 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   € åÜÐ^Ó_Ð_Ý%ä”u—z‘z×.Ñ.Ó0Ó1ò 	/ˆAÜ�K‰K˜ 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±   è ø€ ä!ŸN™N×/Ñ/×6Ñ6×AÑAÐðFØ05Œ�‰×Ñ×"Ñ"Ô-Ü�^‰^×!Ñ!×'Ñ'°ÀÐTXÐejÐ'Ó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Î   è ø€ ä!ŸN™N×/Ñ/×6Ñ6×AÑAÐØ—N‘N€Mð'Ø04Œ�‰×Ñ×"Ñ"Ô-Ø'ˆŒÜ�^‰^×!Ñ!×'Ñ'°ÀÐTXÐeiÐ'Ó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¸   è ø€ ô "ŸN™N×/Ñ/×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¬   ù€ à�M‰Mœ#˜i¨°eÔ<Ô=Ü—:‘:×/Ñ/Ó1Ò=°gˆDØ˜6Ñ!ØÒO¤§¡¨f°XÑ.>Ó!?×!DÑ!DÈÑ!N�Ø˜&Ñ ØÒU ¨¡´U·]±]ÄEÇOÁOÐ4TÐ!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   è ø€ Ò8˜A”�Q—Ñ8ùr%   r&   )r   rÁ   rÂ   ÚfindÚlenr*   r)   )ÚpositionÚversion_strs     r   Ú_get_magma_versionr)  x  s�   € Ø”e×&Ñ&×+Ñ+Ó-Ñ-ØÜ×Ñ×$Ñ$Ó&×+Ñ+¨HÓ5€HÜ×"Ñ"×'Ñ'Ó)¨(´S¸³]Ñ*BÐ*CÐD×JÑJÈ4ÓPÐQRÑS€KÜÑ8 ×!2Ñ!2°3Ó!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>ƒ  ó   è ø€ Ò9˜A”�Q—Ñ9ùr%   r&   )r   r   r/   Ústrr)   r*   )Úcuda_versions    r   r‰   r‰     sE   € Ü‡}�}×ÑÐ!ØÜ”u—}‘}×)Ñ)Ó*€LÜÑ9 ×!3Ñ!3°CÓ!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Ø×%Ñ% c°AÐ%Ó6°qÑ9€LÜÑ9 ×!3Ñ!3°CÓ!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   € ÝØðô —h‘h×8Ñ8Ó:ˆØÐ +°Ò"2ØØ˜uÑ$ˆØ˜uÑ$¨Ñ,ˆØ˜cÑ!ˆØ�u˜eÐ$Ð$øÜœ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   € ÜÒ;œeŸm™m×/Ñ/°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  € ô —(‘(×%Ñ%¤e§h¡h§o¡o°a¸Ó&;¼U¿X¹X¿_¹_ÈQÐPQÓ=RÓS×VÑVÐ^dÐVÓe€KÜ—(‘(×%Ñ%¤e§h¡h§o¡o°a¸Ó&;¼U¿X¹X¿_¹_ÈQÐPQÓ=RÓS×VÑVÐ^dÐVÓe€KÜ	�‰‹ñ Ü˜×.Ñ.Ó0°+×2HÑ2HÓ2JÐSWÔXò 	‰DˆAˆqØ�G‰G�A�Jñ	÷ð �Cˆ[€FØÐ#Ø�‰Ð&Ô'Ù  ×!7Ñ!7Ó!9ÑD¸VÑD€KÙ  ×!7Ñ!7Ó!9ÑD¸VÑD€Kà˜ [°+Ð=Ð=÷ð ú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â   € Ü�[‰[˜ u°VÔ<¼e¿k¹kÈ&ÐX]ÐflÔ>mÐnÜ�[‰[˜ u°VÔ<¼e¿k¹kÈ&ÐX]ÐflÔ>mÐnÜ�[‰[˜ u°VÔ<¼e¿k¹kÈ&ÐX]ÐflÔ>mÐnÜ�[‰[˜ u°VÔ<¼e¿k¹kÈ&ÐX]ÐflÔ>mÐnðp€Dô
 �h‰h×ÑÓ ×#Ñ# FÓ+€Gà€Iä,Ø nÐGWôà	ˆw˜	Ð"ñ#ð #r   c                 ó<   — t         s| S t        j                  | «      S r   )ÚIS_SM89ÚunittestÚexpectedFailure©Úfuncs    r   ÚxfailIfSM89ra  È  ó   € Ýˆ4ÐB¤H×$<Ñ$<¸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ÐC¤X×%=Ñ%=¸dÓ%CÐCr   c                 óH   — t         st        s| S t        j                  | «      S r   )r   Ú	IS_JETSONr]  r^  r_  s    r   ÚxfailIfDistributedNotSupportedrr  å  s   € Ý ¥Iˆ4ÐR´H×4LÑ4LÈTÓ4RÐ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Ð)˜$Ò) &Ò)ô ˜u W×%5Ñ%5Ð$6Ð6ZÐ[Ó\Ð\øô ò 	JØ�L‰LÐ0°×1AÑ1AÐ0BÀ"ÀQÀCÐ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 )zÕMark 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)gñhãˆµøä>){Ú__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ã Û Û ß × Ñ Û Û Û 	Û ð &+§Z¡Z×%>Ñ%>Ó%@Ð "ð Ò<˜eŸj™j×5Ñ5Ó7¸1Ñ<€Ù(1ˆlˆe�l‰l˜8Ô$°t€áÙÑ*Ó+�JáÑwÓx€JáÑ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ñ +2Ñ2fÓ*gÐ ! 4Ó gÙ,3Ñ4lÓ,mÐ # TÓ mÙ*1Ñ2fÓ*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ñ ,3Ñ3hÓ+iÐ " DÓ iòò&òòò
ñ #*Ñ*VÓ"WÐ ˜4Ó WÙ%Ñ&NÓOÐ �tÓ OÙ%,Ñ-\Ó%]Ð ˜dÓ ]Ù+2Ñ3hÓ+iÐ " DÓ iÙ-4Ñ5lÓ-mÐ $ dÓ máðÛØŸ*™*×1Ñ1Ó3‰ð
 €Oð #Ð ò*ð ×ÑñFó ðFð ×Ñò
'ó ð
'ð ×ÑñFó ðFðX°Dô òFò9ò:ò:òò ò<ñ :Ó;Ð Ù;Ó=Ð ð .4ÀEÇKÁKÇOÁOÐfjó >ð$ !'¨e¯k©kÈ%Ï+É+Ï/É/Ðlpó #òCòCòòHòHò
DòSòò,^ð Ð Ø.˜xŸ™°5·=±=×3DÑ3DÐ/Dò 0dØ38·=±=×3DÑ3DÈÐ3Lò0dá3JÓ3LÐOcÑ3cØ/jÐ/c×/jÑ/jð  mAð  0Bó CÐ ñ *Ø‡z�z× Ñ Ô"ÙÐGÓHÐHð #ð *øð ˜ wÐ/ò ØˆØ‹
ðús   É O ÏOÏO