Ë
    ÄÓ^jª˜  ã                  ó  — d dl mZ d dlZd dlmZmZ d dlZd dlmZ ddl	m
Z
mZ ddlmZmZmZmZmZ ddlmZmZmZmZ dd	lmZmZmZmZmZmZ dd
lm Z  ddl!m"Z" erd dl#m$Z$ ddlm%Z%  ejL                  e'«      Z(ejR                  jT                  Z*ed„ «       Z+ed„ «       Z,ed„ «       Z- ede- e"d«      d¬«      Z.dZ/	  ede+de/z   dz   e/z   dz   ¬«      Z0dZ1 ede,de1z   dz   e1z   d z   ¬«      Z2 eejf                  d!d"e*jf                  jh                  ¬#«      Z5d$„ Z6 ee6d«      Z7 G d%„ d&e«      Z8	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d9d'„Z9d(„ Z:d)„ Z; ee*jf                  «      	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d:d*„«       Z3 ee*jx                  «      d+„ «       Z<d,„ Z= ee*jf                  e=«       	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d;d-„Z>	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d;d.„Z?d/„ Z@ ee@d«      ZAd0„ ZB eeBd«      ZCed1„ «       ZD ed2eD e"d3«      ¬«      ZEed4„ «       ZF ed5eF e"d6«      ¬«      ZG ee*j�                  jh                  «      ZI ee*j�                  jh                  «      	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d<d7„«       ZJd8„ ZK ee*j�                  eK«       y)=é    )ÚannotationsN)ÚTYPE_CHECKINGÚ	TypedDict)ÚCKGroupedConvFwdTemplateé   )ÚconfigÚir)Úadd_layout_constraintÚconstrain_to_fx_stridesÚfallback_handlerÚ	loweringsÚregister_lowering)Úautotune_select_algorithmÚExternKernelChoiceÚSymbolicGridFnÚTritonTemplate)Úis_onesÚis_zerosÚpad_listlikeÚsympy_productÚuse_ck_conv_templateÚuse_triton_template)ÚVé   )Úload_kernel_template)ÚSequence)Ú	TensorBoxc               óF   —  || |z  |z  |d   «       |||d   «      |d   fS ©NÚBLOCK_MÚBLOCK_NÚGROUPS© )ÚnÚcÚhÚwÚmetaÚcdivs         úf/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/kernel/conv.pyÚconv2d_gridr+   /   s9   € ñ 	ˆQ�‰U�Q‰Y˜˜Y™Ó(ÙˆQ��Y‘Ó ØˆX‰ðð ó    c               óL   —  || |z  |z  |z  |d   «       |||d   «      |d   fS r   r#   )r$   r%   Údr&   r'   r(   r)   s          r*   Úconv3d_gridr/   8   s=   € ñ 	ˆQ�‰U�Q‰Y˜‰]˜D ™OÓ,ÙˆQ��Y‘Ó ØˆX‰ðð r,   c               óH   —  || |d   «       |||d   «       |||d   «      fS )Nr!   ÚBLOCK_LÚBLOCK_Cr#   )r$   r%   Úlr(   r)   s        r*   Údepthwise_conv1d_gridr4   H   s8   € ñ 	ˆQ��Y‘Ó ÙˆQ��Y‘Ó ÙˆQ��Y‘Ó ðð r,   Údepthwise_conv1dÚtriton_depthwise_convT)ÚnameÚgridÚsourceÚ"cache_codegen_enabled_for_templateaÓ  
        idx_x_h = i - PADDING_H + idx_y_h * STRIDE_H
        idx_x_w = j - PADDING_W + idx_y_w * STRIDE_W
        idx_x_c = tl.arange(0, BLOCK_K) + k

        x_ptrs = x_base + (
            (idx_x_h * stride_xh)[:, None]
            + (idx_x_w * stride_xw)[:, None]
            + (idx_x_c * stride_xc)[None, :]
        )
        mask_x = (
            (idx_n < BATCH)[:, None]
            & (idx_x_h >= 0)[:, None]
            & (idx_x_h < IN_H)[:, None]
            & (idx_x_w >= 0)[:, None]
            & (idx_x_w < IN_W)[:, None]
            & (idx_x_c < GROUP_IN_C)[None, :]
        )
        matrix_x = tl.load(x_ptrs, mask=mask_x, other=0.0)

        w_ptrs = w_base + (
            (idx_x_c * stride_wc_in)[:, None] + (i * stride_wh) + (j * stride_ww)
        )
        mask_w = (idx_x_c[:, None] < GROUP_IN_C) & (idx_y_c[None, :] < GROUP_OUT_C)
        matrix_w = tl.load(w_ptrs, mask=mask_w, other=0.0)
        acc += tl.dot(matrix_x, matrix_w, allow_tf32=ALLOW_TF32)
Úconvolution2da—  
{{def_kernel("X", "W")}}
    # Tensor dimensions
    BATCH = {{size("X", 0)}}
    IN_C = {{size("X", 1)}}
    IN_H = {{size("X", 2)}}
    IN_W = {{size("X", 3)}}
    OUT_C = {{size(None, 1)}}
    OUT_H = {{size(None, 2)}}
    OUT_W = {{size(None, 3)}}

    # Strides:
    stride_xn = {{stride("X", 0)}}
    stride_xc = {{stride("X", 1)}}
    stride_xh = {{stride("X", 2)}}
    stride_xw = {{stride("X", 3)}}
    stride_wc_out = {{stride("W", 0)}}
    stride_wc_in = {{stride("W", 1)}}
    stride_wh = {{stride("W", 2)}}
    stride_ww = {{stride("W", 3)}}

    nhw = tl.program_id(0).to(INDEX_DTYPE) * BLOCK_M + tl.arange(0, BLOCK_M)
    idx_y_w = nhw % OUT_W
    nh = nhw // OUT_W
    idx_y_h = nh % OUT_H
    idx_n = nh // OUT_H
    idx_y_c = tl.program_id(1).to(INDEX_DTYPE) * BLOCK_N + tl.arange(0, BLOCK_N)

{% if GROUPS == 1 %}
    group = 0
    GROUP_IN_C = IN_C
    GROUP_OUT_C = OUT_C
{% else %}
    group = tl.program_id(2).to(INDEX_DTYPE)
    GROUP_IN_C = IN_C // GROUPS
    GROUP_OUT_C = OUT_C // GROUPS
{% endif %}

    x_base = X + (group * stride_xc * GROUP_IN_C + idx_n * stride_xn)[:, None]
    w_base = (
        W + (group * stride_wc_out * GROUP_OUT_C + idx_y_c * stride_wc_out)[None, :]
    )

    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

{% if UNROLL %}
{% for i in range(KERNEL_H) %}
{% for j in range(KERNEL_W) %}
    i = {{i}}
    j = {{j}}
    for k in range(0, GROUP_IN_C, BLOCK_K):
        aÐ  
{% endfor %}
{% endfor %}
{% else %}
    # Could be simplified, but slightly slower:
    # for i in range(KERNEL_H):
    #     for j in range(KERNEL_W):
    #         for k in range(0, GROUP_IN_C, BLOCK_K):
    BLOCK_K_COUNT = (GROUP_IN_C + BLOCK_K - 1) // BLOCK_K
    for ijk in range(KERNEL_H * KERNEL_W * BLOCK_K_COUNT):
        k = (ijk % BLOCK_K_COUNT) * BLOCK_K
        ij = ijk // BLOCK_K_COUNT
        i = ij // KERNEL_W
        j = ij % KERNEL_W
        aÎ  
{% endif %}

    mask = (
        (idx_n < BATCH)[:, None]
        & (idx_y_h < OUT_H)[:, None]
        & (idx_y_w < OUT_W)[:, None]
        & (idx_y_c < GROUP_OUT_C)[None, :]
    )
    idx_n = idx_n[:, None]
    idx_c = idx_y_c[None, :] + group * GROUP_OUT_C
    idx_h = idx_y_h[:, None]
    idx_w = idx_y_w[:, None]

    # inductor generates a suffix
    {{store_output(("idx_n", "idx_c", "idx_h", "idx_w"), "acc", "mask", val_shape=("BLOCK_M", "BLOCK_N"))}}
)r7   r8   r9   a¡  
        idx_x_d = d - PADDING_D + idx_y_d * STRIDE_D
        idx_x_h = i - PADDING_H + idx_y_h * STRIDE_H
        idx_x_w = j - PADDING_W + idx_y_w * STRIDE_W
        idx_x_c = tl.arange(0, BLOCK_K) + k

        x_ptrs = x_base + (
            (idx_x_d * stride_xd)[:, None]
            + (idx_x_h * stride_xh)[:, None]
            + (idx_x_w * stride_xw)[:, None]
            + (idx_x_c * stride_xc)[None, :]
        )
        mask_x = (
            (idx_n < BATCH)[:, None]
            & (idx_x_d >= 0)[:, None]
            & (idx_x_d < IN_D)[:, None]
            & (idx_x_h >= 0)[:, None]
            & (idx_x_h < IN_H)[:, None]
            & (idx_x_w >= 0)[:, None]
            & (idx_x_w < IN_W)[:, None]
            & (idx_x_c < GROUP_IN_C)[None, :]
        )
        matrix_x = tl.load(x_ptrs, mask=mask_x, other=0.0)

        w_ptrs = w_base + (
            (idx_x_c * stride_wc_in)[:, None] +
            (d * stride_wd) + (i * stride_wh) + (j * stride_ww)
        )
        mask_w = (idx_x_c[:, None] < GROUP_IN_C) & (idx_y_c[None, :] < GROUP_OUT_C)
        matrix_w = tl.load(w_ptrs, mask=mask_w, other=0.0)
        acc += tl.dot(matrix_x, matrix_w, allow_tf32=ALLOW_TF32)
Úconvolution3dax  
{{def_kernel("X", "W")}}
    # Tensor dimensions
    BATCH = {{size("X", 0)}}
    IN_C = {{size("X", 1)}}
    IN_D = {{size("X", 2)}}
    IN_H = {{size("X", 3)}}
    IN_W = {{size("X", 4)}}
    OUT_C = {{size(None, 1)}}
    OUT_D = {{size(None, 2)}}
    OUT_H = {{size(None, 3)}}
    OUT_W = {{size(None, 4)}}

    # Strides:
    stride_xn = {{stride("X", 0)}}
    stride_xc = {{stride("X", 1)}}
    stride_xd = {{stride("X", 2)}}
    stride_xh = {{stride("X", 3)}}
    stride_xw = {{stride("X", 4)}}
    stride_wc_out = {{stride("W", 0)}}
    stride_wc_in = {{stride("W", 1)}}
    stride_wd = {{stride("W", 2)}}
    stride_wh = {{stride("W", 3)}}
    stride_ww = {{stride("W", 4)}}

    ndhw = tl.program_id(0).to(INDEX_DTYPE) * BLOCK_M + tl.arange(0, BLOCK_M)
    idx_y_w = ndhw % OUT_W
    ndh = ndhw // OUT_W
    idx_y_h = ndh % OUT_H
    nd = ndh // OUT_H
    idx_y_d = nd % OUT_D
    idx_n = nd // OUT_D
    idx_y_c = tl.program_id(1).to(INDEX_DTYPE) * BLOCK_N + tl.arange(0, BLOCK_N)

{% if GROUPS == 1 %}
    group = 0
    GROUP_IN_C = IN_C
    GROUP_OUT_C = OUT_C
{% else %}
    group = tl.program_id(2).to(INDEX_DTYPE)
    GROUP_IN_C = IN_C // GROUPS
    GROUP_OUT_C = OUT_C // GROUPS
{% endif %}

    x_base = X + (group * stride_xc * GROUP_IN_C + idx_n * stride_xn)[:, None]
    w_base = (
        W + (group * stride_wc_out * GROUP_OUT_C + idx_y_c * stride_wc_out)[None, :]
    )

    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

{% if UNROLL %}
{% for d in range(KERNEL_D) %}
{% for i in range(KERNEL_H) %}
{% for j in range(KERNEL_W) %}
    d = {{d}}
    i = {{i}}
    j = {{j}}
    for k in range(0, GROUP_IN_C, BLOCK_K):
        aF  
{% endfor %}
{% endfor %}
{% endfor %}
{% else %}
    # Could be simplified, but slightly slower:
    # for d in range(KERNEL_D):
    #   for i in range(KERNEL_H):
    #     for j in range(KERNEL_W):
    #         for k in range(0, GROUP_IN_C, BLOCK_K):
    BLOCK_K_COUNT = (GROUP_IN_C + BLOCK_K - 1) // BLOCK_K
    for dijk in range(KERNEL_D * KERNEL_H * KERNEL_W * BLOCK_K_COUNT):
        k = (dijk % BLOCK_K_COUNT) * BLOCK_K
        dij = dijk // BLOCK_K_COUNT
        j = dij % KERNEL_W
        di = dij // KERNEL_W
        i = di % KERNEL_H
        d = di // KERNEL_H
        a  
{% endif %}

    mask = (
        (idx_n < BATCH)[:, None]
        & (idx_y_d < OUT_D)[:, None]
        & (idx_y_h < OUT_H)[:, None]
        & (idx_y_w < OUT_W)[:, None]
        & (idx_y_c < GROUP_OUT_C)[None, :]
    )
    idx_n = idx_n[:, None]
    idx_c = idx_y_c[None, :] + group * GROUP_OUT_C
    idx_d = idx_y_d[:, None]
    idx_h = idx_y_h[:, None]
    idx_w = idx_y_w[:, None]

    # inductor generates a suffix
    {{store_output(("idx_n", "idx_c", "idx_d", "idx_h", "idx_w"), "acc", "mask", val_shape=("BLOCK_M", "BLOCK_N"))}}
zat::convolutionF)Úhas_out_variantÚop_overloadc          
     óî   — t        j                  t        j                  |d«      d«      }t        j                  | j                  dddd«      |j                  dd«      |j                  dddd«      ¬«      S )Néÿÿÿÿr   r   é   r   )Úout)ÚtorchÚsqueezeÚmatmulÚpermute)Úxr'   rB   s      r*   Úconv1x1_via_mmrH   f  s]   € Ü�‰”e—m‘m A rÓ*¨BÓ/€AÜ�<‰<Ø	�	‰	�!�Q˜˜1Ó˜qŸy™y¨¨A›°C·K±KÀÀ1ÀaÈÓ4Kôð r,   c                  óJ   — e Zd ZU ded<   ded<   ded<   ded<   ded<   ded	<   y
)ÚConvLayoutParamsútuple[int, ...]ÚstrideÚpaddingÚdilationÚboolÚ
transposedÚoutput_paddingÚintÚgroupsN)Ú__name__Ú
__module__Ú__qualname__Ú__annotations__r#   r,   r*   rJ   rJ   p  s%   … ØÓØÓØÓØÓØ#Ó#Ø„Kr,   rJ   c	                ó   — t         j                  j                  j                  }	t         j                  j                  5  t
        j                  j                  j                  t        j                  | «      t        j                  |«      t        j                  |«       |	|«       |	|«       |	|«      | |	|«      |«	      }
t        j                  |
j                  «       «      }t        j                  |
j                  «       «      }ddd«       t        j                  | j                  «       | j!                  «       |«      S # 1 sw Y   Œ=xY w)z)Determine output layout for a convolutionN)r   ÚgraphÚsizevarsÚguard_int_seqÚ	fake_moderC   ÚopsÚatenÚconvolutionr	   Úir_node_to_tensorÚconvert_shape_to_inductorÚsizerL   ÚFixedLayoutÚget_device_or_errorÚ	get_dtype)rG   ÚweightÚbiasrL   rM   rN   rP   rQ   rS   ÚguardÚoutputÚsizess               r*   Úconv_layoutrk   y  sü   € ô& �G‰G×Ñ×*Ñ*€EÜ	
�‰×	Ñ	ñ ?Ü—‘—‘×+Ñ+Ü× Ñ  Ó#Ü× Ñ  Ó(Ü× Ñ  Ó&Ù�&‹MÙ�'‹NÙ�(‹OØÙ�.Ó!Øó

ˆô ×,Ñ,¨V¯[©[«]Ó;ˆÜ×-Ñ-¨f¯m©m«oÓ>ˆ÷?ô �>‰>Ø	×ÑÓØ	�‰‹ØØó	ð ÷?ð ?ús   ¿C	EÅEc                ó‚   — t        t        t        | «      «      «      }|j                  d|j	                  d«      «       |S )Nr   r@   )ÚlistÚreversedÚrangeÚinsertÚpop)ÚrankÚorders     r*   Úchannels_last_orderrt   ¤  s0   € Ü”œ% ›+Ó&Ó'€EØ	‡L�L��E—I‘I˜b“MÔ"Ø€Lr,   c                ó   — t        |j                  «       «      }t        |dz
  «      D ]   }t        t        j
                     |d¬«      }Œ" t        t        j                     |ddg«      }t        j                  j                  | t        |«      «      } t        t        |«      «      }|j                  |j                  d«      «       t        t        j                     | |«      } | j                  «       �^ }}t        t        j                     | t        |«      |g«      } |€t        t        j                      | |«      }nt        t        j"                     || |«      }t        t        j                     |g |¢d‘«      }t        t        |«      «      }	|	j%                  d|	j                  d«      «       t        t        j                     ||	«      S )Nr   r@   ©Údimr   r   )ÚlenÚget_sizero   ÚLr^   rD   rF   r	   ÚExternKernelÚrequire_stride_orderrt   rm   Úappendrq   Úreshaper   ÚmmÚaddmmrp   )
rG   rf   rg   rr   Ú_Ú	x_permuterj   Úin_chanÚresultÚresult_permutes
             r*   Úconvert_1x1_conv_to_mmr†   ª  sa  € äˆv�‰Ó Ó!€DÜ�4˜!‘8‹_ò 1ˆÜ”4—<‘<‘ ¨RÔ0‰ð1äŒt�|‰|‰_˜V a¨ VÓ,€Fä
�‰×,Ñ,¨QÔ0CÀDÓ0IÓJ€AÜ”U˜4“[Ó!€IØ×Ñ�Y—]‘] 1Ó%Ô&Ü	Œ$�,‰,‰˜˜9Ó%€AØ—j‘j“l�O€UˆGÜ	Œ$�,‰,‰˜œM¨%Ó0°'Ð:Ó;€AØ€|Ü”4—7‘7‘˜A˜vÓ&‰ä”4—:‘:‘˜t Q¨Ó/ˆÜŒt�|‰|‰_˜V \ u \¨b \Ó2€FÜœ% ›+Ó&€NØ×Ñ˜!˜^×/Ñ/°Ó3Ô4ÜŒT�\‰\‰?˜6 >Ó2Ð2r,   c	                ób  ‡ ‡‡‡— t        |«      }t        |«      }t        |«      }t        |«      }t        |t        «      s)t        j                  j
                  j                  |«      }t        |t        «      sJ ‚t        t        j                  j
                  j                  |«      «      }t        t        j                  j
                  j                  |«      «      }||||||dœŠt        j                  ‰ «      }	t        ‰ j                  «       «      t        ‰j                  «       «      dz
  k(  rVt        t        j                     t        t        t        j                      ‰ dg‰ j                  «       ¢«      ‰|fi ‰¤Žd¬«      S t        j                  j
                  j                  ‰j                  «       «      ^}
}}t        ‰ j                  «       «      dk(  r›t        |«      dk(  r�|	dk(  rˆ‰j#                  d|z   d|z   d|z   d|z   d	œ«       t        t        j$                     ‰ d
¬«      Š t        t        j$                     ‰d
¬«      Št        t        j                     t        ‰ ‰|fi ‰¤Žd
¬«      S t        |«      Št'        |‰«      }t'        |‰«      }t'        |‰«      }t'        |‰«      }ˆˆˆˆ fd„}t(        j*                  xs t(        j,                  }t(        j.                  s	|r“ |«       rŒt1        |«      r�t1        |«      rvt3        |«      rkt1        |«      r`|s^t3        |«      rS|dk(  rNt        j                  j
                  j5                  t7        ‰ j                  «       «      d«      rt9        ‰ ‰|«      S |�¡|	dk7  rœt        ‰ ‰dfi ‰¤Ž}t        j                  j
                  j;                  |j                  «       d   d«      r|S t        t        j<                     |t        t        j>                     ||j                  «       d   g‰dgz  z   «      «      S ‰ jA                  «        ‰jA                  «        t        j                  jB                  ru‰d
k(  rpt        j                  xjD                  dz  c_"        t        jF                  jI                  ‰ «      Š t        jF                  jI                  ‰«      ŠtK        ‰ ‰dfi ‰¤Ž}n”tK        ‰ ‰dfi ‰¤Ž}t        jL                  t        j                  j
                  jO                  |jP                  «      «      }t        jF                  jS                  ‰ |«      Š t        jF                  jS                  ‰|«      Šg d¢}|€‰ ‰g}d‰d<   |jU                  dd«       not        jF                  jW                  |«      }|€J ‚‰ ‰|g}|jY                  «        t        j                  j
                  j                  |j                  «       «       g }tZ        j\                  j^                  ja                  d«      rtc        jd                  |||fi ‰¤Žg}tZ        j\                  j^                  ja                  d«      �rútg        |«      �rît1        |«      �râ|�sßt3        |«      �rÓt        j                  j
                  j;                  ||z  ‰ j                  «       d   «      �r”t1        |«      r@t1        |«      r5t3        |«      r*|dk(  r%|ji                  tj        je                  ||«      «       |dkD  xr |dk(  xr |
|k(  }|rr‰dk(  rmt        jl                  jo                  |	«      }|D ]I  }tq        jr                  |f‰ ‰f||d   |d   |d   |jt                  |jv                  dœ|jx                  ¤Ž ŒK t        jl                  j{                  |	«      }‰ j}                  «       j~                  } |t7        ‰ j                  «       d   g‰ j                  «       d
d ¢«      |
||¬«      D �]K  }t1        |«      }|r|jv                  nt�        |jv                  d«      }‰d
k(  rrtƒ        jr                  |f‰ ‰f||d   |d   |d   |d   |d   |d   ||tZ        j„                  j†                  jˆ                  dk(  |jt                  |dœ|jx                  ¤Ž Œª‰dk(  sŒ°t‹        jr                  |fi d‰ ‰f“d|“d|d   “d|d   “d|d
   “d|d   “d|d   “d|d
   “d|d   “d|d   “d |d
   “d!|“d"|“d#tZ        j„                  j†                  jˆ                  dk(  “d$|jt                  “d%|“|jx                  ¤Ž �ŒN t�        |«      r/t�        j�                  ||‰ ‰f|�|fn	t        «       z   ||||‰¬&«       t“        d'|||«      \  }}|S )(zGLower aten.convolution using Inductor convolution kernels or fallbacks.©rL   rM   rN   rP   rQ   rS   r   r   rv   rA   Úxpu)r   )r   )rL   rM   rN   rQ   r   c                 ó  •— t         j                  j                  r‰dk(  ryt        ‰‰d fi ‰¤Ž} t	        j
                  t         j                  j                  j                  | j                  «      «      }|t        j                  k(  S )Nr   T)
r   rY   Ú
layout_optrk   r	   Úget_stride_orderrZ   Úguarding_hints_or_throwrL   ÚNHWC_STRIDE_ORDER)ÚlayoutÚreq_stride_orderÚkwargsÚndimrf   rG   s     €€€€r*   Úchannels_last_convz'convolution.<locals>.channels_last_conv  sn   ø€ Ü�7‰7×Ò $¨!¢)Øä˜Q ¨Ñ7°Ñ7ˆô ×.Ñ.Ü�G‰G×Ñ×4Ñ4°V·]±]ÓCó
Ðð  ¤2×#7Ñ#7Ñ7Ð7r,   NÚcpurg   ÚATENÚTRITON)Úinput_nodesr�   ÚKERNEL_SIZEÚCONV_STRIDEÚPADDINGÚ
num_stagesÚ	num_warps©Ú
dtype_sizeé   Útf32)r—   r�   ÚKERNEL_HÚKERNEL_WÚSTRIDE_HÚSTRIDE_WÚ	PADDING_HÚ	PADDING_Wr"   ÚUNROLLÚ
ALLOW_TF32r›   rœ   r—   r�   ÚKERNEL_Dr¡   r¢   ÚSTRIDE_Dr£   r¤   Ú	PADDING_Dr¥   r¦   r"   r§   r¨   r›   rœ   )r—   rL   rM   rN   rS   Ún_spatial_dimensionsr_   )JÚtupleÚ
isinstancerR   r   rY   rZ   Ú	guard_intr[   r	   Úget_device_typerx   ry   rz   r^   rD   r_   ÚexpandÚupdateÚ	unsqueezer   r   Úmax_autotuneÚmax_autotune_gemmÚconv_1x1_as_mmr   r   Ústatically_known_gtr   r†   Ústatically_known_equalsÚaddÚviewÚrealizer‹   Únum_channels_last_convr{   Úrequire_channels_lastrk   rŒ   r�   rL   r|   rp   Úrealize_inputÚfreeze_layoutrC   Ú	_inductorÚutilsÚ_use_conv_autotune_backendÚaten_convolutionÚbindr   r}   Úaten_conv1x1_via_mmÚchoicesÚget_depthwise_conv_configsÚdepthwise_conv1d_templateÚmaybe_append_choicer›   rœ   r‘   Úget_conv_configsre   ÚitemsizeÚminÚconv2d_templateÚbackendsÚcudnnÚfp32_precisionÚconv3d_templater   r   Úadd_ck_conv_choicesr   ) rG   rf   rg   rL   rM   rN   rP   rQ   rS   Údevice_typeÚout_chanrƒ   Úkernel_shaper“   Úautotuning_gemmr„   r�   r�   Úordered_kwargs_for_cpp_kernelÚargsrÆ   Úis_depthwiseÚdepthwise_configsÚcfgÚconv_configsrž   Úunrollrœ   Únoder�   r‘   r’   s    ``                            @@r*   r_   r_   Á  sÑ  û€ ô �6‹]€FÜ�G‹n€GÜ�X‹€HÜ˜>Ó*€NÜ�fœcÔ"Ü—‘×!Ñ!×+Ñ+¨FÓ3ˆÜ�fœcÔ"Ð"Ð"ô ”1—7‘7×#Ñ#×1Ñ1°&Ó9Ó:€FÜ”A—G‘G×$Ñ$×2Ñ2°7Ó;Ó<€Gð ØØØ Ø(Øñ €Fô ×$Ñ$ QÓ'€Kä
ˆ1�:‰:‹<ÓœC §¡Ó 1Ó2°QÑ6Ò6ä”—‘‰Üœœ$Ÿ+™+™ q¨1Ð*<¨q¯z©z«|Ð*<Ó=¸vÀtÑVÈvÑVØô
ð 	
ô
 ()§w¡w×'7Ñ'7×'EÑ'EÀfÇoÁoÓFWÓ'XÐ$€Hˆg˜ô
 ˆ1�:‰:‹<Ó˜AÒ¤# lÓ"3°qÒ"8¸[ÈEÒ=QØ�‰à ™-Ø '™>Ø  8™OØ"&¨Ñ"7ñ	ô	
ô Œd�n‰nÑ˜a QÔ'ˆÜ”4—>‘>Ñ" 6¨qÔ1ˆä”—‘‰Ü˜˜6 4Ñ2¨6Ñ2Øô
ð 	
ô
 ˆ|Ó€DÜ˜& $Ó'€FÜ˜7 DÓ)€GÜ˜H dÓ+€HÜ! .°$Ó7€N÷
8ô ×)Ñ)ÒE¬V×-EÑ-E€Oô 
×	Ò	¡?Ñ7IÔ7KÜ�LÔ!Ü�FŒOÜ�WÔÜ�HÔÙÜ�^Ô$Ø�aŠKÜ�G‰G×Ñ×0Ñ0´¸q¿z¹z»|Ó1LÈaÔPä% a¨°Ó6Ð6àÐ˜K¨5Ò0ä˜Q ¨Ñ7°Ñ7ˆÜ�7‰7×Ñ×3Ñ3°F·O±OÓ4EÀaÑ4HÈ!ÔLàˆMÜ”—‘‰{Ø”A”d—i‘i‘L ¨¯©Ó(9¸!Ñ(<Ð'=ÀÈÀsÁ
Ñ'JÓKó
ð 	
ð ‡I�I„KØ
‡N�NÔô
 	‡w�w×Ò˜d ašiÜ	�‰×&Ò&¨!Ñ+Õ&Ü�O‰O×1Ñ1°!Ó4ˆô —‘×6Ñ6°vÓ>ˆÜ˜Q ¨Ñ7°Ñ7‰ä˜Q ¨Ñ7°Ñ7ˆô ×.Ñ.Ü�G‰G×Ñ×4Ñ4°V·]±]ÓCó
Ðô �O‰O×0Ñ0°Ð4DÓEˆÜ—‘×5Ñ5°fÐ>NÓOˆò%Ð!ð €|Ø�6ˆ{ˆØˆˆv‰Ø%×,Ñ,¨Q°Õ7ä�‰×,Ñ,¨TÓ2ˆØÐÐÐØ�6˜4Ð ˆØ×ÑÔÜ	�‰×Ñ×&Ñ& t§}¡}£Ô7à€GÜ‡�×Ñ×7Ñ7¸Ô?ä×!Ñ!ØØØ-ñð ñ	ð
ˆô 	�‰×Ñ×8Ñ8¸ÕBÜ Õ'ä�HÕÚÜ�^Õ$ä�G‰G×Ñ×4Ñ4°W¸vÑ5EÀqÇzÁzÃ|ÐTUÁÕWô �LÔ!Ü˜”Ü˜Ô!Ø˜!’à�N‰NÔ.×3Ñ3°D¸&ÓAÔBà ‘zÒI g°¡lÒI°xÀ6Ñ7IˆÙ˜D AšIÜ !§	¡	× DÑ DÀ[Ó QÐØ(ò �Ü)×=Ñ=Øð
à!" F Ø!Ø ,¨Q¡Ø & q¡	Ø# A™JØ"Ÿ~™~Ø!Ÿm™mñ
ð —j‘jó
ðô —y‘y×1Ñ1°+Ó>ˆà—[‘[“]×+Ñ+ˆ
ÙÜ˜1Ÿ:™:›<¨™?Ð>¨Q¯Z©Z«\¸!¸"Ð-=Ð>Ó?ØØØ!ô	
ó 5	ˆCô ˜\Ó*ˆFñ *0˜Ÿš´S¸¿¹ÈÓ5JˆIà�qŠyÜ×3Ñ3Øðà!" F Ø!Ø)¨!™_Ø)¨!™_Ø# A™YØ# A™YØ% a™jØ% a™jØ!ð "Ü$Ÿ~™~×3Ñ3×BÑBÀfÑLØ"Ÿ~™~Ø'ñ!ð" —j‘jó#ð& ˜“Ü×3Ñ3Øòà!" F¡ðñ "ðð *¨!š_ð	ð
 *¨!š_ðð *¨!š_ðð $ AšYðð $ AšYðð $ AšYðð & ašjðð & ašjðð & ašjðñ "ðñ  "ð!ô"  %Ÿ~™~×3Ñ3×BÑBÀfÒLð#ð$  #Ÿ~š~ð%ñ& (Ø—j‘jô)ðA5	ôl ˜FÔ#Ü ×4Ñ4ØØØ˜F˜°$Ð2B¨¡wÌËÑPØØØØØ!%õ		
ô (¨°wÀÀfÓM�G€Dˆ!Ø€Kr,   c                ó(   — t        | ||||||||«	      S ©N)r_   )rG   rf   rg   rL   rM   rN   rP   rQ   rS   Ú	benchmarkÚdeterministicÚcudnn_enabledÚ
allow_tf32s                r*   Ú_convolutionrå   Ò  s%   € ô  Ø	ˆ6�4˜ ¨(°JÀÐPVóð r,   c                óÔ   — | j                   t        j                  j                  j                  j
                  u sJ ‚t        j                  j                  r||fS t        | g|¢­i |¤ŽS rà   )
ÚtargetrC   r]   r^   r_   Údefaultr   rY   r‹   r   ©Úfx_noderØ   r‘   s      r*   Úconstrain_conv_to_fx_stridesrë   ç  sT   € Ø�>‰>œUŸY™YŸ^™^×7Ñ7×?Ñ?Ñ?Ð?Ð?Ü‡w�w×ÒØ�Vˆ|Ðä& wÐ@°Ò@¸Ñ@Ð@r,   c	                ó¸  — t         j                  j                  j                  }	t         j                  j                  5  t        j                  | «      }
t        j                  |«      }t        j                  |«      }t        j                  j                  j                  |
||d  |	|«       |	|«       |	|«      | |	|«      |d«      \  }}}t        j                  |j                  «       «      }t        j                  |j                  «       «      }d d d «       t        j                  |j                  «       |j!                  «       «      S # 1 sw Y   Œ=xY w)N©TFF©r   rY   rZ   r[   r\   r	   r`   rC   r]   r^   Úconvolution_backwardra   rb   rL   rc   rd   re   )Úgrad_outÚinputrf   rL   rM   rN   rP   rQ   rS   rh   ÚgorG   r'   Údxr�   rj   Ústride_s                    r*   Úconv_bwd_input_layoutrõ   ò  s  € ô �G‰G×Ñ×*Ñ*€EÜ	
�‰×	Ñ	ñ <Ü×!Ñ! (Ó+ˆÜ× Ñ  Ó'ˆÜ× Ñ  Ó(ˆä—9‘9—>‘>×6Ñ6ØØØØÙ�&‹MÙ�'‹NÙ�(‹OØÙ�.Ó!ØØ ó
‰ˆˆAˆqô ×,Ñ,¨R¯W©W«YÓ7ˆÜ×.Ñ.¨r¯y©y«{Ó;ˆ÷'<ô* �>‰>Ø×!Ñ!Ó#Ø�‰ÓØØó	ð ÷+<ð <úó   ¿CEÅEc	                ó¸  — t         j                  j                  j                  }	t         j                  j                  5  t        j                  | «      }
t        j                  |«      }t        j                  |«      }t        j                  j                  j                  |
||d  |	|«       |	|«       |	|«      | |	|«      |d«      \  }}}t        j                  |j                  «       «      }t        j                  |j                  «       «      }d d d «       t        j                  |j                  «       |j!                  «       «      S # 1 sw Y   Œ=xY w)N©FTFrî   )rð   rñ   rf   rL   rM   rN   rP   rQ   rS   rh   rò   rG   r'   r�   Údwrj   rô   s                    r*   Úconv_bwd_weight_layoutrú     s  € ô �G‰G×Ñ×*Ñ*€EÜ	
�‰×	Ñ	ñ <Ü×!Ñ! (Ó+ˆÜ× Ñ  Ó'ˆÜ× Ñ  Ó(ˆä—9‘9—>‘>×6Ñ6ØØØØÙ�&‹MÙ�'‹NÙ�(‹OØÙ�.Ó!ØØ ó
‰ˆˆ2ˆqô ×,Ñ,¨R¯W©W«YÓ7ˆÜ×.Ñ.¨r¯y©y«{Ó;ˆ÷'<ô* �>‰>Ø×"Ñ"Ó$Ø×ÑÓØØó	ð ÷+<ð <úrö   c               ód  — | j                  t        j                  ¬«      rt        j                  }
nt        j                  }
t        j                  ||	j
                  | j                  |
¬«      }t        j                  j                  j                  j                  d |	d || |d ||||||d¬«       |	S )N©Úmemory_format©ÚdtypeÚdevicerý   rø   ©Úout1Úout2Úout3Úgrad_outputrñ   rf   Ú
bias_sizesrL   rM   rN   rP   rQ   rS   Úoutput_mask©Úis_contiguousrC   Úchannels_lastÚcontiguous_formatÚemptyrÿ   r   r]   r^   rï   rB   )Úx_tÚgo_tÚw_shaperL   rM   rN   rP   rQ   rS   rB   Ú
memory_fmtÚdummy_weights               r*   Úcall_aten_dwr  D  s¤   € ð ×Ñ¤u×':Ñ':ÐÔ;Ü×(Ñ(‰
ä×,Ñ,ˆ
ä—;‘;Ø�s—y‘y¨¯©À:ô€Lô 
‡I�I‡N�N×'Ñ'×+Ñ+ØØØØØØØØØØØØ%ØØ(ð ,ô ð  €Jr,   c               ód  — | j                  t        j                  ¬«      rt        j                  }
nt        j                  }
t        j                  ||	j
                  | j                  |
¬«      }t        j                  j                  j                  j                  |	d d | ||d ||||||d¬«       |	S )Nrü   rþ   rí   r  r  )r  Úw_tÚx_shaperL   rM   rN   rP   rQ   rS   rB   r  Údummy_inputs               r*   Úcall_aten_dxr  p  s¤   € ð ×Ñ¬×(;Ñ(;ÐÔ<Ü×(Ñ(‰
ä×,Ñ,ˆ
ä—+‘+Ø�s—y‘y¨¯©ÀJô€Kô 
‡I�I‡N�N×'Ñ'×+Ñ+ØØØØØØØØØØØØ%ØØ(ð ,ô ð  €Jr,   c               óP   — |d   } |||z  |z  |d   «       || |z  |d   «      |fS ©Nr"   r    r!   r#   )ÚcoutÚcinÚkhÚkwr(   r)   Úgs          r*   Úconv2d_bwd_weight_gridr  œ  s@   € àˆX‰€AáˆS�2‰X˜‰]˜D ™OÓ,ÙˆT�Q‰Y˜˜Y™Ó(Ø	ðð r,   Úconvolution2d_bwd_weightÚtriton_conv2d_bwd_weightc               óP   — |d   } || |z  |z  |d   «       |||z  |d   «      |fS r  r#   )r$   r  r&   r'   r(   r)   r  s          r*   Úconv2d_bwd_input_gridr#  ­  s@   € àˆX‰€AáˆQ�‰U�Q‰Y˜˜Y™Ó(ÙˆS�A‰X�t˜I‘Ó'Ø	ðð r,   Úconvolution2d_bwd_inputÚtriton_conv2d_bwd_inputc                ó¦  — t        |«      }t        |«      }t        |«      }t        |«      }t        |	t        «      s)t        j                  j
                  j                  |	«      }	t        j                  j
                  j                  |j                  «       «      ^}}}t        t        j                  j
                  j                  |«      «      }t        t        j                  j
                  j                  |«      «      }t        t        j                  j
                  j                  |«      «      }|j                  «        |j                  «        | j                  «        ||||||	dœ}t        |«      }t        ||«      }t        ||«      }t        ||«      }t        ||«      }t        j                  |«      }t        j                  j                  |«      }|j!                  «       j"                  }d}d}g }g }t%        | ||fi |¤Ž}|
d   �r#t        j                  j&                  ru|dk(  rpt        j                  xj(                  dz  c_        t        j*                  j-                  |«      }t        j*                  j-                  | «      } t%        | ||fi |¤Ž}n‰t        j                  j
                  j                  }t        j.                   ||j0                  «      «      }t        j*                  j3                  ||«      }t        j*                  j3                  | |«      } || g}t4        j6                  j8                  j;                  d«      rÞt=        |«      rÓ|sÑt?        |«      rÆ |tA        |j                  «       d   g|j                  «       dd ¢«      |||¬«      D ]‰  }|dk(  sŒ	d	}tC        jD                  |f|| f||d   |d   |d   |d   |d   |d   |d   |d   |	t4        jF                  jH                  jJ                  |jL                  |jN                  d
œ|jP                  ¤Ž Œ‹ d}d}g }g }tS        | ||fi |¤Ž}|
d   �r#t        j                  j&                  ru|dk(  rpt        j                  xj(                  dz  c_        t        j*                  j-                  | «      } t        j*                  j-                  |«      }tS        | ||fi |¤Ž}n‰t        j                  j
                  j                  }t        j.                   ||j0                  «      «      }t        j*                  j3                  | |«      } t        j*                  j3                  ||«      }| |g}t4        j6                  j8                  jU                  d«      rÞt=        |«      rÓ|sÑt?        |«      rÆ |tA        |j                  «       d   g|j                  «       dd ¢«      |||¬«      D ]‰  }|dk(  sŒ	d	}tW        jD                  |f| |f||d   |d   |d   |d   |d   |d   |d   |d   |	t4        jF                  jH                  jJ                  |jL                  |jN                  d
œ|jP                  ¤Ž Œ‹ |s|stY        | |||||||||	|
«      S |
d   rzt4        j6                  j8                  j;                  d«      s|s>|j[                  t\        j_                  ||g d¢|j                  «       ||||||	¬«
      «       ta        d|||«      \  }} |
d   rzt4        j6                  j8                  jU                  d«      s|s>|j[                  tb        j_                  ||g d¢|j                  «       ||||||	¬«
      «       ta        d|||«      \  }} d}!|
d   r:|�8te        tf        jh                     | dgtk        tm        d|dz   «      «      z   ¬«      }!|||!fS )a¤  
    Lowering function for backward convolution operator.

    TRITON kernels are only registered for supported configurations (currently 2D convolutions).
    For unsupported dimensions or configurations, the choices list remains empty,
    triggering an automatic fallback to the ATen reference implementation.
    This ensures correctness for all cases while enabling TRITON optimizations only where implemented.
    rˆ   FNr   r   r–   r   r�   T)r—   r�   r¡   r¢   r¥   r¦   r£   r¤   Ú
DILATION_HÚ
DILATION_Wr"   r¨   r›   rœ   r•   )r  rL   rM   rN   rP   rQ   rS   )
r—   r�   r×   r  rL   rM   rN   rP   rQ   rS   Úconvolution_bwd_weight)r  rL   rM   rN   rP   rQ   rS   )
r—   r�   r×   r  rL   rM   rN   rP   rQ   rS   Úconvolution_bwd_input)Úaxis)7r­   r®   rR   r   rY   rZ   r¯   r[   ry   r»   rx   r   r	   r°   rÆ   rÊ   re   rË   rú   r‹   r¼   r{   r½   rŒ   rL   r|   rC   rÀ   rÁ   Ú%_use_conv_bwd_weight_autotune_backendr   r   r   Úconv2d_bwd_weight_templaterÉ   rÎ   rÏ   rä   r›   rœ   r‘   rõ   Ú$_use_conv_bwd_input_autotune_backendÚconv2d_bwd_input_templateÚ"aten_convolution_backward_fallbackr}   Úext_kn_aten_dwrÄ   r   Úext_kn_aten_dxrz   r^   Úsumrm   ro   )"rð   rñ   rf   r  rL   rM   rN   rP   rQ   rS   r  rÔ   rƒ   rÕ   r‘   r’   rÓ   rÜ   rž   Úhas_triton_dw_choicesrù   Ú
choices_dwÚargs_wÚ	layout_dwrh   Ústride_orderrÛ   Úhas_triton_dx_choicesró   Ú
choices_dxÚargs_xÚ	layout_dxr�   Údbs"                                     r*   Úconvolution_backward_loweringr>  Á  s  € ô, �6‹]€FÜ�G‹n€GÜ�X‹€HÜ˜>Ó*€NÜ�fœcÔ"Ü—‘×!Ñ!×+Ñ+¨FÓ3ˆä'(§w¡w×'7Ñ'7×'EÑ'EÀfÇoÁoÓFWÓ'XÐ$€Hˆg˜ô ”1—7‘7×#Ñ#×1Ñ1°&Ó9Ó:€FÜ”A—G‘G×$Ñ$×2Ñ2°7Ó;Ó<€GÜ”Q—W‘W×%Ñ%×3Ñ3°HÓ=Ó>€Hà	‡M�M„OØ
‡N�NÔØ×ÑÔð ØØØ Ø(Øñ €Fô ˆ|Ó€DÜ˜& $Ó'€FÜ˜7 DÓ)€GÜ˜H dÓ+€HÜ! .°$Ó7€Nä×$Ñ$ UÓ+€Kä—9‘9×-Ñ-¨kÓ:€LØ—‘Ó"×+Ñ+€Jà!ÐØ	€BØ€JØ€FÜ& x°¸ÑIÀ&ÑI€IØ�1ƒ~Ü�7‰7×Ò $¨!¢)Ü�G‰G×*Ò*¨aÑ/Õ*Ü—O‘O×9Ñ9¸%Ó@ˆEÜ—‘×<Ñ<¸XÓFˆHÜ.¨x¸ÀÑQÈ&ÑQ‰Iä—G‘G×$Ñ$×2Ñ2ˆEÜ×.Ñ.©u°Y×5EÑ5EÓ/FÓGˆLÜ—O‘O×8Ñ8¸ÀÓMˆEÜ—‘×;Ñ;¸HÀlÓSˆHà˜Ð"ˆô �O‰O×!Ñ!×GÑGÈÔQÜ# IÔ.ÙÜ˜Ô(á#Ü˜uŸ~™~Ó/°Ñ2ÐJ°U·^±^Ó5EÀaÀbÐ5IÐJÓKØØØ%ô	ò �ð ˜1“9Ø,0Ð)Ü.×BÑBØ"ðà%*¨HÐ$5Ø(Ø!-¨a¡Ø!-¨a¡Ø")¨!¡*Ø")¨!¡*Ø!'¨¡Ø!'¨¡Ø#+¨A¡;Ø#+¨A¡;Ø%Ü#(§>¡>×#7Ñ#7×#BÑ#BØ#&§>¡>Ø"%§-¡-ñð  Ÿ*™*ó!ðð: "ÐØ	€BØ€JØ€FÜ% h°°vÑHÀÑH€IØ�1ƒ~Ü�7‰7×Ò $¨!¢)Ü�G‰G×*Ò*¨aÑ/Õ*Ü—‘×<Ñ<¸XÓFˆHÜ—_‘_×:Ñ:¸6ÓBˆFÜ-¨h¸¸vÑPÈÑP‰Iä—G‘G×$Ñ$×2Ñ2ˆEÜ×.Ñ.©u°Y×5EÑ5EÓ/FÓGˆLÜ—‘×;Ñ;¸HÀlÓSˆHÜ—_‘_×9Ñ9¸&À,ÓOˆFà˜FÐ#ˆô �O‰O×!Ñ!×FÑFÀxÔPÜ# IÔ.ÙÜ˜Ô(ñ $Ü˜uŸ~™~Ó/°Ñ2ÐJ°U·^±^Ó5EÀaÀbÐ5IÐJÓKØØØ%ô	ò �ð ˜1“9Ø,0Ð)Ü-×AÑAØ"ðà%-¨vÐ$6Ø(Ø!-¨a¡Ø!-¨a¡Ø")¨!¡*Ø")¨!¡*Ø!'¨¡Ø!'¨¡Ø#+¨A¡;Ø#+¨A¡;Ø%Ü#(§>¡>×#7Ñ#7×#BÑ#BØ#&§>¡>Ø"%§-¡-ñð  Ÿ*™*ó!ðñ< !Ñ)>Ü1ØØØØØØØØØØØó
ð 	
ð �1‚~ä�O‰O×!Ñ!×GÑGÈÔOÙ(à×ÑÜ×#Ñ#Ø &Ø$ò3ð #ŸO™OÓ-Ø!Ø#Ø%Ø)Ø#1Ø!ð% $ó ôô2 *Ø$ j°&¸)ó
‰ˆˆAð �1‚~ä�O‰O×!Ñ!×FÑFÀvÔNÙ(à×ÑÜ×#Ñ#Ø &Ø$ò3ð "ŸN™NÓ,Ø!Ø#Ø%Ø)Ø#1Ø!ð% $ó ôô2 *Ø# Z°¸ó
‰ˆˆAð 
€BØ�1‚~˜*Ð0ÜŒt�x‰x‰[˜¨¨¬d´5¸¸DÀ1¹HÓ3EÓ.FÑ(FÔGˆà��Bˆ<Ðr,   c                óÖ   — | j                   t        j                  j                  j                  j
                  k(  sJ ‚t        j                  j                  r||fS t        | g|¢­i |¤ŽS rà   )
rç   rC   r]   r^   rï   rè   r   rY   r‹   r   ré   s      r*   Ú constrain_conv_bwd_to_fx_stridesr@  È  sT   € Ø�>‰>œUŸY™YŸ^™^×@Ñ@×HÑHÒHÐHÐHÜ‡w�w×ÒØ�Vˆ|Ðä& wÐ@°Ò@¸Ñ@Ð@r,   )rG   r   rf   r   rg   úTensorBox | NonerL   úSequence[int]rM   rK   rN   rK   rP   rO   rQ   rK   rS   rR   Úreturnú	ir.Layout)rG   r   rf   r   rg   rA  rL   rB  rM   rB  rN   rB  rP   rO   rQ   rB  rS   rR   )rð   r   rñ   r   rf   r   rL   rB  rM   rK   rN   rK   rP   rO   rQ   rK   rS   rR   rC  rD  )rð   r   rñ   r   rf   r   r  zSequence[int] | NonerL   rB  rM   rB  rN   rB  rP   rO   rQ   rB  rS   rR   r  zSequence[bool])LÚ
__future__r   ÚloggingÚtypingr   r   rC   Ú-torch._inductor.codegen.rocm.ck_conv_templater   Ú r   r	   Úloweringr
   r   r   r   rz   r   Úselect_algorithmr   r   r   r   rÁ   r   r   r   r   r   r   Úvirtualizedr   Ú	mm_commonr   Úcollections.abcr   r   Ú	getLoggerrT   Úlogr]   r^   r+   r/   r4   rÈ   ÚLOOP_BODY_2DrÍ   ÚLOOP_BODY_3DrÑ   r_   rè   rÃ   rH   rÅ   rJ   rk   rt   r†   rå   rë   rõ   rú   r  r1  r  r2  r  r-  r#  r/  rï   r0  r>  r@  r#   r,   r*   ú<module>rS     s  ðå "ã ß +ã Ý Rç ÷õ ÷ó ÷÷ õ Ý +ñ Ý(åà€g×Ñ˜Ó!€ð ‡y�y‡~�~€ð ñó ðð ñó ðð ñó ðñ +Ø	Ø	ÙÐ 7Ó8Ø'+ô	Ð ð€ð8ñ
 !Ø	Ø	ð3ðh ñi4ðjñkCðH ñIDðJñKUôY€ðv€ñB !Ø	Ø	ð;ðx ñy<ðzñ{Oð` ñaPðbñccôg€ñR &Ø	×ÑØØØ× Ñ ×(Ñ(ô	Ð òñ )¨¸Ó>Ð ô�yô ð(Øð(àð(ð ð(ð ð	(ð
 ð(ð ð(ð ð(ð $ð(ð ð(ð ó(òVò3ñ. �4×#Ñ#Ó$ðMØðMàðMð ðMð ð	Mð
 ðMð ðMð ðMð "ðMð òMó %ðMñ` �4×$Ñ$Ó%ñó &ðò(Añ �d×&Ñ&Ð(DÔ Eð&Øð&àð&ð ð&ð ð	&ð
 ð&ð ð&ð ð&ð $ð&ð ð&ð ó&ðR&Øð&àð&ð ð&ð ð	&ð
 ð&ð ð&ð ð&ð $ð&ð ð&ð ó&òR&ñR $ L°$Ó7€ò&ñR $ L°$Ó7€ð ñó ðñ ,Ø	#Ø	ÙÐ :Ó;ôÐ ð ñó ðñ +Ø	"Ø	ÙÐ 9Ó:ôÐ ñ &6°d×6OÑ6O×6WÑ6WÓ%XÐ "ñ �4×,Ñ,×4Ñ4Ó5ðCØðCàðCð ðCð %ð	Cð
 ðCð ðCð ðCð ðCð "ðCð ðCð  òCó 6ðCòLAñ �d×/Ñ/Ð1QÕ Rr,   