Ë
    g^(hX  ã            	      óø  — d dl mZ d dlZd dlmZmZ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 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#m$Z$ erd dl%m&Z& ddlm'Z'  ejP                  e)«      Z*ejV                  jX                  Z,ed„ «       Z-ed„ «       Z.dddœdddœdddœdddœdddœdddœdddœgZ/ e0d„ e/D «       «      Z1ejd                  jf                  r"ejh                  jk                  «       r e#e1«      Z1d„ Z6d„ Z7dZ8	  ede-de8z   d z   e8z   d!z   ¬"«      Z9d#Z: ed$e.d%e:z   d&z   e:z   d'z   ¬"«      Z; eejx                  d(d)e,jx                  jz                  ¬*«      Z>d+„ Z? ee?d«      Z@ G d,„ d-e«      ZA	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d4d.„ZBd/„ ZCd0„ ZD ee,jx                  «      	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d5d1„«       Z< ee,jŠ                  «      d2„ «       ZEd3„ ZF ee,jx                  eF«       y)6é    )ÚannotationsN)ÚcastÚOptionalÚTYPE_CHECKINGÚ	TypedDict)ÚCKGroupedConvFwdTemplateé   )ÚconfigÚir)Úadd_layout_constraintÚconstrain_to_fx_stridesÚ	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é   )Úbuild_rocm_gemm_configsÚfiltered_configs)ÚSequence)Ú	TensorBoxc               óF   —  || |z  |z  |d   «       |||d   «      |d   fS ©NÚBLOCK_MÚBLOCK_NÚGROUPS© )ÚnÚcÚhÚwÚmetaÚcdivs         úY/var/www/skyplay_api_hub/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_gridr1   7   s=   € ñ 	ˆQ�‰U�Q‰Y˜‰]˜D ™OÓ,ÙˆQ��Y‘Ó ØˆX‰ðð r.   )é@   é   é   r	   é   T)r
   Úcond)r3   r2   r4   r	   r5   )i   r4   r4   r   é   )é€   r8   é    r	   r7   )r2   r2   r9   r	   r5   )r2   r3   r9   r	   r7   )r3   r2   r9   r	   r7   c           	   #  ó„   K  — | ]8  }|d    r1t        t        t        t        t        t        t        f   |d   «      –— Œ: y­w)r6   r
   N)r   ÚtupleÚint)Ú.0r
   s     r,   ú	<genexpr>r>   N   s;   è ø€ ò àØˆf‚~ô 	ŒŒs”Cœœc¤3Ð&Ñ	'¨°Ñ)9×:ñùs   ‚>A c                ó8   — | dkD  s
|dkD  s|dkD  ry| |z  |z  dkD  S )Nr3   Ti   r%   )Úmr&   Úks      r,   Ú_is_large_block_for_cpurB   Y   s+   € àˆ3‚w�!�c’'˜Q šWØØˆq‰5�1‰9�uÑÐr.   c               ód   — |dk(  rt        | ||t        dt        ¬«      S t        | ||t        ¬«      S )NÚcpug      à?)ÚconfigsÚscaleÚexclude)rE   )r   Úplatform_configsrB   )r@   r&   rA   Údevice_typeÚkwargss        r,   Úconv_configsrK   `   s=   € Ø�eÒÜØØØÜ$ØÜ+ô
ð 	
ô ˜A˜q !Ô-=Ô>Ð>r.   aÓ  
        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)
Úconvolution2dag  
{{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) * 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) * 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)
    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")}}
)ÚnameÚgridÚsourcea¡  
        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)
Úconvolution3daH  
{{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) * 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) * 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)
    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")}}
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)   rV   s      r,   Úconv1x1_via_mmr\   {  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_paddingr<   ÚgroupsN)Ú__name__Ú
__module__Ú__qualname__Ú__annotations__r%   r.   r,   r^   r^   …  s%   … ØÓØÓØÓØÓØ#Ó#Ø„Kr.   r^   c	                ól  — t         j                  j                  5  t        j                  j
                  j                  t        j                  | d¬«      t        j                  |d¬«      t        j                  |d¬«      t         j                  j                  j                  |«      t         j                  j                  j                  |«      t         j                  j                  j                  |«      |t         j                  j                  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 convolutionT)Úguard_shapeN)r   ÚgraphÚ	fake_moderW   ÚopsÚatenÚconvolutionr   Úir_node_to_tensorÚsizevarsÚ
size_hintsÚconvert_shape_to_inductorÚsizer`   ÚFixedLayoutÚget_device_or_errorÚ	get_dtype)r[   ÚweightÚbiasr`   ra   rb   rd   re   rf   ÚoutputÚsizess              r,   Úconv_layoutr~   Ž  s0  € ô 
�‰×	Ñ	ñ ?Ü—‘—‘×+Ñ+Ü× Ñ  °Ô5Ü× Ñ  °TÔ:Ü× Ñ  °4Ô8Ü�G‰G×Ñ×'Ñ'¨Ó/Ü�G‰G×Ñ×'Ñ'¨Ó0Ü�G‰G×Ñ×'Ñ'¨Ó1ØÜ�G‰G×Ñ×'Ñ'¨Ó7Øó

ˆô ×,Ñ,¨V¯[©[«]Ó;ˆÜ×-Ñ-¨f¯m©m«oÓ>ˆ÷?ô �>‰>Ø	×ÑÓØ	�‰‹ØØó	ð ÷?ð ?ús   ›EF*Æ*F3c                ó‚   — t        t        t        | «      «      «      }|j                  d|j	                  d«      «       |S )Nr   rT   )ÚlistÚreversedÚrangeÚinsertÚpop)ÚrankÚorders     r,   Úchannels_last_orderr‡   ±  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	   rT   ©Údimr   r   )ÚlenÚget_sizer‚   ÚLrp   rX   rZ   r   ÚExternKernelÚrequire_stride_orderr‡   r€   Úappendr„   Úreshaper   ÚmmÚaddmmrƒ   )
r[   rz   r{   r…   Ú_Ú	x_permuter}   Ú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                  «       «      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 t        j                   ‰ «      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 |�wt        j                   ‰ «      dk7  r_t        ‰ ‰d fi ‰¤Ž}t        t        j:                     |t        t        j<                     ||j                  «       d   g‰dgz  z   «      «      S ‰ j?                  «        ‰j?                  «        t        j                  j@                  ru‰d
k(  rpt        j                  xjB                  dz  c_!        t        jD                  jG                  ‰ «      Š t        jD                  jG                  ‰«      ŠtI        ‰ ‰d fi ‰¤Ž}n”tI        ‰ ‰d fi ‰¤Ž}t        jJ                  t        j                  j
                  jM                  |jN                  «      «      }t        jD                  jQ                  ‰ |«      Š t        jD                  jQ                  ‰|«      Šg d¢}|€‰ ‰g}d ‰d<   |jS                  dd«       n\‰ ‰|g}|j?                  «        |jU                  «        t        j                  j
                  j                  |j                  «       «       g }tV        jX                  jZ                  j]                  d«      rt_        j`                  |||fi ‰¤Žg}tV        jX                  jZ                  j]                  d«      �r>tc        |«      �r2t1        |«      �r&|�s#t3        |«      �rt        j                  j
                  je                  |
‰ j                  «       d   «      �rÛt1        |«      r@t1        |«      r5t3        |«      r*|dk(  r%|jg                  th        ja                  ||«      «       tk        t7        ‰ j                  «       d   g‰ j                  «       d
d  ¢«      |	|
t        j                   ‰ «      ¬«      D �]<  }‰d
k(  r‚tm        jn                  |f‰ ‰f||d   |d   |d   |d   |d   |d   |t1        |«      tV        jp                  jr                  jt                  |jv                  |jx                  dœ|jz                  ¤Ž Œ‹‰dk(  sŒ‘t}        jn                  |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t1        |«      “dtV        jp                  jr                  jt                  “d |jv                  “d!|jx                  “|jz                  ¤Ž �Œ? t        |«      r/t�        j‚                  ||‰ ‰f|�|fn	t        «       z   ||||‰¬"«       t…        d#|||«      S )$N)r`   ra   rb   rd   re   rf   r   r   r‰   rU   Úxpu)r   )r   )r`   ra   rb   re   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   rm   Ú
layout_optr~   r   Úget_stride_orderrs   rt   r`   ÚNHWC_STRIDE_ORDER)ÚlayoutÚreq_stride_orderrJ   Úndimrz   r[   s     €€€€r,   Úchannels_last_convz'convolution.<locals>.channels_last_conv  sl   ø€ Ü�7‰7×Ò $¨!¢)Øä˜Q ¨Ñ7°Ñ7ˆÜ×.Ñ.Ü�G‰G×Ñ×'Ñ'¨¯©Ó6ó
Ðð  ¤2×#7Ñ#7Ñ7Ð7r.   rD   r{   ÚATENÚTRITON)rI   )Úinput_nodesr    ÚKERNEL_HÚKERNEL_WÚSTRIDE_HÚSTRIDE_WÚ	PADDING_HÚ	PADDING_Wr$   ÚUNROLLÚ
ALLOW_TF32Ú
num_stagesÚ	num_warpsr¦   r    ÚKERNEL_Dr§   r¨   ÚSTRIDE_Dr©   rª   Ú	PADDING_Dr«   r¬   r$   r­   r®   r¯   r°   )r¦   r`   ra   rb   rf   Ún_spatial_dimensionsrq   )Cr;   Ú
isinstancer<   r   rm   rs   Úevaluate_static_shapeÚevaluate_static_shapesr‹   rŒ   r�   rp   rX   rq   Úexpandr   Úget_device_typeÚupdateÚ	unsqueezer   r
   Úmax_autotuneÚmax_autotune_gemmÚconv_1x1_as_mmr   r   Ústatically_known_gtr   r™   ÚaddÚviewÚrealizer�   Únum_channels_last_convrŽ   Úrequire_channels_lastr~   rž   rt   r`   r�   rƒ   Úfreeze_layoutrW   Ú	_inductorÚutilsÚ_use_conv_autotune_backendÚaten_convolutionÚbindr   Ústatically_known_equalsr�   Úaten_conv1x1_via_mmrK   Úconv2d_templateÚmaybe_append_choiceÚbackendsÚcudnnÚ
allow_tf32r¯   r°   rJ   Úconv3d_templater   r   Úadd_ck_conv_choicesr   )r[   rz   r{   r`   ra   rb   rd   re   rf   Úout_chanr–   Úkernel_shaper£   Úautotuning_gemmr—   r    r¡   Úordered_kwargs_for_cpp_kernelÚargsÚchoicesÚcfgrJ   r¢   s   ``                   @@r,   rq   rq   Î  s·  û€ ô �6‹]€FÜ�G‹n€GÜ�X‹€HÜ˜>Ó*€NÜ�fœcÔ"Ü—‘×!Ñ!×7Ñ7¸Ó?ˆÜ�fœcÔ"Ð"Ð"ô ”1—7‘7×#Ñ#×:Ñ:¸6ÓBÓC€FÜ”A—G‘G×$Ñ$×;Ñ;¸GÓDÓE€Gð ØØØ Ø(Øñ €Fô ˆ1�:‰:‹<ÓœC §¡Ó 1Ó2°QÑ6Ò6ä”—‘‰Üœœ$Ÿ+™+™ q¨1Ð*<¨q¯z©z«|Ð*<Ó=¸vÀtÑVÈvÑVØô
ð 	
ô
 ()§w¡w×'7Ñ'7×'NÑ'NØ�‰Óó(Ð$€Hˆg˜ô 	ˆA�J‰J‹LÓ˜QÒÜ�Ó Ò"Ü×Ñ˜qÓ! UÒ*à�‰à ™-Ø '™>Ø  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àÐœB×.Ñ.¨qÓ1°UÒ:ä˜Q ¨Ñ7°Ñ7ˆÜ”—‘‰{Ø”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×Ñ×'Ñ'¨¯©Ó6ó
Ðô �O‰O×0Ñ0°Ð4DÓEˆÜ—‘×5Ñ5°fÐ>NÓOˆò%Ð!ð €|Ø�6ˆ{ˆØˆˆv‰Ø%×,Ñ,¨Q°Õ7à�6˜4Ð ˆØ�‰ŒØ×ÑÔÜ	�‰×Ñ×/Ñ/°·±³Ô@à€GÜ‡�×Ñ×7Ñ7¸Ô?ä×!Ñ!ØØØ-ñð ñ	ð
ˆô 	�‰×Ñ×8Ñ8¸ÕBÜ Õ'ä�HÕÚÜ�^Õ$ä�G‰G×Ñ×4Ñ4°W¸a¿j¹j»lÈ1¹oÕNô �LÔ!Ü˜”Ü˜Ô!Ø˜!’à�N‰NÔ.×3Ñ3°D¸&ÓAÔBäÜ˜1Ÿ:™:›<¨™?Ð>¨Q¯Z©Z«\¸!¸"Ð-=Ð>Ó?ØØÜ×*Ñ*¨1Ó-ô	
ó 0	ˆCð �qŠyÜ×3Ñ3Øðà!" F Ø!Ø)¨!™_Ø)¨!™_Ø# A™YØ# A™YØ% a™jØ% a™jØ!ô # <Ó0Ü$Ÿ~™~×3Ñ3×>Ñ>Ø"Ÿ~™~Ø!Ÿm™mñ!ð" —j‘jó#ð& ˜“Ü×3Ñ3Øòà!" F¡ðñ "ðð *¨!š_ð	ð
 *¨!š_ðð *¨!š_ðð $ AšYðð $ AšYðð $ AšYðð & ašjðð & ašjðð & ašjðñ "ðô  # <Ô0ð!ô"  %Ÿ~™~×3Ñ3×>Ò>ð#ð$  #Ÿ~š~ð%ð& "ŸmšmØ—j‘jô)ð70	ôb ˜FÔ#Ü ×4Ñ4ØØØ˜F˜°$Ð2B¨¡wÌËÑPØØØØØ!%õ		
ô % ]°G¸TÀ6ÓJÐJr.   c                ó(   — t        | ||||||||«	      S ©N)rq   )r[   rz   r{   r`   ra   rb   rd   re   rf   Ú	benchmarkÚdeterministicÚcudnn_enabledrÑ   s                r,   Ú_convolutionrà   Á  s%   € ô  Ø	ˆ6�4˜ ¨(°JÀÐPVóð r.   c                óÖ   — | j                   t        j                  j                  j                  j
                  k(  sJ ‚t        j                  j                  r||fS t        | g|¢­i |¤ŽS rÜ   )
ÚtargetrW   ro   rp   rq   Údefaultr   rm   r�   r   )Úfx_noderØ   rJ   s      r,   Úconstrain_conv_to_fx_stridesrå   Ö  sT   € Ø�>‰>œUŸY™YŸ^™^×7Ñ7×?Ñ?Ò?Ð?Ð?Ü‡w�w×ÒØ�Vˆ|Ðä& wÐ@°Ò@¸Ñ@Ð@r.   )r[   r   rz   r   r{   úOptional[TensorBox]r`   úSequence[int]ra   r_   rb   r_   rd   rc   re   r_   rf   r<   Úreturnz	ir.Layout)r[   r   rz   r   r{   ræ   r`   rç   ra   rç   rb   rç   rd   rc   re   rç   rf   r<   )GÚ
__future__r   ÚloggingÚtypingr   r   r   r   rW   Ú-torch._inductor.codegen.rocm.ck_conv_templater   Ú r
   r   Úloweringr   r   r   r�   r   Úselect_algorithmr   r   r   r   rÇ   r   r   r   r   r   r   Úvirtualizedr   Ú	mm_commonr   r   Úcollections.abcr   r   Ú	getLoggerrg   Úlogro   rp   r-   r1   Úkernel_configsr;   rH   ÚversionÚhipÚcudaÚis_availablerB   rK   ÚLOOP_BODY_2DrÍ   ÚLOOP_BODY_3DrÒ   rq   rã   rÉ   r\   rÌ   r^   r~   r‡   r™   rà   rå   r%   r.   r,   ú<module>rü      sN  ðå "ã ß ;Ó ;ã Ý Rç ÷ó ÷ó ÷÷ õ ß @ñ Ý(åà€g×Ñ˜Ó!€ð ‡y�y‡~�~€ð ñó ðð ñó ðð #¨DÑ1Ø"¨DÑ1Ø#¨TÑ2Ø#¨TÑ2Ø!¨4Ñ0Ø"¨DÑ1Ø"¨DÑ1ð	€ñ ñ à ôó Ð ð 	‡=�=×Ò˜Ÿ™×0Ñ0Ô2Ù.Ð/?Ó@Ðòò
?ð€ð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ô ð Øð àð ð ð ð ð	 ð
 ð ð ð ð ð ð $ð ð ð ð ó òFò3ñ. �4×#Ñ#Ó$ðoKØðoKàðoKð ðoKð ð	oKð
 ðoKð ðoKð ðoKð "ðoKð òoKó %ðoKñd �4×$Ñ$Ó%ñó &ðò(Añ �d×&Ñ&Ð(DÕ Er.   