
    qi-                        d dl Z d dlmZ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 ddl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"m#Z# ddl$m%Z%m&Z&m'Z'm(Z( erddlm)Z) ddlm*Z*  e jV                  e,      Z-ejD                  j\                  Z.ed        Z/ ede/dd      Z0 eejb                  de.jb                  jd                        Z3 eejb                  dde.jb                  jh                        Z5 eejl                  de.jl                  jd                        Z7 ejp                  e.jb                        dddd       Z9 ejp                  e.jl                        ddddd       Z:y)     N)TYPE_CHECKINGUnion)counters)CKGemmTemplate   )configirlowering)MMKernelInputs)	loweringsmake_pointwisemake_reductiontransform_args)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)opsV   )_is_static_problemis_batch_stride_largest_or_zeromm_argsuse_native_matmul)ChoiceCaller)KernelTemplatec                :     |||d          |||d         z  | dfS )NBLOCK_MBLOCK_Nr    )bmnmetacdivs        `/home/ubuntu/crypto_trading_bot/.venv/lib/python3.12/site-packages/torch/_inductor/kernel/bmm.pybmm_gridr-   +   s*    DO$tAtI'??AFF    bmmaV	  
{{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).to(INDEX_DTYPE)
    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).to(INDEX_DTYPE)  # 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).to(INDEX_DTYPE)  # 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", val_shape=("BLOCK_M", "BLOCK_N"))}}
T)namegridsource"cache_codegen_enabled_for_templatezat::bmm_out)op_overloadzat::_bmm_out_dtype_cuda	bmm_dtype)r0   r4   zat::baddbmm_out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         |      rt        t        j                      d       t        t        j                     |d      }t         |gi d
dd      \  }}	t        j                  j                   rO j"                  t$        j&                  t$        j(                  fv r# fd}
|D cg c]  } t+        |
      |       }} t+        t,        j.                        | } t1        d      |d      }|S t3         |||      \  }}}} }d}t5         |g|      } j                         d	   }t6        d   d| d| d| d| xx   dz  cc<   t8        j;                  d|||| j=                         |j=                         |       t>        }i }|r. jA                         jB                  dk(  sJ d       tD        }d|i}g }g }i }tG               r |jI                  |       |||jJ                  <   tM        |d      r*|| j=                         k(  r|jI                  tN               |jQ                  t        jR                  jU                  ||||             tW        |      \  }}tY         ||      }|rB|r@t[        ||||      r2t]        |      r'ddl/m0} |jc                  |||je                                tg        | |      r'ddl4m5} |jm                  |||je                                to        ||||      r%tq        jr                  |||je                                tu        |||je                         |      S c c}w )z`
    Lowering for autotuning aten.bmm with different backends (Aten, Triton, CUTLASS, etc.)
    c              3   V   K   | ]!  }|j                         j                  d k(   # yw)cpuN)
get_devicetype).0xs     r,   	<genexpr>ztuned_bmm.<locals>.<genexpr>   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_r7   s      r,   is_valid_to_require_contiguousz1tuned_bmm.<locals>.is_valid_to_require_contiguous   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   r&   )sizesstridess     r,    is_preferred_layout_as_bmm_inputz3tuned_bmm.<locals>.is_preferred_layout_as_bmm_input   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)rH   meta_trM   rN   rO   s       r,   may_require_contiguousz)tuned_bmm.<locals>.may_require_contiguous   sT    KK&++-Ekk%(//1G3E7COO66q9Hr.   r   TNF)argskwargs	broadcasttype_promotion_kindconvert_input_to_boolc                 H    t        j                  | j                  d      S )NF)use_compute_types)r   to_dtypedtype)r>   mat1s    r,   	_to_dtypeztuned_bmm.<locals>._to_dtype   s    ||AtzzUKKr.   dot)r7   	out_dtyper/   )rd   aten_mm_infoz	aten.bmm_rI   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 CUDArd   check_max_autotune)kwarg_overrides)CUTLASS3xGemmTemplate)CppBmmTemplate);allget_sizeL	unsqueezesum_mulr   graphcurrent_noderX   r    r   atenr   inductor_configtritoncodegen_upcast_to_fp32r`   torchfloat16bfloat16r   r   rc   r   r   r   r   loginfo	get_dtypeaten_bmmr;   r<   aten_bmm_dtyper   appenduidr   bmm_templateextendchoicesget_template_configsr   r   r   r   codegen.cuda.gemm_templaterj   add_cutlass_gemm_choicesnodesr   codegen.cpp_bmm_templaterk   add_choicesr   r   add_ck_gemm_choicesr   )ra   mat2rd   r7   rJ   rW   	meta_mat1	meta_mat2rX   rY   rb   r>   mul_pointwisedot_reductionr(   r)   kr0   kernel_inputs
batch_sizeaten_handleraten_extra_kwargsr   templates_to_useri   rI   
is_nonzerobatch_stride_largest_or_zerorj   rk   rO   s   `                             @r,   	tuned_bmmr      s   
 
>$
>>==?1"dmmoa&8A&=;;tR(D;;tQ'D66!%%d+!44	9	U	 *$/,,11!4I)$	:D)$/,,11!4I)$	:Dt$(r2(q1% $"'
f !!88TZZMMNNL
 >

L ;??Q-N9-a0?D?/sww/6-u-mQ? #*d6Y#Aq!VT4 D #D$<9EM #J^yAaS!AaSABaGBHHd				 (0L %%/W1WW/%()4"$G IKO-,=(()6e<Y$..*:: 	- NN			&&+	 	' 	
 'v.MAz#B4v#V $ Aq1%F66V]002	
 FD$/=""!	
 FAq!,**7FM<O<O<QR$T7M4G4G4I6RRq @s   ?Q)alphabetar7   c                   t        ||      r|dk(  rd}nt        t        j                     ||       }|dk(  rd}n8t        t        j                     |t        t        j                     ||            }t        t        j
                     ||      S t        ||| |      \  }}	}
}}}} t        | ||gt        ||            }|j                         d   }t        d   d| d| d|	 d|
 xx   dz  cc<   t        j                  d	|||	|
|j                         |j                         | j                         |	       d
}g }g }t               r|j                  t                t#        |d      r|j                  t$               |j'                  t(        j*                  j-                  |||             t/        |||j1                         |      S )z_
    Lowering for autotuning aten.mm with different backends (Aten, Triton, CUTLASS, etc.)
    r   r6   )r   r   )scalarsre   zaten.baddbmm_rI   r   zkTuned aten.baddbmm: batch_size=%s, m=%s, n=%s, k=%s, mat1_dtype=%s, mat2_dtype=%s, inp=%s, output_layout=%sbaddbmmFrg   )r    r   rt   rq   r/   addr   r   dictrm   r   r{   r|   r}   r   r   aten_baddbmmr   r   r   r   r   r   r   r   )inpra   r   r   r   r7   arg1arg2r(   r)   r   r   r   r0   r   r   s                   r,   tuned_baddbmmr     s   
 t$19DTXX&tS1DA:DTXX&ui.A$.MND"4.. (/tT3v'N$Aq!VT4 #	dD4e$#?M
 #J^}ZL!AaS!EF!KFHHu			
 D"$G IK-6e<- NN			&&}6FM %T7M4G4G4I6RRr.   )N);loggingtypingr   r   rx   torch._dynamo.utilsr   7torch._inductor.codegen.rocm.ck_universal_gemm_templater    r   ru   r	   r
   rn   r   r   r   r   r   r   select_algorithmr   r   r   r   utilsr   r   r   r   r   r   virtualizedr   r   	mm_commonr   r   r   r    r!   r"   	getLogger__name__r{   rt   r-   r   r/   outr~   	dtype_outr   r   r   register_loweringr   r   r&   r.   r,   <module>r      sj    '  ( R ; ; * P P   !  !1g!yy~~ G G 		CH (,OHT eiiDHHLLQ#	II	""	 "	MM$$,,2B2B
 TXXTSD TS TSn T\\",-Ad 8S #8Sr.   