Ë
    g^(h$!  ã                   ó.  — d dl Z d dl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mZmZ ddlmZmZmZmZmZ ddlmZ d	d
lmZmZmZmZmZmZ  e j<                  e«      Z ejB                  jD                  Z"ed„ «       Z#d„ Z$d„ Z% ede#d¬«      Z& eejN                  d«      Z( eejR                  de"jR                  jT                  ¬«      Z+ e	jX                  e"jN                  «      ddœd„«       Z- e	jX                  e"jR                  «      d	d	ddœd„«       Z.y)é    N)Úcounters)ÚCKGemmTemplateé   )ÚirÚlowering)Úautotune_select_algorithmÚExternKernelChoiceÚSymbolicGridFnÚTritonTemplate)Úuse_aten_gemm_kernelsÚuse_ck_gemm_templateÚuse_cpp_bmm_templateÚuse_cutlass_templateÚuse_triton_template)ÚVé   )Ú_is_static_problemÚaddmm_epilogueÚmm_argsÚ
mm_configsÚ
mm_optionsÚshould_fallback_to_atenc                ó:   —  |||d   «       |||d   «      z  | dfS )NÚBLOCK_MÚBLOCK_Nr   © )ÚbÚmÚnÚmetaÚcdivs        úX/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/torch/_inductor/kernel/bmm.pyÚbmm_gridr#   %   s*   € á��D˜‘OÓ$¡t¨A¨t°I©Ó'?Ñ?ÀÀAÐFÐFó    c                 ó2   — | dkD  s
|dkD  s|dkD  ry| |z  dkD  S )Né€   Ti   r   )r   r   Úks      r"   Ú_is_large_block_for_cpur(   *   s&   € àˆ3‚w�!�c’'˜Q šWØØˆq‰5�5‰=Ðr$   c                óN   — |dk(  rt        | ||dt        ¬«      S t        | ||«      S )NÚcpug      à?)ÚscaleÚexclude)r   r(   )r   r   r'   Údevice_types       r"   Úbmm_configsr.   1   s-   € Ø�eÒÜ˜!˜Q ¨Ô6MÔNÐNÜ�a˜˜AÓÐr$   ÚbmmaÐ  
{{def_kernel("A", "B")}}
    M = {{size("A", -2)}}
    N = {{size("B", -1)}}
    K = {{size("A", -1)}}

    stride_aq = {{stride("A", 0)}}
    stride_am = {{stride("A", 1)}}
    stride_ak = {{stride("A", 2)}}

    stride_bq = {{stride("B", 0)}}
    stride_bk = {{stride("B", 1)}}
    stride_bn = {{stride("B", 2)}}

    # 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):
        ram = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
    else:
        ram = rm % M
    if (stride_bk == 1 and stride_bn == K) or (stride_bk == N and stride_bn == 1):
        rbn = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
    else:
        rbn = rn % N

    rk = tl.arange(0, BLOCK_K)

    idx_q = tl.program_id(1)  # batch dimension for BMM
    A = A + (ram[:, None] * stride_am + rk[None, :] * stride_ak + idx_q*stride_aq)
    B = B + (rk[:, None] * stride_bk + rbn[None, :] * stride_bn + idx_q*stride_bq)

    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=ACC_TYPE)
    for k in range(K, 0, -BLOCK_K):
        if EVEN_K:
            a = tl.load(A)
            b = tl.load(B)
        else:
            a = tl.load(A, mask=rk[None, :] < k, other=0.)
            b = tl.load(B, mask=rk[:, None] < k, other=0.)
        acc += tl.dot(a, b, allow_tf32=ALLOW_TF32)
        A += BLOCK_K * stride_ak
        B += BLOCK_K * stride_bk

    # 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_q = tl.program_id(1)  # batch dimension for BMM
    idx_m = rm[:, None]
    idx_n = rn[None, :]
    mask = (idx_m < M) & (idx_n < N)

    # inductor generates a suffix
    {{store_output(("idx_q", "idx_m", "idx_n"), "acc", "mask")}}
)ÚnameÚgridÚsourcezat::bmm_outzat::baddbmm_out)Úop_overload©Úlayoutc                ó�  ‡— t        d„ | |fD «       «      rþ| j                  «       d   dk(  s|j                  «       d   dk(  rWt        j                  | d«      } t        j                  |d«      }t        j                  t        j
                  | |«      d¬«      S d„ }d„ Šˆfd„} || «      r0t        j                  j                  j                  d	   } || |«      }  ||«      r0t        j                  j                  j                  d   } |||«      }t        | ||¬
«      \  }}}	}} }t        d   d|› d|› d|	› �xx   dz  cc<   t        j                  d|||	| j                  «       |j                  «       |«       t        «       rt         j#                  | |f|«      gng }
t%        |«      rOt'        |||	t)        j*                  | «      ¬«      D ]*  }t-        j.                  |
f| |f|dœt1        ||||	|«      ¤Ž Œ, t3        |«      \  }}|r+|r)t5        ||||	«      rddlm} |j;                  |
|| |g«       t=        || |«      rddlm } |jC                  |
|| |g«       tE        ||||	«      rtG        jH                  |
|| |g«       tK        |
«      r'|
jM                  t         j#                  | |f|«      «       tO        d|
| |g|«      S )Nc              3   óV   K  — | ]!  }|j                  «       j                  d k(  –— Œ# y­w)r*   N)Ú
get_deviceÚtype)Ú.0Úxs     r"   ú	<genexpr>ztuned_bmm.<locals>.<genexpr>†   s!   è ø€ Ò
>¨Aˆ1�<‰<‹>×Ñ %Õ'Ñ
>ùs   ‚')r   r   éÿÿÿÿ)Úaxisc                 ó–   — t        j                  | «      syt        j                  | d¬«      \  }}t        |t         j                  «      S )NTF)Úfreeze)r   Úis_storage_and_layoutÚas_storage_and_layoutÚ
isinstanceÚFlexibleLayout)ÚtÚ_r5   s      r"   Úis_valid_to_require_contiguousz1tuned_bmm.<locals>.is_valid_to_require_contiguous�   s<   € Ü×+Ñ+¨AÔ.ØÜ×0Ñ0°¸5ÔA‰IˆAˆvÜ˜f¤b×&7Ñ&7Ó8Ð8r$   c                 ó‚   — |d   dk(  xr | d   dk(  xs |d   | d   k\  xs |d   dk(  xr | d   dk(  xs |d   | d   k\  S )Nr=   r   éþÿÿÿr   )ÚsizesÚstridess     r"   Ú is_preferred_layout_as_bmm_inputz3tuned_bmm.<locals>.is_preferred_layout_as_bmm_input“   sf   € ð ˜‘˜qÑ ÒQ e¨B¡i°1¡nÒ&P¸À¹ÀuÈRÁyÑ8PòUà˜"‘+ Ñ"ÒS¨¨b©	°Q©Ò(R¸'À"¹+ÈÈrÉÑ:RðUr$   c                 óÌ   •— |j                   d   j                  «       }|j                   d   j                  «       } ‰||«      st        j                  j                  | «      } | S )NÚval)r    ÚsizeÚstrider   ÚExternKernelÚrequire_contiguous)rE   Úmeta_trJ   rK   rL   s       €r"   Úmay_require_contiguousz)tuned_bmm.<locals>.may_require_contiguous�   sT   ø€ Ø—K‘K Ñ&×+Ñ+Ó-ˆEØ—k‘k %Ñ(×/Ñ/Ó1ˆGÙ3°E¸7ÔCÜ—O‘O×6Ñ6°qÓ9�ØˆHr$   r   r4   Úaten_mm_infoz	aten.bmm_rF   zPTuned aten.bmm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%s©r-   ©Úinput_nodesr5   )ÚCUTLASS3xGemmTemplate)ÚCppBmmTemplater/   )(ÚallÚget_sizeÚLÚ	unsqueezeÚsum_Úmulr   ÚgraphÚcurrent_nodeÚargsr   r   ÚlogÚinfoÚ	get_dtyper   Úaten_bmmÚbindr   r.   r   Úget_device_typeÚbmm_templateÚmaybe_append_choicer   r   r   Úcodegen.cuda.gemm_templaterY   Úadd_cutlass_gemm_choicesr   Úcodegen.cpp_bmm_templaterZ   Úadd_choicesr   r   Úadd_ck_gemm_choicesr   Úappendr   )Úmat1Úmat2r5   rG   rT   Ú	meta_mat1Ú	meta_mat2r   r   r'   ÚchoicesÚconfigÚstatic_shapeÚ
is_nonzerorY   rZ   rL   s                   @r"   Ú	tuned_bmmrz   „   s£  ø€ ä
Ñ
>°$¸°Ô
>Ô>à�=‰=‹?˜1Ñ Ò" d§m¡m£o°aÑ&8¸AÒ&=Ü—;‘;˜t RÓ(ˆDÜ—;‘;˜t QÓ'ˆDÜ—6‘6œ!Ÿ%™%  dÓ+°!Ô4Ð4ò	9ò	Uô	ñ *¨$Ô/ÜŸ™×,Ñ,×1Ñ1°!Ñ4ˆIÙ)¨$°	Ó:ˆDÙ)¨$Ô/ÜŸ™×,Ñ,×1Ñ1°!Ñ4ˆIÙ)¨$°	Ó:ˆDä")¨$°¸VÔ"DÑ€A€qˆ!ˆV�T˜4ô ˆ^Ñ˜y¨¨¨1¨Q¨C¨q°°Ð4Ó5¸Ñ:Ó5Ü‡H�HØZØ	Ø	Ø	Ø�‰ÓØ�‰ÓØôô 8MÔ7NŒx�}‰}˜d D˜\¨6Ó2Ñ3ÐTV€GÜ˜6Ô"Ü! ! Q¨´r×7IÑ7IÈ$Ó7OÔPò 	ˆFÜ×,Ñ,Øðà! 4˜LØñô ˜V Q¨¨1¨fÓ5ó	ð	ô  2°&Ó9Ñ€L�*Ù™
Ô';¸FÀAÀqÈ!Ô'LÝFà×6Ñ6°wÀÈÈtÈÔUä˜F D¨$Ô/Ý=à×"Ñ"ØØØ�4ˆLô	
ô ˜F A q¨!Ô,Ü×*Ñ*¨7°F¸TÀ4¸LÔIä˜wÔ'Ø�‰”x—}‘} d¨D \°6Ó:Ô;ä$ U¨G°d¸D°\À6ÓJÐJr$   )ÚalphaÚbetar5   c                óT  — t        ||| |¬«      \  }}}}}}} t        d   d|› d|› d|› �xx   dz  cc<   t        j                  d||||j	                  «       |j	                  «       | j	                  «       |«       t        «       rt        j                  | ||f|||¬«      gng }	t        |«      rjt        |||t        j                  |«      ¬«      D ]E  }
t        j                  |	f| ||f|d	œt        |
||||«      ¤dt        |j                   ||«      d
œ¤Ž ŒG t#        d|	| ||g|«      S )Nr4   rU   zaten.baddbmm_rF   r   z\Tuned aten.baddbmm: m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, inp=%s, output_layout=%s)r{   r|   rV   rW   )Úprefix_argsÚepilogue_fnÚbaddbmm)r   r   rd   re   rf   r   Úaten_baddbmmrh   r   r.   r   ri   rj   rk   r   r   Údtyper   )Úinprr   rs   r{   r|   r5   r   r   r'   rv   rw   s              r"   Útuned_baddbmmr„   Û   sS  € ä'.¨t°T¸3ÀvÔ'NÑ$€A€qˆ!ˆV�T˜4 ô ˆ^Ñ˜}¨Q¨C¨q°°°1°Q°CÐ8Ó9¸QÑ>Ó9Ü‡H�HØfØ	Ø	Ø	Ø�‰ÓØ�‰ÓØ�‰‹Øô	ô !Ô"ô 
×	Ñ	˜C  tÐ,¨f¸EÈÐ	Ó	MÑNàð ô
 ˜6Ô"Ü! ! Q¨´r×7IÑ7IÈ$Ó7OÔPò 	ˆFÜ×,Ñ,Øðà  $¨Ð-Øñô ˜V Q¨¨1¨fÓ5ð	ð
 Ü*¨6¯<©<¸ÀÓEõð	ô % Y°¸#¸tÀTÐ9JÈFÓSÐSr$   )/ÚloggingÚtorchÚtorch._dynamo.utilsr   Ú7torch._inductor.codegen.rocm.ck_universal_gemm_templater   Ú r   r   r]   Úselect_algorithmr   r	   r
   r   Úutilsr   r   r   r   r   Úvirtualizedr   Ú	mm_commonr   r   r   r   r   r   Ú	getLoggerÚ__name__rd   ÚopsÚatenr#   r(   r.   rj   r/   rg   r€   Úoutr�   Úregister_loweringrz   r„   r   r$   r"   ú<module>r”      s  ðã ã Ý (Ý Rç  ÷ó ÷õ õ ÷÷ ð €g×Ñ˜Ó!€Ø‡y�y‡~�~€ð ñGó ðGòòñ Ø	Ø	ðAôE€ñN ˜eŸi™i¨Ó7€Ù!Ø	‡M�MÐ$°$·,±,×2BÑ2Bô€ð
 €×Ñ�T—X‘XÓØ$(ó SKó ðSKðl €×Ñ�T—\‘\Ó"Ø,-°A¸dó !Tó #ñ!Tr$   