Ë
    g^(hLs  ã                   ó(  — d dl Z d dlZd dlZd dlmZ d dlZd dlmZ d dlm	Z	 d dl
mZmZmZmZ d dlmZ d dlmZ dd	lmZmZ dd
lmZmZ ddlmZ ddlmZ ddlmZmZ ddl m!Z! ddl"m#Z#m$Z$m%Z% ddl&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z,m-Z-m.Z.m/Z/ ddl0m1Z1m2Z2m3Z3m4Z4m5Z5m6Z6m7Z7m8Z8m9Z9m:Z:m;Z;m<Z<m=Z= de>dee?e@df      fd„ZA	 d dlBZB eAeBj†                  «      ZDdZEeD�eD\  ZFZGnd ZFd ZG ej’                  eJ«      ZKej˜                  jš                  ZM e%de7ejœ                  jž                  �eEreFdk\  reGdk\  rdnd¬«      ZP e%de:d¬«      ZQ e j¤                  d«      d „ «       ZS e$ej¨                  d!«      ZU e$ej¬                  d"eMj¬                  j®                  ¬#«      ZX e$ej²                  d$«      ZZ e$ej¶                  d%d¬&«      Z\d'„ Z]d(„ Z^d)„ Z_dddd*œd+„Z` e$e`d«      Za e!eMj¨                  d¬,«      dd-œd.„«       Zb e!eMj²                  d¬,«      dd-œd/„«       Zc e!eMj¬                  d¬,«      dddd0œd1„«       Zd e!eMj¶                  d¬,«      ddd2œd3„«       Ze e j¤                  d«      d4ee@   deffd5„«       Zgd6„ Zhd7„ Zi	 	 d<d8ee@   fd9„Zjd:„ Zkd;„ Zly# eH$ r dZDdZEd ZFd ZGY �Œ¸w xY w)=é    N)ÚOptional)Úcounters)ÚAutoHeuristicSelectAlgorithm)Ú	AHContextÚcontext_add_stridesÚcontext_add_using_tf32Úmm_operations)ÚCppGemmTemplate)ÚVé   )ÚconfigÚir)ÚCUTLASS2xGemmTemplateÚCUTLASS3xGemmTemplate)ÚCKGemmTemplate)ÚPythonWrapperCodegen)ÚFlexibleLayoutÚ	is_triton)Úregister_lowering)Úautotune_select_algorithmÚExternKernelChoiceÚTritonTemplate)	Úget_gpu_shared_memoryÚget_tma_workspace_argÚuse_aten_gemm_kernelsÚuse_ck_gemm_templateÚuse_cpp_gemm_templateÚuse_cutlass_templateÚuse_max_autotuneÚuse_triton_templateÚuse_triton_tma_templateé   )Ú_is_static_problemÚaddmm_epilogueÚextra_mm_configsÚint8_mm_configsÚmm_argsÚ
mm_configsÚmm_gridÚ
mm_optionsÚpersistent_mm_configsÚpersistent_mm_gridÚpersistent_mm_optionsÚshould_fallback_to_atenÚtriton_configÚversion_stringÚreturn.c                 óx   — d}t        j                  || «      }|r t        d„ |j                  «       D «       «      S y )Nz(\d+)\.(\d+)?c              3   ó2   K  — | ]  }t        |«      –— Œ y ­w©N)Úint)Ú.0Úgroups     úW/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/torch/_inductor/kernel/mm.pyú	<genexpr>z parse_version.<locals>.<genexpr>?   s   è ø€ Ò< E”S˜—ZÑ<ùs   ‚)ÚreÚmatchÚtupleÚgroups)r0   Úpatternr;   s      r8   Úparse_versionr?   :   s4   € Ø€GÜ�H‰H�W˜nÓ-€EáÜÑ<¨U¯\©\«^Ô<Ó<Ð<àó    TFÚmmé   aX	  
{{def_kernel("A", "B")}}
    M = {{size("A", 0)}}
    N = {{size("B", 1)}}
    K = {{size("A", 1)}}
    if M * N == 0:
        # early exit due to zero-size input(s)
        return
    stride_am = {{stride("A", 0)}}
    stride_ak = {{stride("A", 1)}}
    stride_bk = {{stride("B", 0)}}
    stride_bn = {{stride("B", 1)}}

    # based on triton.ops.matmul
    pid = tl.program_id(0)
    grid_m = (M + BLOCK_M - 1) // BLOCK_M
    grid_n = (N + BLOCK_N - 1) // BLOCK_N

    # re-order program ID for better L2 performance
    width = GROUP_M * grid_n
    group_id = pid // width
    group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
    pid_m = group_id * GROUP_M + (pid % group_size)
    pid_n = (pid % width) // (group_size)

    rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    if ((stride_am == 1 and stride_ak == M) or (stride_am == K and stride_ak == 1)) and M >= BLOCK_M:
        offs_a_m = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
    else:
        offs_a_m = rm % M
    if ((stride_bk == 1 and stride_bn == K) or (stride_bk == N and stride_bn == 1)) and N >= BLOCK_N:
        offs_b_n = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
    else:
        offs_b_n = rn % N
    offs_k = tl.arange(0, BLOCK_K)
    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=ACC_TYPE)

    for k_idx in range(0, tl.cdiv(K, BLOCK_K)):
        {% if not EVEN_K %}
        a_mask = offs_k[None, :] < (K - k_idx * BLOCK_K)
        b_mask = offs_k[:, None] < (K - k_idx * BLOCK_K)
        {% endif %}
        a_k_idx_vals = offs_k[None, :] + (k_idx * BLOCK_K)
        b_k_idx_vals = offs_k[:, None] + (k_idx * BLOCK_K)

        idx_m = offs_a_m[:, None]
        idx_n = a_k_idx_vals
        {{load_input("A", "a", ("idx_m", "idx_n"), mask=None if EVEN_K else "a_mask", indent_width=8)}}

        idx_m = b_k_idx_vals
        idx_n = offs_b_n[None, :]
        {{load_input("B", "b", ("idx_m", "idx_n"), mask=None if EVEN_K else "b_mask", indent_width=8)}}
        acc += tl.dot(a, b, allow_tf32=ALLOW_TF32)

    # rematerialize rm and rn to save registers
    rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    idx_m = rm[:, None]
    idx_n = rn[None, :]
    mask = (idx_m < M) & (idx_n < N)

    # inductor generates a suffix
    {{store_output(("idx_m", "idx_n"), "acc", "mask")}}
a2	  
{{def_kernel("A", "B")}}
    M = {{size("A", 0)}}
    N = {{size("B", 1)}}
    K = {{size("A", 1)}}
    if M * N == 0:
        # early exit due to zero-size input(s)
        return
    stride_am = {{stride("A", 0)}}
    stride_ak = {{stride("A", 1)}}
    stride_bk = {{stride("B", 0)}}
    stride_bn = {{stride("B", 1)}}

    # based on triton.ops.matmul
    pid = tl.program_id(0)
    grid_m = (M + BLOCK_M - 1) // BLOCK_M
    grid_n = (N + BLOCK_N - 1) // BLOCK_N

    # re-order program ID for better L2 performance
    width = GROUP_M * grid_n
    group_id = pid // width
    group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
    pid_m = group_id * GROUP_M + (pid % group_size)
    pid_n = (pid % width) // (group_size)

    rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    if (stride_am == 1 and stride_ak == M) or (stride_am == K and stride_ak == 1):
        offs_a_m = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
    else:
        offs_a_m = rm % M
    if (stride_bk == 1 and stride_bn == K) or (stride_bk == N and stride_bn == 1):
        offs_b_n = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
    else:
        offs_b_n = rn % N
    offs_k = tl.arange(0, BLOCK_K)
    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=ACC_TYPE)

    for k_idx in range(0, tl.cdiv(K, BLOCK_K)):
        {% if not EVEN_K %}
        a_mask = offs_k[None, :] < (K - k_idx * BLOCK_K)
        b_mask = offs_k[:, None] < (K - k_idx * BLOCK_K)
        {% endif %}
        a_k_idx_vals = offs_k[None, :] + (k_idx * BLOCK_K)
        b_k_idx_vals = offs_k[:, None] + (k_idx * BLOCK_K)

        idx_m = offs_a_m[:, None]
        idx_n = a_k_idx_vals
        {{load_input("A", "a", ("idx_m", "idx_n"), mask=None if EVEN_K else "a_mask", indent_width=8)}}

        idx_m = b_k_idx_vals
        idx_n = offs_b_n[None, :]
        {{load_input("B", "b", ("idx_m", "idx_n"), mask=None if EVEN_K else "b_mask", indent_width=8)}}
        acc += tl.dot(a, b, allow_tf32=ALLOW_TF32)

    # rematerialize rm and rn to save registers
    rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    idx_m = rm[:, None]
    idx_n = rn[None, :]
    mask = (idx_m < M) & (idx_n < N)

    # inductor generates a suffix
    {{store_output(("idx_m", "idx_n"), "acc", "mask")}}
)ÚnameÚgridÚsourceÚmm_persistent_tmaal  
{{def_kernel("A", "B")}}
    M = {{size("A", 0)}}
    N = {{size("B", 1)}}
    K = {{size("A", 1)}}
    if M * N == 0:
        # early exit due to zero-size input(s)
        return

    start_pid = tl.program_id(0)
    grid_m = tl.cdiv(M, BLOCK_M)
    grid_n = tl.cdiv(N, BLOCK_N)
    k_tiles = tl.cdiv(K, BLOCK_K)
    num_tiles = grid_m * grid_n
    tiles_per_SM = num_tiles // NUM_SMS
    if start_pid < num_tiles % NUM_SMS:
        tiles_per_SM += 1

    tile_id = start_pid - NUM_SMS
    ki = -1

    width = GROUP_M * grid_n
    rk_for_mask = tl.arange(0, BLOCK_K)
    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=ACC_TYPE)

    workspace_base = ws_ptr + start_pid * 2 * TMA_SIZE
    a_desc_ptr = workspace_base
    b_desc_ptr = workspace_base + TMA_SIZE

    triton.language.extra.cuda.experimental_device_tensormap_create2d(
        desc_ptr=a_desc_ptr,
        global_address=A,
        load_size=[BLOCK_M, BLOCK_K] if A_ROW_MAJOR else [BLOCK_K, BLOCK_M],
        global_size=[M, K] if A_ROW_MAJOR else [K, M],
        element_ty=A.dtype.element_ty,
    )
    triton.language.extra.cuda.experimental_device_tensormap_create2d(
        desc_ptr=b_desc_ptr,
        global_address=B,
        load_size=[BLOCK_K, BLOCK_N] if B_ROW_MAJOR else [BLOCK_N, BLOCK_K],
        global_size=[K, N] if B_ROW_MAJOR else [N, K],
        element_ty=B.dtype.element_ty,
    )

    tl.extra.cuda.experimental_tensormap_fenceproxy_acquire(a_desc_ptr)
    tl.extra.cuda.experimental_tensormap_fenceproxy_acquire(b_desc_ptr)

    pid_m = 0
    pid_n = 0
    rm = 0
    rn = 0

    for _ in range(0, k_tiles * tiles_per_SM):
        ki = tl.where(ki == k_tiles - 1, 0, ki + 1)
        if ki == 0:
            tile_id += NUM_SMS
            # re-order program ID for better L2 performance
            group_id = tile_id // width
            group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
            pid_m = group_id * GROUP_M + (tile_id % group_size)
            pid_n = (tile_id % width) // (group_size)

            rm = pid_m * BLOCK_M
            rn = pid_n * BLOCK_N

        rk = ki * BLOCK_K

        a = tl._experimental_descriptor_load(
            a_desc_ptr,
            [rm, rk] if A_ROW_MAJOR else [rk, rm],
            [BLOCK_M, BLOCK_K] if A_ROW_MAJOR else [BLOCK_K, BLOCK_M],
            A.dtype.element_ty,
        )
        b = tl._experimental_descriptor_load(
            b_desc_ptr,
            [rk, rn] if B_ROW_MAJOR else [rn, rk],
            [BLOCK_K, BLOCK_N] if B_ROW_MAJOR else [BLOCK_N, BLOCK_K],
            B.dtype.element_ty,
        )
        acc += tl.dot(
            a if A_ROW_MAJOR else a.T,
            b if B_ROW_MAJOR else b.T,
            allow_tf32=ALLOW_TF32,
        )

        if ki == k_tiles - 1:
            # rematerialize rm and rn to save registers
            rcm = rm + tl.arange(0, BLOCK_M)
            rcn = rn + tl.arange(0, BLOCK_N)
            idx_m = rcm[:, None]
            idx_n = rcn[None, :]
            mask = (idx_m < M) & (idx_n < N)

            # inductor generates a suffix
            {{store_output(("idx_m", "idx_n"), "acc", "mask", indent_width=12)}}
            acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=ACC_TYPE)
c                 ó   — t        | «      S r4   )r   )Úfns    r8   Úlazy_register_extern_choicerI   M  s   € ä˜bÓ!Ð!r@   z
at::mm_outzat::addmm_out)Úop_overloadzat::_int_mm_outzat::_sparse_semi_structured_mm)Úhas_out_variantc                 ób   — | j                  «       t        j                  t        j                  fv S r4   )Ú	get_dtypeÚtorchÚint8Úuint8)Úmats    r8   Ú_is_int8_matrR   a  s    € Ø�=‰=‹?œuŸz™z¬5¯;©;Ð7Ð7Ð7r@   c                 ó   — | |z  dkD  S )Ni    © )ÚmÚnÚks      r8   Ú_is_large_block_for_cpurX   e  s   € àˆq‰5�5‰=Ðr@   c                 ó"   — | dk(  r	dt         dœS i S )NÚcpug      à?)ÚscaleÚexclude)rX   )Údevices    r8   Úmm_config_kwargsr^   j  s    € Ø�‚àÜ.ñ
ð 	
ð €Ir@   ©ÚoutÚalphaÚbetac                óÄ   — | j                  d«      dk(  s| j                  d«      dk(  rt        j                  | d   |||||¬«      S t        j                  | |||||¬«      S )z¬
    Giving torch.addmm a 1D tensor calls a different (faster) cublasLt
    kernel under the hood.  There are a few shapes where this is slower,
    but they are rare.
    r   r"   r_   )ÚstrideÚsizerN   Úaddmm)ÚinpÚmat1Úmat2r`   ra   rb   s         r8   Ú
bias_addmmrj   s  sY   € ð ‡z�z�!ƒ}˜Ò˜SŸX™X a›[¨AÒ-Ü�{‰{˜3˜q™6 4¨°3¸eÈ$ÔOÐOÜ�;‰;�s˜D $¨C°uÀ4ÔHÐHr@   )Útype_promotion_kind©Úlayoutc                óö  — t        | ||¬«      \  }}}}} }d}t        d   d|› d|› d|› �xx   dz  cc<   t        j                  d|||| j	                  «       |j	                  «       |«       |}t        «       s,t        |j                  |j                  |j                  ¬«      }t        «       rt        j                  | |f|«      gng }t        |«      \  }	}
|
rët        |«      ràt        |||fi t!        t#        j$                  | «      «      ¤ŽD ]*  }t'        j(                  |f| |f|d	œt+        |||||«      ¤Ž Œ, t-        | |«      r}t/        |||fi t!        t#        j$                  | «      «      ¤ŽD ]P  }t1        j(                  |f| |f|t3        d
| j5                  «       ¬«      dœt+        |||||«      ¤t7        | |«      ¤Ž ŒR |
r't9        ||||«      rt;        j<                  ||| |g«       |
r't?        ||||«      rtA        jB                  ||| |g«       tE        || |«      rtG        jH                  ||| |g«       | |g}|
�r1t        |«      �r%tJ        jL                  jN                  jQ                  |«      rütS        | «      rñg }t        «       r|jU                  d«       tW        |«      }tY        |||fi t!        t#        j$                  | «      «      ¤ŽD ]*  }t'        j(                  |f| |f|d	œt+        |||||«      ¤Ž Œ, t[        | |||||||t]        «       d d|¬«      }tJ        jL                  jN                  j_                  |«      s*|�#tW        |«      dkD  r|D �cg c]	  }||v sŒ|‘Œ }}n|d | }t`        jb                  D ].  }|jU                  te        |«      j                  | |f|«      «       Œ0 tg        |«      r&t        j                  | |f|«      ji                  «       S tk        ||| |g|«      S c c}w )Nrl   rA   Úaten_mm_infozaten.mm_Ú_r"   zOTuned aten.mm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%s©r]   Údtypere   ©Úinput_nodesrm   r   ©Únum_tma_descriptorsr]   ©rt   rm   Úworkspace_argÚ	extern_mmé
   )Útop_kÚalways_includedr   )6r'   r   ÚlogÚinforM   r   r   r]   rr   re   r   Úaten_mmÚbindr#   r    r(   r^   r   Úget_device_typeÚmm_templateÚmaybe_append_choicer*   r!   r+   Úpersistent_tma_mm_templater   Ú
get_devicer-   r   r   Úadd_cutlass_gemm_choicesr   r   Úadd_ck_gemm_choicesr   r
   Úadd_choicesrN   Ú	_inductorr   Úrun_autoheuristicr   ÚappendÚlenr%   Úmm_autoheuristicr	   Úcollect_autoheuristicÚinductor_configÚexternal_matmulrI   r.   Úoutput_noder   )rh   ri   rm   rU   rV   rW   rC   Úaten_layoutÚchoicesÚstatic_shapeÚ
is_nonzeror   rt   r|   Ú num_choices_before_extra_configsÚ
ah_choicesÚchoices                    r8   Útuned_mmr™   �  s)  € ä")¨$°¸VÔ"DÑ€A€qˆ!ˆV�T˜4Ø€Dô ˆ^Ñ˜x¨ s¨!¨A¨3¨a°¨sÐ3Ó4¸Ñ9Ó4Ü‡H�HØYØ	Ø	Ø	Ø�‰ÓØ�‰ÓØôð €KÜÔÜ$Ø—=‘=¨¯©¸6¿;¹;ô
ˆô 6KÔ5LŒ�‰�t˜T�l KÓ	0Ñ1ÐRTð ô  2°&Ó9Ñ€L�*ÙÔ)¨&Ô1Ü   A qÑWÔ,<¼R×=OÑ=OÐPTÓ=UÓ,VÑWò 	ˆFÜ×+Ñ+Øðà! 4˜LØñô ˜V Q¨¨1¨fÓ5ó	ð	ô # 4¨Ô.Ü/Ø�1�añÜ+¬B×,>Ñ,>¸tÓ,DÓEñò �ô +×>Ñ>Øð
à!% t Ø!Ü"7Ø,-Ø#Ÿ™Ó0ô#ñ	
ô ! ¨¨A¨q°&Ó9ð
ô ,¨D°$Ó7ó
ðñ Ô*¨6°1°a¸Ô;Ü×6Ñ6°wÀÈÈtÈÔUáÔ*¨6°1°a¸Ô;Ü×*Ñ*¨7°F¸TÀ4¸LÔIä˜V T¨4Ô0Ü×#Ñ#ØØØ�4ˆLô	
ð ˜�,€KâÜ Õ'Ü�O‰O×"Ñ"×4Ñ4°TÔ:Ü�dŒOàˆÜ Ô"Ø×"Ñ" ;Ô/Ü+.¨w«<Ð(Ü&Øˆq�!ñ
Ü'¬×(:Ñ(:¸4Ó(@ÓAñ
ò 	ˆFô ×+Ñ+Øðà! 4˜LØñô ˜V Q¨¨1¨fÓ5ó	ð	ô &ØØØØØØØØÜ‹OØØØ+ô
ˆ
ô �‰×%Ñ%×;Ñ;¸DÔAàÐ%¬#¨j«/¸AÒ*=ð
 18ÖP f¸6ÀZÒ;Oš6ÐP�ÑPà!Ð"CÐ#CÐD�ä×,Ñ,ò RˆØ�‰Ô2°1Ó5×:Ñ:¸DÀ$¸<ÈÓPÕQðRô ˜wÔ'Ü�|‰|˜T 4˜L¨+Ó6×BÑBÓDÐDä$ T¨7°T¸4°LÀ&ÓIÐIùò Qs   Í	O6Í(O6c                ó  — t        | ||t        j                  ¬«      \  }}}}} }t        d   d|› d|› d|› �xx   dz  cc<   t        j                  d|||| j                  «       |j                  «       |«       t        |«      \  }}|xr |xr t        ||||«      }t        «       rt        j                  | |f|«      gng }	|rt        j                  |	|| |gdd¬«       |rdt        |d¬	«      rWt        |||fi t!        t#        j$                  | «      «      ¤ŽD ]*  }
t'        j(                  |	f| |f|d
œt+        |
||||«      ¤Ž Œ, t-        |	«      r&t        j                  | |f|«      j/                  «       S t1        d|	| |g|«      S )N)rm   Ú	out_dtypero   zaten._int_mm_rp   r"   zTTuned aten._int_mm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sT©ÚfuseableÚnon_fuseable)Úenable_int32rs   Úint_mm)r'   rN   Úint32r   r}   r~   rM   r#   r   r   Úaten__int_mmr€   r   r†   r    r&   r^   r   r�   r‚   rƒ   r*   r.   r‘   r   )rh   ri   rm   rU   rV   rW   r”   r•   Úuse_cutlassr“   r   s              r8   Útuned_int_mmr¤   ú  s°  € ä")Øˆd˜6¬U¯[©[ô#Ñ€A€qˆ!ˆV�T˜4ô
 ˆ^Ñ˜}¨Q¨C¨q°°°1°Q°CÐ8Ó9¸QÑ>Ó9Ü‡H�HØ^Ø	Ø	Ø	Ø�‰ÓØ�‰ÓØôô  2°&Ó9Ñ€L�*ØÒW :ÒWÔ2FÀvÈqÐRSÐUVÓ2W€Kô 6KÔ5LŒ×	Ñ	˜D $˜<¨Ó	0Ñ1ÐRTð ñ Ü×6Ñ6Ø�V˜d D˜\°DÀtõ	
ñ Ô)¨&¸tÕDÜ%Øˆq�!ñ
Ü'¬×(:Ñ(:¸4Ó(@ÓAñ
ò 	ˆFô ×+Ñ+Øðà! 4˜LØñô ˜V Q¨¨1¨fÓ5ó	ð	ô ˜wÔ'Ü× Ñ  $¨ ¨vÓ6×BÑBÓDÐDä$ X¨w¸¸t¸ÀfÓMÐMr@   )ra   rb   rm   c                óš  — d}t        ||| |¬«      \  }}}	}}}}
t        |«      \  }}t        d   d|› d|› d|	› �xx   dz  cc<   t        j	                  d|||	|j                  «       |j                  «       |«       |r
t        «       swdd	lm}m	} t        ||«      r) ||j                  |j                  |j                  ¬
«      }t        «       rt        j!                  | ||f|||¬«      gng }t#        d|| ||g|«      S t        «       rt        j!                  |
||f|||¬«      gng }t        «       ry|
j%                  «       d   dk(  rc|
j'                  «       j(                  dk(  rFt*        j,                  j.                  r,|j1                  dt2        j!                  |
||f|||¬«      «       |�r"t5        |«      �rt7        |||	fi t9        t;        j<                  |«      «      ¤ŽD ]E  }t?        j@                  |f|
||f|dœtC        ||||	|«      ¤dtE        |j                  ||«      dœ¤Ž ŒG tG        ||«      r˜tI        |||	fi t9        t;        j<                  |«      «      ¤ŽD ]k  }tK        j@                  |f|
||f|tM        d|j'                  «       ¬«      dœtC        ||||	|«      ¤tO        ||«      ¤dtE        |j                  ||«      dœ¤Ž Œm |r_|r]tQ        ||||	«      rOtS        jT                  |
jV                  jX                  d   «      dk7  r t[        j\                  |||||
g||g d¢¬«       |r.t_        ||||	«      r ta        jb                  |||||
g||g d¢¬«       te        |||«      rtg        jh                  |||
||g||d¬«       tk        |«      r¥|jm                  t        j!                  |
||f||||¬«      «       |
j%                  «       d   dk(  rc|
j'                  «       j(                  dk(  rFt*        j,                  j.                  r,|j1                  dt2        j!                  |
||f|||¬«      «       t#        d||
||g|«      S )N)rb   ra   rl   ro   zaten.addmm_rp   r"   zRTuned aten.addmm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%sr   )ÚFixedLayoutr   rq   )ra   rb   rf   Úcudars   )Úprefix_argsÚepilogue_fnr   ru   rw   éÿÿÿÿ)r   r   r"   )ra   rb   Úinput_reorderT)ra   rb   Úhas_bias)7r'   r#   r   r}   r~   rM   r   Útorch._inductor.irr¦   r   Ú
isinstancer]   rr   re   r   Ú
aten_addmmr€   r   Ú
get_strider…   Útyper�   ÚtritonÚautotune_cublasLtÚinsertÚaten_bias_addmmr    r(   r^   r   r�   r‚   rƒ   r*   r$   r!   r+   r„   r   r-   r   r   Ústatically_known_int_or_nonerm   rd   r   r†   r   r   r‡   r   r
   rˆ   r.   r‹   )rg   rh   ri   ra   rb   rm   Úordered_kwargs_for_cpp_kernelrU   rV   rW   Úinp_expandedr”   r•   r¦   r   r“   r   s                    r8   Útuned_addmmr¹   (  sÈ  € à$5Ð!Ü07¸¸dÀCÐPVÔ0WÑ-€A€qˆ!ˆV�T˜4 Ü1°&Ó9Ñ€L�*ô ˆ^Ñ˜{¨1¨#¨Q¨q¨c°°1°#Ð6Ó7¸1Ñ<Ó7Ü‡H�HØ\Ø	Ø	Ø	Ø�‰ÓØ�‰ÓØôñ Ô 0Ô 2÷ 	Cä�f˜kÔ*Ù#Ø—}‘}¨F¯L©L¸v¿{¹{ôˆFô %Ô&ô —‘Ø˜$ Ð%ØØØð	  ó ñð ð 	ô )¨°'¸CÀÀtÐ;LÈfÓUÐUô !Ô"ô �O‰OØ˜t TÐ*ØØØð	 ó ñ	
ð ð ô 	ÔØ×#Ñ#Ó% aÑ(¨AÒ-Ø×#Ñ#Ó%×*Ñ*¨fÒ4Ü×"Ñ"×4Ò4ð 	�‰ØÜ× Ñ Ø˜t TÐ*¨F¸%Àdð !ó ô	
ò Ô)¨&Õ1Ü   A qÑWÔ,<¼R×=OÑ=OÐPTÓ=UÓ,VÑWò 	ˆFÜ×+Ñ+Øðà)¨4°Ð6Øñô ˜V Q¨¨1¨fÓ5ð	ð
 Ü*¨6¯<©<¸ÀÓEõð	ô # 4¨Ô.Ü/Ø�1�añÜ+¬B×,>Ñ,>¸tÓ,DÓEñò �ô +×>Ñ>Øðà!-¨t°TÐ :Ø!Ü"7Ø,-Ø#Ÿ™Ó0ô#ñ	ô ! ¨¨A¨q°&Ó9ðô ,¨D°$Ó7ðð !"Ü .¨v¯|©|¸UÀDÓ Iõðñ" ™
Ô';¸FÀAÀqÈ!Ô'Lô
 !×=Ñ=Ø×#Ñ#×*Ñ*¨2Ñ.óð òô
 "×:Ñ:ØØØ�t˜\Ð*ØØÚ'õñ Ô*¨6°1°a¸Ô;Ü×*Ñ*ØØØ�4˜Ð&ØØÚ#õ	
ô ˜V T¨4Ô0Ü×#Ñ#ØØØ˜4 Ð&ØØØõ	
ô ˜wÔ'Ø�‰Ü�O‰OØ˜t TÐ*ØØ-ØØð ó ô	
ð ×#Ñ#Ó% aÑ(¨AÒ-Ø×'Ñ'Ó)×.Ñ.°&Ò8Ü×&Ñ&×8Ò8ð �N‰NØÜ×$Ñ$Ø! 4¨Ð.°¸eÈ$ð %ó ôô %Ø�˜<¨¨tÐ4°fóð r@   )r›   rm   c                ó¤  — ddl m}  || ||«      \  } }}| j                  «       \  }}|j                  «       \  }}	|j                  «       \  }
}t        j                  j
                  j                  ||«      }t        j                  j
                  j                  d|z  |
«      }|€6ddlm}  ||j                  «       |r|n|j                  «       ||g|dg«      }n	|�J d«       ‚t        «       rt        j                  | ||f||¬«      gng }||z  dk7  r+t        ||||«      rt        j                   ||| ||gdd¬	«       t#        d
|| ||g|«      S )Nr   )Úrealize_inputsr   )r¦   r"   z,out_dtype is ignored if layout is specified.)r›   Trœ   Úsparse_semi_structured_mm)Ú torch._inductor.select_algorithmr»   Úget_sizer   ÚgraphÚsizevarsÚguard_equalsr­   r¦   r…   rM   r   Úaten__sparse_semi_structured_mmr€   r   r   r†   r   )rh   Ú	mat1_metari   r›   rm   r»   Úm1Úk1Úm2rp   Úk2rV   rU   rW   r¦   r“   s                   r8   Útuned_sparse_semi_structured_mmrÈ   Í  sj  € õ @á*¨4°¸DÓAÑ€Dˆ)�TØ�]‰]‹_�F€BˆØ×ÑÓ �E€BˆØ�M‰M‹O�E€BˆÜ	�‰×Ñ×%Ñ% b¨"Ó-€AÜ	�‰×Ñ×%Ñ% a¨"¡f¨bÓ1€Aà€~Ý2áØ�O‰OÓÙ"‰I¨¯©Ó(8Ø�ˆFØ�ˆFó	
‰ð Ð ÐPÐ"PÓPÐ ô !Ô"ô	 ,×0Ñ0Ø�y $Ð'¨¸9ð 1ó ñ	
ð ð ð 	ˆ1�u�‚zÔ*¨6°1°a¸Ô;Ü×6Ñ6Ø�V˜d D¨)Ð4¸tÐRVõ	
ô %Ø# W¨t°YÀÐ.EÀvóð r@   Úindexc                 óf   — t         j                  j                  | xs d«      }|j                  dk  S )Nr   é   )rN   r§   Úget_device_propertiesÚmajor)rÉ   Úpropss     r8   Ú_is_sm7x_or_older_gpurÏ   ú  s)   € ä�J‰J×,Ñ,¨UªZ°aÓ8€EØ�;‰;˜!ÑÐr@   c                 ó&   — t        d„ | D «       «      S )Nc              3   ó<   K  — | ]  }t        |t        «      –— Œ y ­wr4   )r®   r5   )r6   Údims     r8   r9   zdims_are_int.<locals>.<genexpr>  s   è ø€ Ò4¨Œz˜#œs×#Ñ4ùs   ‚)Úall)Údimss    r8   Údims_are_intrÕ      s   € ÜÑ4¨tÔ4Ó4Ð4r@   c                 óÜ  — t        ||| ||«      \  } }}t        | ||g«      sy |j                  t        j                  k7  ry t        j
                  j                  «       dk\  rt        «       dk7  ry | dk(  r|dz  dk7  s|dz  dk7  ry | dk  r|dk\  r|dk\  rt        dddd	d
¬«      S | dkD  r| dk  r|dk\  r|dk\  rt        dddd	d
¬«      S | dkD  r| dk  r|dk\  r|dk\  rt        dddd	d
¬«      S y )N)é   r   i Œ r"   é   r   i   é@   é€   é   é   )ÚBLOCK_MÚBLOCK_NÚBLOCK_KÚ
num_stagesÚ	num_warpsé    )	Úget_size_hintsrÕ   rr   rN   Úfloat16r§   Úget_device_capabilityr   r/   )rU   rV   rW   r“   rh   ri   Ú
mat2_dtyperm   s           r8   Útry_heuristicrç     s#  € Ü˜T 4¨¨A¨qÓ1�G€A€qˆ!Ü˜˜A˜q˜	Ô"Øà‡z�z”U—]‘]Ò"Øô �J‰J×,Ñ,Ó.°&Ò8Ü	Ó	  FÒ	*ØàˆA‚v�1�r‘6˜Q’; ! b¡&¨A¢+ØàˆB‚w�1˜’9  d¢ÜØØØØØô
ð 	
ð 
ˆRŠ�A˜’G  T¢	¨a°4ªiÜØØØØØô
ð 	
ð 
ˆRŠ�A˜’G  T¢	¨a°4ªiÜØØØØØô
ð 	
ð r@   r{   c           	      ó   ‡— t        | ||||«      \  }}}t        |||g«      sy t        | |«      \  }}ˆfd„}d„ } ||||| |||«      }t        ||||‰||	¬«      }|
�|j	                  |
|¬«      S |j                  «       S )Nc                 óV  •— t        «       }|j                  d| «       |j                  d|«       |j                  d|«       |j                  d|j                  j                  d¬«       |j                  d|j                  j                  d¬«       t	        |d|«       t	        |d	|«       |j                  d
|j                  j                  «       d¬«       |j                  d|j                  j                  «       d¬«       ‰dk(  r t        ||j                  j                  «       |S )NrU   rW   rV   Ú
mat1_dtypeT)Úis_categoricalræ   rh   ri   Úmat1_iscontigÚmat2_iscontigrA   )r   Úadd_featurerm   rr   r   Úis_contiguousr   )	rU   rW   rV   rh   ri   Úmat1_strideÚmat2_strideÚcontextrC   s	           €r8   Úget_contextz%mm_autoheuristic.<locals>.get_contextE  s  ø€ Ü“+ˆØ×Ñ˜C Ô#Ø×Ñ˜C Ô#Ø×Ñ˜C Ô#Ø×Ñ˜L¨$¯+©+×*;Ñ*;ÈDÐÔQØ×Ñ˜L¨$¯+©+×*;Ñ*;ÈDÐÔQÜ˜G V¨[Ô9Ü˜G V¨[Ô9Ø×ÑØ˜TŸ[™[×6Ñ6Ó8Èð 	ô 	
ð 	×ÑØ˜TŸ[™[×6Ñ6Ó8Èð 	ô 	
ð �4Š<ä" 7¨D¯K©K×,=Ñ,=Ô>Øˆr@   c                   ó   — y r4   rT   rT   r@   r8   Úfallbackz"mm_autoheuristic.<locals>.fallbackY  s   € Ør@   )rõ   r“   rt   rò   rC   Úaugment_contextÚprecondition)r|   )rã   rÕ   Úget_size_hints_stridesr   Úget_top_k_choices_callerÚget_choice_caller)rh   ri   rU   rV   rW   r“   rC   rt   Úopsr÷   r{   r|   rð   rñ   ró   rõ   rò   Úautoheuristics         `           r8   r�   r�   2  s¶   ø€ ô ˜T 4¨¨A¨qÓ1�G€A€qˆ!Ü˜˜A˜q˜	Ô"ØÜ5°d¸DÓAÑ€K�ôò(ñ ˜!˜Q  4¨¨{¸KÓH€GÜ0ØØØØØØØ!ô€Mð Ðà×5Ñ5Ø ?ð 6ó 
ð 	
ð ×*Ñ*Ó,Ð,r@   c                 ó  — t        |t        «      rt        |t        «      s^t        j                  j                  j                  | j                  «       t        j                  j                  j                  ¬«      \  }}t        |t        «      rt        |t        «      s^t        j                  j                  j                  |j                  «       t        j                  j                  j                  ¬«      \  }}|||fS )N©rõ   )r®   r5   r   r¿   rÀ   Ú
size_hintsr¾   rN   r‰   r   Úunbacked_symint_fallback)rh   ri   rU   rV   rW   s        r8   rã   rã   p  sµ   € Ü�aœÔ¤Z°´3Ô%7Ü—‘×!Ñ!×,Ñ,Ø�M‰M‹OÜ—_‘_×+Ñ+×DÑDð -ó 
‰ˆˆAô
 �aœÔ¤Z°´3Ô%7Ü—‘×!Ñ!×,Ñ,Ø�M‰M‹OÜ—_‘_×+Ñ+×DÑDð -ó 
‰ˆˆAð ˆa�ˆ7€Nr@   c                 ód  — | j                   j                  }|j                   j                  }||g}g }|D ]p  }t        |t        «      sMt        j
                  j                  j                  |t        j                  j                  j                  ¬«      }|j                  |«       Œr |d   |d   fS )Nrþ   r   r"   )rm   rd   r®   r5   r   r¿   rÀ   rÿ   rN   r‰   r   r   r‹   )rh   ri   rð   rñ   ÚstridesÚstrides_hintsrd   s          r8   rø   rø     s¤   € Ø—+‘+×$Ñ$€KØ—+‘+×$Ñ$€KØ˜KÐ(€GØ€MØò %ˆÜ˜&¤#Ô&Ü—W‘W×%Ñ%×0Ñ0ØÜŸ™×/Ñ/×HÑHð 1ó ˆFð 	×Ñ˜VÕ$ð%ð ˜Ñ˜]¨1Ñ-Ð-Ð-r@   )NN)mÚ	functoolsÚloggingr:   Útypingr   rN   Útorch._dynamo.utilsr   Ú+torch._inductor.autoheuristic.autoheuristicr   Ú1torch._inductor.autoheuristic.autoheuristic_utilsr   r   r   r	   Ú)torch._inductor.codegen.cpp_gemm_templater
   Útorch._inductor.virtualizedr   Ú r   r�   r   Úcodegen.cuda.gemm_templater   r   Ú'codegen.rocm.ck_universal_gemm_templater   Úcodegen.wrapperr   r   r   Úloweringr   Úselect_algorithmr   r   r   Úutilsr   r   r   r   r   r   r   r    r!   Ú	mm_commonr#   r$   r%   r&   r'   r(   r)   r*   r+   r,   r-   r.   r/   Ústrr<   r5   r?   r²   Ú__version__Útriton_versionÚ
has_tritonÚtriton_majorÚtriton_minorÚImportErrorÚ	getLoggerÚ__name__r}   rû   ÚatenÚversionÚhipr‚   r„   Ú	lru_cacherI   rA   r   rf   Údefaultr¯   Ú_int_mmr¢   Ú_sparse_semi_structured_mmrÂ   rR   rX   r^   rj   rµ   r™   r¤   r¹   rÈ   ÚboolrÏ   rÕ   rç   r�   rã   rø   rT   r@   r8   ú<module>r%     s  ðã Û Û 	Ý ã Ý (Ý T÷ó õ FÝ )ç ,ß UÝ DÝ 2ß *Ý (÷ñ ÷

÷ 
õ 
÷÷ ÷ õ ð" #ð ¨(°5¸¸c¸±?Ñ*Có ðÛá" 6×#5Ñ#5Ó6€NØ€JØÐ!Ø%3Ñ"ˆ‘làˆØˆð €g×Ñ˜Ó!€Ø‡y�y‡~�~€áØ	Ø	ðF �M‰M×ÑÐ%Ù˜<¨1Ò,°ÀÒ1BñE@	ðL@ôUL€ñ\ ,Ø	Ø	ð`ôdÐ ðP €×Ñ�TÓñ"ó ð"ñ ˜UŸX™X |Ó
4€áØ	‡K�K�¨d¯j©j×.@Ñ.@ô€
ñ " %§-¡-Ð1BÓC€á"4Ø	×$Ñ$Ø$Øô#Ð ò8òò
ð (,°1¸1ô Iñ % Z°Ó6€ñ �4—7‘7°Ô5Ø#'ó uJó 6ðuJñp �4—<‘<°TÔ:Ø'+ó *Nó ;ð*NñZ �4—:‘:°4Ô8Ø*+°!¸Dó aó 9ðañH �4×2Ñ2ÈÔMà(,°Tó)ó Nð)ðX €×Ñ�TÓð ¨#¡ð °4ò ó ðò
5ò+ðr  Øñ;-ð �C‰=ó;-ò|ó.øðc ò Ø€NØ€JØ€LØƒLð	ús   Â4$J  Ê JÊJ