o i @sJddlZddlmZddlmZejdejdejdejddfddZdS) N)scale_and_clamp BLOCK_SIZE HAS_X_SCALE HAS_O_SCALEreturnc Cstjddtj} tjdd} ||| } ||| } | |td|}||k}tj| ||dtj}tj| |||dtj}|rUt|}||9}||9}t|||}| rmt|}t|||j j }tj | |||ddS)aSigmoid Linear Unit and Multiplication Kernel Args: o_ptr: Pointer to the 2D output tensor. o_stride: Output tensor stride. o_scale_ptr: The optional, known scale of the output activations. x_ptr: Pointer to the 2D input tensor. x_stride: Input tensor stride. x_scale_ptr: The optional, known scale of the input tensor. d: The number of elements along the second dimension. BLOCK_SIZE: Tunable block size to process in each kernel. Operating on a 2D grid, computes the following: ``` out[i, j] = sigmoid(x[i, j]) * x[i, j] * x[i, j + d] ``` If scales are provided, the input and output tensors are scaled. r)axis)maskN) tl program_idtoint64arangeloadfloat32sigmoidrdtype element_tystore)o_ptro_stride o_scale_ptrx_ptrx_stride x_scale_ptrdrrrij o_row_ptr x_row_ptroffsetsr abx_scaleresulto_scaler&Z/home/ubuntu/vllm_env/lib/python3.10/site-packages/flashinfer/triton/kernels/activation.pysilu_and_mul_kernels""     r() tritontriton.languagelanguager flashinfer.triton.kernels.quantrjit constexprr(r&r&r&r's