§
     ’Þjú<  ã                   ó°  — d Z ddlZddlZddlZddlmZ ddlmZ ddl	m
Z
  e
ddd¦  «        ZdZej        Zd	„ Zd
„ ZdZd,d„Z ej        g d¢¬¦  «        dej        dej        dej        dej        dej        dej        dej        dej        dej        dej        dej        dej        fd„¦   «         ZdZ ej        dg¬¦  «        dej        dej        d ej        fd!„¦   «         Zd"„ Zd#„ Zd-d$„Zd%„ Z ej        g d&¢¬¦  «        dej        dej        d'ej        dej        d(ej        dej        dej        fd)„¦   «         Zd,d*„Zdd+lmZ  e¦   «          dS ).a~  Triton implicit-GEMM 3D / 2D convolution on channels-last 16-bit activations ("fast" kernel precision), shared by the Wan VAEs and the
streamed LTX VAE: FP32 accumulation, the bias and an optional residual added as PyTorch rounds them, optionally the nearest x2 upsampling
of the input read in place (Wan decoder). 1.1-2x faster than cuDNN on RTX 50 GPUs, where it is enabled; only the accumulation order
differs. One binary per channel count, kernel shape, mode and dtype; sizes and paddings are not specialized. Launch configurations start
from a static choice and are tuned once per identity and GPU on scratch outputs of a large call.é    N)Údispatch)ÚLaunchTunerÚwan2gp_conv3d_tritoné   ÚVAEi   c                 ód   — t          j        t          j        | ¦  «        t          j        ¦  «        S )z‚NVIDIA GPUs with compute capability 8.0+ (the 16-bit kernels' floor, shared.kernels.dispatch); other devices use the PyTorch path.)r   Údevice_supportedÚtorchÚdeviceÚbfloat16©r   s    úDC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\conv3d_triton.pyÚ	supportedr      s!   € åÔ$¥U¤\°&Ñ%9Ô%9½5¼>ÑJÔJÐJó    c                 óŠ   — t          | ¦  «        o4t          j        t          j        | ¦  «        j        ¦  «        d         dk    S )zÊTriton convolutions: measured faster than cuDNN on compute capability 12.x GPUs only (GeForce RTX 50; 0.7-1.3x slower on 8.9), so
    other GPUs keep cuDNN. A performance choice, not an eligibility one.r   é   )r   r   Ú
capabilityr
   r   Úindexr   s    r   Úconv_supportedr      s;   € õ �VÑÔÐY¥Ô!4µU´\À&Ñ5IÔ5IÔ5OÑ!PÔ!PÐQRÔ!SÐWYÒ!YÐYr   i    Fc                 ó*  — |j         j        dd…         \  }}|r-| j        d         | j        d         z  | j        d         z  dz  }nj| j        d         | j        d         |j        d         z
  dz   z  | j        d         |j        d         z
  dz   z  | j        d         |j        d         z
  dz   z  }| j        t          j        k    oN|dz  dk    oE|dz  dk    o<t          d„ |j        D ¦   «         ¦  «        o|t          k    ot          | j
        ¦  «        S )	zwTriton convolution for conv (stride 1, channels multiple of 32) on the 16-bit channels-last input x of a validated GPU.Né   r   é   é   r   é    c              3   ó"   K  — | ]
}|d k    V — ŒdS )r   N© )Ú.0Úss     r   ú	<genexpr>z conv_eligible.<locals>.<genexpr>*   s*   è è € ÐPmÐPmÐ\]ÐQRÐVWÒQWÐPmÐPmÐPmÐPmÐPmÐPmr   )ÚweightÚshapeÚkernel_sizeÚdtyper
   Úfloat32ÚallÚstrideÚ_CONV_MIN_PIXELSr   r   )ÚxÚconvÚupsampleÚcoutÚcinÚpixelss         r   Úconv_eligibler.   #   sk  € à”Ô! " 1 "Ô%�I€Dˆ#Øð WØ”˜”˜aœg aœjÑ(¨1¬7°1¬:Ñ5¸Ñ9ˆˆà”˜”˜qœw qœz¨DÔ,<¸QÔ,?Ñ?À!ÑCÑDÈÌÐPQÌ
ÐUYÔUeÐfgÔUhÑHhÐklÑHlÑmÐqrÔqxÐyzÔq{ð  Cô  Oð  PQô  Rñ  rRð  UVñ  rVñ  WˆØŒ7•e”mÒ#ð  j¨¨b©°Aªð  j¸$À¹)Àqº.ð  jÍSÐPmÐPmÐaeÔalÐPmÑPmÔPmÑMmÔMmð  jÐrxõ  }Mò  sMð  jõ  R`ð  abô  aiñ  Rjô  Rjð  jr   )ÚMÚTÚHÚWÚTPÚHPÚWP)Údo_not_specializeÚCINÚCOUTÚKTÚKHÚKWÚUPSAMPLEÚRESIDUALÚBLOCK_MÚBN1ÚBN2ÚBLOCK_KÚGROUP_Mc                 óœ  — ||z   }t          j        d¦  «        }t          j        ||¦  «        }||z  }|||z  z  |z  }t          j        ||z
  |¦  «        }||||z  z  |z  z   }|||z  z  |z  }|                     t           j        ¦  «        |z  t          j        d|¦  «        z   } | |k     }!| |z  }"| |z  }#|#|z  }$|#|z  }#|#|z  }%|#|z  }&||z  t          j        d|¦  «        z   }'||z  |z   t          j        d|dk    r|n|¦  «        z   }(t          j        d|¦  «        })t          j        ||ft           j        ¬¦  «        }*t          j        ||dk    r|n|ft           j        ¬¦  «        }+|�r‰t          j	        d¦  «        D �]q},|$|,z   dz
  }-t          j	        d¦  «        D �]P}.|"|.z   dz
  }/|!|-dk    z  |-|k     z  |/dk    z  |/|k     z  }0|&|
z  |-dz  z   |z  |/dz  z   |z  }1t          d||¦  «        D ]þ}2t          j        | |1dd…df         z   |2|)z   ddd…f         z   |0dd…df         d¬¦  «        }3||,dz  |.z   |z  |2z   |)z   ddd…f         z   }4t          j        |3t          j        t          j        |4|'d	z  |z  dd…df         z   ¦  «        ¦  «        |*¦  «        }*|dk    rMt          j        |3t          j        t          j        |4|(d	z  |z  dd…df         z   ¦  «        ¦  «        |+¦  «        }+Œÿ�ŒR�Œs�n›|&|	z  |%z   |
z  |$z   |z  |"z   |z  }5t          j	        |¦  «        D �]n}6t          j	        |¦  «        D �]U},|5|6|
z  |,z   |                     t           j        ¦  «        z  |z  z   }1t          d||z  |¦  «        D �]}2t          j        | |1dd…df         z   |2|)z   ddd…f         z   |!dd…df         d¬¦  «        }3||6|z  |,z   |z  |z  |2z   |)z   ddd…f         z   }4t          j        |3t          j        t          j        |4|'||z  |z  |z  z  dd…df         z   ¦  «        ¦  «        |*¦  «        }*|dk    rSt          j        |3t          j        t          j        |4|(||z  |z  |z  z  dd…df         z   ¦  «        ¦  «        |+¦  «        }+�Œ�ŒW�Œp|j        j        }7|*                     |7¦  «                             t           j        ¦  «        t          j        ||'z   ¦  «                             t           j        ¦  «        ddd…f         z   }8|rŽ|8                     |7¦  «                             t           j        ¦  «        t          j        || |z  dd…df         z   |'ddd…f         z   |!dd…df         d¬¦  «                             t           j        ¦  «        z   }8t          j        || |z  dd…df         z   |'ddd…f         z   |8                     |7¦  «        |!dd…df         ¦  «         |dk    �rS|+                     |7¦  «                             t           j        ¦  «        t          j        ||(z   ¦  «                             t           j        ¦  «        ddd…f         z   }8|rŽ|8                     |7¦  «                             t           j        ¦  «        t          j        || |z  dd…df         z   |(ddd…f         z   |!dd…df         d¬¦  «                             t           j        ¦  «        z   }8t          j        || |z  dd…df         z   |(ddd…f         z   |8                     |7¦  «        |!dd…df         ¦  «         dS dS )
a&  Implicit-GEMM convolution with FP32 accumulation of BLOCK_M output pixels x BN1 + BN2 output channels (two products when the
    channels per program are no power of two) of the channels-last output (batch, T, H, W, COUT). Weights WT channels last (COUT, KT, KH,
    KW, CIN): rows of K contiguous per output channel, as cuDNN's channels-last convolution converts them.
    Not UPSAMPLE: X is the padded input (batch, TP, HP, WP, CIN) of a valid convolution; the KW taps of an input row and their channels are
    one contiguous K range. UPSAMPLE: X (batch, HP, WP, CIN), T = TP = 1, a 3x3 convolution with zero padding 1 of its nearest x2
    upsampling (H = 2 HP, W = 2 WP), read in place. As PyTorch: the convolution rounded to the output dtype, plus the bias, rounded, plus
    RES (RESIDUAL), rounded.r   ©r#   r   r   r   Nç        ©Úotheré	   )ÚtlÚ
program_idÚcdivÚminimumÚtoÚint64ÚarangeÚzerosr$   Ústatic_rangeÚrangeÚloadÚdotÚtransr#   Ú
element_tyÚstore)9ÚXÚWTÚBIASÚRESÚYr/   r0   r1   r2   r3   r4   r5   r7   r8   r9   r:   r;   r<   r=   r>   r?   r@   rA   rB   ÚBNÚpidÚnum_mÚnum_nÚfirst_mÚsize_mÚpid_mÚpid_nÚmÚrowsÚwÚrÚhÚtÚbÚn1Ún2ÚkÚacc1Úacc2ÚkhÚuhÚkwÚuwÚinsideÚrowÚk0ÚaÚweightsÚbaseÚktr#   Úys9                                                            r   Ú_conv_kernelr}   -   sì  € ð ˜S‘y€BÝ
Œ-˜Ñ
Ô
€CÝŒG�A�wÑÔ€EØ "™*€EØ�w ‘Ñ'¨7Ñ2€GÝŒZ˜ ™¨Ñ1Ô1€FØ�s˜g¨™oÑ.°&Ñ8Ñ8€EØ�G˜e‘OÑ$¨Ñ/€EØ�Š•”ÑÔ˜WÑ$¥r¤y°°GÑ'<Ô'<Ñ<€AØˆqŠ5€DØ	ˆA‰€AØ	ˆQ‰€AØ	ˆA‰€AØ	ˆQ‰€AØ	ˆA‰€AØ	ˆQ‰€AØ	�‰•b”i  3Ñ'Ô'Ñ	'€BØ	�‰�cÑ	�BœI a°°a²°¨¨¸SÑAÔAÑ	A€BÝ
Œ	�!�WÑÔ€AÝŒ8�W˜c�N­"¬*Ð5Ñ5Ô5€DÝŒ8�W S¨1¢W W˜c˜c°#Ð6½b¼jÐIÑIÔI€DØñ rÝ”/ !Ñ$Ô$ð 	eñ 	eˆBØ�R‘˜!‘ˆBÝ”o aÑ(Ô(ð 	eñ 	e�Ø˜‘V˜a‘Z�Ø  q¢Ñ)¨R°!ªVÑ4¸¸aºÑ@ÀBÈÂFÑK�Ø˜B™  q¡Ñ(¨BÑ.°°q±Ñ8¸CÑ?�Ý  3¨Ñ0Ô0ð eð e�BÝœ  C¨¨¨¨4¨¤LÑ 0°B¸±F¸DÀ!À!À!¸GÔ3DÑ DÀfÈQÈQÈQÐPTÈWÄoÐ]`ÐaÑaÔa�AØ  R¨!¡V¨b¡[°CÑ$7¸"Ñ$<¸qÑ$@À$ÈÈÈÀ'Ô#JÑJ�GÝœ6 !¥R¤X­b¬g°gÀÀaÁÈ#ÁÈqÈqÈqÐRVÈwÔ@WÑ6WÑ.XÔ.XÑ%YÔ%YÐ[_Ñ`Ô`�DØ˜Q’w�wÝ!œv a­¬µ"´'¸'ÀRÈ!ÁVÈcÁ\ÐSTÐSTÐSTÐVZÐSZÔD[Ñ:[Ñ2\Ô2\Ñ)]Ô)]Ð_cÑdÔd˜øñeñ		eñ	eð �b‘&˜1‘* Ñ" QÑ&¨"Ñ,¨qÑ0°CÑ7ˆÝ”/ "Ñ%Ô%ð 	rñ 	rˆBÝ”o bÑ)Ô)ð rñ r�Ø˜b 2™g¨™l¨b¯eªeµB´H©o¬oÑ=ÀÑCÑC�Ý  2¨¡8¨WÑ5Ô5ð rñ r�BÝœ  C¨¨¨¨4¨¤LÑ 0°B¸±F¸DÀ!À!À!¸GÔ3DÑ DÀdÈ1È1È1ÈdÈ7ÄmÐ[^Ð_Ñ_Ô_�AØ  R¨"¡W¨r¡\°RÑ$7¸#Ñ$=ÀÑ$BÀQÑ$FÈÈaÈaÈaÈÔ#PÑP�GÝœ6 !¥R¤X­b¬g°gÀÀrÈBÁwÐQSÁ|ÐVYÑGYÑAZÐ\]Ð\]Ð\]Ð_cÐ\cÔ@dÑ6dÑ.eÔ.eÑ%fÔ%fÐhlÑmÔm�DØ˜Q’w�wÝ!œv a­¬µ"´'¸'ÀRÈ2ÐPRÉ7ÐUWÉ<ÐZ]ÑK]ÑE^Ð`aÐ`aÐ`aÐcgÐ`gÔDhÑ:hÑ2iÔ2iÑ)jÔ)jÐlpÑqÔq˜ùñrñrð ŒGÔ€EØ�Š�‰Œ×Ò�"œ*Ñ%Ô%­¬°°r±	Ñ(:Ô(:×(=Ò(=½b¼jÑ(IÔ(IÈ$ÐPQÐPQÐPQÈ'Ô(RÑR€AØð CØ�DŠD�‰KŒK�NŠN�2œ:Ñ&Ô&­¬°¸¸D¹À!À!À!ÀTÀ'Ô7JÑ1JÈRÐPTÐVWÐVWÐVWÐPWÌ[Ñ1XÐZ^Ð_`Ð_`Ð_`ÐbfÐ_fÔZgÐorÐ)sÑ)sÔ)s×)vÒ)vÕwyô  xBñ  *Cô  *Cñ  CˆÝ„HˆQ�!�d‘(˜A˜A˜A˜t˜GÔ$Ñ$ r¨$°°°¨'¤{Ñ2°A·D²D¸±K´KÀÀaÀaÀaÈÀgÄÑOÔOÐOØ
ˆQ‚w�wØ�GŠG�E‰NŒN×Ò�bœjÑ)Ô)­B¬G°D¸2±IÑ,>Ô,>×,AÒ,AÅ"Ä*Ñ,MÔ,MÈdÐTUÐTUÐTUÈgÔ,VÑVˆØð 	GØ—’�U‘”—’�rœzÑ*Ô*­R¬W°S¸AÀ¹HÀaÀaÀaÈÀgÔ;NÑ5NÐQSÐTXÐZ[ÐZ[ÐZ[ÐT[ÔQ\Ñ5\Ð^bÐcdÐcdÐcdÐfjÐcjÔ^kÐsvÐ-wÑ-wÔ-w×-zÒ-zÕ{}ô  |Fñ  .Gô  .Gñ  GˆAÝ
Œ��a˜$‘h    4 Ô(Ñ(¨2¨d°A°A°A¨g¬;Ñ6¸¿º¸U¹¼ÀTÈ!È!È!ÈTÈ'Ä]ÑSÔSÐSÐSÐSð	 €wr   i  @ ÚTAPSÚBLOCK_CÚBLOCK_Tc           
      ó  — t          j        d¦  «                             t           j        ¦  «        |z  |z  }t          j        d¦  «        |z  t          j        d|¦  «        z   }t          j        d|¦  «        }||k     dd…df         ||k     ddd…f         z  }	t          j        ||z   ||z  ddd…f         z   |dd…df         z   t          j        | |z   ||z  dd…df         z   |ddd…f         z   |	¦  «        |	¦  «         dS )zBLOCK_C input channels of one output channel of a convolution weight (cout, cin, taps) written channels last (cout, taps, cin).r   r   N)rI   rJ   rM   rN   rO   rW   rS   )
r2   ÚOUTr7   r~   r   r€   rz   Úcrj   Úmasks
             r   Ú_weight_kernelr…   t   s  € õ Œ=˜ÑÔ×Ò�rœxÑ(Ô(¨3Ñ.°Ñ5€DÝ
Œ�aÑÔ˜7Ñ"¥R¤Y¨q°'Ñ%:Ô%:Ñ:€AÝ
Œ	�!�WÑÔ€AØ�ŠG�Q�Q�Q˜�WÔ  T¢¨4°°°¨7Ô 3Ñ3€DÝ„HˆS�4‰Z˜1˜s™7 D¨!¨!¨! GÔ,Ñ,¨q°°°°D°¬zÑ9½2¼7À1ÀtÁ8ÈqÐSWÉxÐYZÐYZÐYZÐ\`ÐY`ÔNaÑCaÐdeÐfjÐlmÐlmÐlmÐfmÔdnÑCnÐptÑ;uÔ;uÐw{Ñ|Ô|Ð|Ð|Ð|r   c           
      óv  — | d                               ¦   «         }|                      ¦   «         dk    rt          nt          j        }|dk    sK|                      ¦   «         dk    s3|                       ¦   «         t
          k     s|                      |¬¦  «        r|                      |¬¦  «        S t          j        | |¬¦  «        }t          | j
        d         t          j        | j
        d         d¦  «        f         |                      ¦   «         || j
        d         |dt          j        |¦  «        d¬¦  «         |S )	a  weight (cout, cin, *kernel) in the channels-last layout of the channels-last convolutions: large 3D weights by a tiled transpose,
    2.5x faster than PyTorch's strided copy (tiled decodes convert every weight once per tile and chunk); small ones are launch-bound.)r   r   é   r   ©Úmemory_formatr   é@   r   )Ú	num_warps)ÚnumelÚdimÚ_CLr
   Úchannels_lastÚ_WEIGHT_KERNEL_ELEMENTSÚis_contiguousÚ
contiguousÚ
empty_liker…   r!   ÚtritonrK   Únext_power_of_2)r    ÚtapsÚlayoutÚouts       r   Úchannels_last_weightr™   ~   s8  € ð �$Œ<×ÒÑÔ€DØ—J’J‘L”L AÒ%Ð%�SˆS­5Ô+>€FØˆq‚y€y�F—J’J‘L”L AÒ%Ð%¨¯ª©¬Õ:QÒ)QÐ)QÐU[×UiÒUiÐx~ÐUiÑUÔUÐ)QØ× Ò ¨vÐ Ñ6Ô6Ð6Ý
Ô
˜6°Ð
8Ñ
8Ô
8€CÝ�F”L ”O¥V¤[°´¸a´À"Ñ%EÔ%EÐFÔGÈ×HYÒHYÑH[ÔH[Ð]`ÐbhÔbnÐopÔbqÐswÐy{õ  ~Dô  ~Tð  UYñ  ~Zô  ~Zð  fgð  hñ  hô  hð  hØ€Jr   c                 óè   — | dz  dk    rdn| dz  dk    rdn| dz  dk    rdnd}d	g|¢d
‘d‘d‘R d	g|¢d
‘d‘d‘R dg|¢d
‘d‘d‘R g}| d	z  dk    r|                      d¦  «         t          |¦  «        S )zÅLaunch configurations (BLOCK_M, BN1, BN2, BLOCK_K, warps, stages), the first being the static choice: output channels by 64
    (96 = 64 + 32: no power-of-two padding), and by 128 when they divide.rŠ   r   )rŠ   r   é`   )rŠ   r   é    )é€   r   )r   r   r�   r   r   r   )r�   r�   r   r   r   r   )ÚappendÚtuple)r+   ÚsplitÚ
candidatess      r   Ú_conv_candidatesr¢   Š   sÍ   € ð ˜b‘y A’~�~ˆGˆG°t¸b±yÀA²~°~¨8¨8ÐX\Ð_bÑXbÐfgÒXgÐXgÈ9È9Ðmt€EØÐ)˜Ð) Ð) AÐ) qÐ)Ð)¨CÐ+B°%Ð+B¸Ð+B¸QÐ+BÀÐ+BÐ+BÀRÐDZÈ%ÐDZÐQSÐDZÐUVÐDZÐXYÐDZÐDZÐ[€JØˆc�z�Q‚€Ø×ÒÐ1Ñ2Ô2Ð2Ý�ÑÔÐr   c                 ó¤  ‡ ‡‡‡‡‡‡‡‡‡‡‡‡‡‡‡‡‡— |j         dd…         \  ŠŠ‰                      ‰rt          j        nt          ¬¦  «        Š ‰€dn‰                     t          ¬¦  «        Š‰rT‰ j         \  }}ŠŠdddddd‰z  d‰z  f\  ŠŠŠŠŠŠŠt          j        |‰‰‰f‰ j        ‰ j        t          j        ¬¦  «        }ne‰ j         \  }}ŠŠŠ|j         dd…         \  ŠŠŠ‰‰z
  dz   ‰‰z
  dz   ‰‰z
  dz   cŠŠŠt          j        |‰‰‰‰f‰ j        ‰ j        t          ¬¦  «        }|‰z  ‰z  ‰z  Št          |¦  «        Š‰fˆˆˆˆˆˆˆˆˆˆˆˆˆˆˆ fd„	Šˆˆˆˆˆ fd„}‰› d	‰› d	‰› d
‰› d
‰› d	‰rdnd› d	‰�dnd› d	‰ j        › �}	t          ‰¦  «        }
 ‰|‰t                               d|	‰ j        |
|‰‰z  t          k    ¦  «        ¦  «         |S )af  Convolution (stride 1) of the padded channels-last input x (batch, cin, TP, HP, WP) by weight (cout, cin, kt, kh, kw), or with
    upsample of the nearest x2 upsampling of x (frames, cin, h, w) by a 3x3 weight with zero padding 1, plus bias (and residual, the
    output's shape and layout), as the PyTorch path rounds them. Returns the channels-last output.Nr   rˆ   r   r   ©r#   r   r‰   c                 ó¾   •— |\  }}}}}}	t          t          j        ||¦  «        ‰||z   z  z  f         ‰‰‰
|€| n|| |‰‰‰‰‰‰‰‰‰‰‰‰|d u||||d||	¬¦  «         d S ©Né   )r‹   Ú
num_stages)r}   r”   rK   )ÚtargetÚcountÚconfigÚresÚblock_mÚbn1Úbn2Úblock_kÚwarpsÚstagesÚbiasr,   r+   ri   Úhprq   r{   rs   rj   Útpr*   rg   ÚwpÚwtr(   s             €€€€€€€€€€€€€€€r   Úlaunchzconv.<locals>.launch§   sÝ   ø€ Ø4:Ñ1ˆ��c˜7 E¨6Ý•f”k %¨Ñ1Ô1°T¸cÀC¹iÑ5HÑIÐKÔLÈQÐPRÐTXÐdgÐdoÐZ`ÐZ`Ðuxð  {Að  CHð  JKð  MNð  PQØMOÐQSÐUWÐY\Ð^bÐdfÐhjÐlnÐpxÐz}ð  FJð  {JØMTÐVYÐ[^Ð`gÐijÐv{ð  IOð	Pñ 	Pô 	Pð 	Pð 	Pð 	Pr   c                  ó€   •‡ ‡— t          ‰d¦  «        Š t          j        ‰ ‰z  ‰j        ‰j        ¬¦  «        Šˆ ˆˆˆfd„S )Né   ©r#   r   c                 ó    •—  ‰‰‰| ‰¦  «        S ©Nr   )r«   rª   r¸   Úresidualr©   s    €€€€r   ú<lambda>z'conv.<locals>.scratch.<locals>.<lambda>°   s   ø€ ˜f˜f V¨U°F¸HÑEÔE€ r   ©Úminr
   Úemptyr#   r   )rª   r©   r+   r¸   r-   r¾   r(   s   @@€€€€€r   Úscratchzconv.<locals>.scratch­   sJ   øøø€ Ý�F˜GÑ$Ô$ˆÝ”˜U T™\°´ÀÄÐJÑJÔJˆØEÐEÐEÐEÐEÐEÐEÐEr   ú|r(   r*   Úpaddedr¾   Ú r)   )r!   r’   r
   r�   rŽ   rÂ   r#   r   r™   r¢   Ú_TUNERr«   Ú_TUNE_ELEMENTS)r(   r    r³   r¾   r*   ÚbatchÚ_r˜   rÃ   Úidentityr¡   r,   r+   ri   r´   rq   r{   rs   r¸   r-   rj   rµ   rg   r¶   r·   s   ` ```      @@@@@@@@@@@@@@r   r)   r)   ”   s  øøøøøøøøøøøøøøøøøø€ ð ”˜R˜a˜RÔ �I€Dˆ#Ø	�Š¸(Ð#K¥5Ô#6Ð#6ÍˆÑLÔL€AØÐ'ˆtˆt¨X×-@Ò-@ÍsÐ-@Ñ-SÔ-S€HØð eØœ7Ñˆˆq�"�bØ"# Q¨¨1¨a°°R±¸¸R¹Ð"?ÑˆˆB��B˜˜1˜aÝŒk˜5 $¨¨1Ð-°Q´WÀQÄXÕ]bÔ]pÐqÑqÔqˆˆà œwÑˆˆq�"�b˜"Ø”\ ! " "Ô%‰
ˆˆB�Ø�r‘'˜A‘+˜r B™w¨™{¨B°©G°a©Kˆˆˆ1ˆaÝŒk˜5 $¨¨1¨aÐ0¸¼ÈÌÕ`cÐdÑdÔdˆØ�Q‰Y˜‰]˜QÑ€FÝ	˜fÑ	%Ô	%€Bà*2ð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð PðFð Fð Fð Fð Fð Fð Fð Fð Fð
 ð  Lð  L˜ð  Lð  L˜rð  Lð  L Bð  Lð  L¨ð  Lð  L¸HÐ.R¨j¨jÈ(ð  Lð  LÐckÐcwÐU_ÐU_Ð}ð  Lð  Lð  CDô  CJð  Lð  L€HÝ! $Ñ'Ô'€JØ
€Fˆ3��Ÿš f¨h¸¼À*ÈgÐW]Ð`dÑWdÕhvÒWvÑwÔwÑxÔxÐxØ€Jr   c                 ó‚  — | j         t          j        k    oª| j        o£|j        dk    o˜|j        dk    o�|j        dk    o‚|j        dk    ow|j        dk    ol|j	        dk    oa|j
        dz  dk    oS|j        dz  dk    oE| j        d         | j        d         z  | j        d	         z  t          k    ot          | j        ¦  «        S )
z»Triton 2D convolution for conv (3x3, padding 1 with zeros, stride 1, input channels multiple of 32, output channels of 64) on the
    16-bit input x (batch, cin, h, w) of a validated GPU.)r   r   )r   r   r   rP   r   r   rŠ   r   r   )r#   r
   r$   Úis_cudar"   Úpaddingr&   ÚdilationÚgroupsÚpadding_modeÚin_channelsÚout_channelsr!   r'   r   r   )r(   r)   s     r   Úconv2d_eligiblerÔ   ¼   s	  € ð ŒG•u”}Ò$ð F¨¬ð F°tÔ7GÈ6Ò7Qð FÐVZÔVbÐflÒVlð FÐquÔq|ð  AGò  rGð FØ” Ò'ðFØ,0¬K¸1Ò,<ðFØAEÔARÐV]ÒA]ðFØbfÔbrÐuwÑbwÐ{|Òb|ðFàÔ! BÑ&¨!Ò+ðFà01´¸´
¸Q¼WÀQ¼ZÑ0GÈ!Ì'ÐRSÌ*Ñ0TÕXhÒ0hðFåm{Ð|}ô  }Eñ  nFô  nFðGr   )r/   r1   r2   ÚRELUr]   c                 ó>  — t          j        d¦  «        }t          j        ||
¦  «        }||z  }|||z  z  |z  }t          j        ||z
  |¦  «        }||||z  z  |z  z   }|||z  z  |z  }|                     t           j        ¦  «        |
z  t          j        d|
¦  «        z   }||k     }||z  }||z  }||z  }||z  }||z  t          j        d|¦  «        z   }t          j        d|¦  «        }t          j        |
|ft           j        ¬¦  «        }t          j	        d¦  «        D �]}||z   dz
  }t          j	        d¦  «        D ]ô} || z   dz
  }!||dk    z  ||k     z  |!dk    z  |!|k     z  }"||z  |z   |z  |!z   |z  }#t          d||¦  «        D ]©}$t          j        | |#dd…df         z   |$|z   ddd…f         z   |"dd…df         d¬¦  «        }%t          j        |%t          j        t          j        ||dz  |z  dd…df         z   |dz  | z   |z  |$z   |z   ddd…f         z   ¦  «        ¦  «        |¦  «        }ŒªŒõ�Œ|j        j        }&|                     |&¦  «                             t           j        ¦  «        t          j        ||z   ¦  «                             t           j        ¦  «        ddd…f         z                        |&¦  «        }'|	rt          j        |'d¦  «        }'t          j        |||z  dd…df         z   |ddd…f         z   |'|dd…df         ¦  «         dS )	a
  Implicit-GEMM 3x3 convolution, zero padding 1, FP32 accumulation, of BLOCK_M output pixels x BN output channels: X channels last
    (batch, H, W, CIN), WT (COUT, 3, 3, CIN). As PyTorch: the convolution rounded to the output dtype, plus the bias, rounded, then ReLU.r   rD   r   r   NrE   rF   rH   )rI   rJ   rK   rL   rM   rN   rO   rP   r$   rQ   rR   rS   rT   rU   r#   rV   ÚmaximumrW   )(rX   rY   rZ   r\   r/   r1   r2   r7   r8   rÕ   r>   r]   rA   rB   r^   r_   r`   ra   rb   rc   rd   re   rf   rg   rh   ri   rk   Únrn   Úaccrq   rr   rs   rt   ru   rv   rw   rx   r#   r|   s(                                           r   Ú_conv2d_kernelrÚ   Ä   sm  € õ
 Œ-˜Ñ
Ô
€CÝŒG�A�wÑÔ€EØ "™*€EØ�w ‘Ñ'¨7Ñ2€GÝŒZ˜ ™¨Ñ1Ô1€FØ�s˜g¨™oÑ.°&Ñ8Ñ8€EØ�G˜e‘OÑ$¨Ñ/€EØ�Š•”ÑÔ˜WÑ$¥r¤y°°GÑ'<Ô'<Ñ<€AØˆqŠ5€DØ	ˆA‰€AØ	ˆQ‰€AØ	ˆA‰€AØ	ˆQ‰€AØ�‰
•R”Y˜q "Ñ%Ô%Ñ%€AÝ
Œ	�!�WÑÔ€AÝ
Œ(�G˜R�=­¬
Ð
3Ñ
3Ô
3€CÝŒo˜aÑ Ô ð ñ ˆØ�‰V�a‰ZˆÝ”/ !Ñ$Ô$ð 	ð 	ˆBØ�R‘˜!‘ˆBØ˜R 1šWÑ%¨¨aªÑ0°B¸!²GÑ<ÀÀQÂÑGˆFØ˜‘E˜B‘J !Ñ# bÑ(¨CÑ/ˆCÝ˜A˜s GÑ,Ô,ð ð �Ý”G˜A  A A A t G¤Ñ,°°Q±¸¸a¸a¸a¸Ô/@Ñ@À&ÈÈÈÈDÈÄ/ÐY\Ð]Ñ]Ô]�Ý”f˜Q¥¤­¬°°q¸1±u¸s±{ÀAÀAÀAÀtÀGÔ6LÑ1LÐQSÐVWÑQWÐZ\ÑQ\Ð`cÑPcÐfhÑPhÐklÑPlÐnrÐtuÐtuÐtuÐnuÔOvÑ1vÑ)wÔ)wÑ xÔ xÐz}Ñ~Ô~��ðñ		ð ŒGÔ€EØ	�Š�‰Œ×	Ò	�"œ*Ñ	%Ô	%­¬°°q±Ñ(9Ô(9×(<Ò(<½R¼ZÑ(HÔ(HÈÈqÈqÈqÈÔ(QÑ	Q×UÒUÐV[Ñ\Ô\€AØð ÝŒJ�q˜#ÑÔˆÝ„HˆQ�!�d‘(˜A˜A˜A˜t˜GÔ$Ñ$ q¨¨q¨q¨q¨¤zÑ1°1°d¸1¸1¸1¸d¸7´mÑDÔDÐDÐDÐDr   c                 ór  ‡ ‡‡‡	‡
‡‡‡‡‡— ‰ j         \  }Š	ŠŠ|j         d         Š
‰                      t          j        ¬¦  «        Š ‰€!t          j        ‰
‰ j        ‰ j        ¬¦  «        n‰Št          j        |‰
‰‰f‰ j        ‰ j        t          j        ¬¦  «        }|‰z  ‰z  |                     t          j        ¬¦  «        cŠŠˆˆ	ˆ
ˆˆˆˆˆ fd„Šˆ
ˆˆˆ fd„}d‰
d	z  dk    rd
ndz   }‰	› d‰
› d‰rdnd› d‰ j        › �} ‰|‰t           	                    d|‰ j        ||‰‰
z  t          k    ¦  «        ¦  «         |S )zË3x3 convolution with zero padding 1 of x (batch, cin, h, w) by weight (cout, cin, 3, 3) plus bias (None: no bias), then ReLU (relu),
    as the PyTorch path rounds them. Returns the channels-last output.r   rˆ   Nr»   r¤   c                 ó–   •— |\  }}}}}t          t          j        ||¦  «        ‰
|z  z  f         ‰‰‰| |‰‰‰	‰
‰|||d||¬¦  «         d S r¦   )rÚ   r”   rK   )r©   rª   r«   r­   Úbnr°   r±   r²   r³   r,   r+   ri   Úrelurg   r·   r(   s           €€€€€€€€r   r¸   zconv2d.<locals>.launchó   s”   ø€ Ø.4Ñ+ˆ��W˜e VÝ�œ E¨7Ñ3Ô3°t¸r±zÑBÐDÔEÀaÈÈTÐSYÐ[`ÐbcÐefÐhkÐmqÐswð  zAð  CEð  GNð  PQØOTÐagð	iñ 	iô 	ið 	ið 	ið 	ir   c                  ó~   •‡ ‡— t          ‰d¦  «        Š t          j        ‰ ‰z  ‰j        ‰j        ¬¦  «        Šˆ ˆˆfd„S )Nrº   r»   c                 ó   •—  ‰‰‰| ¦  «        S r½   r   )r«   rª   r¸   r©   s    €€€r   r¿   z)conv2d.<locals>.scratch.<locals>.<lambda>û   s   ø€ ˜f˜f V¨U°FÑ;Ô;€ r   rÀ   )rª   r©   r+   r¸   r-   r(   s   @@€€€€r   rÃ   zconv2d.<locals>.scratchø   sG   øøø€ Ý�F˜GÑ$Ô$ˆÝ”˜U T™\°´ÀÄÐJÑJÔJˆØ;Ð;Ð;Ð;Ð;Ð;Ð;r   ))r�   rŠ   r   r   r   )r�   rŠ   rŠ   r   r   )rŠ   rŠ   r   r   r   r�   ))r�   r�   r   r§   r   r   rÄ   z|3x3|rÞ   rÆ   Úconv2d)r!   r’   r
   r�   rP   r#   r   rÂ   rÇ   r«   rÈ   )r(   r    r³   rÞ   rÉ   r˜   rÃ   r¡   rË   r,   r+   ri   r¸   r-   rg   r·   s   ` ``     @@@@@@@r   rá   rá   é   s½  øøøøøøøøøø€ ð ”wÑ€Eˆ3��1ØŒ<˜Œ?€DØ	�Š¥5Ô#6ˆÑ7Ô7€AØ@DÀ�5Œ;�t 1¤7°1´8Ð<Ñ<Ô<Ð<ÐRV€DÝ
Œ+�u˜d A qÐ)°´ÀÄÕY^ÔYlÐ
mÑ
mÔ
m€CØ˜‘˜Q‘ × 1Ò 1ÅÔ@SÐ 1Ñ TÔ T€J€FˆBðið ið ið ið ið ið ið ið ið ið ið ið
<ð <ð <ð <ð <ð <ð <ð <ð
 PÐnrÐuxÑnxÐ|}Òn}Ðn}ÐSjÐSjð  DFñ  G€JØÐEÐE˜ÐEÐE¨TÐ#9 6 6°rÐEÐE¸A¼GÐEÐE€HØ
€Fˆ3��Ÿš h°¸!¼(ÀJÐPWÐY_ÐbfÑYfÕjxÒYxÑyÔyÑzÔzÐzØ€Jr   )Ú!install_triton_compilation_logger)F)NF) Ú__doc__Ú	functoolsr
   r”   Útriton.languageÚlanguagerI   Úshared.kernelsr   Úshared.kernels.triton_tunerr   rÇ   rÈ   Úchannels_last_3drŽ   r   r   r'   r.   ÚjitÚ	constexprr}   r�   r…   r™   r¢   r)   rÔ   rÚ   rá   Ú%shared.kernels.triton_compilation_logrâ   r   r   r   ú<module>rí      s  ððdð dð
 Ð Ð Ð à €€€Ø €€€Ø Ð Ð Ð Ð Ð à #Ð #Ð #Ð #Ð #Ð #Ø 3Ð 3Ð 3Ð 3Ð 3Ð 3à	ˆÐ+¨Q°Ñ	6Ô	6€Ø€ØÔ€ðKð Kð Kð
Zð Zð Zð Ð ðjð jð jð jð €„ÐDÐDÐDÐEÑEÔEð@TÀ2Ä<ð @TÐWYÔWcð @TÐikÔiuð @TÐ{}ô  |Hð @TØ”\ð@TØ-/¬\ð@TØEGÄ\ð@TØ\^Ô\hð@TØoqÔo{ð@Tà”lð@Tà-/¬\ð@TàDFÄLð@Tð @Tð @Tñ FÔEð@TðF "Ð ð €„˜u˜gÐ&Ñ&Ô&ð} b¤lð }¸R¼\ð }ÐTVÔT`ð }ð }ð }ñ 'Ô&ð}ð	ð 	ð 	ðð ð ð!ð !ð !ð !ðPGð Gð Gð €„˜o˜o˜oÐ.Ñ.Ô.ð!E°´ð !EÀRÄ\ð !EÐY[ÔYeð !EÐprÔp|ð !EØ”|ð!EØ.0¬lð!EØEGÄ\ð!Eð !Eð !Eñ /Ô.ð!EðHð ð ð ð4 TÐ SÐ SÐ SÐ SÐ Sà !Ð !Ñ #Ô #Ð #Ð #Ð #r   