
    8hP&                     x   S SK r S SKrS SKJr  S SKJr  SSKJrJr	  SSK
JrJrJrJr  SSKJrJrJrJrJrJr  SSKJr  S	S
KJrJrJrJrJrJr  \ R>                  " \ 5      r!\RD                  RF                  r#\S 5       r$S r%\" S\$SSS9r&\" \RN                  S5      r(\" \RN                  SS\#RN                  RR                  S9r*\" \RV                  S\#RV                  RX                  S9r-\	R\                  " \#RN                  5      SSS.S jj5       r/\	R\                  " \#RV                  5      S	S	SS.S j5       r0g)    N)counters)CKGemmTemplate   )irlowering)autotune_select_algorithmExternKernelChoiceSymbolicGridFnTritonTemplate)_use_cutlass_for_opuse_aten_gemm_kernelsuse_ck_gemm_templateuse_cpp_bmm_templateuse_cutlass_templateuse_triton_template)V   )_is_static_problemaddmm_epilogueis_batch_stride_largestmm_argsmm_config_kwargs
mm_optionsc                6    U" XS   5      U" X#S   5      -  U S4$ )NBLOCK_MBLOCK_Nr    )bmnmetacdivs        T/var/www/fran/franai/venv/lib/python3.13/site-packages/torch/_inductor/kernel/bmm.pybmm_gridr$   &   s&    O$tAI'??AFF    c                 6    U S:  d  US:  d  US:  a  gX-  S:  $ )N   Ti   r   )r   r    ks      r#   _is_large_block_for_cpur)   +   s$    3w!c'QW55=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)
    tl.assume(pid_m >= 0)
    tl.assume(pid_n >= 0)

    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")}}
T)namegridsource"cache_codegen_enabled_for_templatezat::bmm_outzat::_bmm_out_dtype_cuda	bmm_dtype)r+   op_overloadzat::baddbmm_out)r0   layoutc                  ^ [        S X4 5       5      (       Ga	  U R                  5       S   S:X  d  UR                  5       S   S:X  aW  [        R                  " U S5      n [        R                  " US5      n[        R                  " [        R
                  " X5      SS9$ S nS mU4S jnU" U 5      (       a/  [        R                  R                  R                  S	   nU" X5      n U" U5      (       a/  [        R                  R                  R                  S   nU" X5      n[        XX2S
9u  pppU R                  5       S	   n[        S   SU SU SU	 SU
 3==   S-  ss'   [        R                  SUUU	U
U R                  5       UR                  5       U5        U(       a;  U R                  5       R                   S:X  d   S5       e["        R%                  X4X2S9nO[&        R%                  X4U5      n[)        5       (       a  U/O/ n[*        R,                  " U 5      n[        R.                  R1                  U5      nU R                  5       n[3        U5      (       a[  Ub   S5       eU" UU	U
40 [5        U[6        UR8                  5      D6 H)  n[:        R<                  " U4X4US.[?        UXX5      D6  M+     [A        U5      u  nn[C        XU5      nU(       aA  U(       a:  [E        X8X5      (       a)  [G        S5      (       a  SSK$J%n  URM                  XX/5        [O        X0U5      (       a  SSK(J)n  URU                  UUX/5        [W        X8X5      (       a  [X        RZ                  " XX/5        []        SXU/U5      $ )zX
Lowering for autotuning aten.bmm with different backends (Aten, Triton, CUTLASS, etc.)
c              3   Z   #    U  H!  oR                  5       R                  S :H  v   M#     g7f)cpuN)
get_devicetype).0xs     r#   	<genexpr>tuned_bmm.<locals>.<genexpr>   s     
>A<<>%'s   )+r   r   )axisc                     [         R                  " U 5      (       d  g[         R                  " U SS9u  p[        U[         R                  5      $ )NTF)freeze)r   is_storage_and_layoutas_storage_and_layout
isinstanceFlexibleLayout)t_r2   s      r#   is_valid_to_require_contiguous1tuned_bmm.<locals>.is_valid_to_require_contiguous   s=    ++A..005AIAfb&7&788r%   c                     US   S:H  =(       a    U S   S:H  =(       d    US   U S   :  =(       d)    US   S:H  =(       a    U S   S:H  =(       d    US   U S   :  $ )Nr<   r   r   )sizesstridess     r#    is_preferred_layout_as_bmm_input3tuned_bmm.<locals>.is_preferred_layout_as_bmm_input   sf     q QeBi1n&PuRy8PU"+"Sb	Q(R'"+r:RUr%   c                    > UR                   S   R                  5       nUR                   S   R                  5       nT" X#5      (       d  [        R                  R                  U 5      n U $ )Nval)r!   sizestrider   ExternKernelrequire_contiguous)rD   meta_trJ   rK   rL   s       r#   may_require_contiguous)tuned_bmm.<locals>.may_require_contiguous   sU    KK&++-Ekk%(//1G3ECCOO66q9Hr%   r   )r2   	out_dtypeaten_mm_infoz	aten.bmm_rE   zZTuned aten.bmm: batch=%s, m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, output_layout=%scudaz$out_dtype is only supported for CUDA)rW   z%out_dtype is not supported for Tritoninput_nodesr2   r*   )CUTLASS3xGemmTemplate)CppBmmTemplate)/allget_sizeL	unsqueezesum_mulr   graphcurrent_nodeargsr   r   loginfo	get_dtyper6   r7   aten_bmm_dtypebindaten_bmmr   r   get_device_typechoicesget_base_mm_configsr   r   r)   itemsizebmm_templatemaybe_append_choicer   r   r   r   r   codegen.cuda.gemm_templater\   add_cutlass_gemm_choicesr   codegen.cpp_bmm_templater]   add_choicesr   r   add_ck_gemm_choicesr   )mat1mat2rW   r2   rF   rU   	meta_mat1	meta_mat2r   r    r(   
batch_size	aten_funcrn   device_typebmm_configsdtypeconfigrE   
is_nonzerobatch_stride_largestr\   r]   rL   s                          @r#   	tuned_bmmr      s.   
 
>$
>>>==?1"dmmoa&8A&=;;tR(D;;tQ'D66!%%+!44	9	U	 *$//,,11!4I)$:D)$//,,11!4I)$:D")6#A!T
 #J^yAaS!AaSABaGBHHd				  %%/W1WW/"''f'R	MM4,7	 344yk"G$$T*K))//<KNNE6"" I"II !
 {,CU^^T	
F ,,!L VQ15	
 'v.MAz24vF A11&&F66wUF$//=""L	
 Fq,,**7TLI$UGD\6JJr%   )alphabetar2   c                   [        XXS9u  pgppn UR                  5       S   n	[        S   SU	 SU SU SU 3==   S-  ss'   [        R	                  SU	UUUUR                  5       UR                  5       U R                  5       U5	        [        5       (       a  [        R                  XU4XSUS9/O/ n
[        R                  " U5      n[        R                  R                  U5      n[        U5      (       av  U" XgU40 [        U[         5      D6 HZ  n["        R$                  " U
4XU4US	.['        XXxU5      DS[)        UR*                  X45      [-        S
UR*                  X4/5      S.D6  M\     [/        SXX/U5      $ )Nr1   r   rX   zaten.baddbmm_rE   r   zkTuned aten.baddbmm: batch_size=%s, m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, inp=%s, output_layout=%s)r   r   rZ   r   )prefix_argsepilogue_fnepilogue_fn_hashbaddbmm)r   r_   r   rg   rh   ri   r   aten_baddbmmrk   r   rm   r   rn   ro   r   r   r)   rq   rr   r   r   r   strr   )inprx   ry   r   r   r2   r   r    r(   r|   rn   r~   r   r   s                 r#   tuned_baddbmmr      s   '.t3'N$A!T #J^}ZL!AaS!EF!KFHHu			
 !"" 
		Ct,f		MN  $$T*K))//<K6""!!
'5LM
F ,, - Vf5	
 *6<<E!$&6e%R!S
 %Yt9JFSSr%   )N)1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   r   virtualizedr   	mm_commonr   r   r   r   r   r   	getLogger__name__rg   opsatenr$   r)   rq   r*   rl   	dtype_outrj   r   outr   register_loweringr   r   r   r%   r#   <module>r      sG     ( R       !yy~~ G G 		CH (,OHT eii7#	II	""	 "	MM$$,,2B2B
 TXXoKD oK oKd T\\",-Ad *T #*Tr%   