
    sg                        d dl Z d dlZddlmZmZ ddlmZmZ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mZmZ  e j2                  e      Zej8                  j:                  Zd
 Z e	ded      Z eej@                  d      Z! eejD                  d      Z# ejH                  ej@                        ddd       Z%dddddZ&y)    N   )irlowering)autotune_select_algorithmExternKernelChoiceTritonTemplate)ceildivuse_aten_gemm_kernelsuse_cutlass_templateuse_triton_template)V   )_is_static_problem)addmm_epiloguemm_args
mm_configs
mm_optionsc                 F    t        ||d         t        ||d         z  | dfS )NBLOCK_MBLOCK_Nr   )cdiv)bmnmetas       M/var/www/html/venv/lib/python3.12/site-packages/torch/_inductor/kernel/bmm.pybmm_gridr      s*    DO$tAtI'??AFF    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layoutc                V   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               rt        j                  | |f|      gng }
t        |      r:t        |||	      D ]*  }t!        j"                  |
f| |f|dt%        ||||	|       , t'        | |g|      \  }}|r+|r)t)        ||||	      rddlm} |j/                  |
|| |g       t1        |
      d	k(  r<t2        j5                  d       |
j7                  t        j                  | |f|             t9        d|
| |g|      S )Nc              3   V   K   | ]!  }|j                         j                  d k(   # yw)cpuN)
get_devicetype).0xs     r   	<genexpr>ztuned_bmm.<locals>.<genexpr>l   s!     
>A1<<>%'
>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_r$   s      r   is_valid_to_require_contiguousz1tuned_bmm.<locals>.is_valid_to_require_contiguouss   s<    ++A.005AIAvfb&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    )sizesstridess     r    is_preferred_layout_as_bmm_inputz3tuned_bmm.<locals>.is_preferred_layout_as_bmm_inputy   sf     q QeBi1n&PuRy8P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)r5   meta_tr;   r<   r=   s       r   may_require_contiguousz)tuned_bmm.<locals>.may_require_contiguous   sT    KK&++-Ekk%(//1G3E7COO66q9Hr   r   r#   input_nodesr$   )CUTLASS3xGemmTemplatez3No choices for GEMM, using ATen backend as fallbackr   )allget_sizeL	unsqueezesum_mulr   graphcurrent_nodeargsr   r
   aten_bmmbindr   r   bmm_templatemaybe_append_choicer   r   r   codegen.cuda.gemm_templaterH   add_cutlass_gemm_choiceslenlogwarningappendr   )mat1mat2r$   r7   rE   	meta_mat1	meta_mat2r   r   kchoicesconfigstatic_shape
is_nonzerorH   r=   s                  @r   	tuned_bmmre   j   s   

>$
>>==?1"dmmoa&8A&=;;tR(D;;tQ'D66!%%d+!44	9	U	 *$/,,11!4I)$	:D)$/,,11!4I)$	:D")$V"DAq!VT4 8M7Nx}}dD\623TVG6" Aq) 	F,,!4L VQ1f5		  24,GL*
';FAq!'LF66wtU
7|qIJx}}dD\6:;$UGdD\6JJr   )alphabetar$   c                d   t        ||| |      \  }}}}}}} t               rt        j                  | ||f|||      gng }	t	        |      rUt        |||      D ]E  }
t        j                  |	f| ||f|dt        |
||||      dt        |j                  ||      d G t        d|	| ||g|      S )Nr#   )rf   rg   rF   r   )prefix_argsepilogue_fnbaddbmm)r   r
   aten_baddbmmrS   r   r   rT   rU   r   r   dtyper   )inpr\   r]   rf   rg   r$   r   r   r`   ra   rb   s              r   tuned_baddbmmro      s    '.tT3v'N$Aq!VT4
 !" 
		Ct,fE		MN 
 6" Aq) 	F,, $- VQ1f5	
 *6<<E	 %Y#tT9JFSSr   )'loggingtorch r   r   rK   select_algorithmr   r   r   utilsr	   r   r
   r   r   virtualizedr   mmr   	mm_commonr   r   r   r   	getLogger__name__rY   opsatenr   rT   r   rR   rk   rl   register_loweringre   ro   r:   r   r   <module>r}      s        
   " F F g!yy~~G 		AEN eii7!%--1BC TXX$( <K <KB -.Ad Tr   