o iiQ@s&ddlmZddlZddlZddlmZddlmZm Z ddl m Z m Z m Z mZmZmZddZddZejejd d d d d dddejd d d d d dddejd dd d d dddejdd d d d dddejd d d d d dddejd dd d d dddejdd d d d dddejd d d d d dddejdd d d d dddejd d d d d dddejd d d d d dddejd dd d d dddejdd d d d dddejd d d d d dddejd ddd d dddejdd dd d dddejd d dd d dddejdd dd d dddgegdee dddedddiejdejdejdejdejd ejdejd!ejd"ejd#ejd$ejd%ejfd&d'Z  ( )d9d*ejd+ejd,eejd-ed.ed/ejf d0d1Zejejd d d d d dddejd d d d d dddejd dd d d dddejdd d d d dddejd d d d d dddejd dd d d dddejdd d d d dddejd d d d d dddejdd d d d dddejd d d d d dddejd d d d d dddejd dd d d dddejdd d d d dddejd d d d d dddejd ddd d dddejdd dd d dddejd d dd d dddejdd dd d dddgegdee dddedd2diejdejdejdejdejd ejdejd%ejfd3d4Z ( d:d5ejd+ejd-ed6eejd/ejf d7d8ZdS);)OptionalN)early_config_pruneestimate_matmul_time)gelu gelu_approxgelu_approx_grad gelu_grad squared_relusquared_relu_gradcs fddS)Ncs |S)N)zero_)nargsnameR/home/ubuntu/vllm_env/lib/python3.10/site-packages/flash_attn/ops/triton/linear.pys zinit_to_zero..rr rr r init_to_zeros rc Cs`g}dD])}dD]$}dD]}dD]}|dkrdnd}|tj|||dd ||d qq qq|S) N)) )r@)rrrrrBLOCK_MBLOCK_NBLOCK_KSPLIT_K num_stages num_warps)appendtritonConfig)configsr$block_mblock_kblock_nr%rrrget_configs_io_bounds*r-rrrrrrr#rrrr) CACHE_KEY_M CACHE_KEY_N CACHE_KEY_K )r perf_modeltop_k)r)keyprune_configs_byEVEN_KcC|d|d|ddkSNKr!r"rrargsrrrrrrGROUP_Mr r!r" A_ROWMAJOR B_COLMAJORBIASSAVE_ACT_INPUT ACTIVATIONc.Cstjdd}||d|}||d|}||}||}t||||} |||| }!||| }"|!|td|}#|"|td|}$tt|#|||}%tt|$|||}&td|}'|r~||%dddf| |'dddf}n||%dddf| |'dddf| }|r||'dddf|&dddf|}n||'dddf||&dddf|}tj||ftjd}(t|d| D]S})|rt |}*t |}+n tj ||'dddf|)kdd}*tj ||'dddf|)kdd}+|(t |*|+7}(|r||7}n||| 7}|r||7}q|||7}q|rCtj ||$|$|kdd tj}|(|dddf7}(|r`||%dddf| |&dddf},t |,|(|dkrjt |(}(n|d krtt|(}(n |d kr}t|(}(|!|td|}#|"|td|}$||#dddf| |$dddf}|#|kdddf|$|kdddf@}-t ||(dS) aW Kernel for computing Out = activation(A x W + C) - Input has shape (M, K) - Weight has shape (K, N) - Bias has shape (N,) - Output has shape (M, N) - ActInputs (optional) has shape (M, N) 'ActInputs' optionally saves the A x W + C intermediate for backward computations This kernel will consolidate over K raxisrNdtypemaskotherrrr )tl program_idminarangemax_contiguous multiple_ofzerosfloat32rangeloaddottostorerrr ).C ACT_INPUTABbiasMNr:r/r0r1 stride_cm stride_am stride_ak stride_bn stride_bkrr>r r!r"r7r?r@rArBrCpidgrid_mgrid_nwidthgroup_id group_sizepid_mpid_nrmrnramrbnrkacckab act_in_ptrsrJrrr kernel_fwd5sd ~  *,*,        (      ((rwidFxweightr] activationsave_act_inputreturnc s`|dvsJ|jdd|jd}}|}|||}|ddkr/|ddkr/|}|ddkrA|ddkrA|}|durI|nd}|j|jks]Jd|jd|j|durs|j|jkssJd|jd|j|jd|jdksJd |jd |j|dus|jd|jdksJd |j\} |j\} tjf|j|jd } |rt | nd} fd d} t | | | |||dur|n|| dd| df | d|d|d|d|d|du|||ddk|ddkdd |s| jg|| jdRS| jg|| jdR| jg|| jdRfS)a Compute e = activation(x @ weight.T + bias). This wrapper kicks the `kernel_fwd` Triton kernel :param x: input tensor :param weight: weight matrix :param bias: an optional bias tensor :param activation: Activation name. Needs to be a Triton kernel. :param act_input: an optional tensor to save the activation inputs (for backward) :return: result tensor rxrrr Nrrz/Input and weight must have the same dtype, got  and z-Input and bias must have the same dtype, got Incompatible dimensions:  - z2Incompatible dimensions in between weight and biasdevicerGc"t|dt|dfSNrr r'cdivMETAr^r_rrr9"z#triton_linear_act..rr.) r`rarbrdrcrArBrCr?r@r>) shapenumelreshapestride contiguousrGtorchemptyr empty_likerw) ryrzr]r{r| batch_shapen batch_dim x_reshapedr:output act_inputgridrrrtriton_linear_actsr       rcCr8r9rr;rrrrr=c*Cstjdd}||d|}||d|}||}||}t||||}||||}|||}||td|}||td|}tt||||} tt||||}!td|}"|| dddf| |"dddf| }||"dddf| |!dddf|}tj||ftjd}#t|d| D]C}$|rt |}%t |}&n tj ||"dddf|$kdd}%tj ||"dddf|$kdd}&|#t |%|&7}#||| 7}||| 7}q|dkr || dddf| |!dddf}'t |' |#j }(|d kr|#t |(9}#n|d kr"|#t|(9}#n |d kr-|#t|(9}#||td|}||td|}||dddf| |dddf}||kdddf||kdddf@})tj||#|)d dS) a= Kernel for computing Out = activation(A x W + C) - Input has shape (M, K) - Weight has shape (K, N) - Output has shape (M, N) - ActInputs (optional) has shape (M, N) 'ActInputs' optionally saves the A x W + C intermediate for backward computations This kernel will consolidate over K rrDrNrFrHrIrxrrr )rJ)rLrMrNrOrPrQrRrSrTrUrVrWrGrrr rX)*rYrZr[r\r^r_r:r/r0r1r`rarbrdrcrr>r r!r"r7rCrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrrJrrr kernel_bwd^sN x  ,,      (    ((r grad_outputrc s|dvsJ|jdd|jd}}|}|||}|ddkr/|ddkr/|}|ddkrA|ddkrA|}|j|jksSJd|jd|j|jd|jdksiJd|jd |j|d krx|dusxJd ||j\}|j\}tjf|j|jd } fd d} t | | ||||dd|d| d|d|d|d|d|dd| jg|| jdRS)ap Compute e = activation(grad_output @ weight + bias). This wrapper kicks the `kernel_fwd` Triton kernel :param grad_output: input tensor :param weight: weight matrix :param activation: Activation name. Needs to be a Triton kernel. :param act_input: an optional tensor to save the activation inputs (for backward) :return: result tensor r~Nrrrz5grad_output and weight must have the same dtype, got rrrrxz%act_input is required for activation rcrrrrrrrr;rz"triton_dgrad_act..rr.)r`rarbrdrcrCr>) rrrrrrGrrrr) rrzr{rrrrgrad_output_reshapedr: grad_inputrrrrtriton_dgrad_actsP    r)NrxF)rxN)typingrrr'triton.languagelanguagerLtriton.ops.matmul_perf_modelrr#flash_attn.ops.triton.k_activationsrrrrr r rr-autotuner( heuristicsjit constexprrwTensorstrboolrrrrrrrs       ?@H !" \   ?@Hh