diff --git a/benchmark/kernels/bench_silu_and_mul.py b/benchmark/kernels/bench_silu_and_mul.py
new file mode 100644
index 000000000000..776c04102f91
--- /dev/null
+++ b/benchmark/kernels/bench_silu_and_mul.py
@@ -0,0 +1,97 @@
+from itertools import product
+
+import torch
+from flag_gems import silu_and_mul as flag_gems_silu_and_mul
+from flashinfer.activation import silu_and_mul as flashinfer_silu_and_mul
+from torch.utils.benchmark import Timer
+from vllm import _custom_ops as ops
+
+
+def forward_vllm(x: torch.Tensor) -> torch.Tensor:
+    d = x.shape[-1] // 2
+    output_shape = x.shape[:-1] + (d,)
+    out = torch.empty(output_shape, dtype=torch.float16, device=x.device)
+    ops.silu_and_mul(out, x)
+    return out
+
+
+def forward_flashinfer(x: torch.Tensor) -> torch.Tensor:
+    d = x.shape[-1] // 2
+    out = torch.empty((*x.shape[:-1], d), dtype=torch.float16, device=x.device)
+    flashinfer_silu_and_mul(out, x)
+    return out
+
+
+def forward_flag_gems(x: torch.Tensor) -> torch.Tensor:
+    d = x.shape[-1] // 2
+    return flag_gems_silu_and_mul(x[..., :d], x[..., d:])
+
+
+def test_consistency():
+    x = torch.randn(2, 4, 2 * d, dtype=torch.float16, device=device)
+    out_vllm = forward_vllm(x)
+    out_flashinfer = forward_flashinfer(x)
+    out_flag_gems = forward_flag_gems(x)
+    assert torch.allclose(out_vllm, out_flashinfer, atol=1e-3, rtol=1e-3)
+    assert torch.allclose(out_vllm, out_flag_gems, atol=1e-3, rtol=1e-3)
+    assert torch.allclose(out_flashinfer, out_flag_gems, atol=1e-3, rtol=1e-3)
+    print("Consistency test passed!")
+
+
+device = torch.device("cuda")
+d = 4096
+
+test_consistency()
+
+results = []
+sizes = [2, 8, 32, 128, 512]
+
+for batch_size, seq_length in product(sizes, sizes):
+    label = "SiLU and Mul"
+    sub_label = f"[{batch_size}, {seq_length}]"
+
+    input_tensor = torch.randn(
+        batch_size, seq_length, 2 * d, dtype=torch.float16, device=device
+    )
+
+    min_run_time = max(0.1, min(1, batch_size * seq_length / 1e6))
+
+    for num_threads in [1, 4, 16, 32]:
+        results.append(
+            Timer(
+                stmt="forward_vllm(input_tensor)",
+                setup="from __main__ import forward_vllm",
+                globals={"input_tensor": input_tensor},
+                num_threads=num_threads,
+                label=label,
+                sub_label=sub_label,
+                description="vLLM",
+            ).blocked_autorange(min_run_time=min_run_time)
+        )
+
+        results.append(
+            Timer(
+                stmt="forward_flashinfer(input_tensor)",
+                setup="from __main__ import forward_flashinfer",
+                globals={"input_tensor": input_tensor},
+                num_threads=num_threads,
+                label=label,
+                sub_label=sub_label,
+                description="FlashInfer",
+            ).blocked_autorange(min_run_time=min_run_time)
+        )
+
+        results.append(
+            Timer(
+                stmt="forward_flag_gems(input_tensor)",
+                setup="from __main__ import forward_flag_gems",
+                globals={"input_tensor": input_tensor},
+                num_threads=num_threads,
+                label=label,
+                sub_label=sub_label,
+                description="Flag_gems",
+            ).blocked_autorange(min_run_time=min_run_time)
+        )
+
+compare = torch.utils.benchmark.Compare(results)
+compare.print()
diff --git a/python/sglang/bench_latency.py b/python/sglang/bench_latency.py
index ffd6b24f09e7..50ec8a67f287 100644
--- a/python/sglang/bench_latency.py
+++ b/python/sglang/bench_latency.py
@@ -238,7 +238,7 @@ def correctness_test(
 
     # Decode
     output_ids = [input_ids[i] + [next_token_ids[i]] for i in range(len(input_ids))]
-    for _ in range(bench_args.output_len):
+    for _ in range(bench_args.output_len[0]):
         next_token_ids, _ = decode(next_token_ids, batch, model_runner)
         for i in range(len(reqs)):
             output_ids[i].append(next_token_ids[i])
diff --git a/python/sglang/srt/kernels/silu_and_mul.py b/python/sglang/srt/kernels/silu_and_mul.py
new file mode 100644
index 000000000000..020829afd0f0
--- /dev/null
+++ b/python/sglang/srt/kernels/silu_and_mul.py
@@ -0,0 +1,38 @@
+"""
+Copyright 2023-2024 SGLang Team
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+"""
+
+import torch
+import triton
+import triton.language as tl
+
+from sglang.srt.kernels.utils.pointwise_dynamic import pointwise_dynamic
+
+
+@pointwise_dynamic(promotion_methods=[(0, 1, "DEFAULT")])
+@triton.jit
+def silu_and_mul_kernel(x, y):
+    x_fp32 = x.to(tl.float32)
+    x_silu = tl.fdiv(x_fp32, (1.0 + tl.exp(-x_fp32)))
+    return x_silu * y
+
+
+class SiluAndMul(torch.autograd.Function):
+    @staticmethod
+    def forward(ctx, A, B):
+        return silu_and_mul_kernel(A, B)
+
+
+def silu_and_mul(A, B):
+    return SiluAndMul.apply(A, B)
diff --git a/python/sglang/srt/kernels/utils/pointwise_dynamic.py b/python/sglang/srt/kernels/utils/pointwise_dynamic.py
new file mode 100644
index 000000000000..ac5113d4605a
--- /dev/null
+++ b/python/sglang/srt/kernels/utils/pointwise_dynamic.py
@@ -0,0 +1,869 @@
+"""
+Copyright 2023-2024 SGLang Team
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+"""
+
+# Adapted from https://github.com/FlagOpen/FlagGems/blob/4f6ea4cc2fc004d806d48a7d8f21dcb6a79a6dc5/src/flag_gems/utils/pointwise_dynamic.py
+
+import importlib
+import os
+import threading
+from typing import Any, Callable, List, Mapping, Optional, Tuple
+
+import torch
+import torch._prims_common as utils
+import triton
+from triton import language as tl
+from triton.runtime.jit import JITFunction
+
+from sglang.srt.kernels.utils.triton_utils import (
+    IndentedBuffer,
+    NameSpace,
+    broadcast_shapes,
+    cache_dir,
+)
+
+
+# ------------------ Operation Description ---------------------------
+def _type_name(type) -> str:
+    "Render typename as string, work for both (bool, int, float, str) and torch.dtype object"
+    if type in (bool, int, float, str):
+        return type.__name__
+    if isinstance(type, torch.dtype):
+        return str(type)
+    return str(type)
+
+
+def _check_typed_list(container, type):
+    for item in container:
+        assert isinstance(item, type)
+
+
+def _check_sized_list(container, size):
+    assert len(container) == size
+
+
+class OPDesc:
+    _num_inputs: int
+    _is_tensor: List[bool]
+    _dtypes: List[Optional[type]]
+
+    _num_input_tensors: int
+    _num_non_tensor_inputs: int
+
+    _num_outputs: int
+    _promotion_methods: List[Tuple[int, ...]]
+
+    def __init__(
+        self,
+        *,
+        num_inputs: Optional[int] = None,
+        is_tensor: Optional[List[bool]] = None,
+        dtypes: Optional[List[Optional[type]]] = None,
+        num_outputs: Optional[int] = None,
+        promotion_methods: Optional[List[Tuple[int, ...]]] = None,
+    ):
+        if is_tensor is not None:
+            _check_typed_list(is_tensor, bool)
+        if dtypes is not None:
+            _check_typed_list(dtypes, (type, type(None)))
+        if promotion_methods is None:
+            raise ValueError(
+                "No type promotion method provided! You must provide type promotion method for each output!"
+            )
+        else:
+            self._promotion_methods = promotion_methods
+
+        if num_inputs is not None:
+            self._num_inputs = num_inputs
+            if is_tensor is not None:
+                _check_sized_list(is_tensor, num_inputs)
+                self._is_tensor = is_tensor
+            else:
+                self._is_tensor = [True] * num_inputs
+
+            if dtypes is not None:
+                _check_sized_list(dtypes, num_inputs)
+                self._dtypes = dtypes
+            else:
+                self._dtypes = [None] * num_inputs
+        elif is_tensor is not None:
+            self._num_inputs = len(is_tensor)
+            self._is_tensor = is_tensor
+            if dtypes is not None:
+                _check_sized_list(dtypes, self._num_inputs)
+                self._dtypes = dtypes
+            else:
+                self._dtypes = [None] * self._num_inputs
+        elif dtypes is not None:
+            self._num_inputs = len(dtypes)
+            self._dtypes = dtypes
+            if is_tensor is not None:
+                _check_sized_list(is_tensor, self._num_inputs)
+                self._is_tensor = is_tensor
+            else:
+                self._is_tensor = [item is None for item in dtypes]
+        else:
+            raise ValueError(
+                "Cannot make OPDesc when none of (num_inputs, is_tensor, dtypes) is specified."
+            )
+
+        if num_outputs is not None:
+            self._num_outputs = num_outputs
+            _check_sized_list(promotion_methods, num_outputs)
+        else:
+            self._num_outputs = len(promotion_methods)
+
+        assert self._num_inputs >= 1
+        assert self._num_outputs >= 1
+
+        self._num_input_tensors = sum(self._is_tensor)
+        self._num_non_tensor_inputs = self._num_inputs - self._num_input_tensors
+
+    def num_inputs(self):
+        # num of arguments, outputs not included
+        return self._num_inputs
+
+    def num_outputs(self):
+        return self._num_outputs
+
+    def is_tensor(self, arg_id: int) -> bool:
+        return self._is_tensor[arg_id]
+
+    def input_type(self, arg_id) -> Optional[type]:
+        return self._dtypes[arg_id]
+
+    def num_input_tensors(self) -> int:
+        return self._num_input_tensors
+
+    def num_output_tensors(self) -> int:
+        return self._num_outputs
+
+    def num_non_tensor_args(self) -> int:
+        return self._num_non_tensor_inputs
+
+    def type_promotion_methods(self) -> List[Tuple[int, ...]]:
+        return self._promotion_methods
+
+    def _match_enum_by_string(
+        self, input_str: str
+    ) -> utils.ELEMENTWISE_TYPE_PROMOTION_KIND:
+        for kind in utils.ELEMENTWISE_TYPE_PROMOTION_KIND:
+            if input_str.lower() == kind.name.lower():
+                return kind
+        raise ValueError(f"No matching enum member found for input: {input_str}")
+
+    def ith_type_promotion_args(self, i) -> List[int]:
+        return self._promotion_methods[i][:-1]
+
+    def ith_type_promotion_kind(self, i) -> utils.ELEMENTWISE_TYPE_PROMOTION_KIND:
+        return self._match_enum_by_string(self._promotion_methods[i][-1])
+
+    def signature(self, outputs_in_arg: bool = False):
+        input_types = []
+        for is_tensor, dtype in zip(self._is_tensor, self._dtypes):
+            if is_tensor:
+                input_types.append("Tensor")
+            else:
+                if dtype is None:
+                    input_types.append("scalar")
+                else:
+                    input_types.append(_type_name(dtype))
+
+        output_types = []
+        for _ in range(self.num_outputs()):
+            output_types.append("Tensor")
+        if outputs_in_arg:
+            input_types.extend(output_types)
+        sig = f'Pointwise: ({", ".join(input_types)}) -> ({", ".join(output_types)})'
+        return sig
+
+    def __str__(self) -> str:
+        return self.signature(outputs_in_arg=False)
+
+
+# --------------------------- pointwise wrapper genration -----------------------------------
+def parameter_for_wrapper(op_desc: OPDesc, include_outputs: bool = False) -> str:
+    """Generate parameter declaration with type annotation for wrapper function.
+    Example: in0: torch.Tensor, val0: float, out0: torch.Tensor
+    """
+    parameters: List[str] = []
+
+    input_tensor_index = 0
+    non_tensor_index = 0
+    for i in range(op_desc.num_inputs()):
+        if op_desc._is_tensor[i]:
+            parameters.append(f"in{input_tensor_index}: torch.Tensor")
+            input_tensor_index += 1
+        else:
+            if op_desc.input_type(i) is not None:
+                parameters.append(
+                    f"val{non_tensor_index}: {_type_name(op_desc.input_type(i))}"
+                )
+            else:
+                parameters.append(f"val{non_tensor_index}")
+            non_tensor_index += 1
+
+    if include_outputs:
+        output_tensor_index = 0
+        for i in range(op_desc.num_outputs()):
+            parameters.append(f"out{output_tensor_index}: torch.Tensor")
+            output_tensor_index += 1
+
+    parameters.append("**kwargs")
+
+    return ", ".join(parameters)
+
+
+def ith_parameter_for_type_promotion(op_desc: OPDesc, ith: int) -> str:
+    """Generate parameter reference for i-th type promotion rule
+    Example: in0, val0, out0
+    """
+    parameters: List[str] = []
+
+    input_tensor_index = 0
+    non_tensor_index = 0
+    for i in range(op_desc.num_inputs()):
+        if i not in op_desc.ith_type_promotion_args(ith):
+            if op_desc._is_tensor[i]:
+                input_tensor_index += 1
+            else:
+                non_tensor_index += 1
+            continue
+        if op_desc._is_tensor[i]:
+            parameters.append(f"in{input_tensor_index}")
+            input_tensor_index += 1
+        else:
+            parameters.append(f"val{non_tensor_index}")
+            non_tensor_index += 1
+
+    return ", ".join(parameters)
+
+
+def parameter_ref_for_wrapper(
+    op_desc: OPDesc,
+    include_outputs: bool = False,
+    include_offset: bool = False,
+    include_kwargs: bool = False,
+) -> str:
+    """Generate parameter reference for wrapper function.
+    Example: in0, val0, out0, out0_offset
+    """
+    parameters: List[str] = []
+
+    input_tensor_index = 0
+    non_tensor_index = 0
+    for i in range(op_desc.num_inputs()):
+        if op_desc._is_tensor[i]:
+            parameters.append(f"in{input_tensor_index}")
+            input_tensor_index += 1
+        else:
+            parameters.append(f"val{non_tensor_index}")
+            non_tensor_index += 1
+
+    if include_outputs:
+        output_tensor_index = 0
+        for i in range(op_desc.num_outputs()):
+            parameters.append(f"out{output_tensor_index}")
+            if include_offset:
+                parameters.append(f"out{output_tensor_index}_offset")
+            output_tensor_index += 1
+
+    if include_kwargs:
+        parameters.append("**kwargs")
+
+    return ", ".join(parameters)
+
+
+def output_ref_for_wrapper(op_desc: OPDesc) -> str:
+    """Generate output variable refernece for wrapper function.
+    Example: out0, out1
+    """
+    parameters: List[str] = [f"out{i}" for i in range(op_desc.num_outputs())]
+    return ", ".join(parameters)
+
+
+def docstring_for_functional_wrapper(op_desc: OPDesc):
+    doc = f'"""Generated wrapper function with {str(op_desc)}"""'
+    return doc
+
+
+def docstring_for_destination_passing_wrapper(op_desc: OPDesc):
+    doc = f'"""Generated wrapper function with {op_desc.signature(outputs_in_arg=True)}"""'
+    return doc
+
+
+def generate_imports(code: IndentedBuffer) -> IndentedBuffer:
+    code.writeline("import math")
+    code.writeline("import torch")
+    code.writeline("import triton")
+    code.writeline("from triton import language as tl")
+    code.newline()
+    code.writeline("from sglang.srt.kernels.utils.triton_utils import (")
+    code.writeline("    broadcast_shapes,")
+    code.writeline("    broadcasted_stride,")
+    code.writeline("    c_contiguous_stride,")
+    code.writeline("    volume,")
+    code.writeline("    libentry,")
+    code.writeline("    type_promotion,")
+    code.writeline(")")
+    code.writeline("import torch._prims_common as utils")
+    code.newline()
+    code.newline()
+    return code
+
+
+def generate_functional_pointwise_wrapper(
+    op_desc: OPDesc,
+    wrapper_name: str,
+    destination_passing_func_name: str,
+    code: IndentedBuffer,
+) -> IndentedBuffer:
+    # wrapper signature
+    parameters: str = parameter_for_wrapper(op_desc, include_outputs=False)
+    wrapper_signature: str = f"def {wrapper_name}({parameters}):"
+    code.writeline(wrapper_signature)
+
+    with code.indent():
+        # docstring
+        wrapper_docstring = docstring_for_functional_wrapper(op_desc)
+        code.writeline(wrapper_docstring)
+
+        shapes_str = ", ".join(
+            f"in{i}.shape" for i in range(op_desc.num_input_tensors())
+        )
+        code.writeline(f"shape = broadcast_shapes([{shapes_str}])")
+
+        # output allocation
+        num_output_tensor_index = 0
+        for i in range(op_desc.num_outputs()):
+            type_promotion_args = ith_parameter_for_type_promotion(op_desc, i)
+            k_type_promotion = op_desc.ith_type_promotion_kind(i)
+            code.writeline(
+                (
+                    f"out{num_output_tensor_index} = "
+                    f"torch.empty(shape, dtype=type_promotion"
+                    f"({type_promotion_args}, type_promotion=utils.{k_type_promotion})[1], "
+                    f"device=in0.device)"
+                )
+            )
+            num_output_tensor_index += 1
+
+        # call destination_passing_func
+        output_names: str = output_ref_for_wrapper(op_desc)
+        call_str = (
+            f"{output_names} = {destination_passing_func_name}"
+            f"({parameter_ref_for_wrapper(op_desc, include_outputs=True, include_offset=False, include_kwargs=True)})"
+        )
+        code.writeline(call_str)
+
+        return_str = f"return {output_names}"
+        code.writeline(return_str)
+        code.newline()
+        code.newline()
+    return code
+
+
+def generate_destination_passing_pointwise_wrapper(
+    op_desc: OPDesc,
+    rank: int,
+    wrapper_name: str,
+    kernel_name: str,
+    code: IndentedBuffer,
+) -> IndentedBuffer:
+    # wrapper signature
+    parameters: str = parameter_for_wrapper(op_desc, include_outputs=True)
+    wrapper_signature: str = f"def {wrapper_name}({parameters}):"
+    code.writeline(wrapper_signature)
+
+    with code.indent():
+        # docstring
+        wrapper_docstring = docstring_for_destination_passing_wrapper(op_desc)
+        code.writeline(wrapper_docstring)
+
+        if rank > 0:
+            code.writeline("shape = out0.shape")
+            code.writeline("num_tasks = volume(shape)")
+
+        if rank > 0:
+            code.writeline("tile_size = min(512, triton.next_power_of_2(num_tasks))")
+            code.writeline("num_warps = 4")
+            code.writeline("num_ctas = min(65535, triton.cdiv(num_tasks, tile_size))")
+            code.writeline(
+                "tiles_per_cta = triton.cdiv(num_tasks, tile_size * num_ctas)"
+            )
+        else:
+            code.writeline("num_warps = 1")
+            code.writeline("num_ctas = 1")
+        code.writeline("grid = (num_ctas, 1, 1)")
+        code.newline()
+
+        # input strides for each input tensor w.r.t. the task index space
+        if rank > 0:
+            code.writeline("# strides of each tensor argument w.r.t the task space")
+            for i in range(op_desc.num_input_tensors()):
+                code.writeline(
+                    f"in{i}_strides = broadcasted_stride(in{i}.shape, in{i}.stride(), shape)"
+                )
+            for i in range(op_desc.num_output_tensors()):
+                code.writeline(f"if 'out{i}_offset' in kwargs:")
+                with code.indent():
+                    code.writeline(f"out{i}_offset = kwargs['out{i}_offset']")
+                code.writeline("else:")
+                with code.indent():
+                    code.writeline(f"out{i}_offset = 0")
+
+                code.writeline(f"if 'out{i}_strides' in kwargs:")
+                with code.indent():
+                    code.writeline(f"out{i}_strides = kwargs['out{i}_strides']")
+                code.writeline("else:")
+                with code.indent():
+                    code.writeline(f"out{i}_strides = out{i}.stride()")
+        else:
+            for i in range(op_desc.num_output_tensors()):
+                code.writeline(f"out{i}_offset = 0")
+        code.newline()
+
+        # grid
+        code.writeline("# kernel launch")
+
+        # launch kernel
+        code.writeline("with torch.cuda.device(in0.device.index):")
+        with code.indent():
+            kernel_launch: str = f"{kernel_name}[grid]("
+            code.writeline(kernel_launch)
+
+            with code.indent():
+                code.writeline(
+                    "{},".format(
+                        parameter_ref_for_wrapper(
+                            op_desc,
+                            include_outputs=True,
+                            include_offset=True,
+                            include_kwargs=False,
+                        )
+                    )
+                )
+
+                if rank > 0:
+                    for i in range(op_desc.num_input_tensors()):
+                        s = ", ".join(f"in{i}_strides[{j}]" for j in range(rank))
+                        code.writeline(f"{s}, # stride for in{i}")
+
+                    for i in range(op_desc.num_output_tensors()):
+                        s = ", ".join(f"out{i}_strides[{j}]" for j in range(rank))
+                        code.writeline(f"{s}, # stride for out{i}")
+
+                    shape_args: str = ", ".join(f"shape[{i}]" for i in range(rank))
+                    code.writeline(f"{shape_args}, # task indexing space")
+                    code.writeline("num_tasks, # num tasks")
+                    code.writeline("tiles_per_cta=tiles_per_cta, # tiles_per_cta")
+                    code.writeline("tile_size=tile_size,")
+                    code.writeline("one_tile_per_cta=tiles_per_cta==1,")
+                code.writeline("num_warps=num_warps,")
+            code.writeline(")")
+
+        # return
+        code.writeline(f"return {output_ref_for_wrapper(op_desc)}")
+        code.newline()
+        code.newline()
+    return code
+
+
+def generate_pointwise_kernel(
+    op_desc: OPDesc,
+    scalar_fn: JITFunction,
+    rank: int,
+    kernel_name: str,
+    code: IndentedBuffer,
+) -> IndentedBuffer:
+    # make the inlined function visible in the context
+    fn_name = scalar_fn.__name__
+    code.writeline(f"from {scalar_fn.__module__} import {fn_name}")
+    code.writeline(f"inlined_f = {fn_name}._scalar_fn")
+    code.newline()
+
+    # the decorators
+    code.writeline("@libentry()")
+    if op_desc.num_non_tensor_args() > 0:
+        # we do not specialize non tensor args since they are passed into the inlined function
+        # which means that their values may not deserve specialization
+        non_specialize_arg_names = [
+            f"val{i}" for i in range(op_desc.num_non_tensor_args())
+        ]
+        code.writeline(f"@triton.jit(do_not_specialize={non_specialize_arg_names})")
+    else:
+        code.writeline("@triton.jit")
+
+    # signature
+    code.writeline(f"def {kernel_name}(")
+    function_ns = NameSpace()
+    with code.indent():
+        input_tensor_index = 0
+        non_tensor_index = 0
+        output_tensor_index = 0
+        # signature: inputs ptrs & non tensor inputs
+        for i in range(op_desc.num_inputs()):
+            if op_desc.is_tensor(i):
+                code.writeline(
+                    f"in{input_tensor_index}_ptr: tl.tensor, # of tl.pointer_type"
+                )
+                function_ns.create_name(f"in{input_tensor_index}_ptr")
+                input_tensor_index += 1
+            else:
+                if op_desc.input_type(i) is not None:
+                    code.writeline(
+                        f"val{non_tensor_index}: {_type_name(op_desc.input_type(i))},"
+                    )
+                else:
+                    code.writeline(f"val{non_tensor_index},")
+                function_ns.create_name(f"val{non_tensor_index}")
+                non_tensor_index += 1
+
+        # signature: output ptrs
+        for i in range(op_desc.num_outputs()):
+            code.writeline(
+                f"out{output_tensor_index}_ptr: tl.tensor, # of tl.pointer_type"
+            )
+            code.writeline(f"out{output_tensor_index}_offset: int,")
+            function_ns.create_name(f"out{output_tensor_index}_ptr")
+            function_ns.create_name(f"out{output_tensor_index}_offset")
+            output_tensor_index += 1
+
+        # signature: strides, for each tensor arguments
+        # only add this arguments when rank > 0
+        if rank > 0:
+            # strides for inputs
+            for i in range(op_desc.num_input_tensors()):
+                for j in range(rank):
+                    function_ns.create_name(f"in{i}_stride{j}")
+                stride_args = ", ".join(f"in{i}_stride{j}: int" for j in range(rank))
+                code.writeline(f"{stride_args}, # strides for in{i}")
+
+            # strides for outputs
+            for i in range(op_desc.num_output_tensors()):
+                for j in range(rank):
+                    function_ns.create_name(f"out{i}_stride{j}")
+                stride_args = ", ".join(f"out{i}_stride{j}: int" for j in range(rank))
+                code.writeline(f"{stride_args}, # strides for out{i}")
+
+            # task space, used to reconstruct multi index
+            task_space_args = ", ".join(f"s{i}: int" for i in range(rank))
+            for i in range(rank):
+                function_ns.create_name(f"s{i}")
+            code.writeline(f"{task_space_args}, # task_space")
+
+            # number of tasks, used to compute mask
+            code.writeline("num_tasks: int,")
+            function_ns.create_name("num_tasks")
+
+        # tile size & tiles_per_cta, gsl style
+        if rank > 0:
+            code.writeline("tiles_per_cta,")
+            function_ns.create_name("tiles_per_cta")
+
+            code.writeline("tile_size: tl.constexpr,")
+            function_ns.create_name("tile_size")
+
+            code.writeline("one_tile_per_cta: tl.constexpr,")
+            function_ns.create_name("one_tile_per_cta")
+    code.writeline("):")
+
+    # input & output names
+    inputs_to_scalar_fn = []
+    input_tensor_index = 0
+    non_tensor_index = 0
+    for i in range(op_desc.num_inputs()):
+        if op_desc.is_tensor(i):
+            inputs_to_scalar_fn.append(f"in{input_tensor_index}")
+            input_tensor_index += 1
+        else:
+            inputs_to_scalar_fn.append(f"val{non_tensor_index}")
+            non_tensor_index += 1
+    inputs_to_scalar_fn: str = ", ".join(inputs_to_scalar_fn)
+
+    outputs_to_scalar_fn = [f"out{i}" for i in range(op_desc.num_outputs())]
+    outputs_to_scalar_fn: str = ", ".join(outputs_to_scalar_fn)
+
+    # function body for rank-0
+    if rank == 0:
+        with code.indent():
+            code.writeline("# loads")
+            for i in range(op_desc.num_input_tensors()):
+                ptrs_expr: str = f"in{i}_ptr"
+                load_stmt: str = f"in{i} = tl.load({ptrs_expr})"
+                function_ns.create_name(f"in{i}")  # add to the namespace
+                code.writeline(load_stmt)
+            code.newline()
+
+            code.writeline("# compute")
+            code.writeline(f"{outputs_to_scalar_fn} = inlined_f({inputs_to_scalar_fn})")
+            code.newline()
+
+            code.writeline("# stores")
+            for i in range(op_desc.num_output_tensors()):
+                ptrs_expr: str = f"out{i}_ptr + out{i}_offset"
+                store_stmt: str = f"tl.store({ptrs_expr}, out{i})"
+                code.writeline(store_stmt)
+            code.newline()
+            return code
+
+    with code.indent():
+        # get pid
+        code.writeline("# task id & masking")
+        pid_stmt = "pid = tl.program_id(0)"
+        code.writeline(pid_stmt)
+        function_ns.create_name("pid")
+
+        code.writeline("num_ctas = tl.num_programs(0)")
+        function_ns.create_name("num_ctas")
+
+        # get tid (a.k.a task id)
+        tid_stmt = "init_tid = pid * tile_size + tl.arange(0, tile_size)"
+        code.writeline(tid_stmt)
+        function_ns.create_name("init_tid")
+
+        # one-tile-per-cta, monolithic kernel style
+        code.writeline("if one_tile_per_cta: # monolitic kernel style")
+        with code.indent():
+            tid_stmt = "tid = init_tid"
+            code.writeline(tid_stmt)
+            function_ns.create_name("tid")
+
+            # only apply masking when rank > 0
+            # since we only load a value instead of a block of values when the rank is 0
+            mask_stmt: str = "mask = tid < num_tasks"
+            code.writeline(mask_stmt)
+            function_ns.create_name("mask")
+            code.newline()
+
+            # reconstruct multi index
+            code.writeline("# multi index recontruction")
+            for i in reversed(range(rank)):
+                if i > 0:
+                    code.writeline(f"i{i} = tid % s{i}")
+                    code.writeline(f"tid //= s{i}")
+                else:
+                    code.writeline(f"i{i} = tid")
+                function_ns.create_name(f"{i}")
+            code.newline()
+
+            # loads
+            code.writeline("# loads")
+            for i in range(op_desc.num_input_tensors()):
+                ptrs_expr: str = " + ".join(
+                    f"i{j} * in{i}_stride{j}" for j in range(rank)
+                )
+                ptrs_expr: str = f"in{i}_ptr + {ptrs_expr}"
+                load_stmt: str = f"in{i} = tl.load({ptrs_expr}, mask=mask)"
+                function_ns.create_name(f"in{i}")  # add to the namespace
+                code.writeline(load_stmt)
+            code.newline()
+
+            # compute
+            code.writeline("# compute")
+            code.writeline(f"{outputs_to_scalar_fn} = inlined_f({inputs_to_scalar_fn})")
+            code.newline()
+
+            # stores
+            code.writeline("# stores")
+            for i in range(op_desc.num_output_tensors()):
+                ptrs_expr: str = " + ".join(
+                    f"i{j} * out{i}_stride{j}" for j in range(rank)
+                )
+                ptrs_expr: str = f"out{i}_ptr + out{i}_offset + {ptrs_expr}"
+                store_stmt: str = f"tl.store({ptrs_expr}, out{i}, mask=mask)"
+                code.writeline(store_stmt)
+
+        # https://developer.nvidia.com/blog/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/
+        code.writeline("else: # grid-stride-loop style kernel")
+        with code.indent():
+            code.writeline("for j in range(0, tiles_per_cta):")
+            function_ns.create_name("j")
+            with code.indent():
+                tid_stmt = "tid = init_tid + j * tile_size * num_ctas"
+                code.writeline(tid_stmt)
+                function_ns.create_name("tid")
+
+                # only apply masking when rank > 0
+                # since we only load a value instead of a block of values when the rank is 0
+                mask_stmt: str = "mask = tid < num_tasks"
+                code.writeline(mask_stmt)
+                function_ns.create_name("mask")
+                code.newline()
+
+                # reconstruct multi index
+                code.writeline("# multi index recontruction")
+                for i in reversed(range(rank)):
+                    if i > 0:
+                        code.writeline(f"i{i} = tid % s{i}")
+                        code.writeline(f"tid //= s{i}")
+                    else:
+                        code.writeline(f"i{i} = tid")
+                    function_ns.create_name(f"{i}")
+                code.newline()
+
+                # loads
+                code.writeline("# loads")
+                for i in range(op_desc.num_input_tensors()):
+                    ptrs_expr: str = " + ".join(
+                        f"i{j} * in{i}_stride{j}" for j in range(rank)
+                    )
+                    ptrs_expr: str = f"in{i}_ptr + {ptrs_expr}"
+                    load_stmt: str = f"in{i} = tl.load({ptrs_expr}, mask=mask)"
+                    function_ns.create_name(f"in{i}")  # add to the namespace
+                    code.writeline(load_stmt)
+                code.newline()
+
+                # compute
+                code.writeline("# compute")
+                code.writeline(
+                    f"{outputs_to_scalar_fn} = inlined_f({inputs_to_scalar_fn})"
+                )
+                code.newline()
+
+                # stores
+                code.writeline("# stores")
+                for i in range(op_desc.num_output_tensors()):
+                    ptrs_expr: str = " + ".join(
+                        f"i{j} * out{i}_stride{j}" for j in range(rank)
+                    )
+                    ptrs_expr: str = f"out{i}_ptr + out{i}_offset + {ptrs_expr}"
+                    store_stmt: str = f"tl.store({ptrs_expr}, out{i}, mask=mask)"
+                    code.writeline(store_stmt)
+                code.newline()
+    return code
+
+
+def generate_code(
+    op_desc: OPDesc,
+    scalar_fn: JITFunction,
+    inputs: Tuple[Any],
+    wrapper_name: str,
+    destination_passing_func_name: str,
+    kernel_name: str,
+    code: IndentedBuffer,
+) -> IndentedBuffer:
+    assert (
+        len(inputs) == op_desc.num_inputs()
+    ), "the number of inputs does not match {str(op_desc)}"
+    input_tensor_ids = [i for i in range(op_desc.num_inputs()) if op_desc.is_tensor(i)]
+    tensor_shapes = [inputs[i].shape for i in input_tensor_ids]
+    shape = broadcast_shapes(tensor_shapes)
+    rank = len(shape)
+
+    # the only runtime determined factor is the rank of the task space
+    code = generate_imports(code)
+    code = generate_functional_pointwise_wrapper(
+        op_desc, wrapper_name, destination_passing_func_name, code
+    )
+    code = generate_destination_passing_pointwise_wrapper(
+        op_desc, rank, destination_passing_func_name, kernel_name, code
+    )
+    code = generate_pointwise_kernel(op_desc, scalar_fn, rank, kernel_name, code)
+    return code
+
+
+class PointwiseDynamicFunction:
+    """Utility to generate function for general pointwise operation. It generate wrapper & JITFunction
+    which are specialized according to the rank of the task space(the broadcasted shape of all input tensors).
+    The generated code are written out to the cache directory.
+    """
+
+    def __init__(self, op_desc: OPDesc, scalar_fn: JITFunction):
+        self._op_desc = op_desc
+
+        assert isinstance(scalar_fn, JITFunction)
+        self._scalar_fn = scalar_fn
+        self._scalar_fn_cache_key = scalar_fn.cache_key
+        self.pid = os.getpid()
+        self.lock = threading.Lock()
+
+        # instantiated & cached overloads
+        self.overloads: Mapping[str, Callable] = {}
+
+    def __call__(self, *args, **kwargs):
+        # note: kwargs should not be used in JITFunction directly
+        key = f"{self.arg_key(*args)}"
+        cache = self.overloads
+        lock = self.lock
+
+        while key not in cache:
+            # generate file & import it
+            with lock:
+                if key in cache:
+                    break
+                code = IndentedBuffer()
+                code = generate_code(
+                    self._op_desc,
+                    self._scalar_fn,
+                    args,
+                    "_wrapper",
+                    "_wrapper_out",
+                    "_jit_function",
+                    code,
+                )
+
+                file_name = f"pointwise_dynamic_{self._scalar_fn_cache_key}_rank_{key}_pid_{self.pid}.py"
+
+                with open(
+                    os.path.join(cache_dir(), file_name), "wt", encoding="utf-8"
+                ) as f:
+                    f.write(code.getvalue())
+
+                # load
+                spec = importlib.util.spec_from_file_location(
+                    f"_gen_module_{self._scalar_fn_cache_key}_rank_{key}_pid_{self.pid}",
+                    f.name,
+                )
+                m = importlib.util.module_from_spec(spec)
+                # do not expose it to sys.modules
+                # sys.modules["_add_module"] = m
+                spec.loader.exec_module(m)
+                overload = getattr(m, "_wrapper")
+                cache[key] = overload
+
+        overload = self.overloads[key]
+        return overload(*args, **kwargs)
+
+    def arg_key(self, *args):
+        tensors = [item for item in args if torch.is_tensor(item)]
+        max_rank = max(item.ndim for item in tensors)
+        return max_rank
+
+
+def pointwise_dynamic(
+    f: Optional[JITFunction] = None,
+    *,
+    num_inputs: Optional[int] = None,
+    is_tensor: Optional[List[bool]] = None,
+    dtypes: Optional[List[Optional[type]]] = None,
+    num_outputs: Optional[int] = None,
+    promotion_methods: Optional[Tuple[int, ...]] = None,
+):
+    def decorator(fn):
+        nonlocal num_inputs
+        if (num_inputs is None) and (is_tensor is None) and (dtypes is None):
+            num_inputs = len(fn.arg_names)
+        op_desc = OPDesc(
+            num_inputs=num_inputs,
+            is_tensor=is_tensor,
+            dtypes=dtypes,
+            num_outputs=num_outputs,
+            promotion_methods=promotion_methods,
+        )
+        return PointwiseDynamicFunction(op_desc, fn)
+
+    if f is not None:
+        return decorator(f)
+    return decorator
diff --git a/python/sglang/srt/kernels/utils/triton_utils.py b/python/sglang/srt/kernels/utils/triton_utils.py
new file mode 100644
index 000000000000..a70ab76e8d05
--- /dev/null
+++ b/python/sglang/srt/kernels/utils/triton_utils.py
@@ -0,0 +1,396 @@
+"""
+Copyright 2023-2024 SGLang Team
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+"""
+
+import builtins
+import contextlib
+import functools
+import inspect
+import keyword
+import operator
+import os
+import re
+import threading
+from collections import defaultdict
+from io import StringIO
+from pathlib import Path
+from typing import Dict, Iterable, Set, Tuple
+
+import torch
+import torch._prims_common as utils
+import triton
+from triton.runtime.cache import default_cache_dir
+
+
+def type_promotion(*args, type_promotion: utils.ELEMENTWISE_TYPE_PROMOTION_KIND):
+    computation_dtype, result_dtype = utils.elementwise_dtypes(
+        *args,
+        type_promotion_kind=type_promotion,
+    )
+    return computation_dtype, result_dtype
+
+
+def broadcast(s1: Tuple[int], s2: Tuple[int]) -> Tuple[int]:
+    _s1, _s2 = s1, s2
+    r1 = len(s1)
+    if r1 == 0:
+        return s2
+    r2 = len(s2)
+    if r2 == 0:
+        return s1
+
+    s1, s2 = (s1, s2) if r1 >= r2 else (s2, s1)
+    r1, r2 = (r1, r2) if r1 >= r2 else (r2, r1)
+
+    d = r1 - r2
+    s = list(s1)
+
+    for i in range(r2):
+        if s1[d + i] == 1:
+            s[d + i] = s2[i]
+        elif s2[i] == 1:
+            s[d + i] = s1[d + i]
+        elif s2[i] == s1[d + i]:
+            s[d + i] = s2[i]
+        else:
+            raise ValueError(f"Unbroadcastable {_s1} and {_s2}")
+    s = tuple(s)
+    return s
+
+
+def broadcastable_to(s1: Tuple[int], s2: Tuple[int]) -> bool:
+    r1 = len(s1)
+    if r1 == 0:
+        return True
+    r2 = len(s2)
+    if r2 == 0:  # r1 > 0
+        return False
+
+    if r1 > r2:
+        return False
+
+    d = r2 - r1
+    for i in range(r1):
+        if s1[i] == 1 or s1[i] == s2[d + i]:
+            continue
+        return False
+    return True
+
+
+def broadcast_shapes(shapes: Iterable[Tuple[int]]) -> Tuple[int]:
+    if len(shapes) == 0:
+        return ()
+    shape = shapes[0]
+    for s in shapes[1:]:
+        shape = broadcast(shape, s)
+    return shape
+
+
+def broadcasted_stride(
+    shape: Tuple[int], stride: Tuple[int], new_shape: Tuple[int]
+) -> Tuple[int]:
+    assert broadcastable_to(shape, new_shape)
+    r1 = len(shape)
+    r2 = len(new_shape)
+    d = r2 - r1
+    new_stride = [0 for _ in range(r2)]
+    for i in range(r1):
+        new_stride[d + i] = 0 if (shape[i] == 1 and new_shape[d + i] > 1) else stride[i]
+    return tuple(new_stride)
+
+
+def volume(shape: Tuple[int]) -> int:
+    return functools.reduce(operator.mul, shape, 1)
+
+
+def c_contiguous_stride(shape: Tuple[int]) -> Tuple[int]:
+    strides = []
+    s = 1
+    for size in reversed(shape):
+        strides.append(s)
+        s *= size
+
+    return tuple(reversed(strides))
+
+
+def size_in_bytes(a):
+    return a.numel() * a.element_size()
+
+
+class LibEntry(triton.KernelInterface):
+    def __init__(
+        self,
+        fn,
+    ):
+        self.fn = fn
+        self.arg_names = fn.arg_names
+        self.divisibility = 16
+        self.kernel_cache = tuple(dict() for _ in range(torch.cuda.device_count()))
+
+        fn = self.fn
+        while not isinstance(fn, triton.runtime.JITFunction):
+            fn = fn.fn
+        self.jit_function: triton.runtime.JITFunction = fn
+        self.specialize_indices = [
+            p.num
+            for p in self.jit_function.params
+            if not p.is_constexpr and not p.do_not_specialize
+        ]
+        self.do_not_specialize_indices = [
+            p.num
+            for p in self.jit_function.params
+            if not p.is_constexpr and p.do_not_specialize
+        ]
+        self.lock = threading.Lock()
+
+    def key(self, spec_args, dns_args, const_args):
+        spec_key = [
+            (
+                (arg.dtype, arg.data_ptr() % self.divisibility == 0)
+                if hasattr(arg, "data_ptr")
+                else (type(arg), arg)
+            )
+            for arg in spec_args
+        ]
+        dns_key = [
+            (
+                arg.dtype
+                if hasattr(arg, "data_ptr")
+                else (
+                    type(arg)
+                    if not isinstance(arg, int)
+                    else (
+                        "i32"
+                        if -(2**31) <= arg and arg <= 2**31 - 1
+                        else "u64" if 2**63 <= arg and arg <= 2**64 - 1 else "i64"
+                    )
+                )
+            )
+            for arg in dns_args
+        ]
+        # const args passed by position
+        return tuple(spec_key + dns_key + const_args)
+
+    def run(self, *args, **kwargs):
+        grid = kwargs["grid"]
+
+        # collect all the arguments
+        spec_args = []  # specialize arguments
+        dns_args = []  # do not specialize arguments
+        const_args = []  # constexpr arguments
+        k_args = []  # kernel arguments
+        for i, arg in enumerate(args):
+            if i in self.specialize_indices:
+                k_args.append(arg)
+                spec_args.append(arg)
+            elif i in self.do_not_specialize_indices:
+                k_args.append(arg)
+                dns_args.append(arg)
+            else:
+                const_args.append(arg)
+        for p in self.jit_function.params[len(args) :]:
+            if p.name in kwargs:
+                val = kwargs[p.name]
+            elif p.default is inspect._empty:
+                continue
+            else:
+                val = p.default
+
+            if p.is_constexpr:
+                const_args.append(val)
+            elif p.do_not_specialize:
+                dns_args.append(val)
+                k_args.append(val)
+            else:
+                spec_args.append(val)
+                k_args.append(val)
+
+        entry_key = self.key(spec_args, dns_args, const_args)
+        device = torch.cuda.current_device()
+        cache = self.kernel_cache[device]
+        while entry_key not in cache:
+            # NOTE: we serialize the first run of a jit function regardless of which device to run on
+            # because Triton runtime is currently not threadsafe.
+            with self.lock:
+                if entry_key in cache:
+                    break
+                kernel = self.fn.run(*args, **kwargs)
+                fn = self.fn
+                # collect constexpr arguments for grid computation
+                constexprs = {}
+                while not isinstance(fn, triton.runtime.JITFunction):
+                    if isinstance(fn, triton.runtime.Autotuner):
+                        config = fn.best_config
+                        constexprs["num_warps"] = config.num_warps
+                        constexprs["num_stages"] = config.num_stages
+                        constexprs["num_ctas"] = config.num_ctas
+                        constexprs = {**constexprs, **config.kwargs}
+                    elif isinstance(fn, triton.runtime.Heuristics):
+                        for v, heur in fn.values.items():
+                            constexprs[v] = heur(
+                                {
+                                    **dict(zip(fn.arg_names, args)),
+                                    **kwargs,
+                                    **constexprs,
+                                }
+                            )
+                    else:
+                        raise RuntimeError("Invalid Runtime Function")
+                    fn = fn.fn
+                for p in self.jit_function.params:
+                    if p.is_constexpr and p.name not in constexprs:
+                        constexprs[p.name] = p.default
+                cache[entry_key] = (kernel, constexprs)
+            return kernel, constexprs
+
+        kernel, constexprs = cache[entry_key]
+
+        if callable(grid):
+            # collect all arguments to the grid fn，ie:
+            # 1. args,
+            # 2. kwargs,
+            # 3. all all other captured arguments in CompiledKernel from Autotunner & Heuristics
+            # when kwargs & captured args conflict, captured args have higher priority
+            meta = {**dict(zip(self.arg_names, args)), **kwargs, **constexprs}
+            grid = grid(meta)
+        grid = grid + (1, 1)
+
+        kernel[grid[0:3]](*k_args)
+        return kernel, constexprs
+
+
+def libentry():
+    """
+    Decorator for triton library entries.
+    """
+
+    def decorator(fn):
+        return LibEntry(fn)
+
+    return decorator
+
+
+class IndentedBuffer:
+    tabwidth = 4
+
+    def __init__(self, initial_indent=0):
+        self._lines = []
+        self._indent = initial_indent
+
+    def getvalue(self) -> str:
+        buf = StringIO()
+        for line in self._lines:
+            assert isinstance(line, str)
+            buf.write(line)
+            buf.write("\n")
+        return buf.getvalue()
+
+    def clear(self):
+        self._lines.clear()
+
+    def __bool__(self):
+        return bool(self._lines)
+
+    def prefix(self):
+        return " " * (self._indent * self.tabwidth)
+
+    def newline(self):
+        self.writeline("\n")
+
+    def writeline(self, line):
+        if line.strip():
+            self._lines.append(f"{self.prefix()}{line}")
+        else:
+            self._lines.append("")
+
+    def writelines(self, lines):
+        for line in lines:
+            self.writeline(line)
+
+    def indent(self, offset=1):
+        @contextlib.contextmanager
+        def ctx():
+            self._indent += offset
+            try:
+                yield
+            finally:
+                self._indent -= offset
+
+        return ctx()
+
+
+class NameSpace:
+    def __init__(self):
+        self._used_names: Set[str] = set()
+        self._base_count: Dict[str, int] = defaultdict(int)
+
+        self._illegal_char_regex = re.compile("[^0-9a-zA-Z_]+")
+        self._name_suffix_regex = re.compile(r"(.*)_(\d+)$")
+
+    def create_name(self, candidate: str) -> str:
+        """Create a unique name.
+
+        Arguments:
+            candidate: used as the basis for the unique name, relevant to the user.
+        """
+        # delete all characters that are illegal in a Python identifier
+        candidate = self._illegal_char_regex.sub("_", candidate)
+
+        if not candidate:
+            candidate = "_unnamed"
+
+        if candidate[0].isdigit():
+            candidate = f"_{candidate}"
+
+        match = self._name_suffix_regex.match(candidate)
+        if match is None:
+            base = candidate
+            num = None
+        else:
+            base, num_str = match.group(1, 2)
+            num = int(num_str)
+
+        candidate = base if num is None else f"{base}_{num}"
+        if not num:
+            num = self._base_count[base]
+
+        while candidate in self._used_names or self._is_illegal_name(candidate):
+            num += 1
+            candidate = f"{base}_{num}"
+
+        self._used_names.add(candidate)
+        self._base_count[base] = num
+        return candidate
+
+    def _is_illegal_name(self, name: str) -> bool:
+        # 1. keywords are never allowed as names.
+        if name in keyword.kwlist:
+            return True
+
+        # 2. Can't shadow a builtin name, unless you *are* that builtin.
+        if name in builtins.__dict__:
+            return True
+
+        return False
+
+
+@functools.lru_cache(maxsize=None)
+def cache_dir_path() -> Path:
+    return default_cache_dir()
+
+
+def cache_dir() -> Path:
+    _cache_dir = cache_dir_path()
+    os.makedirs(_cache_dir, exist_ok=True)
+    return _cache_dir
diff --git a/python/sglang/srt/layers/activation.py b/python/sglang/srt/layers/activation.py
new file mode 100644
index 000000000000..a8c81e103303
--- /dev/null
+++ b/python/sglang/srt/layers/activation.py
@@ -0,0 +1,29 @@
+"""
+Copyright 2023-2024 SGLang Team
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+"""
+
+import torch
+import torch.nn as nn
+
+from sglang.srt.kernels.silu_and_mul import silu_and_mul
+
+
+class SiluAndMul(nn.Module):
+    def forward_native(self, x: torch.Tensor) -> torch.Tensor:
+        d = x.shape[-1] // 2
+        return F.silu(x[..., :d]) * x[..., d:]
+
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        d = x.shape[-1] // 2
+        return silu_and_mul(x[..., :d], x[..., d:])
diff --git a/python/sglang/srt/models/internlm2.py b/python/sglang/srt/models/internlm2.py
index 394d005042de..d9bcf7cf312d 100644
--- a/python/sglang/srt/models/internlm2.py
+++ b/python/sglang/srt/models/internlm2.py
@@ -23,7 +23,6 @@
 from transformers import PretrainedConfig
 from vllm.config import CacheConfig
 from vllm.distributed import get_tensor_model_parallel_world_size
-from vllm.model_executor.layers.activation import SiluAndMul
 from vllm.model_executor.layers.layernorm import RMSNorm
 from vllm.model_executor.layers.linear import (
     MergedColumnParallelLinear,
@@ -38,13 +37,13 @@
 )
 from vllm.model_executor.model_loader.weight_utils import default_weight_loader
 
+from sglang.srt.layers.activation import SiluAndMul
 from sglang.srt.layers.logits_processor import LogitsProcessor
 from sglang.srt.layers.radix_attention import RadixAttention
 from sglang.srt.model_executor.forward_batch_info import InputMetadata
 
 
 class InternLM2MLP(nn.Module):
-
     def __init__(
         self,
         hidden_size: int,
@@ -74,7 +73,6 @@ def forward(self, x):
 
 
 class InternLM2Attention(nn.Module):
-
     def __init__(
         self,
         hidden_size: int,
@@ -150,7 +148,6 @@ def forward(
 
 
 class InternLMDecoderLayer(nn.Module):
-
     def __init__(
         self,
         config: PretrainedConfig,
@@ -207,7 +204,6 @@ def forward(
 
 
 class InternLM2Model(nn.Module):
-
     def __init__(
         self,
         config: PretrainedConfig,
@@ -254,7 +250,6 @@ def forward(
 
 
 class InternLM2ForCausalLM(nn.Module):
-
     def __init__(
         self,
         config: PretrainedConfig,