diff --git a/.claude/skills/add-jit-kernel/SKILL.md b/.claude/skills/add-jit-kernel/SKILL.md
index 6b9758041..b71cd9df5 100644
--- a/.claude/skills/add-jit-kernel/SKILL.md
+++ b/.claude/skills/add-jit-kernel/SKILL.md
@@ -18,7 +18,7 @@ Add a new operation that scales each element of a tensor by a scalar factor:
 ## When to use JIT vs AOT (`sgl-kernel`)
 
 - **JIT (`jit_kernel`)**: prefer this first for kernels that do **not** depend on CUTLASS or another large C++ project. It is the default choice for lightweight kernels that benefit from rapid iteration and first-use compilation.
-- **AOT (`sgl-kernel`)**: prefer this when the kernel **does** depend on CUTLASS or another large C++ project, or when it should live in `sgl-kernel/` and participate in the wheel build / torch op registration flow.
+- **AOT (`sgl-kernel`)**: prefer this when the kernel **does** depend on CUTLASS or another large C++ project, or when it should live in `python/sglang/kernels/aot/` and participate in the wheel build / torch op registration flow.
 - **Exception**: kernels that depend on `flashinfer`, or on CUTLASS that is already provided through `flashinfer`, can still be implemented as `jit_kernel`.
 
 ---
diff --git a/.claude/skills/add-sgl-kernel/SKILL.md b/.claude/skills/add-sgl-kernel/SKILL.md
index 7bd17d589..1916a4abb 100644
--- a/.claude/skills/add-sgl-kernel/SKILL.md
+++ b/.claude/skills/add-sgl-kernel/SKILL.md
@@ -14,7 +14,7 @@ Add a new operation that scales each element of a tensor by a scalar factor:
 - Input: tensor `x` (CUDA) and scalar `factor` (float)
 - Output: `x * factor` (element-wise, in-place or into pre-allocated `out`)
 - Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**
-  - Dispatched via `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro (defined in `sgl-kernel/include/utils.h`)
+  - Dispatched via `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro (defined in `python/sglang/kernels/aot/include/utils.h`)
 
 ## Two rules of thumb (must follow)
 
@@ -33,13 +33,13 @@ In addition, every new kernel must ship with:
 
 You will typically touch these files/areas:
 
-- Implementation: `sgl-kernel/csrc/elementwise/scale.cu` (pick the right subdirectory)
-- Public declarations: `sgl-kernel/include/sgl_kernel_ops.h`
-- Torch extension registration: `sgl-kernel/csrc/common_extension.cc`
-- Build: `sgl-kernel/CMakeLists.txt` (`set(SOURCES ...)`)
-- Python API: `sgl-kernel/python/sgl_kernel/` and `sgl-kernel/python/sgl_kernel/__init__.py`
-- Tests: `sgl-kernel/tests/test_scale.py`
-- Benchmarks: `sgl-kernel/benchmark/bench_scale.py`
+- Implementation: `python/sglang/kernels/aot/csrc/elementwise/scale.cu` (pick the right subdirectory)
+- Public declarations: `python/sglang/kernels/aot/include/sgl_kernel_ops.h`
+- Torch extension registration: `python/sglang/kernels/aot/csrc/common_extension.cc`
+- Build: `python/sglang/kernels/aot/CMakeLists.txt` (`set(SOURCES ...)`)
+- Python API: `python/sglang/kernels/aot/python/sgl_kernel/` and `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`
+- Tests: `python/sglang/kernels/aot/tests/test_scale.py`
+- Benchmarks: `python/sglang/kernels/aot/benchmark/bench_scale.py`
 
 ---
 
@@ -50,7 +50,7 @@ Pick the right subdirectory:
 - `csrc/elementwise/` — for element-wise ops (our example)
 - `csrc/gemm/`, `csrc/attention/`, `csrc/moe/` — for other categories
 
-Create `sgl-kernel/csrc/elementwise/scale.cu`:
+Create `python/sglang/kernels/aot/csrc/elementwise/scale.cu`:
 
 ```cpp
 #include <ATen/cuda/CUDAContext.h>
@@ -115,7 +115,7 @@ void scale(at::Tensor& out, const at::Tensor& input, double factor) {
 
 ## Step 2: Add a C++ declaration in `include/sgl_kernel_ops.h`
 
-Edit `sgl-kernel/include/sgl_kernel_ops.h`, add to the elementwise section:
+Edit `python/sglang/kernels/aot/include/sgl_kernel_ops.h`, add to the elementwise section:
 
 ```cpp
 void scale(at::Tensor& out, const at::Tensor& input, double factor);
@@ -125,7 +125,7 @@ void scale(at::Tensor& out, const at::Tensor& input, double factor);
 
 ## Step 3: Register the op in `csrc/common_extension.cc`
 
-Edit `sgl-kernel/csrc/common_extension.cc`, inside `TORCH_LIBRARY_FRAGMENT(sgl_kernel, m)`:
+Edit `python/sglang/kernels/aot/csrc/common_extension.cc`, inside `TORCH_LIBRARY_FRAGMENT(sgl_kernel, m)`:
 
 ```cpp
 // From csrc/elementwise
@@ -143,7 +143,7 @@ m.impl("scale", torch::kCUDA, &scale);
 
 ## Step 4: Add the new source file to `CMakeLists.txt`
 
-Edit `sgl-kernel/CMakeLists.txt`, add to `set(SOURCES ...)`:
+Edit `python/sglang/kernels/aot/CMakeLists.txt`, add to `set(SOURCES ...)`:
 
 ```cmake
 csrc/elementwise/scale.cu
@@ -156,14 +156,14 @@ csrc/elementwise/scale.cu
 
 ---
 
-## Step 5: Expose a Python API under `sgl-kernel/python/sgl_kernel/`
+## Step 5: Expose a Python API under `python/sglang/kernels/aot/python/sgl_kernel/`
 
 Prefer following the existing module organization first. For elementwise kernels, the usual pattern is:
 
-- implement the Python wrapper in `sgl-kernel/python/sgl_kernel/elementwise.py`
-- then re-export it from `sgl-kernel/python/sgl_kernel/__init__.py`
+- implement the Python wrapper in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`
+- then re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py`
 
-For example, in `sgl-kernel/python/sgl_kernel/elementwise.py`, add:
+For example, in `python/sglang/kernels/aot/python/sgl_kernel/elementwise.py`, add:
 
 ```python
 import torch
@@ -190,13 +190,13 @@ def scale(
     return out
 ```
 
-Then re-export it from `sgl-kernel/python/sgl_kernel/__init__.py` following the existing import style used by other kernels.
+Then re-export it from `python/sglang/kernels/aot/python/sgl_kernel/__init__.py` following the existing import style used by other kernels.
 
 ---
 
 ## Step 6: Write tests (required)
 
-Create `sgl-kernel/tests/test_scale.py`:
+Create `python/sglang/kernels/aot/tests/test_scale.py`:
 ```python
 import pytest
 
@@ -241,7 +241,7 @@ if __name__ == "__main__":
 
 ## Step 7: Add a benchmark (required)
 
-Create `sgl-kernel/benchmark/bench_scale.py`:
+Create `python/sglang/kernels/aot/benchmark/bench_scale.py`:
 
 ```python
 import itertools
@@ -306,14 +306,14 @@ if __name__ == "__main__":
 Build:
 
 ```bash
-cd sgl-kernel
+cd python/sglang/kernels/aot
 make build -j16
 ```
 
 If you need to limit host resource usage:
 
 ```bash
-cd sgl-kernel
+cd python/sglang/kernels/aot
 make build -j1 MAX_JOBS=2 CMAKE_ARGS="-DSGL_KERNEL_COMPILE_THREADS=1"
 ```
 
@@ -324,8 +324,8 @@ make build -j1 MAX_JOBS=2 CMAKE_ARGS="-DSGL_KERNEL_COMPILE_THREADS=1"
 After building successfully, run the test and benchmark:
 
 ```bash
-pytest sgl-kernel/tests/test_scale.py -q
-python sgl-kernel/benchmark/bench_scale.py
+pytest python/sglang/kernels/aot/tests/test_scale.py -q
+python python/sglang/kernels/aot/benchmark/bench_scale.py
 ```
 
 PR CI also runs `pr-test-sgl-kernel.yml`, including the B200 job
@@ -339,29 +339,29 @@ Blackwell coverage signal for AOT `sgl-kernel` changes.
 - **Async CUDA errors**: `CUDA_LAUNCH_BLOCKING=1`
 - **Memory errors**: `compute-sanitizer --tool memcheck python ...`
 - **Build is too slow / OOM**: reduce `MAX_JOBS` and `SGL_KERNEL_COMPILE_THREADS`
-- **Binary bloat**: use `sgl-kernel/analyze_whl_kernel_sizes.py`
+- **Binary bloat**: use `python/sglang/kernels/aot/analyze_whl_kernel_sizes.py`
 - **CMake sources list**: if your `.cu` file is missing from `SOURCES`, the symbol will be undefined at link time
 
 ---
 
 ## References
 
-- `sgl-kernel/README.md`
-- `sgl-kernel/include/sgl_kernel_ops.h`
-- `sgl-kernel/csrc/common_extension.cc`
-- `sgl-kernel/CMakeLists.txt`
-- `sgl-kernel/include/utils.h` — `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro and friends
-- `sgl-kernel/csrc/elementwise/activation.cu` — reference for the FP16/BF16/FP32 dispatch pattern
+- `python/sglang/kernels/aot/README.md`
+- `python/sglang/kernels/aot/include/sgl_kernel_ops.h`
+- `python/sglang/kernels/aot/csrc/common_extension.cc`
+- `python/sglang/kernels/aot/CMakeLists.txt`
+- `python/sglang/kernels/aot/include/utils.h` — `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro and friends
+- `python/sglang/kernels/aot/csrc/elementwise/activation.cu` — reference for the FP16/BF16/FP32 dispatch pattern
 
 ## Summary of Files Created/Modified
 
 ```
-sgl-kernel/csrc/elementwise/scale.cu          # NEW: CUDA kernel + launcher
-sgl-kernel/include/sgl_kernel_ops.h           # MODIFIED: C++ declaration
-sgl-kernel/csrc/common_extension.cc           # MODIFIED: schema + dispatch registration
-sgl-kernel/CMakeLists.txt                     # MODIFIED: add source file (alphabetical)
-sgl-kernel/python/sgl_kernel/elementwise.py   # MODIFIED: Python wrapper
-sgl-kernel/python/sgl_kernel/__init__.py      # MODIFIED: re-export Python API
-sgl-kernel/tests/test_scale.py                # NEW: tests
-sgl-kernel/benchmark/bench_scale.py           # NEW: benchmark
+python/sglang/kernels/aot/csrc/elementwise/scale.cu          # NEW: CUDA kernel + launcher
+python/sglang/kernels/aot/include/sgl_kernel_ops.h           # MODIFIED: C++ declaration
+python/sglang/kernels/aot/csrc/common_extension.cc           # MODIFIED: schema + dispatch registration
+python/sglang/kernels/aot/CMakeLists.txt                     # MODIFIED: add source file (alphabetical)
+python/sglang/kernels/aot/python/sgl_kernel/elementwise.py   # MODIFIED: Python wrapper
+python/sglang/kernels/aot/python/sgl_kernel/__init__.py      # MODIFIED: re-export Python API
+python/sglang/kernels/aot/tests/test_scale.py                # NEW: tests
+python/sglang/kernels/aot/benchmark/bench_scale.py           # NEW: benchmark
 ```
diff --git a/.claude/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md b/.claude/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md
index 606b28818..612710180 100644
--- a/.claude/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md
+++ b/.claude/skills/llm-torch-profiler-analysis/references/fuse-overlap-catalog.md
@@ -145,7 +145,7 @@ Stable entries should be folded into the mainline family rows above.
 | PR `#21491` FlashInfer TRTLLM FP8 MoE with fused shared experts | `num_fused_shared_experts`<br>`trtllm_fp8_block_scale_moe` | `PR #21491`<br>`python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py`<br>`python/sglang/srt/models/deepseek_v2.py` | FlashInfer TRTLLM FP8 MoE path can fuse shared experts inside the routed MoE kernel | On FP8 TRTLLM MoE discussions, treat fused shared experts as an upstream pattern that already has a concrete PR. |
 | PR `#22005` fused add + RMSNorm + per-token FP8 quant | `fused_add_rmsnorm_per_token_quant`<br>`per_token_quant_fp8` | `PR #22005`<br>`python/sglang/kernels/jit/csrc/elementwise/fused_add_rmsnorm_per_token_quant.cuh`<br>`python/sglang/kernels/jit/fused_add_rmsnorm_per_token_quant.py` | CUDA JIT kernel keeps normed values in registers and emits BF16 + FP8 outputs plus per-token scales | If FP8 online-quant traces show add+norm followed by per-token quant, treat this as an in-flight upstream CUDA fuse family. |
 | PR `#20667` Qwen3.5 fused QK norm + RoPE + KV cache write | `fused_qk_norm_rope_cache_pts_quant_shuffle`<br>`fused_qk_norm_mrope_3d_cache_pts_quant_shuffle`<br>`rotary_dim` | `PR #20667`<br>`python/sglang/srt/models/qwen3_5.py`<br>`python/sglang/srt/models/utils.py` | ROCm / AITER path fuses Q / K RMSNorm, partial or 3D RoPE, and direct KV cache write for Qwen3.5 attention | Treat split QK-norm + RoPE + cache-store on Qwen3.5 as a concrete in-flight upstream family, not a novel idea. |
-| PR `#22392` CUTLASS FP8 GEMM replacing nvjet | `cutlass_scaled_mm`<br>`fp8_scaled_mm`<br>`nvjet`<br>`cudaMemsetAsync` | `PR #22392`<br>`sgl-kernel/python/sgl_kernel/gemm.py`<br>`python/sglang/srt/layers/quantization/fp8_utils.py` | Runtime replacement swaps nvjet FP8 GEMMs for CUTLASS kernels, removing per-launch memset bubbles and extra output-copy kernels | Treat nvjet GEMM + memset bubble ladders as an in-flight SGLang linear-kernel family before calling them novel. |
+| PR `#22392` CUTLASS FP8 GEMM replacing nvjet | `cutlass_scaled_mm`<br>`fp8_scaled_mm`<br>`nvjet`<br>`cudaMemsetAsync` | `PR #22392`<br>`python/sglang/kernels/aot/python/sgl_kernel/gemm.py`<br>`python/sglang/srt/layers/quantization/fp8_utils.py` | Runtime replacement swaps nvjet FP8 GEMMs for CUTLASS kernels, removing per-launch memset bubbles and extra output-copy kernels | Treat nvjet GEMM + memset bubble ladders as an in-flight SGLang linear-kernel family before calling them novel. |
 | PR `#18612` NVFP4 CUTLASS MoE fused SiLU+Mul+quant | `silu_and_mul_scaled_nvfp4`<br>`nvfp4 expert quant`<br>`cutlass moe` | `PR #18612`<br>`python/sglang/srt/layers/moe/cutlass_w4a8_moe.py`<br>`python/sglang/kernels/ops/quantization/nvfp4_gemm_swiglu_nvfp4_quant.py` | Fuses MoE activation epilogue and NVFP4 expert quantization before the CUTLASS MoE second GEMM | Treat split SiLU+Mul then NVFP4 expert quant in CUTLASS MoE traces as an in-flight upstream SGLang family. |
 | PR `#22918` FlashInfer per-token NVFP4 MoE | `per_token_nvfp4`<br>`trtllm_fp4_block_scale_moe`<br>`FlashInfer MoE` | `PR #22918`<br>`python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py` | Adds FlashInfer-backed per-token NVFP4 MoE execution so expert quant/dequant work can move into the fused MoE backend | Treat standalone per-token NVFP4 MoE support kernels as a candidate missing backend-selection path, not an automatically novel kernel idea. |
 | PR `#22851` NSA top-k backend and FlashInfer / PyTorch top-k split | `nsa topk`<br>`flashinfer_topk`<br>`pytorch_topk`<br>`fast_topk_transform` | `PR #22851`<br>`python/sglang/srt/layers/attention/nsa_backend.py` | Makes NSA top-k backend selection explicit and aligns fused top-k transform with FlashInfer / PyTorch fallbacks | When NSA top-k dominates decode, first classify it as backend selection or fused-transform eligibility work. |
diff --git a/.claude/skills/llm-torch-profiler-analysis/scripts/triage_kernel_helpers.py b/.claude/skills/llm-torch-profiler-analysis/scripts/triage_kernel_helpers.py
index cb410a1d8..81a71cb96 100644
--- a/.claude/skills/llm-torch-profiler-analysis/scripts/triage_kernel_helpers.py
+++ b/.claude/skills/llm-torch-profiler-analysis/scripts/triage_kernel_helpers.py
@@ -840,7 +840,7 @@ FUSION_PATTERN_REGISTRY: Tuple[FusionPatternSpec, ...] = (
         pattern="PR #22392 CUTLASS FP8 scaled MM replacing nvjet",
         candidate_path=(
             "PR #22392"
-            "<br>sgl-kernel/python/sgl_kernel/gemm.py"
+            "<br>python/sglang/kernels/aot/python/sgl_kernel/gemm.py"
             "<br>python/sglang/srt/layers/quantization/fp8_utils.py"
         ),
         active_keywords=("cutlass_scaled_mm", "fp8_scaled_mm"),
@@ -2632,7 +2632,7 @@ def fusion_framework_hints(spec: FusionPatternSpec) -> set[str]:
         hints.add("tokenspeed")
     if "tensorrt_llm/" in text:
         hints.add("trtllm")
-    if any(token in text for token in ("python/sglang/", "sgl-kernel/", "sgl_kernel/")):
+    if any(token in text for token in ("python/sglang/", "sgl_kernel/")):
         hints.add("sglang")
     return hints
 
diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index be8b70ce9..9730ceb3d 100644
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -67,8 +67,8 @@
 /python/sglang/srt/speculative @Ying1123 @merrymercy @hnyls2002 @Qiaolin-Yu
 /python/sglang/srt/utils/hf_transformers @JustinTong0323
 /python/sglang/srt/weight_cache @liusy58 @QiuMike @alexnails
-/sgl-kernel @ispobock @BBuf @yizhang2077 @merrymercy @FlamingoPg @HaiShaw
-/sgl-kernel/csrc/musa @yeahdongcn
+/python/sglang/kernels/aot @ispobock @BBuf @yizhang2077 @merrymercy @FlamingoPg @HaiShaw
+/python/sglang/kernels/aot/csrc/musa @yeahdongcn
 /sgl-model-gateway @slin1237 @CatherineSue
 /sgl-model-gateway/benches @slin1237
 /sgl-model-gateway/bindings/python @CatherineSue @key4ng @slin1237
diff --git a/.github/labeler.yml b/.github/labeler.yml
index 618083f0c..ea7835346 100644
--- a/.github/labeler.yml
+++ b/.github/labeler.yml
@@ -9,12 +9,14 @@ model-gateway:
 # Kernel specific
 sgl-kernel:
   - changed-files:
-    - any-glob-to-any-file: 'sgl-kernel/**/*'
+    - any-glob-to-any-file: 'python/sglang/kernels/aot/**/*'
 
 # JIT kernel specific
 jit-kernel:
   - changed-files:
-    - any-glob-to-any-file: 'python/sglang/kernels/**/*'
+    - any-glob-to-any-file:
+      - 'python/sglang/kernels/!(*.md)'
+      - 'python/sglang/kernels/!(aot)/**/*'
 
 # Documentation
 documentation:
@@ -90,7 +92,7 @@ blackwell:
   - changed-files:
     - any-glob-to-any-file:
       - '**/*nvfp4*'
-      - 'sgl-kernel/csrc/attention/cutlass_sm100_mla/**/*'
+      - 'python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/**/*'
       - 'python/sglang/srt/layers/attention/trtllm_mla_backend.py'
       - 'python/sglang/srt/layers/attention/trtllm_mha_backend.py'
 
diff --git a/.github/workflows/_pr-test-check-changes.yml b/.github/workflows/_pr-test-check-changes.yml
index ae3b060db..fe9263a5d 100644
--- a/.github/workflows/_pr-test-check-changes.yml
+++ b/.github/workflows/_pr-test-check-changes.yml
@@ -82,7 +82,9 @@ jobs:
               - ".github/workflows/pr-gate.yml"
               - ".github/actions/**"
               - "python/pyproject.toml"
-              - "python/sglang/!(multimodal_gen)/**/!(*.md)"
+              - "python/sglang/!(multimodal_gen|kernels)/**/!(*.md)"
+              - "python/sglang/kernels/!(*.md)"
+              - "python/sglang/kernels/!(aot)/**/!(*.md)"
               - "scripts/ci/cuda/*"
               - "scripts/ci/utils/*"
               - "test/**/!(*.md)"
@@ -104,11 +106,12 @@ jobs:
               - "test/registered/kernels/**"
               # sglang.kernels is the migrated kernel namespace (RFC #29630 / #30044); the
               # base-b-kernel suites import it directly, so kernel edits must run them.
-              - "python/sglang/kernels/**"
+              - "python/sglang/kernels/!(*.md)"
+              - "python/sglang/kernels/!(aot)/**"
             sgl_kernel:
               # Intentionally excludes ".github/workflows/pr-test-sgl-kernel.yml" —
               # see API-side detector below for rationale.
-              - "sgl-kernel/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
+              - "python/sglang/kernels/aot/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
 
       - name: Determine full-parallel mode
         id: parallel-mode
diff --git a/.github/workflows/_pr-test-sgl-kernel-build.yml b/.github/workflows/_pr-test-sgl-kernel-build.yml
index 75981dc16..5a72c5514 100644
--- a/.github/workflows/_pr-test-sgl-kernel-build.yml
+++ b/.github/workflows/_pr-test-sgl-kernel-build.yml
@@ -100,19 +100,19 @@ jobs:
 
       - name: Build wheel for Python ${{ matrix.python-version }} and CUDA ${{ matrix.cuda-version }}
         run: |
-          cd sgl-kernel
+          cd python/sglang/kernels/aot
           ./build.sh "${{ matrix.python-version }}" "${{ matrix.cuda-version }}"
         env:
           USE_CCACHE: 1
 
       - name: Verify wheel artifacts
         run: |
-          ls -alh sgl-kernel/dist
-          ls -alh sgl-kernel/dist/*.whl
+          ls -alh python/sglang/kernels/aot/dist
+          ls -alh python/sglang/kernels/aot/dist/*.whl
 
       - name: Upload artifacts
         uses: actions/upload-artifact@v4
         with:
           name: wheel-python${{ matrix.python-version }}-cuda${{ matrix.cuda-version }}${{ inputs.arch_suffix }}
-          path: sgl-kernel/dist/*
+          path: python/sglang/kernels/aot/dist/*
           if-no-files-found: error
diff --git a/.github/workflows/_pr-test-stage.yml b/.github/workflows/_pr-test-stage.yml
index bdf7f70f2..3dd77c072 100644
--- a/.github/workflows/_pr-test-stage.yml
+++ b/.github/workflows/_pr-test-stage.yml
@@ -108,7 +108,7 @@ jobs:
         if: ${{ fromJson(inputs.check_changes).sgl_kernel == 'true' && steps.rc.outputs.artifact_version == 'v4' }}
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -116,7 +116,7 @@ jobs:
         if: ${{ fromJson(inputs.check_changes).sgl_kernel == 'true' && steps.rc.outputs.artifact_version == 'v6' }}
         uses: actions/download-artifact@v6
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
diff --git a/.github/workflows/diffusion-ci-gt-gen.yml b/.github/workflows/diffusion-ci-gt-gen.yml
index ce175a6bc..28ed6305f 100644
--- a/.github/workflows/diffusion-ci-gt-gen.yml
+++ b/.github/workflows/diffusion-ci-gt-gen.yml
@@ -186,17 +186,17 @@ jobs:
             diffusion-ci/consistency_gt/official_generated/case_map.json
           sparse-checkout-cone-mode: false
 
-      - name: Prepare sgl-kernel/dist for prebuilt wheel
+      - name: Prepare AOT dist for prebuilt wheel
         if: inputs.kernel_artifact_run_id != ''
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download prebuilt sgl-kernel wheel
         if: inputs.kernel_artifact_run_id != ''
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           name: wheel-python3.10-cuda13.0
           run-id: ${{ inputs.kernel_artifact_run_id }}
@@ -382,17 +382,17 @@ jobs:
         with:
           ref: ${{ inputs.ref || github.ref }}
 
-      - name: Prepare sgl-kernel/dist for prebuilt wheel
+      - name: Prepare AOT dist for prebuilt wheel
         if: inputs.kernel_artifact_run_id != ''
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download prebuilt sgl-kernel wheel
         if: inputs.kernel_artifact_run_id != ''
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           name: wheel-python3.10-cuda13.0
           run-id: ${{ inputs.kernel_artifact_run_id }}
@@ -457,17 +457,17 @@ jobs:
         with:
           ref: ${{ inputs.ref || github.ref }}
 
-      - name: Prepare sgl-kernel/dist for prebuilt wheel
+      - name: Prepare AOT dist for prebuilt wheel
         if: inputs.kernel_artifact_run_id != ''
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download prebuilt sgl-kernel wheel
         if: inputs.kernel_artifact_run_id != ''
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           name: wheel-python3.10-cuda13.0
           run-id: ${{ inputs.kernel_artifact_run_id }}
@@ -532,17 +532,17 @@ jobs:
         with:
           ref: ${{ inputs.ref || github.ref }}
 
-      - name: Prepare sgl-kernel/dist for prebuilt wheel
+      - name: Prepare AOT dist for prebuilt wheel
         if: inputs.kernel_artifact_run_id != ''
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download prebuilt sgl-kernel wheel
         if: inputs.kernel_artifact_run_id != ''
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           name: wheel-python3.10-cuda13.0
           run-id: ${{ inputs.kernel_artifact_run_id }}
diff --git a/.github/workflows/lint.yml b/.github/workflows/lint.yml
index 7cb53216a..b4c637430 100644
--- a/.github/workflows/lint.yml
+++ b/.github/workflows/lint.yml
@@ -59,7 +59,7 @@ jobs:
       - name: Run sgl-kernel clang-format checks
         uses: DoozyX/clang-format-lint-action@v0.20
         with:
-          source: sgl-kernel
+          source: python/sglang/kernels/aot
           extensions: h,c,cpp,hpp,cu,cuh,cc
           clangFormatVersion: 20
           style: file
diff --git a/.github/workflows/nightly-72-gpu-gb200.yml b/.github/workflows/nightly-72-gpu-gb200.yml
index 43b93d026..d3f404c9d 100644
--- a/.github/workflows/nightly-72-gpu-gb200.yml
+++ b/.github/workflows/nightly-72-gpu-gb200.yml
@@ -132,7 +132,7 @@ jobs:
       image_tag: ${{ steps.build.outputs.image_tag }}
     steps:
       # Self-hosted runners retain the workspace across jobs. Prior `docker buildx`
-      # runs on this node leave root-owned build artifacts (e.g. sgl-kernel/build/)
+      # runs on this node leave root-owned build artifacts (e.g. python/sglang/kernels/aot/build/)
       # that actions/checkout cannot remove, causing EACCES on rmdir. Wipe them
       # via a throwaway root container before checkout recreates the workspace.
       - name: Clean workspace (remove root-owned files from prior runs)
diff --git a/.github/workflows/nightly-test-musa.yml b/.github/workflows/nightly-test-musa.yml
index eafb7f23b..0d9087e22 100644
--- a/.github/workflows/nightly-test-musa.yml
+++ b/.github/workflows/nightly-test-musa.yml
@@ -74,11 +74,11 @@ jobs:
         if: steps.gate.outputs.run_job == 'true'
         timeout-minutes: 30
         run: |
-          pytest sgl-kernel/tests/test_per_token_quant_fp8.py
-          pytest sgl-kernel/tests/speculative/test_eagle_utils.py
-          pytest sgl-kernel/tests/speculative/test_ngram_utils.py
-          pytest sgl-kernel/tests/speculative/test_speculative_sampling.py
-          pytest sgl-kernel/tests/test_torch_defaults_reset.py
+          pytest python/sglang/kernels/aot/tests/test_per_token_quant_fp8.py
+          pytest python/sglang/kernels/aot/tests/speculative/test_eagle_utils.py
+          pytest python/sglang/kernels/aot/tests/speculative/test_ngram_utils.py
+          pytest python/sglang/kernels/aot/tests/speculative/test_speculative_sampling.py
+          pytest python/sglang/kernels/aot/tests/test_torch_defaults_reset.py
 
   # ==================== General: multimodal layer ====================
   nightly-test-musa-general-multimodal-layer:
diff --git a/.github/workflows/pr-test-amd-rocm720.yml b/.github/workflows/pr-test-amd-rocm720.yml
index afdf9662f..d437e13e8 100644
--- a/.github/workflows/pr-test-amd-rocm720.yml
+++ b/.github/workflows/pr-test-amd-rocm720.yml
@@ -15,7 +15,7 @@ on:
   #     - "python/**"
   #     - "scripts/ci/**"
   #     - "test/**"
-  #     - "sgl-kernel/**"
+  #     - "python/sglang/kernels/aot/**"
   #     - ".github/workflows/pr-test-amd-rocm720.yml"
   #     - "docker/rocm.Dockerfile"
   # pull_request:
@@ -23,7 +23,7 @@ on:
   #     - "python/**"
   #     - "scripts/ci/**"
   #     - "test/**"
-  #     - "sgl-kernel/**"
+  #     - "python/sglang/kernels/aot/**"
   #     - ".github/workflows/pr-test-amd-rocm720.yml"
   #     - "docker/rocm.Dockerfile"
   workflow_dispatch:
@@ -187,7 +187,9 @@ jobs:
         with:
           filters: |
             main_package:
-              - "python/sglang/!(multimodal_gen)/**/!(*.md)"
+              - "python/sglang/!(multimodal_gen|kernels)/**/!(*.md)"
+              - "python/sglang/kernels/!(*.md)"
+              - "python/sglang/kernels/!(aot)/**/!(*.md)"
               - "python/pyproject_rocm.toml"
               - "python/pyproject_other.toml"
               - "scripts/ci/amd/*"
@@ -195,10 +197,11 @@ jobs:
               - "test/**/!(*.md)"
               - ".github/workflows/pr-test-amd-rocm720.yml"
             sgl_kernel:
-              - "sgl-kernel/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
+              - "python/sglang/kernels/aot/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
               - ".github/workflows/pr-test-amd-rocm720.yml"
             jit_kernel:
-              - "python/sglang/kernels/**"
+              - "python/sglang/kernels/!(*.md)"
+              - "python/sglang/kernels/!(aot)/**"
               - "test/registered/kernels/**"
               - ".github/workflows/pr-test-amd-rocm720.yml"
             multimodal_gen:
@@ -268,15 +271,15 @@ jobs:
       - name: Run test
         timeout-minutes: 30
         run: |
-          docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_moe_align.py
-          docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_moe_topk_softmax.py
-          docker exec -w /sglang-checkout/sgl-kernel/tests/speculative ci_sglang python3 -m pytest test_eagle_utils.py
-          docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_apply_token_bitmask_inplace.py
-          docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_activation.py
-          docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_topk.py
-          docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_kvcacheio.py
-          docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_moe_topk_sigmoid.py
-          docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_torch_defaults_reset.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_moe_align.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_moe_topk_softmax.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests/speculative ci_sglang python3 -m pytest test_eagle_utils.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_apply_token_bitmask_inplace.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_activation.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_topk.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_kvcacheio.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_moe_topk_sigmoid.py
+          docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_torch_defaults_reset.py
 
   sgl-kernel-unit-test-2-gpu-amd-rocm720:
     needs: [check-changes]
@@ -657,7 +660,7 @@ jobs:
         if: needs.check-changes.outputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda12.9
 
@@ -797,7 +800,7 @@ jobs:
         if: needs.check-changes.outputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda12.9
 
@@ -934,7 +937,7 @@ jobs:
         if: needs.check-changes.outputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda12.9
 
diff --git a/.github/workflows/pr-test-amd.yml b/.github/workflows/pr-test-amd.yml
index 566200d8e..4e5320805 100644
--- a/.github/workflows/pr-test-amd.yml
+++ b/.github/workflows/pr-test-amd.yml
@@ -11,7 +11,7 @@ on:
       - "python/**"
       - "scripts/ci/**"
       - "test/**"
-      - "sgl-kernel/**"
+      - "python/sglang/kernels/aot/**"
       - ".github/workflows/pr-test-amd.yml"
       - "docker/rocm.Dockerfile"
   workflow_dispatch:
@@ -175,7 +175,9 @@ jobs:
         with:
           filters: |
             main_package:
-              - "python/sglang/!(multimodal_gen)/**/!(*.md)"
+              - "python/sglang/!(multimodal_gen|kernels)/**/!(*.md)"
+              - "python/sglang/kernels/!(*.md)"
+              - "python/sglang/kernels/!(aot)/**/!(*.md)"
               - "python/pyproject_rocm.toml"
               - "python/pyproject_other.toml"
               - "scripts/ci/amd/*"
@@ -183,10 +185,11 @@ jobs:
               - "test/**/!(*.md)"
               - ".github/workflows/pr-test-amd.yml"
             sgl_kernel:
-              - "sgl-kernel/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
+              - "python/sglang/kernels/aot/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
               - ".github/workflows/pr-test-amd.yml"
             jit_kernel:
-              - "python/sglang/kernels/**"
+              - "python/sglang/kernels/!(*.md)"
+              - "python/sglang/kernels/!(aot)/**"
               - "test/registered/kernels/**"
               - ".github/workflows/pr-test-amd.yml"
             multimodal_gen:
@@ -266,15 +269,15 @@ jobs:
               "$@"
             fi
           }
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_moe_align.py
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_moe_topk_softmax.py
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests/speculative ci_sglang python3 -m pytest test_eagle_utils.py
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_apply_token_bitmask_inplace.py
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_activation.py
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_topk.py
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_kvcacheio.py
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_moe_topk_sigmoid.py
-          run_pytest docker exec -w /sglang-checkout/sgl-kernel/tests ci_sglang python3 -m pytest test_torch_defaults_reset.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_moe_align.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_moe_topk_softmax.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests/speculative ci_sglang python3 -m pytest test_eagle_utils.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_apply_token_bitmask_inplace.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_activation.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_topk.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_kvcacheio.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_moe_topk_sigmoid.py
+          run_pytest docker exec -w /sglang-checkout/python/sglang/kernels/aot/tests ci_sglang python3 -m pytest test_torch_defaults_reset.py
           exit $failures
 
   sgl-kernel-unit-test-2-gpu-amd:
@@ -681,7 +684,7 @@ jobs:
         if: needs.check-changes.outputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda12.9
 
@@ -821,7 +824,7 @@ jobs:
         if: needs.check-changes.outputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda12.9
 
@@ -960,7 +963,7 @@ jobs:
         if: needs.check-changes.outputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda12.9
 
diff --git a/.github/workflows/pr-test-arm64.yml b/.github/workflows/pr-test-arm64.yml
index d49b45c2a..64bccb6b4 100644
--- a/.github/workflows/pr-test-arm64.yml
+++ b/.github/workflows/pr-test-arm64.yml
@@ -55,7 +55,7 @@ jobs:
               - "python/sglang/!(multimodal_gen)/**/!(*.md)"
               - "python/pyproject_cpu.toml"
               - "test/**/!(*.md)"
-              - "sgl-kernel/**/*.!(md|txt)"
+              - "python/sglang/kernels/aot/**/*.!(md|txt)"
               - ".github/workflows/pr-test-arm64.yml"
               - "docker/arm64.Dockerfile"
 
diff --git a/.github/workflows/pr-test-jit-kernel.yml b/.github/workflows/pr-test-jit-kernel.yml
index 8ef7285b8..3c91479b9 100644
--- a/.github/workflows/pr-test-jit-kernel.yml
+++ b/.github/workflows/pr-test-jit-kernel.yml
@@ -56,14 +56,14 @@ jobs:
       - name: Cleanup
         if: inputs.sgl_kernel == 'true'
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download artifacts
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda13.0
 
@@ -94,14 +94,14 @@ jobs:
       - name: Cleanup
         if: inputs.sgl_kernel == 'true'
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download artifacts
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda13.0
 
@@ -134,14 +134,14 @@ jobs:
       - name: Cleanup
         if: inputs.sgl_kernel == 'true'
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download artifacts
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda13.0
 
@@ -174,14 +174,14 @@ jobs:
       - name: Cleanup
         if: inputs.sgl_kernel == 'true'
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download artifacts
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda13.0
 
diff --git a/.github/workflows/pr-test-mlx.yml b/.github/workflows/pr-test-mlx.yml
index 44eca2238..96d863f93 100644
--- a/.github/workflows/pr-test-mlx.yml
+++ b/.github/workflows/pr-test-mlx.yml
@@ -43,12 +43,14 @@ jobs:
           base: main
           filters: |
             main_package:
-              - "python/sglang/!(multimodal_gen)/**/!(*.md)"
+              - "python/sglang/!(multimodal_gen|kernels)/**/!(*.md)"
+              - "python/sglang/kernels/!(*.md)"
+              - "python/sglang/kernels/!(aot)/**/!(*.md)"
               - "python/pyproject_other.toml"
               - "test/**/!(*.md)"
               - ".github/workflows/pr-test-mlx.yml"
             sgl_kernel:
-              - "sgl-kernel/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
+              - "python/sglang/kernels/aot/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
               - ".github/workflows/pr-test-mlx.yml"
 
   # ==================== PR Gate ==================== #
diff --git a/.github/workflows/pr-test-multimodal-gen.yml b/.github/workflows/pr-test-multimodal-gen.yml
index c2c8957c9..8cf90602a 100644
--- a/.github/workflows/pr-test-multimodal-gen.yml
+++ b/.github/workflows/pr-test-multimodal-gen.yml
@@ -98,7 +98,7 @@ jobs:
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -164,7 +164,7 @@ jobs:
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -227,7 +227,7 @@ jobs:
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -289,7 +289,7 @@ jobs:
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -356,7 +356,7 @@ jobs:
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -401,7 +401,7 @@ jobs:
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -455,7 +455,7 @@ jobs:
         if: inputs.sgl_kernel == 'true'
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
diff --git a/.github/workflows/pr-test-musa.yml b/.github/workflows/pr-test-musa.yml
index ec55ffa32..550136fdc 100644
--- a/.github/workflows/pr-test-musa.yml
+++ b/.github/workflows/pr-test-musa.yml
@@ -83,7 +83,7 @@ jobs:
               - "python/sglang/multimodal_gen/test/server/musa/**"
             sgl_kernel:
               - ".github/workflows/pr-test-musa.yml"
-              - "sgl-kernel/csrc/musa/**"
+              - "python/sglang/kernels/aot/csrc/musa/**"
 
   # ==================== PR Gate ==================== #
   pr-gate:
@@ -217,11 +217,11 @@ jobs:
       - name: Run sgl-kernel test
         timeout-minutes: 20
         run: |
-          pytest sgl-kernel/tests/test_per_token_quant_fp8.py
-          pytest sgl-kernel/tests/speculative/test_eagle_utils.py
-          pytest sgl-kernel/tests/speculative/test_ngram_utils.py
-          pytest sgl-kernel/tests/speculative/test_speculative_sampling.py
-          pytest sgl-kernel/tests/test_torch_defaults_reset.py
+          pytest python/sglang/kernels/aot/tests/test_per_token_quant_fp8.py
+          pytest python/sglang/kernels/aot/tests/speculative/test_eagle_utils.py
+          pytest python/sglang/kernels/aot/tests/speculative/test_ngram_utils.py
+          pytest python/sglang/kernels/aot/tests/speculative/test_speculative_sampling.py
+          pytest python/sglang/kernels/aot/tests/test_torch_defaults_reset.py
 
 
   pr-test-musa-finish:
diff --git a/.github/workflows/pr-test-sgl-kernel.yml b/.github/workflows/pr-test-sgl-kernel.yml
index 272b0951f..dd0ccd7d1 100644
--- a/.github/workflows/pr-test-sgl-kernel.yml
+++ b/.github/workflows/pr-test-sgl-kernel.yml
@@ -44,13 +44,13 @@ jobs:
 
       - name: Cleanup
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download artifacts
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -62,7 +62,7 @@ jobs:
       - name: Run test
         timeout-minutes: 30
         run: |
-          cd sgl-kernel
+          cd python/sglang/kernels/aot
           pytest tests/
 
   sgl-kernel-benchmark-test:
@@ -79,13 +79,13 @@ jobs:
 
       - name: Cleanup
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download artifacts
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -97,7 +97,7 @@ jobs:
       - name: Run benchmark tests
         timeout-minutes: 45
         run: |
-          cd sgl-kernel/benchmark
+          cd python/sglang/kernels/aot/benchmark
           echo "Running sgl-kernel benchmark tests in CI mode..."
 
           echo "CI environment variable: $CI"
@@ -126,13 +126,13 @@ jobs:
 
       - name: Cleanup
         run: |
-          ls -alh sgl-kernel/dist || true
-          rm -rf sgl-kernel/dist/* || true
+          ls -alh python/sglang/kernels/aot/dist || true
+          rm -rf python/sglang/kernels/aot/dist/* || true
 
       - name: Download artifacts
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-python3.10-cuda*
 
@@ -144,7 +144,7 @@ jobs:
       - name: Run sgl-kernel unit tests on B200
         timeout-minutes: 30
         run: |
-          cd sgl-kernel
+          cd python/sglang/kernels/aot
           pytest tests/
 
   # Adding a single CUDA13 build-and-run check for the kernel
@@ -157,13 +157,13 @@ jobs:
 
   #     - name: Cleanup
   #       run: |
-  #         ls -alh sgl-kernel/dist || true
-  #         rm -rf sgl-kernel/dist/* || true
+  #         ls -alh python/sglang/kernels/aot/dist || true
+  #         rm -rf python/sglang/kernels/aot/dist/* || true
 
   #     - name: Download CUDA 13.0 artifacts
   #       uses: actions/download-artifact@v4
   #       with:
-  #         path: sgl-kernel/dist/
+  #         path: python/sglang/kernels/aot/dist/
   #         merge-multiple: true
   #         pattern: wheel-python3.10-cuda*
 
@@ -174,5 +174,5 @@ jobs:
   #     - name: Run kernel unit tests
   #       timeout-minutes: 30
   #       run: |
-  #         cd sgl-kernel
+  #         cd python/sglang/kernels/aot
   #         pytest tests/
diff --git a/.github/workflows/pr-test-xeon.yml b/.github/workflows/pr-test-xeon.yml
index b593a1dcc..8e10e955e 100644
--- a/.github/workflows/pr-test-xeon.yml
+++ b/.github/workflows/pr-test-xeon.yml
@@ -66,7 +66,7 @@ jobs:
               - "python/sglang/!(multimodal_gen)/**/!(*.md)"
               - "python/pyproject_cpu.toml"
               - "test/**/!(*.md)"
-              - "sgl-kernel/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
+              - "python/sglang/kernels/aot/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
               - ".github/workflows/pr-test-xeon.yml"
               - "docker/xeon.Dockerfile"
 
diff --git a/.github/workflows/pr-test-xpu.yml b/.github/workflows/pr-test-xpu.yml
index da3095115..ce94e6dae 100644
--- a/.github/workflows/pr-test-xpu.yml
+++ b/.github/workflows/pr-test-xpu.yml
@@ -57,7 +57,7 @@ jobs:
               - "python/sglang/!(multimodal_gen)/**/!(*.md)"
               - "python/pyproject_xpu.toml"
               - "test/**/!(*.md)"
-              - "sgl-kernel/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
+              - "python/sglang/kernels/aot/**/!(*.md|THIRDPARTYNOTICES.txt|LICENSE)"
               - ".github/workflows/pr-test-xpu.yml"
               - "docker/xpu.Dockerfile"
 
diff --git a/.github/workflows/release-whl-kernel.yml b/.github/workflows/release-whl-kernel.yml
index 98a89df32..50adb14f3 100644
--- a/.github/workflows/release-whl-kernel.yml
+++ b/.github/workflows/release-whl-kernel.yml
@@ -5,7 +5,7 @@ on:
     branches:
       - main
     paths:
-      - sgl-kernel/python/sgl_kernel/version.py
+      - python/sglang/kernels/aot/python/sgl_kernel/version.py
   workflow_dispatch:
     inputs:
       target:
@@ -54,7 +54,7 @@ jobs:
     runs-on: ${{ matrix.runner }}
     steps:
       # Self-hosted build nodes retain the workspace across jobs. Prior builds
-      # leave root-owned artifacts under sgl-kernel/build/ that actions/checkout
+      # leave root-owned artifacts under python/sglang/kernels/aot/build/ that actions/checkout
       # cannot remove, causing EACCES on rmdir. Wipe them via a throwaway root
       # container before checkout recreates the workspace.
       - name: Clean workspace (remove root-owned files from prior runs)
@@ -74,7 +74,7 @@ jobs:
 
       - name: Build wheels
         run: |
-          cd sgl-kernel
+          cd python/sglang/kernels/aot
           chmod +x ./build.sh
           ./build.sh "${{ matrix.python-version }}" "${{ matrix.cuda-version }}" ${{ matrix.arch == 'aarch64' && 'aarch64' || '' }}
         env:
@@ -85,7 +85,7 @@ jobs:
         uses: actions/upload-artifact@v4
         with:
           name: wheel-python${{ matrix.python-version }}-cuda${{ matrix.cuda-version }}${{ matrix.arch == 'aarch64' && '-aarch64' || '' }}
-          path: sgl-kernel/dist/*
+          path: python/sglang/kernels/aot/dist/*
 
   release-cu129:
     needs: build-cu129-matrix
@@ -98,7 +98,7 @@ jobs:
       - name: Download artifacts
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-*
 
@@ -106,7 +106,7 @@ jobs:
         id: set_tag_name
         run: |
           if [ -z "${{ inputs.tag_name }}" ]; then
-            TAG_NAME="v$(cat sgl-kernel/python/sgl_kernel/version.py | cut -d'"' -f2)"
+            TAG_NAME="v$(cat python/sglang/kernels/aot/python/sgl_kernel/version.py | cut -d'"' -f2)"
             echo "tag_name=$TAG_NAME" >> $GITHUB_OUTPUT
           else
             echo "tag_name=${{ inputs.tag_name }}" >> $GITHUB_OUTPUT
@@ -119,7 +119,7 @@ jobs:
           repository: sgl-project/whl
           token: ${{ secrets.GH_PAT_FOR_WHL_RELEASE }}
           files: |
-            sgl-kernel/dist/*
+            python/sglang/kernels/aot/dist/*
 
       - name: Clone wheel index
         run: git clone https://oauth2:${WHL_TOKEN}@github.com/sgl-project/whl.git sgl-whl
@@ -155,7 +155,7 @@ jobs:
     runs-on: ${{ matrix.runner }}
     steps:
       # Self-hosted build nodes retain the workspace across jobs. Prior builds
-      # leave root-owned artifacts under sgl-kernel/build/ that actions/checkout
+      # leave root-owned artifacts under python/sglang/kernels/aot/build/ that actions/checkout
       # cannot remove, causing EACCES on rmdir. Wipe them via a throwaway root
       # container before checkout recreates the workspace.
       - name: Clean workspace (remove root-owned files from prior runs)
@@ -175,7 +175,7 @@ jobs:
 
       - name: Build wheels
         run: |
-          cd sgl-kernel
+          cd python/sglang/kernels/aot
           chmod +x ./build.sh
           ./build.sh "${{ matrix.python-version }}" "${{ matrix.cuda-version }}" ${{ matrix.arch == 'aarch64' && 'aarch64' || '' }}
         env:
@@ -183,7 +183,7 @@ jobs:
           NVCC_THREADS: 8
 
       - name: Strip +cu130 local version for PyPI upload
-        working-directory: sgl-kernel
+        working-directory: python/sglang/kernels/aot
         run: |
           set -eux
           pip install wheel
@@ -208,7 +208,7 @@ jobs:
           ls -lh dist-pypi/
 
       - name: Upload to PyPI
-        working-directory: sgl-kernel
+        working-directory: python/sglang/kernels/aot
         run: |
           pip install twine
           python3 -m twine upload --skip-existing dist-pypi/* -u __token__ -p ${{ secrets.PYPI_TOKEN_SGLANG_KERNEL }}
@@ -217,7 +217,7 @@ jobs:
         uses: actions/upload-artifact@v4
         with:
           name: wheel-python${{ matrix.python-version }}-cuda${{ matrix.cuda-version }}${{ matrix.arch == 'aarch64' && '-aarch64' || '' }}
-          path: sgl-kernel/dist/*
+          path: python/sglang/kernels/aot/dist/*
 
   release-cu130:
     needs: build-cu130-matrix
@@ -230,7 +230,7 @@ jobs:
       - name: Download artifacts
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-*
 
@@ -238,7 +238,7 @@ jobs:
         id: set_tag_name
         run: |
           if [ -z "${{ inputs.tag_name }}" ]; then
-            TAG_NAME="v$(cat sgl-kernel/python/sgl_kernel/version.py | cut -d'"' -f2)"
+            TAG_NAME="v$(cat python/sglang/kernels/aot/python/sgl_kernel/version.py | cut -d'"' -f2)"
             echo "tag_name=$TAG_NAME" >> $GITHUB_OUTPUT
           else
             echo "tag_name=${{ inputs.tag_name }}" >> $GITHUB_OUTPUT
@@ -251,7 +251,7 @@ jobs:
           repository: sgl-project/whl
           token: ${{ secrets.GH_PAT_FOR_WHL_RELEASE }}
           files: |
-            sgl-kernel/dist/*
+            python/sglang/kernels/aot/dist/*
 
       - name: Clone wheel index
         run: git clone https://oauth2:${WHL_TOKEN}@github.com/sgl-project/whl.git sgl-whl
@@ -281,7 +281,7 @@ jobs:
         rocm-version: ["700", "720"]
     steps:
       # Self-hosted build nodes retain the workspace across jobs. Prior builds
-      # leave root-owned artifacts under sgl-kernel/build/ that actions/checkout
+      # leave root-owned artifacts under python/sglang/kernels/aot/build/ that actions/checkout
       # cannot remove, causing EACCES on rmdir. Wipe them via a throwaway root
       # container before checkout recreates the workspace.
       - name: Clean workspace (remove root-owned files from prior runs)
@@ -301,8 +301,8 @@ jobs:
 
       - name: Build wheels
         run: |
-          cp 3rdparty/amd/wheel/sgl-kernel/* sgl-kernel/
-          cd sgl-kernel
+          cp 3rdparty/amd/wheel/sgl-kernel/* python/sglang/kernels/aot/
+          cd python/sglang/kernels/aot
           chmod +x ./build_rocm.sh
           ./build_rocm.sh "${{ matrix.rocm-version }}"
 
@@ -310,7 +310,7 @@ jobs:
         uses: actions/upload-artifact@v4
         with:
           name: wheel-python${{ matrix.python-version }}-rocm${{ matrix.rocm-version }}
-          path: sgl-kernel/dist/*
+          path: python/sglang/kernels/aot/dist/*
 
   release-rocm700:
     needs: build-rocm-matrix
@@ -323,7 +323,7 @@ jobs:
       - name: Download artifacts
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-*-rocm700
 
@@ -331,7 +331,7 @@ jobs:
         id: set_tag_name
         run: |
           if [ -z "${{ inputs.tag_name }}" ]; then
-            TAG_NAME="v$(cat sgl-kernel/python/sgl_kernel/version.py | cut -d'"' -f2)"
+            TAG_NAME="v$(cat python/sglang/kernels/aot/python/sgl_kernel/version.py | cut -d'"' -f2)"
             echo "tag_name=$TAG_NAME" >> $GITHUB_OUTPUT
           else
             echo "tag_name=${{ inputs.tag_name }}" >> $GITHUB_OUTPUT
@@ -344,7 +344,7 @@ jobs:
           repository: sgl-project/whl
           token: ${{ secrets.GH_PAT_FOR_WHL_RELEASE }}
           files: |
-            sgl-kernel/dist/*
+            python/sglang/kernels/aot/dist/*
 
       - name: Clone wheel index
         run: git clone https://oauth2:${WHL_TOKEN}@github.com/sgl-project/whl.git sgl-whl
@@ -374,7 +374,7 @@ jobs:
       - name: Download artifacts
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-*-rocm720
 
@@ -382,7 +382,7 @@ jobs:
         id: set_tag_name
         run: |
           if [ -z "${{ inputs.tag_name }}" ]; then
-            TAG_NAME="v$(cat sgl-kernel/python/sgl_kernel/version.py | cut -d'"' -f2)"
+            TAG_NAME="v$(cat python/sglang/kernels/aot/python/sgl_kernel/version.py | cut -d'"' -f2)"
             echo "tag_name=$TAG_NAME" >> $GITHUB_OUTPUT
           else
             echo "tag_name=${{ inputs.tag_name }}" >> $GITHUB_OUTPUT
@@ -395,7 +395,7 @@ jobs:
           repository: sgl-project/whl
           token: ${{ secrets.GH_PAT_FOR_WHL_RELEASE }}
           files: |
-            sgl-kernel/dist/*
+            python/sglang/kernels/aot/dist/*
 
       - name: Clone wheel index
         run: git clone https://oauth2:${WHL_TOKEN}@github.com/sgl-project/whl.git sgl-whl
@@ -425,7 +425,7 @@ jobs:
         musa-version: ["43"]
     steps:
       # Self-hosted build nodes retain the workspace across jobs. Prior builds
-      # leave root-owned artifacts under sgl-kernel/build/ that actions/checkout
+      # leave root-owned artifacts under python/sglang/kernels/aot/build/ that actions/checkout
       # cannot remove, causing EACCES on rmdir. Wipe them via a throwaway root
       # container before checkout recreates the workspace.
       - name: Clean workspace (remove root-owned files from prior runs)
@@ -451,19 +451,19 @@ jobs:
 
       - name: Build wheels
         run: |
-          cd sgl-kernel
+          cd python/sglang/kernels/aot
           mv pyproject_musa.toml pyproject.toml
           python setup_musa.py sdist bdist_wheel
 
       - name: Rename MUSA wheels
         run: |
-          bash scripts/ci/musa/rename_wheels_musa.sh ${{ matrix.musa-version }} sgl-kernel/dist
+          bash scripts/ci/musa/rename_wheels_musa.sh ${{ matrix.musa-version }} python/sglang/kernels/aot/dist
 
       - name: Upload artifacts
         uses: actions/upload-artifact@v4
         with:
           name: wheel-python${{ matrix.python-version }}-musa${{ matrix.musa-version }}
-          path: sgl-kernel/dist/*
+          path: python/sglang/kernels/aot/dist/*
 
   release-musa43:
     needs: build-musa43
@@ -474,7 +474,7 @@ jobs:
       - name: Download artifacts
         uses: actions/download-artifact@v4
         with:
-          path: sgl-kernel/dist/
+          path: python/sglang/kernels/aot/dist/
           merge-multiple: true
           pattern: wheel-*
 
@@ -482,7 +482,7 @@ jobs:
         id: set_tag_name
         run: |
           if [ -z "${{ inputs.tag_name }}" ]; then
-            TAG_NAME="v$(cat sgl-kernel/python/sgl_kernel/version.py | cut -d'"' -f2)"
+            TAG_NAME="v$(cat python/sglang/kernels/aot/python/sgl_kernel/version.py | cut -d'"' -f2)"
             echo "tag_name=$TAG_NAME" >> $GITHUB_OUTPUT
           else
             echo "tag_name=${{ inputs.tag_name }}" >> $GITHUB_OUTPUT
@@ -495,7 +495,7 @@ jobs:
           repository: sgl-project/whl
           token: ${{ secrets.GH_PAT_FOR_WHL_RELEASE }}
           files: |
-            sgl-kernel/dist/*
+            python/sglang/kernels/aot/dist/*
 
       - name: Clone wheel index
         run: git clone https://oauth2:${WHL_TOKEN}@github.com/sgl-project/whl.git sgl-whl
diff --git a/.gitignore b/.gitignore
index 737f90f43..88a0a12b4 100644
--- a/.gitignore
+++ b/.gitignore
@@ -263,9 +263,9 @@ python/kernel.lock
 
 # MUSA section
 # Generated source files by torchada
-sgl-kernel/csrc_musa/
-sgl-kernel/include_musa/
-sgl-kernel/csrc/**/*_musa/
+python/sglang/kernels/aot/csrc_musa/
+python/sglang/kernels/aot/include_musa/
+python/sglang/kernels/aot/csrc/**/*_musa/
 
 # MUSA core dump files
 *.mudmp
diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index 88444f67b..67cd0178e 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -42,6 +42,7 @@ repos:
           (?x)^(
           .*/__init__\.py$|
           .*\.ipynb$|
+          python/sglang/kernels/aot/.*|
           python/sglang/srt/grpc/.*_pb2\.py$|
           python/sglang/srt/grpc/.*_pb2_grpc\.py$|
           python/sglang/srt/grpc/.*_pb2\.pyi$|
diff --git a/docker/arm64.Dockerfile b/docker/arm64.Dockerfile
index 5173e46be..73a240b74 100644
--- a/docker/arm64.Dockerfile
+++ b/docker/arm64.Dockerfile
@@ -42,7 +42,7 @@ RUN source $HOME/.local/bin/env && \
     cd python && \
     cp pyproject_cpu.toml pyproject.toml && \
     uv pip install . && \
-    cd ../sgl-kernel && \
+    cd sglang/kernels/aot && \
     cp pyproject_cpu.toml pyproject.toml && \
     uv pip install .
 
diff --git a/docker/rocm.Dockerfile b/docker/rocm.Dockerfile
index 1bee164d0..76348ec6f 100644
--- a/docker/rocm.Dockerfile
+++ b/docker/rocm.Dockerfile
@@ -317,11 +317,11 @@ RUN if [ "$BRANCH_TYPE" = "local" ]; then \
        fi \
     && rm -rf /tmp/local_src \
     && cd sglang \
-    && cd sgl-kernel \
+    && cd python/sglang/kernels/aot \
     && rm -f pyproject.toml \
     && mv pyproject_rocm.toml pyproject.toml \
     && AMDGPU_TARGET=$GPU_ARCH_LIST python setup_rocm.py install \
-    && cd .. \
+    && cd ../../../.. \
     && rm -rf python/pyproject.toml && mv python/pyproject_other.toml python/pyproject.toml \
     && if [ "$BUILD_TYPE" = "srt" ]; then \
          export SETUPTOOLS_SCM_PRETEND_VERSION="${SETUPTOOLS_SCM_PRETEND_VERSION}" && python -m pip --no-cache-dir install -e "python[srt_hip,diffusion_hip]"; \
diff --git a/docker/xeon.Dockerfile b/docker/xeon.Dockerfile
index 7850e9f8e..c01a87902 100644
--- a/docker/xeon.Dockerfile
+++ b/docker/xeon.Dockerfile
@@ -40,7 +40,7 @@ RUN source /opt/.venv/bin/activate && \
     cd python && \
     cp pyproject_cpu.toml pyproject.toml && \
     uv pip install . && \
-    cd ../sgl-kernel && \
+    cd sglang/kernels/aot && \
     cp pyproject_cpu.toml pyproject.toml && \
     uv pip install .
 
diff --git a/docs_new/docs/developer_guide/contribution_guide.mdx b/docs_new/docs/developer_guide/contribution_guide.mdx
index 1339257fa..300962e25 100644
--- a/docs_new/docs/developer_guide/contribution_guide.mdx
+++ b/docs_new/docs/developer_guide/contribution_guide.mdx
@@ -164,7 +164,7 @@ Users listed in [CI_PERMISSIONS.json](https://github.com/sgl-project/sglang/blob
 
 ## How to update sgl-kernel
 Since sglang and the `sglang-kernel` (prior `sgl-kernel`) distribution are separate Python packages, our current GitHub CI infrastructure does not support updating a kernel and using it immediately within the same pull request (PR).
-To add a new kernel or modify an existing one in the `sgl-kernel/` source tree, you must use multiple PRs.
+To add a new kernel or modify an existing one in the `python/sglang/kernels/aot/` source tree, you must use multiple PRs.
 
 Follow these steps:
 
diff --git a/docs_new/docs/hardware-platforms/amd_gpu.mdx b/docs_new/docs/hardware-platforms/amd_gpu.mdx
index 644f4fdfe..00cb4645c 100644
--- a/docs_new/docs/hardware-platforms/amd_gpu.mdx
+++ b/docs_new/docs/hardware-platforms/amd_gpu.mdx
@@ -50,11 +50,11 @@ cd sglang
 
 # Compile sgl-kernel
 pip install --upgrade pip
-cd sgl-kernel
+cd python/sglang/kernels/aot
 python setup_rocm.py install
 
 # Install sglang python package along with diffusion support
-cd ..
+cd ../../../..
 rm -rf python/pyproject.toml && mv python/pyproject_other.toml python/pyproject.toml
 pip install -e "python[all_hip]"
 ```
diff --git a/docs_new/docs/hardware-platforms/apple_metal.mdx b/docs_new/docs/hardware-platforms/apple_metal.mdx
index f1eb7de59..188bebe27 100644
--- a/docs_new/docs/hardware-platforms/apple_metal.mdx
+++ b/docs_new/docs/hardware-platforms/apple_metal.mdx
@@ -36,7 +36,7 @@ source sglang-metal/bin/activate
 
 # (Optional) Compile sgl-kernel
 uv pip install --upgrade pip
-uv run sgl-kernel/setup_metal.py install
+uv run python/sglang/kernels/aot/setup_metal.py install
 
 # Install sglang python package along with diffusion support
 rm -f python/pyproject.toml && mv python/pyproject_other.toml python/pyproject.toml
diff --git a/docs_new/docs/hardware-platforms/cpu_server.mdx b/docs_new/docs/hardware-platforms/cpu_server.mdx
index 9de015e73..6ac568698 100644
--- a/docs_new/docs/hardware-platforms/cpu_server.mdx
+++ b/docs_new/docs/hardware-platforms/cpu_server.mdx
@@ -126,7 +126,7 @@ uv pip install --upgrade pip setuptools
 uv pip install .
 
 # Build the CPU backend kernels
-cd ../sgl-kernel
+cd sglang/kernels/aot
 cp pyproject_cpu.toml pyproject.toml
 uv pip install .
 ```
diff --git a/docs_new/docs/hardware-platforms/mthreads_gpu.mdx b/docs_new/docs/hardware-platforms/mthreads_gpu.mdx
index bc9263bbe..ce31173c2 100644
--- a/docs_new/docs/hardware-platforms/mthreads_gpu.mdx
+++ b/docs_new/docs/hardware-platforms/mthreads_gpu.mdx
@@ -19,11 +19,11 @@ cd sglang
 
 # Compile sgl-kernel
 pip install --upgrade pip
-cd sgl-kernel
+cd python/sglang/kernels/aot
 python setup_musa.py install
 
 # Install sglang python package along with diffusion support
-cd ..
+cd ../../../..
 rm -f python/pyproject.toml && mv python/pyproject_other.toml python/pyproject.toml
 pip install -e "python[all_musa]"
 ```
diff --git a/python/MANIFEST.in b/python/MANIFEST.in
new file mode 100644
index 000000000..cc9b3e3cc
--- /dev/null
+++ b/python/MANIFEST.in
@@ -0,0 +1 @@
+prune sglang/kernels/aot
diff --git a/python/pyproject.toml b/python/pyproject.toml
index efc6cf46c..89267f397 100755
--- a/python/pyproject.toml
+++ b/python/pyproject.toml
@@ -199,6 +199,12 @@ killall_sglang = "sglang.cli.killall:main"
   "multimodal_gen/apps/realtime_webui/**/*"
 ]
 
+[tool.setuptools.exclude-package-data]
+"sglang" = [
+  "kernels/aot/*",
+  "kernels/aot/**/*",
+]
+
 [tool.setuptools.packages.find]
 exclude = [
   "assets*",
@@ -207,6 +213,7 @@ exclude = [
   "dist*",
   "playground*",
   "scripts*",
+  "sglang.kernels.aot*",
   "tests*",
 ]
 
@@ -218,6 +225,7 @@ exclude = [
   "dist*",
   "playground*",
   "scripts*",
+  "sglang/kernels/aot*",
   "tests*",
 ]
 
diff --git a/sgl-kernel/.clang-format b/python/sglang/kernels/aot/.clang-format
similarity index 100%
rename from sgl-kernel/.clang-format
rename to python/sglang/kernels/aot/.clang-format
diff --git a/sgl-kernel/CMakeLists.txt b/python/sglang/kernels/aot/CMakeLists.txt
similarity index 100%
rename from sgl-kernel/CMakeLists.txt
rename to python/sglang/kernels/aot/CMakeLists.txt
diff --git a/sgl-kernel/Dockerfile b/python/sglang/kernels/aot/Dockerfile
similarity index 100%
rename from sgl-kernel/Dockerfile
rename to python/sglang/kernels/aot/Dockerfile
diff --git a/sgl-kernel/LICENSE b/python/sglang/kernels/aot/LICENSE
similarity index 100%
rename from sgl-kernel/LICENSE
rename to python/sglang/kernels/aot/LICENSE
diff --git a/sgl-kernel/Makefile b/python/sglang/kernels/aot/Makefile
similarity index 100%
rename from sgl-kernel/Makefile
rename to python/sglang/kernels/aot/Makefile
diff --git a/sgl-kernel/README.md b/python/sglang/kernels/aot/README.md
similarity index 85%
rename from sgl-kernel/README.md
rename to python/sglang/kernels/aot/README.md
index de3bdf05d..0227c4b17 100644
--- a/sgl-kernel/README.md
+++ b/python/sglang/kernels/aot/README.md
@@ -1,6 +1,6 @@
 # sglang-kernel (prior sgl-kernel)
 
-[Kernel Library](https://github.com/sgl-project/sglang/tree/main/sgl-kernel) for LLM inference engines
+[Kernel Library](https://github.com/sgl-project/sglang/tree/main/python/sglang/kernels/aot) for LLM inference engines
 
 <div align="center">
 
@@ -9,7 +9,7 @@
 
 </div>
 
-`sglang-kernel` provides optimized compute primitives for LLM inference engines, enabling efficient inference for large language models and vision-language models through custom kernel operations. The source tree remains under the `sgl-kernel/` directory and the Python import path remains `sgl_kernel`.
+`sglang-kernel` provides optimized compute primitives for LLM inference engines, enabling efficient inference for large language models and vision-language models through custom kernel operations. The source tree lives under the `python/sglang/kernels/aot/` directory and the Python import path remains `sgl_kernel`.
 
 ## Installation
 Requires torch == 2.11.0
@@ -48,11 +48,11 @@ make build MAX_JOBS=2 CMAKE_ARGS="-DSGL_KERNEL_COMPILE_THREADS=1"
 
 ### Steps to add a new kernel:
 
-1. Implement the kernel in [csrc](https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc)
-2. Expose the interface in [include/sgl_kernel_ops.h](https://github.com/sgl-project/sglang/blob/main/sgl-kernel/include/sgl_kernel_ops.h)
-3. Create torch extension in [csrc/common_extension.cc](https://github.com/sgl-project/sglang/blob/main/sgl-kernel/csrc/common_extension.cc)
-4. Update [CMakeLists.txt](https://github.com/sgl-project/sglang/blob/main/sgl-kernel/CMakeLists.txt) to include new CUDA source
-5. Expose Python interface in [python](https://github.com/sgl-project/sglang/blob/main/sgl-kernel/python/sgl_kernel)
+1. Implement the kernel in [csrc](https://github.com/sgl-project/sglang/tree/main/python/sglang/kernels/aot/csrc)
+2. Expose the interface in [include/sgl_kernel_ops.h](https://github.com/sgl-project/sglang/blob/main/python/sglang/kernels/aot/include/sgl_kernel_ops.h)
+3. Create torch extension in [csrc/common_extension.cc](https://github.com/sgl-project/sglang/blob/main/python/sglang/kernels/aot/csrc/common_extension.cc)
+4. Update [CMakeLists.txt](https://github.com/sgl-project/sglang/blob/main/python/sglang/kernels/aot/CMakeLists.txt) to include new CUDA source
+5. Expose Python interface in [python](https://github.com/sgl-project/sglang/blob/main/python/sglang/kernels/aot/python/sgl_kernel)
 6. Add test and benchmark
 
 ### Development Tips
@@ -95,7 +95,7 @@ m.impl("fwd", torch::kCUDA, make_pytorch_shim(&mha_fwd));
 
 ### Testing & Benchmarking
 
-1. Add pytest tests in [tests/](https://github.com/sgl-project/sglang/tree/main/sgl-kernel/tests), if you need to skip some test, please use `@pytest.mark.skipif`
+1. Add pytest tests in [tests/](https://github.com/sgl-project/sglang/tree/main/python/sglang/kernels/aot/tests), if you need to skip some test, please use `@pytest.mark.skipif`
 
 ```python
 @pytest.mark.skipif(
@@ -103,7 +103,7 @@ m.impl("fwd", torch::kCUDA, make_pytorch_shim(&mha_fwd));
 )
 ```
 
-2. Add benchmarks using [triton benchmark](https://triton-lang.org/main/python-api/generated/triton.testing.Benchmark.html) in [benchmark/](https://github.com/sgl-project/sglang/tree/main/sgl-kernel/benchmark)
+2. Add benchmarks using [triton benchmark](https://triton-lang.org/main/python-api/generated/triton.testing.Benchmark.html) in [benchmark/](https://github.com/sgl-project/sglang/tree/main/python/sglang/kernels/aot/benchmark)
 
    **We recommend using `triton.testing.do_bench_cudagraph` for kernel benchmarking**:
 
diff --git a/sgl-kernel/THIRDPARTYNOTICES.txt b/python/sglang/kernels/aot/THIRDPARTYNOTICES.txt
similarity index 100%
rename from sgl-kernel/THIRDPARTYNOTICES.txt
rename to python/sglang/kernels/aot/THIRDPARTYNOTICES.txt
diff --git a/sgl-kernel/analyze_whl_kernel_sizes.py b/python/sglang/kernels/aot/analyze_whl_kernel_sizes.py
similarity index 100%
rename from sgl-kernel/analyze_whl_kernel_sizes.py
rename to python/sglang/kernels/aot/analyze_whl_kernel_sizes.py
diff --git a/sgl-kernel/benchmark/bench_activation.py b/python/sglang/kernels/aot/benchmark/bench_activation.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_activation.py
rename to python/sglang/kernels/aot/benchmark/bench_activation.py
diff --git a/sgl-kernel/benchmark/bench_amd_deterministic_allreduce.py b/python/sglang/kernels/aot/benchmark/bench_amd_deterministic_allreduce.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_amd_deterministic_allreduce.py
rename to python/sglang/kernels/aot/benchmark/bench_amd_deterministic_allreduce.py
diff --git a/sgl-kernel/benchmark/bench_awq_dequant.py b/python/sglang/kernels/aot/benchmark/bench_awq_dequant.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_awq_dequant.py
rename to python/sglang/kernels/aot/benchmark/bench_awq_dequant.py
diff --git a/sgl-kernel/benchmark/bench_cutlass_mla.py b/python/sglang/kernels/aot/benchmark/bench_cutlass_mla.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_cutlass_mla.py
rename to python/sglang/kernels/aot/benchmark/bench_cutlass_mla.py
diff --git a/sgl-kernel/benchmark/bench_dsv4_norm_rope.py b/python/sglang/kernels/aot/benchmark/bench_dsv4_norm_rope.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_dsv4_norm_rope.py
rename to python/sglang/kernels/aot/benchmark/bench_dsv4_norm_rope.py
diff --git a/sgl-kernel/benchmark/bench_es_fp8_blockwise_grouped_gemm.py b/python/sglang/kernels/aot/benchmark/bench_es_fp8_blockwise_grouped_gemm.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_es_fp8_blockwise_grouped_gemm.py
rename to python/sglang/kernels/aot/benchmark/bench_es_fp8_blockwise_grouped_gemm.py
diff --git a/sgl-kernel/benchmark/bench_fp4_gemm.py b/python/sglang/kernels/aot/benchmark/bench_fp4_gemm.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_fp4_gemm.py
rename to python/sglang/kernels/aot/benchmark/bench_fp4_gemm.py
diff --git a/sgl-kernel/benchmark/bench_fp8_blockwise_group_gemm.py b/python/sglang/kernels/aot/benchmark/bench_fp8_blockwise_group_gemm.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_fp8_blockwise_group_gemm.py
rename to python/sglang/kernels/aot/benchmark/bench_fp8_blockwise_group_gemm.py
diff --git a/sgl-kernel/benchmark/bench_fp8_gemm.py b/python/sglang/kernels/aot/benchmark/bench_fp8_gemm.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_fp8_gemm.py
rename to python/sglang/kernels/aot/benchmark/bench_fp8_gemm.py
diff --git a/sgl-kernel/benchmark/bench_fp8_gemm_swap_ab.py b/python/sglang/kernels/aot/benchmark/bench_fp8_gemm_swap_ab.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_fp8_gemm_swap_ab.py
rename to python/sglang/kernels/aot/benchmark/bench_fp8_gemm_swap_ab.py
diff --git a/sgl-kernel/benchmark/bench_int8_gemm.py b/python/sglang/kernels/aot/benchmark/bench_int8_gemm.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_int8_gemm.py
rename to python/sglang/kernels/aot/benchmark/bench_int8_gemm.py
diff --git a/sgl-kernel/benchmark/bench_moe_align_block_size.py b/python/sglang/kernels/aot/benchmark/bench_moe_align_block_size.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_moe_align_block_size.py
rename to python/sglang/kernels/aot/benchmark/bench_moe_align_block_size.py
diff --git a/sgl-kernel/benchmark/bench_moe_ep_post_reorder.py b/python/sglang/kernels/aot/benchmark/bench_moe_ep_post_reorder.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_moe_ep_post_reorder.py
rename to python/sglang/kernels/aot/benchmark/bench_moe_ep_post_reorder.py
diff --git a/sgl-kernel/benchmark/bench_moe_topk_sigmoid.py b/python/sglang/kernels/aot/benchmark/bench_moe_topk_sigmoid.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_moe_topk_sigmoid.py
rename to python/sglang/kernels/aot/benchmark/bench_moe_topk_sigmoid.py
diff --git a/sgl-kernel/benchmark/bench_moe_topk_softmax.py b/python/sglang/kernels/aot/benchmark/bench_moe_topk_softmax.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_moe_topk_softmax.py
rename to python/sglang/kernels/aot/benchmark/bench_moe_topk_softmax.py
diff --git a/sgl-kernel/benchmark/bench_mrope.py b/python/sglang/kernels/aot/benchmark/bench_mrope.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_mrope.py
rename to python/sglang/kernels/aot/benchmark/bench_mrope.py
diff --git a/sgl-kernel/benchmark/bench_per_tensor_quant_fp8.py b/python/sglang/kernels/aot/benchmark/bench_per_tensor_quant_fp8.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_per_tensor_quant_fp8.py
rename to python/sglang/kernels/aot/benchmark/bench_per_tensor_quant_fp8.py
diff --git a/sgl-kernel/benchmark/bench_per_token_group_quant_8bit.py b/python/sglang/kernels/aot/benchmark/bench_per_token_group_quant_8bit.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_per_token_group_quant_8bit.py
rename to python/sglang/kernels/aot/benchmark/bench_per_token_group_quant_8bit.py
diff --git a/sgl-kernel/benchmark/bench_per_token_quant_fp8.py b/python/sglang/kernels/aot/benchmark/bench_per_token_quant_fp8.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_per_token_quant_fp8.py
rename to python/sglang/kernels/aot/benchmark/bench_per_token_quant_fp8.py
diff --git a/sgl-kernel/benchmark/bench_rmsnorm.py b/python/sglang/kernels/aot/benchmark/bench_rmsnorm.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_rmsnorm.py
rename to python/sglang/kernels/aot/benchmark/bench_rmsnorm.py
diff --git a/sgl-kernel/benchmark/bench_rotary_embedding.py b/python/sglang/kernels/aot/benchmark/bench_rotary_embedding.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_rotary_embedding.py
rename to python/sglang/kernels/aot/benchmark/bench_rotary_embedding.py
diff --git a/sgl-kernel/benchmark/bench_sum_scale.py b/python/sglang/kernels/aot/benchmark/bench_sum_scale.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_sum_scale.py
rename to python/sglang/kernels/aot/benchmark/bench_sum_scale.py
diff --git a/sgl-kernel/benchmark/bench_top_k_top_p_sampling.py b/python/sglang/kernels/aot/benchmark/bench_top_k_top_p_sampling.py
similarity index 100%
rename from sgl-kernel/benchmark/bench_top_k_top_p_sampling.py
rename to python/sglang/kernels/aot/benchmark/bench_top_k_top_p_sampling.py
diff --git a/sgl-kernel/build.sh b/python/sglang/kernels/aot/build.sh
similarity index 100%
rename from sgl-kernel/build.sh
rename to python/sglang/kernels/aot/build.sh
diff --git a/sgl-kernel/cmake/flashmla.cmake b/python/sglang/kernels/aot/cmake/flashmla.cmake
similarity index 100%
rename from sgl-kernel/cmake/flashmla.cmake
rename to python/sglang/kernels/aot/cmake/flashmla.cmake
diff --git a/sgl-kernel/cmake/utils.cmake b/python/sglang/kernels/aot/cmake/utils.cmake
similarity index 100%
rename from sgl-kernel/cmake/utils.cmake
rename to python/sglang/kernels/aot/cmake/utils.cmake
diff --git a/sgl-kernel/csrc/allreduce/custom_all_reduce.cu b/python/sglang/kernels/aot/csrc/allreduce/custom_all_reduce.cu
similarity index 100%
rename from sgl-kernel/csrc/allreduce/custom_all_reduce.cu
rename to python/sglang/kernels/aot/csrc/allreduce/custom_all_reduce.cu
diff --git a/sgl-kernel/csrc/allreduce/custom_all_reduce.cuh b/python/sglang/kernels/aot/csrc/allreduce/custom_all_reduce.cuh
similarity index 100%
rename from sgl-kernel/csrc/allreduce/custom_all_reduce.cuh
rename to python/sglang/kernels/aot/csrc/allreduce/custom_all_reduce.cuh
diff --git a/sgl-kernel/csrc/allreduce/custom_all_reduce.hip b/python/sglang/kernels/aot/csrc/allreduce/custom_all_reduce.hip
similarity index 100%
rename from sgl-kernel/csrc/allreduce/custom_all_reduce.hip
rename to python/sglang/kernels/aot/csrc/allreduce/custom_all_reduce.hip
diff --git a/sgl-kernel/csrc/allreduce/custom_all_reduce_hip.cuh b/python/sglang/kernels/aot/csrc/allreduce/custom_all_reduce_hip.cuh
similarity index 100%
rename from sgl-kernel/csrc/allreduce/custom_all_reduce_hip.cuh
rename to python/sglang/kernels/aot/csrc/allreduce/custom_all_reduce_hip.cuh
diff --git a/sgl-kernel/csrc/allreduce/deterministic_all_reduce.hip b/python/sglang/kernels/aot/csrc/allreduce/deterministic_all_reduce.hip
similarity index 100%
rename from sgl-kernel/csrc/allreduce/deterministic_all_reduce.hip
rename to python/sglang/kernels/aot/csrc/allreduce/deterministic_all_reduce.hip
diff --git a/sgl-kernel/csrc/allreduce/quick_all_reduce.cu b/python/sglang/kernels/aot/csrc/allreduce/quick_all_reduce.cu
similarity index 100%
rename from sgl-kernel/csrc/allreduce/quick_all_reduce.cu
rename to python/sglang/kernels/aot/csrc/allreduce/quick_all_reduce.cu
diff --git a/sgl-kernel/csrc/allreduce/quick_all_reduce.cuh b/python/sglang/kernels/aot/csrc/allreduce/quick_all_reduce.cuh
similarity index 100%
rename from sgl-kernel/csrc/allreduce/quick_all_reduce.cuh
rename to python/sglang/kernels/aot/csrc/allreduce/quick_all_reduce.cuh
diff --git a/sgl-kernel/csrc/allreduce/quick_all_reduce.h b/python/sglang/kernels/aot/csrc/allreduce/quick_all_reduce.h
similarity index 100%
rename from sgl-kernel/csrc/allreduce/quick_all_reduce.h
rename to python/sglang/kernels/aot/csrc/allreduce/quick_all_reduce.h
diff --git a/sgl-kernel/csrc/allreduce/quick_all_reduce_base.h b/python/sglang/kernels/aot/csrc/allreduce/quick_all_reduce_base.h
similarity index 100%
rename from sgl-kernel/csrc/allreduce/quick_all_reduce_base.h
rename to python/sglang/kernels/aot/csrc/allreduce/quick_all_reduce_base.h
diff --git a/sgl-kernel/csrc/attention/cutlass_mla_kernel.cu b/python/sglang/kernels/aot/csrc/attention/cutlass_mla_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/attention/cutlass_mla_kernel.cu
rename to python/sglang/kernels/aot/csrc/attention/cutlass_mla_kernel.cu
diff --git a/sgl-kernel/csrc/attention/cutlass_sm100_mla/device/sm100_mla.hpp b/python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/device/sm100_mla.hpp
similarity index 100%
rename from sgl-kernel/csrc/attention/cutlass_sm100_mla/device/sm100_mla.hpp
rename to python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/device/sm100_mla.hpp
diff --git a/sgl-kernel/csrc/attention/cutlass_sm100_mla/kernel/sm100_fmha_mla_reduction.hpp b/python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/kernel/sm100_fmha_mla_reduction.hpp
similarity index 100%
rename from sgl-kernel/csrc/attention/cutlass_sm100_mla/kernel/sm100_fmha_mla_reduction.hpp
rename to python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/kernel/sm100_fmha_mla_reduction.hpp
diff --git a/sgl-kernel/csrc/attention/cutlass_sm100_mla/kernel/sm100_fmha_mla_tma_warpspecialized.hpp b/python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/kernel/sm100_fmha_mla_tma_warpspecialized.hpp
similarity index 100%
rename from sgl-kernel/csrc/attention/cutlass_sm100_mla/kernel/sm100_fmha_mla_tma_warpspecialized.hpp
rename to python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/kernel/sm100_fmha_mla_tma_warpspecialized.hpp
diff --git a/sgl-kernel/csrc/attention/cutlass_sm100_mla/kernel/sm100_mla_tile_scheduler.hpp b/python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/kernel/sm100_mla_tile_scheduler.hpp
similarity index 100%
rename from sgl-kernel/csrc/attention/cutlass_sm100_mla/kernel/sm100_mla_tile_scheduler.hpp
rename to python/sglang/kernels/aot/csrc/attention/cutlass_sm100_mla/kernel/sm100_mla_tile_scheduler.hpp
diff --git a/sgl-kernel/csrc/attention/merge_attn_states.cu b/python/sglang/kernels/aot/csrc/attention/merge_attn_states.cu
similarity index 100%
rename from sgl-kernel/csrc/attention/merge_attn_states.cu
rename to python/sglang/kernels/aot/csrc/attention/merge_attn_states.cu
diff --git a/sgl-kernel/csrc/attention/vertical_slash_index.cu b/python/sglang/kernels/aot/csrc/attention/vertical_slash_index.cu
similarity index 100%
rename from sgl-kernel/csrc/attention/vertical_slash_index.cu
rename to python/sglang/kernels/aot/csrc/attention/vertical_slash_index.cu
diff --git a/sgl-kernel/csrc/common_extension.cc b/python/sglang/kernels/aot/csrc/common_extension.cc
similarity index 100%
rename from sgl-kernel/csrc/common_extension.cc
rename to python/sglang/kernels/aot/csrc/common_extension.cc
diff --git a/sgl-kernel/csrc/common_extension_musa.cc b/python/sglang/kernels/aot/csrc/common_extension_musa.cc
similarity index 100%
rename from sgl-kernel/csrc/common_extension_musa.cc
rename to python/sglang/kernels/aot/csrc/common_extension_musa.cc
diff --git a/sgl-kernel/csrc/common_extension_rocm.cc b/python/sglang/kernels/aot/csrc/common_extension_rocm.cc
similarity index 100%
rename from sgl-kernel/csrc/common_extension_rocm.cc
rename to python/sglang/kernels/aot/csrc/common_extension_rocm.cc
diff --git a/sgl-kernel/csrc/cpu/CMakeLists.txt b/python/sglang/kernels/aot/csrc/cpu/CMakeLists.txt
similarity index 100%
rename from sgl-kernel/csrc/cpu/CMakeLists.txt
rename to python/sglang/kernels/aot/csrc/cpu/CMakeLists.txt
diff --git a/sgl-kernel/csrc/cpu/aarch64/gemm_int8.cpp b/python/sglang/kernels/aot/csrc/cpu/aarch64/gemm_int8.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/aarch64/gemm_int8.cpp
rename to python/sglang/kernels/aot/csrc/cpu/aarch64/gemm_int8.cpp
diff --git a/sgl-kernel/csrc/cpu/aarch64/moe.cpp b/python/sglang/kernels/aot/csrc/cpu/aarch64/moe.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/aarch64/moe.cpp
rename to python/sglang/kernels/aot/csrc/cpu/aarch64/moe.cpp
diff --git a/sgl-kernel/csrc/cpu/aarch64/op.h b/python/sglang/kernels/aot/csrc/cpu/aarch64/op.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/aarch64/op.h
rename to python/sglang/kernels/aot/csrc/cpu/aarch64/op.h
diff --git a/sgl-kernel/csrc/cpu/aarch64/shm.h b/python/sglang/kernels/aot/csrc/cpu/aarch64/shm.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/aarch64/shm.h
rename to python/sglang/kernels/aot/csrc/cpu/aarch64/shm.h
diff --git a/sgl-kernel/csrc/cpu/activation.cpp b/python/sglang/kernels/aot/csrc/cpu/activation.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/activation.cpp
rename to python/sglang/kernels/aot/csrc/cpu/activation.cpp
diff --git a/sgl-kernel/csrc/cpu/bmm.cpp b/python/sglang/kernels/aot/csrc/cpu/bmm.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/bmm.cpp
rename to python/sglang/kernels/aot/csrc/cpu/bmm.cpp
diff --git a/sgl-kernel/csrc/cpu/common.h b/python/sglang/kernels/aot/csrc/cpu/common.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/common.h
rename to python/sglang/kernels/aot/csrc/cpu/common.h
diff --git a/sgl-kernel/csrc/cpu/conv3d.cpp b/python/sglang/kernels/aot/csrc/cpu/conv3d.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/conv3d.cpp
rename to python/sglang/kernels/aot/csrc/cpu/conv3d.cpp
diff --git a/sgl-kernel/csrc/cpu/decode.cpp b/python/sglang/kernels/aot/csrc/cpu/decode.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/decode.cpp
rename to python/sglang/kernels/aot/csrc/cpu/decode.cpp
diff --git a/sgl-kernel/csrc/cpu/extend.cpp b/python/sglang/kernels/aot/csrc/cpu/extend.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/extend.cpp
rename to python/sglang/kernels/aot/csrc/cpu/extend.cpp
diff --git a/sgl-kernel/csrc/cpu/flash_attn.cpp b/python/sglang/kernels/aot/csrc/cpu/flash_attn.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/flash_attn.cpp
rename to python/sglang/kernels/aot/csrc/cpu/flash_attn.cpp
diff --git a/sgl-kernel/csrc/cpu/flash_attn.h b/python/sglang/kernels/aot/csrc/cpu/flash_attn.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/flash_attn.h
rename to python/sglang/kernels/aot/csrc/cpu/flash_attn.h
diff --git a/sgl-kernel/csrc/cpu/gemm.cpp b/python/sglang/kernels/aot/csrc/cpu/gemm.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/gemm.cpp
rename to python/sglang/kernels/aot/csrc/cpu/gemm.cpp
diff --git a/sgl-kernel/csrc/cpu/gemm.h b/python/sglang/kernels/aot/csrc/cpu/gemm.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/gemm.h
rename to python/sglang/kernels/aot/csrc/cpu/gemm.h
diff --git a/sgl-kernel/csrc/cpu/gemm_fp8.cpp b/python/sglang/kernels/aot/csrc/cpu/gemm_fp8.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/gemm_fp8.cpp
rename to python/sglang/kernels/aot/csrc/cpu/gemm_fp8.cpp
diff --git a/sgl-kernel/csrc/cpu/gemm_int4.cpp b/python/sglang/kernels/aot/csrc/cpu/gemm_int4.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/gemm_int4.cpp
rename to python/sglang/kernels/aot/csrc/cpu/gemm_int4.cpp
diff --git a/sgl-kernel/csrc/cpu/gemm_int8.cpp b/python/sglang/kernels/aot/csrc/cpu/gemm_int8.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/gemm_int8.cpp
rename to python/sglang/kernels/aot/csrc/cpu/gemm_int8.cpp
diff --git a/sgl-kernel/csrc/cpu/interface.cpp b/python/sglang/kernels/aot/csrc/cpu/interface.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/interface.cpp
rename to python/sglang/kernels/aot/csrc/cpu/interface.cpp
diff --git a/sgl-kernel/csrc/cpu/kvcache.cpp b/python/sglang/kernels/aot/csrc/cpu/kvcache.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/kvcache.cpp
rename to python/sglang/kernels/aot/csrc/cpu/kvcache.cpp
diff --git a/sgl-kernel/csrc/cpu/mamba/conv.cpp b/python/sglang/kernels/aot/csrc/cpu/mamba/conv.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/mamba/conv.cpp
rename to python/sglang/kernels/aot/csrc/cpu/mamba/conv.cpp
diff --git a/sgl-kernel/csrc/cpu/mamba/fla.cpp b/python/sglang/kernels/aot/csrc/cpu/mamba/fla.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/mamba/fla.cpp
rename to python/sglang/kernels/aot/csrc/cpu/mamba/fla.cpp
diff --git a/sgl-kernel/csrc/cpu/model/qwen3.cpp b/python/sglang/kernels/aot/csrc/cpu/model/qwen3.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/model/qwen3.cpp
rename to python/sglang/kernels/aot/csrc/cpu/model/qwen3.cpp
diff --git a/sgl-kernel/csrc/cpu/moe.cpp b/python/sglang/kernels/aot/csrc/cpu/moe.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/moe.cpp
rename to python/sglang/kernels/aot/csrc/cpu/moe.cpp
diff --git a/sgl-kernel/csrc/cpu/moe.h b/python/sglang/kernels/aot/csrc/cpu/moe.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/moe.h
rename to python/sglang/kernels/aot/csrc/cpu/moe.h
diff --git a/sgl-kernel/csrc/cpu/moe_fp8.cpp b/python/sglang/kernels/aot/csrc/cpu/moe_fp8.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/moe_fp8.cpp
rename to python/sglang/kernels/aot/csrc/cpu/moe_fp8.cpp
diff --git a/sgl-kernel/csrc/cpu/moe_int4.cpp b/python/sglang/kernels/aot/csrc/cpu/moe_int4.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/moe_int4.cpp
rename to python/sglang/kernels/aot/csrc/cpu/moe_int4.cpp
diff --git a/sgl-kernel/csrc/cpu/moe_int8.cpp b/python/sglang/kernels/aot/csrc/cpu/moe_int8.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/moe_int8.cpp
rename to python/sglang/kernels/aot/csrc/cpu/moe_int8.cpp
diff --git a/sgl-kernel/csrc/cpu/norm.cpp b/python/sglang/kernels/aot/csrc/cpu/norm.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/norm.cpp
rename to python/sglang/kernels/aot/csrc/cpu/norm.cpp
diff --git a/sgl-kernel/csrc/cpu/numa_utils.cpp b/python/sglang/kernels/aot/csrc/cpu/numa_utils.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/numa_utils.cpp
rename to python/sglang/kernels/aot/csrc/cpu/numa_utils.cpp
diff --git a/sgl-kernel/csrc/cpu/preprocessor.cpp b/python/sglang/kernels/aot/csrc/cpu/preprocessor.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/preprocessor.cpp
rename to python/sglang/kernels/aot/csrc/cpu/preprocessor.cpp
diff --git a/sgl-kernel/csrc/cpu/qkv_proj.cpp b/python/sglang/kernels/aot/csrc/cpu/qkv_proj.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/qkv_proj.cpp
rename to python/sglang/kernels/aot/csrc/cpu/qkv_proj.cpp
diff --git a/sgl-kernel/csrc/cpu/rope.cpp b/python/sglang/kernels/aot/csrc/cpu/rope.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/rope.cpp
rename to python/sglang/kernels/aot/csrc/cpu/rope.cpp
diff --git a/sgl-kernel/csrc/cpu/shm.cpp b/python/sglang/kernels/aot/csrc/cpu/shm.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/shm.cpp
rename to python/sglang/kernels/aot/csrc/cpu/shm.cpp
diff --git a/sgl-kernel/csrc/cpu/shm.h b/python/sglang/kernels/aot/csrc/cpu/shm.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/shm.h
rename to python/sglang/kernels/aot/csrc/cpu/shm.h
diff --git a/sgl-kernel/csrc/cpu/spec.cpp b/python/sglang/kernels/aot/csrc/cpu/spec.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/spec.cpp
rename to python/sglang/kernels/aot/csrc/cpu/spec.cpp
diff --git a/sgl-kernel/csrc/cpu/topk.cpp b/python/sglang/kernels/aot/csrc/cpu/topk.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/topk.cpp
rename to python/sglang/kernels/aot/csrc/cpu/topk.cpp
diff --git a/sgl-kernel/csrc/cpu/torch_extension_cpu.cpp b/python/sglang/kernels/aot/csrc/cpu/torch_extension_cpu.cpp
similarity index 100%
rename from sgl-kernel/csrc/cpu/torch_extension_cpu.cpp
rename to python/sglang/kernels/aot/csrc/cpu/torch_extension_cpu.cpp
diff --git a/sgl-kernel/csrc/cpu/vec.h b/python/sglang/kernels/aot/csrc/cpu/vec.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/vec.h
rename to python/sglang/kernels/aot/csrc/cpu/vec.h
diff --git a/sgl-kernel/csrc/cpu/vec_pack.h b/python/sglang/kernels/aot/csrc/cpu/vec_pack.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/vec_pack.h
rename to python/sglang/kernels/aot/csrc/cpu/vec_pack.h
diff --git a/sgl-kernel/csrc/cpu/x86_64/shm.h b/python/sglang/kernels/aot/csrc/cpu/x86_64/shm.h
similarity index 100%
rename from sgl-kernel/csrc/cpu/x86_64/shm.h
rename to python/sglang/kernels/aot/csrc/cpu/x86_64/shm.h
diff --git a/sgl-kernel/csrc/cutlass_extensions/common.hpp b/python/sglang/kernels/aot/csrc/cutlass_extensions/common.hpp
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/common.hpp
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/common.hpp
diff --git a/sgl-kernel/csrc/cutlass_extensions/detail/collective/mixed_input_utils.hpp b/python/sglang/kernels/aot/csrc/cutlass_extensions/detail/collective/mixed_input_utils.hpp
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/detail/collective/mixed_input_utils.hpp
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/detail/collective/mixed_input_utils.hpp
diff --git a/sgl-kernel/csrc/cutlass_extensions/epilogue/broadcast_load_epilogue_c3x.hpp b/python/sglang/kernels/aot/csrc/cutlass_extensions/epilogue/broadcast_load_epilogue_c3x.hpp
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/epilogue/broadcast_load_epilogue_c3x.hpp
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/epilogue/broadcast_load_epilogue_c3x.hpp
diff --git a/sgl-kernel/csrc/cutlass_extensions/epilogue/epilogue_per_row_per_col_scale.h b/python/sglang/kernels/aot/csrc/cutlass_extensions/epilogue/epilogue_per_row_per_col_scale.h
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/epilogue/epilogue_per_row_per_col_scale.h
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/epilogue/epilogue_per_row_per_col_scale.h
diff --git a/sgl-kernel/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp b/python/sglang/kernels/aot/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/collective/builders/sm90_gmma_builder_mixed_input.inl b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/collective/builders/sm90_gmma_builder_mixed_input.inl
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/collective/builders/sm90_gmma_builder_mixed_input.inl
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/collective/builders/sm90_gmma_builder_mixed_input.inl
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/collective/collective_builder_mixed_input.hpp b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/collective/collective_builder_mixed_input.hpp
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/collective/collective_builder_mixed_input.hpp
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/collective/collective_builder_mixed_input.hpp
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/collective/collective_mma_array_mixed_input.hpp b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/collective/collective_mma_array_mixed_input.hpp
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/collective/collective_mma_array_mixed_input.hpp
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/collective/collective_mma_array_mixed_input.hpp
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/cutlass_gemm_caller.cuh b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/cutlass_gemm_caller.cuh
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/cutlass_gemm_caller.cuh
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/cutlass_gemm_caller.cuh
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/dispatch_policy.hpp b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/dispatch_policy.hpp
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/dispatch_policy.hpp
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/dispatch_policy.hpp
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/fp8_gemm_sm90_dispatch.cuh b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/fp8_gemm_sm90_dispatch.cuh
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/fp8_gemm_sm90_dispatch.cuh
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/fp8_gemm_sm90_dispatch.cuh
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/gemm_universal_base_compat.h b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/gemm_universal_base_compat.h
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/gemm_universal_base_compat.h
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/gemm_universal_base_compat.h
diff --git a/sgl-kernel/csrc/cutlass_extensions/gemm/gemm_with_epilogue_visitor.h b/python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/gemm_with_epilogue_visitor.h
similarity index 100%
rename from sgl-kernel/csrc/cutlass_extensions/gemm/gemm_with_epilogue_visitor.h
rename to python/sglang/kernels/aot/csrc/cutlass_extensions/gemm/gemm_with_epilogue_visitor.h
diff --git a/sgl-kernel/csrc/elementwise/activation.cu b/python/sglang/kernels/aot/csrc/elementwise/activation.cu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/activation.cu
rename to python/sglang/kernels/aot/csrc/elementwise/activation.cu
diff --git a/sgl-kernel/csrc/elementwise/concat_mla.cu b/python/sglang/kernels/aot/csrc/elementwise/concat_mla.cu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/concat_mla.cu
rename to python/sglang/kernels/aot/csrc/elementwise/concat_mla.cu
diff --git a/sgl-kernel/csrc/elementwise/copy.cu b/python/sglang/kernels/aot/csrc/elementwise/copy.cu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/copy.cu
rename to python/sglang/kernels/aot/csrc/elementwise/copy.cu
diff --git a/sgl-kernel/csrc/elementwise/deepseek_v4_topk.cu b/python/sglang/kernels/aot/csrc/elementwise/deepseek_v4_topk.cu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/deepseek_v4_topk.cu
rename to python/sglang/kernels/aot/csrc/elementwise/deepseek_v4_topk.cu
diff --git a/sgl-kernel/csrc/elementwise/dsv4_norm_rope.cu b/python/sglang/kernels/aot/csrc/elementwise/dsv4_norm_rope.cu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/dsv4_norm_rope.cu
rename to python/sglang/kernels/aot/csrc/elementwise/dsv4_norm_rope.cu
diff --git a/sgl-kernel/csrc/elementwise/fused_add_rms_norm_kernel.cu b/python/sglang/kernels/aot/csrc/elementwise/fused_add_rms_norm_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/fused_add_rms_norm_kernel.cu
rename to python/sglang/kernels/aot/csrc/elementwise/fused_add_rms_norm_kernel.cu
diff --git a/sgl-kernel/csrc/elementwise/fused_add_rms_norm_kernel.mu b/python/sglang/kernels/aot/csrc/elementwise/fused_add_rms_norm_kernel.mu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/fused_add_rms_norm_kernel.mu
rename to python/sglang/kernels/aot/csrc/elementwise/fused_add_rms_norm_kernel.mu
diff --git a/sgl-kernel/csrc/elementwise/pos_enc.cu b/python/sglang/kernels/aot/csrc/elementwise/pos_enc.cu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/pos_enc.cu
rename to python/sglang/kernels/aot/csrc/elementwise/pos_enc.cu
diff --git a/sgl-kernel/csrc/elementwise/pos_enc.cuh b/python/sglang/kernels/aot/csrc/elementwise/pos_enc.cuh
similarity index 100%
rename from sgl-kernel/csrc/elementwise/pos_enc.cuh
rename to python/sglang/kernels/aot/csrc/elementwise/pos_enc.cuh
diff --git a/sgl-kernel/csrc/elementwise/topk.cu b/python/sglang/kernels/aot/csrc/elementwise/topk.cu
similarity index 100%
rename from sgl-kernel/csrc/elementwise/topk.cu
rename to python/sglang/kernels/aot/csrc/elementwise/topk.cu
diff --git a/sgl-kernel/csrc/elementwise/utils.cuh b/python/sglang/kernels/aot/csrc/elementwise/utils.cuh
similarity index 100%
rename from sgl-kernel/csrc/elementwise/utils.cuh
rename to python/sglang/kernels/aot/csrc/elementwise/utils.cuh
diff --git a/sgl-kernel/csrc/expert_specialization/es_fp8_blockwise.cu b/python/sglang/kernels/aot/csrc/expert_specialization/es_fp8_blockwise.cu
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_fp8_blockwise.cu
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_fp8_blockwise.cu
diff --git a/sgl-kernel/csrc/expert_specialization/es_fp8_blockwise_functor.cuh b/python/sglang/kernels/aot/csrc/expert_specialization/es_fp8_blockwise_functor.cuh
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_fp8_blockwise_functor.cuh
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_fp8_blockwise_functor.cuh
diff --git a/sgl-kernel/csrc/expert_specialization/es_fp8_blockwise_launcher.cuh b/python/sglang/kernels/aot/csrc/expert_specialization/es_fp8_blockwise_launcher.cuh
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_fp8_blockwise_launcher.cuh
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_fp8_blockwise_launcher.cuh
diff --git a/sgl-kernel/csrc/expert_specialization/es_fp8_blockwise_traits.cuh b/python/sglang/kernels/aot/csrc/expert_specialization/es_fp8_blockwise_traits.cuh
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_fp8_blockwise_traits.cuh
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_fp8_blockwise_traits.cuh
diff --git a/sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled.cu b/python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled.cu
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled.cu
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled.cu
diff --git a/sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_functor.cuh b/python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_functor.cuh
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_functor.cuh
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_functor.cuh
diff --git a/sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cu b/python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cu
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cu
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cu
diff --git a/sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cuh b/python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cuh
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cuh
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cuh
diff --git a/sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_launcher.cuh b/python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_launcher.cuh
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_launcher.cuh
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_launcher.cuh
diff --git a/sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_traits.cuh b/python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_traits.cuh
similarity index 100%
rename from sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_traits.cuh
rename to python/sglang/kernels/aot/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_traits.cuh
diff --git a/sgl-kernel/csrc/flash_extension.cc b/python/sglang/kernels/aot/csrc/flash_extension.cc
similarity index 100%
rename from sgl-kernel/csrc/flash_extension.cc
rename to python/sglang/kernels/aot/csrc/flash_extension.cc
diff --git a/sgl-kernel/csrc/flashmla_extension.cc b/python/sglang/kernels/aot/csrc/flashmla_extension.cc
similarity index 100%
rename from sgl-kernel/csrc/flashmla_extension.cc
rename to python/sglang/kernels/aot/csrc/flashmla_extension.cc
diff --git a/sgl-kernel/csrc/gemm/awq_kernel.cu b/python/sglang/kernels/aot/csrc/gemm/awq_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/gemm/awq_kernel.cu
rename to python/sglang/kernels/aot/csrc/gemm/awq_kernel.cu
diff --git a/sgl-kernel/csrc/gemm/fp8_gemm_kernel.cu b/python/sglang/kernels/aot/csrc/gemm/fp8_gemm_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/gemm/fp8_gemm_kernel.cu
rename to python/sglang/kernels/aot/csrc/gemm/fp8_gemm_kernel.cu
diff --git a/sgl-kernel/csrc/gemm/gptq/compat.cuh b/python/sglang/kernels/aot/csrc/gemm/gptq/compat.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/gptq/compat.cuh
rename to python/sglang/kernels/aot/csrc/gemm/gptq/compat.cuh
diff --git a/sgl-kernel/csrc/gemm/gptq/gptq_kernel.cu b/python/sglang/kernels/aot/csrc/gemm/gptq/gptq_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/gemm/gptq/gptq_kernel.cu
rename to python/sglang/kernels/aot/csrc/gemm/gptq/gptq_kernel.cu
diff --git a/sgl-kernel/csrc/gemm/gptq/matrix_view.cuh b/python/sglang/kernels/aot/csrc/gemm/gptq/matrix_view.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/gptq/matrix_view.cuh
rename to python/sglang/kernels/aot/csrc/gemm/gptq/matrix_view.cuh
diff --git a/sgl-kernel/csrc/gemm/gptq/qdq_2.cuh b/python/sglang/kernels/aot/csrc/gemm/gptq/qdq_2.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/gptq/qdq_2.cuh
rename to python/sglang/kernels/aot/csrc/gemm/gptq/qdq_2.cuh
diff --git a/sgl-kernel/csrc/gemm/gptq/qdq_3.cuh b/python/sglang/kernels/aot/csrc/gemm/gptq/qdq_3.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/gptq/qdq_3.cuh
rename to python/sglang/kernels/aot/csrc/gemm/gptq/qdq_3.cuh
diff --git a/sgl-kernel/csrc/gemm/gptq/qdq_4.cuh b/python/sglang/kernels/aot/csrc/gemm/gptq/qdq_4.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/gptq/qdq_4.cuh
rename to python/sglang/kernels/aot/csrc/gemm/gptq/qdq_4.cuh
diff --git a/sgl-kernel/csrc/gemm/gptq/qdq_8.cuh b/python/sglang/kernels/aot/csrc/gemm/gptq/qdq_8.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/gptq/qdq_8.cuh
rename to python/sglang/kernels/aot/csrc/gemm/gptq/qdq_8.cuh
diff --git a/sgl-kernel/csrc/gemm/gptq/qdq_util.cuh b/python/sglang/kernels/aot/csrc/gemm/gptq/qdq_util.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/gptq/qdq_util.cuh
rename to python/sglang/kernels/aot/csrc/gemm/gptq/qdq_util.cuh
diff --git a/sgl-kernel/csrc/gemm/int8_gemm_kernel.cu b/python/sglang/kernels/aot/csrc/gemm/int8_gemm_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/gemm/int8_gemm_kernel.cu
rename to python/sglang/kernels/aot/csrc/gemm/int8_gemm_kernel.cu
diff --git a/sgl-kernel/csrc/gemm/marlin/dequant.h b/python/sglang/kernels/aot/csrc/gemm/marlin/dequant.h
similarity index 100%
rename from sgl-kernel/csrc/gemm/marlin/dequant.h
rename to python/sglang/kernels/aot/csrc/gemm/marlin/dequant.h
diff --git a/sgl-kernel/csrc/gemm/marlin/kernel.h b/python/sglang/kernels/aot/csrc/gemm/marlin/kernel.h
similarity index 100%
rename from sgl-kernel/csrc/gemm/marlin/kernel.h
rename to python/sglang/kernels/aot/csrc/gemm/marlin/kernel.h
diff --git a/sgl-kernel/csrc/gemm/marlin/marlin.cuh b/python/sglang/kernels/aot/csrc/gemm/marlin/marlin.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/marlin/marlin.cuh
rename to python/sglang/kernels/aot/csrc/gemm/marlin/marlin.cuh
diff --git a/sgl-kernel/csrc/gemm/marlin/marlin_dtypes.cuh b/python/sglang/kernels/aot/csrc/gemm/marlin/marlin_dtypes.cuh
similarity index 100%
rename from sgl-kernel/csrc/gemm/marlin/marlin_dtypes.cuh
rename to python/sglang/kernels/aot/csrc/gemm/marlin/marlin_dtypes.cuh
diff --git a/sgl-kernel/csrc/gemm/marlin/marlin_template.h b/python/sglang/kernels/aot/csrc/gemm/marlin/marlin_template.h
similarity index 100%
rename from sgl-kernel/csrc/gemm/marlin/marlin_template.h
rename to python/sglang/kernels/aot/csrc/gemm/marlin/marlin_template.h
diff --git a/sgl-kernel/csrc/gemm/math.hpp b/python/sglang/kernels/aot/csrc/gemm/math.hpp
similarity index 100%
rename from sgl-kernel/csrc/gemm/math.hpp
rename to python/sglang/kernels/aot/csrc/gemm/math.hpp
diff --git a/sgl-kernel/csrc/gemm/per_token_group_quant_8bit.cu b/python/sglang/kernels/aot/csrc/gemm/per_token_group_quant_8bit.cu
similarity index 100%
rename from sgl-kernel/csrc/gemm/per_token_group_quant_8bit.cu
rename to python/sglang/kernels/aot/csrc/gemm/per_token_group_quant_8bit.cu
diff --git a/sgl-kernel/csrc/gemm/per_token_group_quant_8bit_v2.cu b/python/sglang/kernels/aot/csrc/gemm/per_token_group_quant_8bit_v2.cu
similarity index 100%
rename from sgl-kernel/csrc/gemm/per_token_group_quant_8bit_v2.cu
rename to python/sglang/kernels/aot/csrc/gemm/per_token_group_quant_8bit_v2.cu
diff --git a/sgl-kernel/csrc/gemm/per_token_quant_fp8.cu b/python/sglang/kernels/aot/csrc/gemm/per_token_quant_fp8.cu
similarity index 100%
rename from sgl-kernel/csrc/gemm/per_token_quant_fp8.cu
rename to python/sglang/kernels/aot/csrc/gemm/per_token_quant_fp8.cu
diff --git a/sgl-kernel/csrc/grammar/apply_token_bitmask_inplace_cuda.cu b/python/sglang/kernels/aot/csrc/grammar/apply_token_bitmask_inplace_cuda.cu
similarity index 100%
rename from sgl-kernel/csrc/grammar/apply_token_bitmask_inplace_cuda.cu
rename to python/sglang/kernels/aot/csrc/grammar/apply_token_bitmask_inplace_cuda.cu
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/flash_api.cpp b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/flash_api.cpp
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/flash_api.cpp
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/flash_api.cpp
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/block_info.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/block_info.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/block_info.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/block_info.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/dropout.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/dropout.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/dropout.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/dropout.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/flash.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/flash.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_blockmask.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_blockmask.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_blockmask.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_blockmask.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_kernel.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_kernel.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_kernel.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_kernel.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_launch_template.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_launch_template.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_launch_template.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_launch_template.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim128_bf16_causal_sm80.cu b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim128_bf16_causal_sm80.cu
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim128_bf16_causal_sm80.cu
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim128_bf16_causal_sm80.cu
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim128_bf16_sm80.cu b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim128_bf16_sm80.cu
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim128_bf16_sm80.cu
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim128_bf16_sm80.cu
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim64_bf16_causal_sm80.cu b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim64_bf16_causal_sm80.cu
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim64_bf16_causal_sm80.cu
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim64_bf16_causal_sm80.cu
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim64_bf16_sm80.cu b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim64_bf16_sm80.cu
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim64_bf16_sm80.cu
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/flash_fwd_split_hdim64_bf16_sm80.cu
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/hardware_info.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/hardware_info.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/hardware_info.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/hardware_info.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/kernel_traits.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/kernel_traits.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/kernel_traits.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/kernel_traits.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/mask.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/mask.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/mask.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/mask.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/philox.cuh b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/philox.cuh
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/philox.cuh
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/philox.cuh
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/philox_unpack.cuh b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/philox_unpack.cuh
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/philox_unpack.cuh
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/philox_unpack.cuh
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/rotary.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/rotary.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/rotary.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/rotary.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/softmax.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/softmax.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/softmax.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/softmax.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/static_switch.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/static_switch.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/static_switch.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/static_switch.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_attn/src/utils.h b/python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/utils.h
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_attn/src/utils.h
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_attn/src/utils.h
diff --git a/sgl-kernel/csrc/infllm_v2/flash_extension.cc b/python/sglang/kernels/aot/csrc/infllm_v2/flash_extension.cc
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/flash_extension.cc
rename to python/sglang/kernels/aot/csrc/infllm_v2/flash_extension.cc
diff --git a/sgl-kernel/csrc/infllm_v2/max_pooling.cu b/python/sglang/kernels/aot/csrc/infllm_v2/max_pooling.cu
similarity index 100%
rename from sgl-kernel/csrc/infllm_v2/max_pooling.cu
rename to python/sglang/kernels/aot/csrc/infllm_v2/max_pooling.cu
diff --git a/sgl-kernel/csrc/kvcacheio/transfer.cu b/python/sglang/kernels/aot/csrc/kvcacheio/transfer.cu
similarity index 100%
rename from sgl-kernel/csrc/kvcacheio/transfer.cu
rename to python/sglang/kernels/aot/csrc/kvcacheio/transfer.cu
diff --git a/sgl-kernel/csrc/mamba/causal_conv1d.cu b/python/sglang/kernels/aot/csrc/mamba/causal_conv1d.cu
similarity index 100%
rename from sgl-kernel/csrc/mamba/causal_conv1d.cu
rename to python/sglang/kernels/aot/csrc/mamba/causal_conv1d.cu
diff --git a/sgl-kernel/csrc/mamba/causal_conv1d.h b/python/sglang/kernels/aot/csrc/mamba/causal_conv1d.h
similarity index 100%
rename from sgl-kernel/csrc/mamba/causal_conv1d.h
rename to python/sglang/kernels/aot/csrc/mamba/causal_conv1d.h
diff --git a/sgl-kernel/csrc/memory/weak_ref_tensor.cpp b/python/sglang/kernels/aot/csrc/memory/weak_ref_tensor.cpp
similarity index 100%
rename from sgl-kernel/csrc/memory/weak_ref_tensor.cpp
rename to python/sglang/kernels/aot/csrc/memory/weak_ref_tensor.cpp
diff --git a/sgl-kernel/csrc/metal/README.md b/python/sglang/kernels/aot/csrc/metal/README.md
similarity index 100%
rename from sgl-kernel/csrc/metal/README.md
rename to python/sglang/kernels/aot/csrc/metal/README.md
diff --git a/sgl-kernel/csrc/metal/rope_pool_fused.cpp b/python/sglang/kernels/aot/csrc/metal/rope_pool_fused.cpp
similarity index 100%
rename from sgl-kernel/csrc/metal/rope_pool_fused.cpp
rename to python/sglang/kernels/aot/csrc/metal/rope_pool_fused.cpp
diff --git a/sgl-kernel/csrc/metal/rope_pool_fused.metal b/python/sglang/kernels/aot/csrc/metal/rope_pool_fused.metal
similarity index 100%
rename from sgl-kernel/csrc/metal/rope_pool_fused.metal
rename to python/sglang/kernels/aot/csrc/metal/rope_pool_fused.metal
diff --git a/sgl-kernel/csrc/moe/cutlass_moe/w4a8/scaled_mm_entry.cu b/python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/scaled_mm_entry.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/cutlass_moe/w4a8/scaled_mm_entry.cu
rename to python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/scaled_mm_entry.cu
diff --git a/sgl-kernel/csrc/moe/cutlass_moe/w4a8/w4a8_get_group_starts.cuh b/python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/w4a8_get_group_starts.cuh
similarity index 100%
rename from sgl-kernel/csrc/moe/cutlass_moe/w4a8/w4a8_get_group_starts.cuh
rename to python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/w4a8_get_group_starts.cuh
diff --git a/sgl-kernel/csrc/moe/cutlass_moe/w4a8/w4a8_grouped_mm_c3x.cu b/python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/w4a8_grouped_mm_c3x.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/cutlass_moe/w4a8/w4a8_grouped_mm_c3x.cu
rename to python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/w4a8_grouped_mm_c3x.cu
diff --git a/sgl-kernel/csrc/moe/cutlass_moe/w4a8/w4a8_grouped_mm_c3x.cuh b/python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/w4a8_grouped_mm_c3x.cuh
similarity index 100%
rename from sgl-kernel/csrc/moe/cutlass_moe/w4a8/w4a8_grouped_mm_c3x.cuh
rename to python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/w4a8_grouped_mm_c3x.cuh
diff --git a/sgl-kernel/csrc/moe/cutlass_moe/w4a8/w4a8_moe_data.cu b/python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/w4a8_moe_data.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/cutlass_moe/w4a8/w4a8_moe_data.cu
rename to python/sglang/kernels/aot/csrc/moe/cutlass_moe/w4a8/w4a8_moe_data.cu
diff --git a/sgl-kernel/csrc/moe/cutlass_moe_helper.cu b/python/sglang/kernels/aot/csrc/moe/cutlass_moe_helper.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/cutlass_moe_helper.cu
rename to python/sglang/kernels/aot/csrc/moe/cutlass_moe_helper.cu
diff --git a/sgl-kernel/csrc/moe/fp8_blockwise_moe_kernel.cu b/python/sglang/kernels/aot/csrc/moe/fp8_blockwise_moe_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/fp8_blockwise_moe_kernel.cu
rename to python/sglang/kernels/aot/csrc/moe/fp8_blockwise_moe_kernel.cu
diff --git a/sgl-kernel/csrc/moe/fused_qknorm_rope_kernel.cu b/python/sglang/kernels/aot/csrc/moe/fused_qknorm_rope_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/fused_qknorm_rope_kernel.cu
rename to python/sglang/kernels/aot/csrc/moe/fused_qknorm_rope_kernel.cu
diff --git a/sgl-kernel/csrc/moe/moe_align_kernel.cu b/python/sglang/kernels/aot/csrc/moe/moe_align_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/moe_align_kernel.cu
rename to python/sglang/kernels/aot/csrc/moe/moe_align_kernel.cu
diff --git a/sgl-kernel/csrc/moe/moe_sum.cu b/python/sglang/kernels/aot/csrc/moe/moe_sum.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/moe_sum.cu
rename to python/sglang/kernels/aot/csrc/moe/moe_sum.cu
diff --git a/sgl-kernel/csrc/moe/moe_sum_reduce.cu b/python/sglang/kernels/aot/csrc/moe/moe_sum_reduce.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/moe_sum_reduce.cu
rename to python/sglang/kernels/aot/csrc/moe/moe_sum_reduce.cu
diff --git a/sgl-kernel/csrc/moe/moe_topk_sigmoid_kernels.cu b/python/sglang/kernels/aot/csrc/moe/moe_topk_sigmoid_kernels.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/moe_topk_sigmoid_kernels.cu
rename to python/sglang/kernels/aot/csrc/moe/moe_topk_sigmoid_kernels.cu
diff --git a/sgl-kernel/csrc/moe/moe_topk_softmax_kernels.cu b/python/sglang/kernels/aot/csrc/moe/moe_topk_softmax_kernels.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/moe_topk_softmax_kernels.cu
rename to python/sglang/kernels/aot/csrc/moe/moe_topk_softmax_kernels.cu
diff --git a/sgl-kernel/csrc/moe/prepare_moe_input.cu b/python/sglang/kernels/aot/csrc/moe/prepare_moe_input.cu
similarity index 100%
rename from sgl-kernel/csrc/moe/prepare_moe_input.cu
rename to python/sglang/kernels/aot/csrc/moe/prepare_moe_input.cu
diff --git a/sgl-kernel/csrc/musa/common.muh b/python/sglang/kernels/aot/csrc/musa/common.muh
similarity index 100%
rename from sgl-kernel/csrc/musa/common.muh
rename to python/sglang/kernels/aot/csrc/musa/common.muh
diff --git a/sgl-kernel/csrc/musa/dtype.muh b/python/sglang/kernels/aot/csrc/musa/dtype.muh
similarity index 100%
rename from sgl-kernel/csrc/musa/dtype.muh
rename to python/sglang/kernels/aot/csrc/musa/dtype.muh
diff --git a/sgl-kernel/csrc/musa/moe_gemv_swiglu.mu b/python/sglang/kernels/aot/csrc/musa/moe_gemv_swiglu.mu
similarity index 100%
rename from sgl-kernel/csrc/musa/moe_gemv_swiglu.mu
rename to python/sglang/kernels/aot/csrc/musa/moe_gemv_swiglu.mu
diff --git a/sgl-kernel/csrc/musa/pos_encoding_contiguous.mu b/python/sglang/kernels/aot/csrc/musa/pos_encoding_contiguous.mu
similarity index 100%
rename from sgl-kernel/csrc/musa/pos_encoding_contiguous.mu
rename to python/sglang/kernels/aot/csrc/musa/pos_encoding_contiguous.mu
diff --git a/sgl-kernel/csrc/musa/ternary.mu b/python/sglang/kernels/aot/csrc/musa/ternary.mu
similarity index 100%
rename from sgl-kernel/csrc/musa/ternary.mu
rename to python/sglang/kernels/aot/csrc/musa/ternary.mu
diff --git a/sgl-kernel/csrc/musa/top_k_top_p_sampling.mu b/python/sglang/kernels/aot/csrc/musa/top_k_top_p_sampling.mu
similarity index 100%
rename from sgl-kernel/csrc/musa/top_k_top_p_sampling.mu
rename to python/sglang/kernels/aot/csrc/musa/top_k_top_p_sampling.mu
diff --git a/sgl-kernel/csrc/quantization/gguf/dequantize.cuh b/python/sglang/kernels/aot/csrc/quantization/gguf/dequantize.cuh
similarity index 100%
rename from sgl-kernel/csrc/quantization/gguf/dequantize.cuh
rename to python/sglang/kernels/aot/csrc/quantization/gguf/dequantize.cuh
diff --git a/sgl-kernel/csrc/quantization/gguf/ggml-common.h b/python/sglang/kernels/aot/csrc/quantization/gguf/ggml-common.h
similarity index 100%
rename from sgl-kernel/csrc/quantization/gguf/ggml-common.h
rename to python/sglang/kernels/aot/csrc/quantization/gguf/ggml-common.h
diff --git a/sgl-kernel/csrc/quantization/gguf/gguf_kernel.cu b/python/sglang/kernels/aot/csrc/quantization/gguf/gguf_kernel.cu
similarity index 100%
rename from sgl-kernel/csrc/quantization/gguf/gguf_kernel.cu
rename to python/sglang/kernels/aot/csrc/quantization/gguf/gguf_kernel.cu
diff --git a/sgl-kernel/csrc/quantization/gguf/mmq.cuh b/python/sglang/kernels/aot/csrc/quantization/gguf/mmq.cuh
similarity index 100%
rename from sgl-kernel/csrc/quantization/gguf/mmq.cuh
rename to python/sglang/kernels/aot/csrc/quantization/gguf/mmq.cuh
diff --git a/sgl-kernel/csrc/quantization/gguf/mmvq.cuh b/python/sglang/kernels/aot/csrc/quantization/gguf/mmvq.cuh
similarity index 100%
rename from sgl-kernel/csrc/quantization/gguf/mmvq.cuh
rename to python/sglang/kernels/aot/csrc/quantization/gguf/mmvq.cuh
diff --git a/sgl-kernel/csrc/quantization/gguf/moe.cuh b/python/sglang/kernels/aot/csrc/quantization/gguf/moe.cuh
similarity index 100%
rename from sgl-kernel/csrc/quantization/gguf/moe.cuh
rename to python/sglang/kernels/aot/csrc/quantization/gguf/moe.cuh
diff --git a/sgl-kernel/csrc/quantization/gguf/moe_vec.cuh b/python/sglang/kernels/aot/csrc/quantization/gguf/moe_vec.cuh
similarity index 100%
rename from sgl-kernel/csrc/quantization/gguf/moe_vec.cuh
rename to python/sglang/kernels/aot/csrc/quantization/gguf/moe_vec.cuh
diff --git a/sgl-kernel/csrc/quantization/gguf/vecdotq.cuh b/python/sglang/kernels/aot/csrc/quantization/gguf/vecdotq.cuh
similarity index 100%
rename from sgl-kernel/csrc/quantization/gguf/vecdotq.cuh
rename to python/sglang/kernels/aot/csrc/quantization/gguf/vecdotq.cuh
diff --git a/sgl-kernel/csrc/spatial/cuda_utils.h b/python/sglang/kernels/aot/csrc/spatial/cuda_utils.h
similarity index 100%
rename from sgl-kernel/csrc/spatial/cuda_utils.h
rename to python/sglang/kernels/aot/csrc/spatial/cuda_utils.h
diff --git a/sgl-kernel/csrc/spatial/greenctx_stream.cu b/python/sglang/kernels/aot/csrc/spatial/greenctx_stream.cu
similarity index 100%
rename from sgl-kernel/csrc/spatial/greenctx_stream.cu
rename to python/sglang/kernels/aot/csrc/spatial/greenctx_stream.cu
diff --git a/sgl-kernel/csrc/spatial/greenctx_stream.h b/python/sglang/kernels/aot/csrc/spatial/greenctx_stream.h
similarity index 100%
rename from sgl-kernel/csrc/spatial/greenctx_stream.h
rename to python/sglang/kernels/aot/csrc/spatial/greenctx_stream.h
diff --git a/sgl-kernel/csrc/spatial_extension.cc b/python/sglang/kernels/aot/csrc/spatial_extension.cc
similarity index 100%
rename from sgl-kernel/csrc/spatial_extension.cc
rename to python/sglang/kernels/aot/csrc/spatial_extension.cc
diff --git a/sgl-kernel/csrc/speculative/eagle_utils.cu b/python/sglang/kernels/aot/csrc/speculative/eagle_utils.cu
similarity index 100%
rename from sgl-kernel/csrc/speculative/eagle_utils.cu
rename to python/sglang/kernels/aot/csrc/speculative/eagle_utils.cu
diff --git a/sgl-kernel/csrc/speculative/ngram_utils.cu b/python/sglang/kernels/aot/csrc/speculative/ngram_utils.cu
similarity index 100%
rename from sgl-kernel/csrc/speculative/ngram_utils.cu
rename to python/sglang/kernels/aot/csrc/speculative/ngram_utils.cu
diff --git a/sgl-kernel/csrc/speculative/packbit.cu b/python/sglang/kernels/aot/csrc/speculative/packbit.cu
similarity index 100%
rename from sgl-kernel/csrc/speculative/packbit.cu
rename to python/sglang/kernels/aot/csrc/speculative/packbit.cu
diff --git a/sgl-kernel/csrc/speculative/speculative_sampling.cu b/python/sglang/kernels/aot/csrc/speculative/speculative_sampling.cu
similarity index 100%
rename from sgl-kernel/csrc/speculative/speculative_sampling.cu
rename to python/sglang/kernels/aot/csrc/speculative/speculative_sampling.cu
diff --git a/sgl-kernel/csrc/speculative/speculative_sampling.cuh b/python/sglang/kernels/aot/csrc/speculative/speculative_sampling.cuh
similarity index 100%
rename from sgl-kernel/csrc/speculative/speculative_sampling.cuh
rename to python/sglang/kernels/aot/csrc/speculative/speculative_sampling.cuh
diff --git a/sgl-kernel/include/hip/hip_act_and_mul.cuh b/python/sglang/kernels/aot/include/hip/hip_act_and_mul.cuh
similarity index 100%
rename from sgl-kernel/include/hip/hip_act_and_mul.cuh
rename to python/sglang/kernels/aot/include/hip/hip_act_and_mul.cuh
diff --git a/sgl-kernel/include/hip/hip_math_def.h b/python/sglang/kernels/aot/include/hip/hip_math_def.h
similarity index 100%
rename from sgl-kernel/include/hip/hip_math_def.h
rename to python/sglang/kernels/aot/include/hip/hip_math_def.h
diff --git a/sgl-kernel/include/hip/hip_vec_dtypes.h b/python/sglang/kernels/aot/include/hip/hip_vec_dtypes.h
similarity index 100%
rename from sgl-kernel/include/hip/hip_vec_dtypes.h
rename to python/sglang/kernels/aot/include/hip/hip_vec_dtypes.h
diff --git a/sgl-kernel/include/hip/impl/hip_vec_bf16_impl.h b/python/sglang/kernels/aot/include/hip/impl/hip_vec_bf16_impl.h
similarity index 100%
rename from sgl-kernel/include/hip/impl/hip_vec_bf16_impl.h
rename to python/sglang/kernels/aot/include/hip/impl/hip_vec_bf16_impl.h
diff --git a/sgl-kernel/include/hip/impl/hip_vec_fp32_impl.h b/python/sglang/kernels/aot/include/hip/impl/hip_vec_fp32_impl.h
similarity index 100%
rename from sgl-kernel/include/hip/impl/hip_vec_fp32_impl.h
rename to python/sglang/kernels/aot/include/hip/impl/hip_vec_fp32_impl.h
diff --git a/sgl-kernel/include/hip/impl/hip_vec_half_impl.h b/python/sglang/kernels/aot/include/hip/impl/hip_vec_half_impl.h
similarity index 100%
rename from sgl-kernel/include/hip/impl/hip_vec_half_impl.h
rename to python/sglang/kernels/aot/include/hip/impl/hip_vec_half_impl.h
diff --git a/sgl-kernel/include/musa/dispatch_utils.h b/python/sglang/kernels/aot/include/musa/dispatch_utils.h
similarity index 100%
rename from sgl-kernel/include/musa/dispatch_utils.h
rename to python/sglang/kernels/aot/include/musa/dispatch_utils.h
diff --git a/sgl-kernel/include/musa/integer_subbyte.h b/python/sglang/kernels/aot/include/musa/integer_subbyte.h
similarity index 100%
rename from sgl-kernel/include/musa/integer_subbyte.h
rename to python/sglang/kernels/aot/include/musa/integer_subbyte.h
diff --git a/sgl-kernel/include/pytorch_extension_utils_rocm.h b/python/sglang/kernels/aot/include/pytorch_extension_utils_rocm.h
similarity index 100%
rename from sgl-kernel/include/pytorch_extension_utils_rocm.h
rename to python/sglang/kernels/aot/include/pytorch_extension_utils_rocm.h
diff --git a/sgl-kernel/include/scalar_type.hpp b/python/sglang/kernels/aot/include/scalar_type.hpp
similarity index 100%
rename from sgl-kernel/include/scalar_type.hpp
rename to python/sglang/kernels/aot/include/scalar_type.hpp
diff --git a/sgl-kernel/include/sgl_flash_kernel_ops.h b/python/sglang/kernels/aot/include/sgl_flash_kernel_ops.h
similarity index 100%
rename from sgl-kernel/include/sgl_flash_kernel_ops.h
rename to python/sglang/kernels/aot/include/sgl_flash_kernel_ops.h
diff --git a/sgl-kernel/include/sgl_kernel_musa_ops.h b/python/sglang/kernels/aot/include/sgl_kernel_musa_ops.h
similarity index 100%
rename from sgl-kernel/include/sgl_kernel_musa_ops.h
rename to python/sglang/kernels/aot/include/sgl_kernel_musa_ops.h
diff --git a/sgl-kernel/include/sgl_kernel_ops.h b/python/sglang/kernels/aot/include/sgl_kernel_ops.h
similarity index 100%
rename from sgl-kernel/include/sgl_kernel_ops.h
rename to python/sglang/kernels/aot/include/sgl_kernel_ops.h
diff --git a/sgl-kernel/include/sgl_kernel_torch_shim.h b/python/sglang/kernels/aot/include/sgl_kernel_torch_shim.h
similarity index 100%
rename from sgl-kernel/include/sgl_kernel_torch_shim.h
rename to python/sglang/kernels/aot/include/sgl_kernel_torch_shim.h
diff --git a/sgl-kernel/include/utils.h b/python/sglang/kernels/aot/include/utils.h
similarity index 100%
rename from sgl-kernel/include/utils.h
rename to python/sglang/kernels/aot/include/utils.h
diff --git a/sgl-kernel/kernel-runner-setup.sh b/python/sglang/kernels/aot/kernel-runner-setup.sh
similarity index 100%
rename from sgl-kernel/kernel-runner-setup.sh
rename to python/sglang/kernels/aot/kernel-runner-setup.sh
diff --git a/sgl-kernel/pyproject.toml b/python/sglang/kernels/aot/pyproject.toml
similarity index 90%
rename from sgl-kernel/pyproject.toml
rename to python/sglang/kernels/aot/pyproject.toml
index aeded1e80..f64edb0af 100644
--- a/sgl-kernel/pyproject.toml
+++ b/python/sglang/kernels/aot/pyproject.toml
@@ -24,7 +24,7 @@ classifiers = [
 dependencies = []
 
 [project.urls]
-"Homepage" = "https://github.com/sgl-project/sglang/tree/main/sgl-kernel"
+"Homepage" = "https://github.com/sgl-project/sglang/tree/main/python/sglang/kernels/aot"
 "Bug Tracker" = "https://github.com/sgl-project/sglang/issues"
 
 [tool.wheel]
diff --git a/sgl-kernel/pyproject_cpu.toml b/python/sglang/kernels/aot/pyproject_cpu.toml
similarity index 89%
rename from sgl-kernel/pyproject_cpu.toml
rename to python/sglang/kernels/aot/pyproject_cpu.toml
index 241832f20..f928eb28b 100644
--- a/sgl-kernel/pyproject_cpu.toml
+++ b/python/sglang/kernels/aot/pyproject_cpu.toml
@@ -20,7 +20,7 @@ classifiers = [
 dependencies = []
 
 [project.urls]
-"Homepage" = "https://github.com/sgl-project/sglang/tree/main/sgl-kernel"
+"Homepage" = "https://github.com/sgl-project/sglang/tree/main/python/sglang/kernels/aot"
 "Bug Tracker" = "https://github.com/sgl-project/sglang/issues"
 
 [tool.wheel]
diff --git a/sgl-kernel/pyproject_musa.toml b/python/sglang/kernels/aot/pyproject_musa.toml
similarity index 87%
rename from sgl-kernel/pyproject_musa.toml
rename to python/sglang/kernels/aot/pyproject_musa.toml
index d7f4538ad..45dc6c065 100644
--- a/sgl-kernel/pyproject_musa.toml
+++ b/python/sglang/kernels/aot/pyproject_musa.toml
@@ -23,7 +23,7 @@ classifiers = [
 dependencies = []
 
 [project.urls]
-"Homepage" = "https://github.com/sgl-project/sglang/tree/main/sgl-kernel"
+"Homepage" = "https://github.com/sgl-project/sglang/tree/main/python/sglang/kernels/aot"
 "Bug Tracker" = "https://github.com/sgl-project/sglang/issues"
 
 [tool.wheel]
diff --git a/sgl-kernel/pyproject_rocm.toml b/python/sglang/kernels/aot/pyproject_rocm.toml
similarity index 87%
rename from sgl-kernel/pyproject_rocm.toml
rename to python/sglang/kernels/aot/pyproject_rocm.toml
index 4c1fda66a..f7e9ba6ed 100644
--- a/sgl-kernel/pyproject_rocm.toml
+++ b/python/sglang/kernels/aot/pyproject_rocm.toml
@@ -22,7 +22,7 @@ classifiers = [
 dependencies = []
 
 [project.urls]
-"Homepage" = "https://github.com/sgl-project/sglang/tree/main/sgl-kernel"
+"Homepage" = "https://github.com/sgl-project/sglang/tree/main/python/sglang/kernels/aot"
 "Bug Tracker" = "https://github.com/sgl-project/sglang/issues"
 
 [tool.wheel]
diff --git a/sgl-kernel/python/sgl_kernel/__init__.py b/python/sglang/kernels/aot/python/sgl_kernel/__init__.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/__init__.py
rename to python/sglang/kernels/aot/python/sgl_kernel/__init__.py
diff --git a/sgl-kernel/python/sgl_kernel/allreduce.py b/python/sglang/kernels/aot/python/sgl_kernel/allreduce.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/allreduce.py
rename to python/sglang/kernels/aot/python/sgl_kernel/allreduce.py
diff --git a/sgl-kernel/python/sgl_kernel/attention.py b/python/sglang/kernels/aot/python/sgl_kernel/attention.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/attention.py
rename to python/sglang/kernels/aot/python/sgl_kernel/attention.py
diff --git a/sgl-kernel/python/sgl_kernel/cutlass_moe.py b/python/sglang/kernels/aot/python/sgl_kernel/cutlass_moe.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/cutlass_moe.py
rename to python/sglang/kernels/aot/python/sgl_kernel/cutlass_moe.py
diff --git a/sgl-kernel/python/sgl_kernel/debug_utils.py b/python/sglang/kernels/aot/python/sgl_kernel/debug_utils.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/debug_utils.py
rename to python/sglang/kernels/aot/python/sgl_kernel/debug_utils.py
diff --git a/sgl-kernel/python/sgl_kernel/elementwise.py b/python/sglang/kernels/aot/python/sgl_kernel/elementwise.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/elementwise.py
rename to python/sglang/kernels/aot/python/sgl_kernel/elementwise.py
diff --git a/sgl-kernel/python/sgl_kernel/expert_specialization.py b/python/sglang/kernels/aot/python/sgl_kernel/expert_specialization.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/expert_specialization.py
rename to python/sglang/kernels/aot/python/sgl_kernel/expert_specialization.py
diff --git a/sgl-kernel/python/sgl_kernel/flash_attn.py b/python/sglang/kernels/aot/python/sgl_kernel/flash_attn.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/flash_attn.py
rename to python/sglang/kernels/aot/python/sgl_kernel/flash_attn.py
diff --git a/sgl-kernel/python/sgl_kernel/flash_mla.py b/python/sglang/kernels/aot/python/sgl_kernel/flash_mla.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/flash_mla.py
rename to python/sglang/kernels/aot/python/sgl_kernel/flash_mla.py
diff --git a/sgl-kernel/python/sgl_kernel/gemm.py b/python/sglang/kernels/aot/python/sgl_kernel/gemm.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/gemm.py
rename to python/sglang/kernels/aot/python/sgl_kernel/gemm.py
diff --git a/sgl-kernel/python/sgl_kernel/grammar.py b/python/sglang/kernels/aot/python/sgl_kernel/grammar.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/grammar.py
rename to python/sglang/kernels/aot/python/sgl_kernel/grammar.py
diff --git a/sgl-kernel/python/sgl_kernel/infllm_v2/__init__.py b/python/sglang/kernels/aot/python/sgl_kernel/infllm_v2/__init__.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/infllm_v2/__init__.py
rename to python/sglang/kernels/aot/python/sgl_kernel/infllm_v2/__init__.py
diff --git a/sgl-kernel/python/sgl_kernel/infllm_v2/_loader.py b/python/sglang/kernels/aot/python/sgl_kernel/infllm_v2/_loader.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/infllm_v2/_loader.py
rename to python/sglang/kernels/aot/python/sgl_kernel/infllm_v2/_loader.py
diff --git a/sgl-kernel/python/sgl_kernel/infllm_v2/attention.py b/python/sglang/kernels/aot/python/sgl_kernel/infllm_v2/attention.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/infllm_v2/attention.py
rename to python/sglang/kernels/aot/python/sgl_kernel/infllm_v2/attention.py
diff --git a/sgl-kernel/python/sgl_kernel/infllm_v2/max_pooling.py b/python/sglang/kernels/aot/python/sgl_kernel/infllm_v2/max_pooling.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/infllm_v2/max_pooling.py
rename to python/sglang/kernels/aot/python/sgl_kernel/infllm_v2/max_pooling.py
diff --git a/sgl-kernel/python/sgl_kernel/kvcacheio.py b/python/sglang/kernels/aot/python/sgl_kernel/kvcacheio.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/kvcacheio.py
rename to python/sglang/kernels/aot/python/sgl_kernel/kvcacheio.py
diff --git a/sgl-kernel/python/sgl_kernel/load_utils.py b/python/sglang/kernels/aot/python/sgl_kernel/load_utils.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/load_utils.py
rename to python/sglang/kernels/aot/python/sgl_kernel/load_utils.py
diff --git a/sgl-kernel/python/sgl_kernel/mamba.py b/python/sglang/kernels/aot/python/sgl_kernel/mamba.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/mamba.py
rename to python/sglang/kernels/aot/python/sgl_kernel/mamba.py
diff --git a/sgl-kernel/python/sgl_kernel/memory.py b/python/sglang/kernels/aot/python/sgl_kernel/memory.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/memory.py
rename to python/sglang/kernels/aot/python/sgl_kernel/memory.py
diff --git a/sgl-kernel/python/sgl_kernel/metal.py b/python/sglang/kernels/aot/python/sgl_kernel/metal.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/metal.py
rename to python/sglang/kernels/aot/python/sgl_kernel/metal.py
diff --git a/sgl-kernel/python/sgl_kernel/moe.py b/python/sglang/kernels/aot/python/sgl_kernel/moe.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/moe.py
rename to python/sglang/kernels/aot/python/sgl_kernel/moe.py
diff --git a/sgl-kernel/python/sgl_kernel/musa.py b/python/sglang/kernels/aot/python/sgl_kernel/musa.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/musa.py
rename to python/sglang/kernels/aot/python/sgl_kernel/musa.py
diff --git a/sgl-kernel/python/sgl_kernel/quantization/__init__.py b/python/sglang/kernels/aot/python/sgl_kernel/quantization/__init__.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/quantization/__init__.py
rename to python/sglang/kernels/aot/python/sgl_kernel/quantization/__init__.py
diff --git a/sgl-kernel/python/sgl_kernel/quantization/gguf.py b/python/sglang/kernels/aot/python/sgl_kernel/quantization/gguf.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/quantization/gguf.py
rename to python/sglang/kernels/aot/python/sgl_kernel/quantization/gguf.py
diff --git a/sgl-kernel/python/sgl_kernel/sampling.py b/python/sglang/kernels/aot/python/sgl_kernel/sampling.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/sampling.py
rename to python/sglang/kernels/aot/python/sgl_kernel/sampling.py
diff --git a/sgl-kernel/python/sgl_kernel/scalar_type.py b/python/sglang/kernels/aot/python/sgl_kernel/scalar_type.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/scalar_type.py
rename to python/sglang/kernels/aot/python/sgl_kernel/scalar_type.py
diff --git a/sgl-kernel/python/sgl_kernel/sparse_flash_attn.py b/python/sglang/kernels/aot/python/sgl_kernel/sparse_flash_attn.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/sparse_flash_attn.py
rename to python/sglang/kernels/aot/python/sgl_kernel/sparse_flash_attn.py
diff --git a/sgl-kernel/python/sgl_kernel/spatial.py b/python/sglang/kernels/aot/python/sgl_kernel/spatial.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/spatial.py
rename to python/sglang/kernels/aot/python/sgl_kernel/spatial.py
diff --git a/sgl-kernel/python/sgl_kernel/speculative.py b/python/sglang/kernels/aot/python/sgl_kernel/speculative.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/speculative.py
rename to python/sglang/kernels/aot/python/sgl_kernel/speculative.py
diff --git a/sgl-kernel/python/sgl_kernel/test_utils.py b/python/sglang/kernels/aot/python/sgl_kernel/test_utils.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/test_utils.py
rename to python/sglang/kernels/aot/python/sgl_kernel/test_utils.py
diff --git a/sgl-kernel/python/sgl_kernel/testing/__init__.py b/python/sglang/kernels/aot/python/sgl_kernel/testing/__init__.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/testing/__init__.py
rename to python/sglang/kernels/aot/python/sgl_kernel/testing/__init__.py
diff --git a/sgl-kernel/python/sgl_kernel/testing/rotary_embedding.py b/python/sglang/kernels/aot/python/sgl_kernel/testing/rotary_embedding.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/testing/rotary_embedding.py
rename to python/sglang/kernels/aot/python/sgl_kernel/testing/rotary_embedding.py
diff --git a/sgl-kernel/python/sgl_kernel/top_k.py b/python/sglang/kernels/aot/python/sgl_kernel/top_k.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/top_k.py
rename to python/sglang/kernels/aot/python/sgl_kernel/top_k.py
diff --git a/sgl-kernel/python/sgl_kernel/utils.py b/python/sglang/kernels/aot/python/sgl_kernel/utils.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/utils.py
rename to python/sglang/kernels/aot/python/sgl_kernel/utils.py
diff --git a/sgl-kernel/python/sgl_kernel/version.py b/python/sglang/kernels/aot/python/sgl_kernel/version.py
similarity index 100%
rename from sgl-kernel/python/sgl_kernel/version.py
rename to python/sglang/kernels/aot/python/sgl_kernel/version.py
diff --git a/sgl-kernel/rename_wheels.sh b/python/sglang/kernels/aot/rename_wheels.sh
similarity index 100%
rename from sgl-kernel/rename_wheels.sh
rename to python/sglang/kernels/aot/rename_wheels.sh
diff --git a/sgl-kernel/setup_metal.py b/python/sglang/kernels/aot/setup_metal.py
similarity index 100%
rename from sgl-kernel/setup_metal.py
rename to python/sglang/kernels/aot/setup_metal.py
diff --git a/sgl-kernel/setup_musa.py b/python/sglang/kernels/aot/setup_musa.py
similarity index 100%
rename from sgl-kernel/setup_musa.py
rename to python/sglang/kernels/aot/setup_musa.py
diff --git a/sgl-kernel/setup_rocm.py b/python/sglang/kernels/aot/setup_rocm.py
similarity index 100%
rename from sgl-kernel/setup_rocm.py
rename to python/sglang/kernels/aot/setup_rocm.py
diff --git a/sgl-kernel/tests/conftest.py b/python/sglang/kernels/aot/tests/conftest.py
similarity index 100%
rename from sgl-kernel/tests/conftest.py
rename to python/sglang/kernels/aot/tests/conftest.py
diff --git a/sgl-kernel/tests/spatial/test_greenctx_stream.py b/python/sglang/kernels/aot/tests/spatial/test_greenctx_stream.py
similarity index 100%
rename from sgl-kernel/tests/spatial/test_greenctx_stream.py
rename to python/sglang/kernels/aot/tests/spatial/test_greenctx_stream.py
diff --git a/sgl-kernel/tests/speculative/test_eagle_utils.py b/python/sglang/kernels/aot/tests/speculative/test_eagle_utils.py
similarity index 100%
rename from sgl-kernel/tests/speculative/test_eagle_utils.py
rename to python/sglang/kernels/aot/tests/speculative/test_eagle_utils.py
diff --git a/sgl-kernel/tests/speculative/test_ngram_utils.py b/python/sglang/kernels/aot/tests/speculative/test_ngram_utils.py
similarity index 100%
rename from sgl-kernel/tests/speculative/test_ngram_utils.py
rename to python/sglang/kernels/aot/tests/speculative/test_ngram_utils.py
diff --git a/sgl-kernel/tests/speculative/test_speculative_sampling.py b/python/sglang/kernels/aot/tests/speculative/test_speculative_sampling.py
similarity index 100%
rename from sgl-kernel/tests/speculative/test_speculative_sampling.py
rename to python/sglang/kernels/aot/tests/speculative/test_speculative_sampling.py
diff --git a/sgl-kernel/tests/test_activation.py b/python/sglang/kernels/aot/tests/test_activation.py
similarity index 100%
rename from sgl-kernel/tests/test_activation.py
rename to python/sglang/kernels/aot/tests/test_activation.py
diff --git a/sgl-kernel/tests/test_apply_token_bitmask_inplace.py b/python/sglang/kernels/aot/tests/test_apply_token_bitmask_inplace.py
similarity index 100%
rename from sgl-kernel/tests/test_apply_token_bitmask_inplace.py
rename to python/sglang/kernels/aot/tests/test_apply_token_bitmask_inplace.py
diff --git a/sgl-kernel/tests/test_awq_dequant.py b/python/sglang/kernels/aot/tests/test_awq_dequant.py
similarity index 100%
rename from sgl-kernel/tests/test_awq_dequant.py
rename to python/sglang/kernels/aot/tests/test_awq_dequant.py
diff --git a/sgl-kernel/tests/test_causal_conv1d.py b/python/sglang/kernels/aot/tests/test_causal_conv1d.py
similarity index 100%
rename from sgl-kernel/tests/test_causal_conv1d.py
rename to python/sglang/kernels/aot/tests/test_causal_conv1d.py
diff --git a/sgl-kernel/tests/test_copy.py b/python/sglang/kernels/aot/tests/test_copy.py
similarity index 100%
rename from sgl-kernel/tests/test_copy.py
rename to python/sglang/kernels/aot/tests/test_copy.py
diff --git a/sgl-kernel/tests/test_custom_allreduce.py b/python/sglang/kernels/aot/tests/test_custom_allreduce.py
similarity index 100%
rename from sgl-kernel/tests/test_custom_allreduce.py
rename to python/sglang/kernels/aot/tests/test_custom_allreduce.py
diff --git a/sgl-kernel/tests/test_cutlass_mla.py b/python/sglang/kernels/aot/tests/test_cutlass_mla.py
similarity index 100%
rename from sgl-kernel/tests/test_cutlass_mla.py
rename to python/sglang/kernels/aot/tests/test_cutlass_mla.py
diff --git a/sgl-kernel/tests/test_cutlass_w4a8_moe_mm.py b/python/sglang/kernels/aot/tests/test_cutlass_w4a8_moe_mm.py
similarity index 100%
rename from sgl-kernel/tests/test_cutlass_w4a8_moe_mm.py
rename to python/sglang/kernels/aot/tests/test_cutlass_w4a8_moe_mm.py
diff --git a/sgl-kernel/tests/test_dsv4_norm_rope.py b/python/sglang/kernels/aot/tests/test_dsv4_norm_rope.py
similarity index 100%
rename from sgl-kernel/tests/test_dsv4_norm_rope.py
rename to python/sglang/kernels/aot/tests/test_dsv4_norm_rope.py
diff --git a/sgl-kernel/tests/test_es_fp8_blockwise_moe.py b/python/sglang/kernels/aot/tests/test_es_fp8_blockwise_moe.py
similarity index 100%
rename from sgl-kernel/tests/test_es_fp8_blockwise_moe.py
rename to python/sglang/kernels/aot/tests/test_es_fp8_blockwise_moe.py
diff --git a/sgl-kernel/tests/test_es_mxfp8_blockscaled_moe.py b/python/sglang/kernels/aot/tests/test_es_mxfp8_blockscaled_moe.py
similarity index 100%
rename from sgl-kernel/tests/test_es_mxfp8_blockscaled_moe.py
rename to python/sglang/kernels/aot/tests/test_es_mxfp8_blockscaled_moe.py
diff --git a/sgl-kernel/tests/test_flash_attention.py b/python/sglang/kernels/aot/tests/test_flash_attention.py
similarity index 100%
rename from sgl-kernel/tests/test_flash_attention.py
rename to python/sglang/kernels/aot/tests/test_flash_attention.py
diff --git a/sgl-kernel/tests/test_flash_attn_sparse.py b/python/sglang/kernels/aot/tests/test_flash_attn_sparse.py
similarity index 100%
rename from sgl-kernel/tests/test_flash_attn_sparse.py
rename to python/sglang/kernels/aot/tests/test_flash_attn_sparse.py
diff --git a/sgl-kernel/tests/test_flashmla.py b/python/sglang/kernels/aot/tests/test_flashmla.py
similarity index 100%
rename from sgl-kernel/tests/test_flashmla.py
rename to python/sglang/kernels/aot/tests/test_flashmla.py
diff --git a/sgl-kernel/tests/test_fp8_blockwise_moe.py b/python/sglang/kernels/aot/tests/test_fp8_blockwise_moe.py
similarity index 100%
rename from sgl-kernel/tests/test_fp8_blockwise_moe.py
rename to python/sglang/kernels/aot/tests/test_fp8_blockwise_moe.py
diff --git a/sgl-kernel/tests/test_fp8_gemm.py b/python/sglang/kernels/aot/tests/test_fp8_gemm.py
similarity index 100%
rename from sgl-kernel/tests/test_fp8_gemm.py
rename to python/sglang/kernels/aot/tests/test_fp8_gemm.py
diff --git a/sgl-kernel/tests/test_fused_qk_norm_rope.py b/python/sglang/kernels/aot/tests/test_fused_qk_norm_rope.py
similarity index 100%
rename from sgl-kernel/tests/test_fused_qk_norm_rope.py
rename to python/sglang/kernels/aot/tests/test_fused_qk_norm_rope.py
diff --git a/sgl-kernel/tests/test_gguf.py b/python/sglang/kernels/aot/tests/test_gguf.py
similarity index 100%
rename from sgl-kernel/tests/test_gguf.py
rename to python/sglang/kernels/aot/tests/test_gguf.py
diff --git a/sgl-kernel/tests/test_gptq_kernel.py b/python/sglang/kernels/aot/tests/test_gptq_kernel.py
similarity index 100%
rename from sgl-kernel/tests/test_gptq_kernel.py
rename to python/sglang/kernels/aot/tests/test_gptq_kernel.py
diff --git a/sgl-kernel/tests/test_infllm_v2_attention.py b/python/sglang/kernels/aot/tests/test_infllm_v2_attention.py
similarity index 100%
rename from sgl-kernel/tests/test_infllm_v2_attention.py
rename to python/sglang/kernels/aot/tests/test_infllm_v2_attention.py
diff --git a/sgl-kernel/tests/test_infllm_v2_max_pooling.py b/python/sglang/kernels/aot/tests/test_infllm_v2_max_pooling.py
similarity index 100%
rename from sgl-kernel/tests/test_infllm_v2_max_pooling.py
rename to python/sglang/kernels/aot/tests/test_infllm_v2_max_pooling.py
diff --git a/sgl-kernel/tests/test_int8_gemm.py b/python/sglang/kernels/aot/tests/test_int8_gemm.py
similarity index 100%
rename from sgl-kernel/tests/test_int8_gemm.py
rename to python/sglang/kernels/aot/tests/test_int8_gemm.py
diff --git a/sgl-kernel/tests/test_kvcacheio.py b/python/sglang/kernels/aot/tests/test_kvcacheio.py
similarity index 100%
rename from sgl-kernel/tests/test_kvcacheio.py
rename to python/sglang/kernels/aot/tests/test_kvcacheio.py
diff --git a/sgl-kernel/tests/test_merge_state_v2.py b/python/sglang/kernels/aot/tests/test_merge_state_v2.py
similarity index 100%
rename from sgl-kernel/tests/test_merge_state_v2.py
rename to python/sglang/kernels/aot/tests/test_merge_state_v2.py
diff --git a/sgl-kernel/tests/test_moe_align.py b/python/sglang/kernels/aot/tests/test_moe_align.py
similarity index 100%
rename from sgl-kernel/tests/test_moe_align.py
rename to python/sglang/kernels/aot/tests/test_moe_align.py
diff --git a/sgl-kernel/tests/test_moe_topk_sigmoid.py b/python/sglang/kernels/aot/tests/test_moe_topk_sigmoid.py
similarity index 100%
rename from sgl-kernel/tests/test_moe_topk_sigmoid.py
rename to python/sglang/kernels/aot/tests/test_moe_topk_sigmoid.py
diff --git a/sgl-kernel/tests/test_moe_topk_softmax.py b/python/sglang/kernels/aot/tests/test_moe_topk_softmax.py
similarity index 100%
rename from sgl-kernel/tests/test_moe_topk_softmax.py
rename to python/sglang/kernels/aot/tests/test_moe_topk_softmax.py
diff --git a/sgl-kernel/tests/test_norm.py b/python/sglang/kernels/aot/tests/test_norm.py
similarity index 100%
rename from sgl-kernel/tests/test_norm.py
rename to python/sglang/kernels/aot/tests/test_norm.py
diff --git a/sgl-kernel/tests/test_per_token_group_quant_8bit.py b/python/sglang/kernels/aot/tests/test_per_token_group_quant_8bit.py
similarity index 100%
rename from sgl-kernel/tests/test_per_token_group_quant_8bit.py
rename to python/sglang/kernels/aot/tests/test_per_token_group_quant_8bit.py
diff --git a/sgl-kernel/tests/test_per_token_quant_fp8.py b/python/sglang/kernels/aot/tests/test_per_token_quant_fp8.py
similarity index 100%
rename from sgl-kernel/tests/test_per_token_quant_fp8.py
rename to python/sglang/kernels/aot/tests/test_per_token_quant_fp8.py
diff --git a/sgl-kernel/tests/test_sampling.py b/python/sglang/kernels/aot/tests/test_sampling.py
similarity index 100%
rename from sgl-kernel/tests/test_sampling.py
rename to python/sglang/kernels/aot/tests/test_sampling.py
diff --git a/sgl-kernel/tests/test_topk.py b/python/sglang/kernels/aot/tests/test_topk.py
similarity index 100%
rename from sgl-kernel/tests/test_topk.py
rename to python/sglang/kernels/aot/tests/test_topk.py
diff --git a/sgl-kernel/tests/test_torch_defaults_reset.py b/python/sglang/kernels/aot/tests/test_torch_defaults_reset.py
similarity index 100%
rename from sgl-kernel/tests/test_torch_defaults_reset.py
rename to python/sglang/kernels/aot/tests/test_torch_defaults_reset.py
diff --git a/sgl-kernel/tests/utils.py b/python/sglang/kernels/aot/tests/utils.py
similarity index 100%
rename from sgl-kernel/tests/utils.py
rename to python/sglang/kernels/aot/tests/utils.py
diff --git a/python/sglang/srt/hardware_backend/mlx/aot.py b/python/sglang/srt/hardware_backend/mlx/aot.py
index 94158b59e..e04b7651f 100644
--- a/python/sglang/srt/hardware_backend/mlx/aot.py
+++ b/python/sglang/srt/hardware_backend/mlx/aot.py
@@ -33,8 +33,8 @@ def _load_metal_rope_pool_fused():
         raise ImportError(
             "sgl_kernel.metal is importable, but the native Metal extension "
             f"or metallib is not available.{reason} Install the Metal kernels "
-            "with `uv run sgl-kernel/setup_metal.py install` from the SGLang "
-            "repo root in the active environment."
+            "with `uv run python/sglang/kernels/aot/setup_metal.py install` "
+            "from the SGLang repo root in the active environment."
         ) from import_error
     return metal.rope_pool_fused
 
diff --git a/python/sglang/srt/mem_cache/cpp_radix_tree/.clang-format b/python/sglang/srt/mem_cache/cpp_radix_tree/.clang-format
index 5a7a8cea7..7d1cf7aa8 120000
--- a/python/sglang/srt/mem_cache/cpp_radix_tree/.clang-format
+++ b/python/sglang/srt/mem_cache/cpp_radix_tree/.clang-format
@@ -1 +1 @@
-../../../../../sgl-kernel/.clang-format
\ No newline at end of file
+../../../kernels/aot/.clang-format
\ No newline at end of file
diff --git a/scripts/ci/amd/amd_ci_install_dependency.sh b/scripts/ci/amd/amd_ci_install_dependency.sh
index 344ba75fb..1f9764373 100755
--- a/scripts/ci/amd/amd_ci_install_dependency.sh
+++ b/scripts/ci/amd/amd_ci_install_dependency.sh
@@ -128,7 +128,7 @@ else
   # Also clear cache in sglang-checkout
   docker exec ci_sglang find /sglang-checkout -name "*.pyc" -delete || true
   docker exec ci_sglang find /sglang-checkout -name "__pycache__" -type d -exec rm -rf {} + || true
-  docker exec -w /sglang-checkout/sgl-kernel ci_sglang bash -c "rm -f pyproject.toml && mv pyproject_rocm.toml pyproject.toml && python3 setup_rocm.py install"
+  docker exec -w /sglang-checkout/python/sglang/kernels/aot ci_sglang bash -c "rm -f pyproject.toml && mv pyproject_rocm.toml pyproject.toml && python3 setup_rocm.py install"
 
   docker exec ci_sglang bash -c 'rm -rf python/pyproject.toml && mv python/pyproject_other.toml python/pyproject.toml'
   install_with_retry docker exec ci_sglang pip install --cache-dir=/sgl-data/pip-cache -e "python[${EXTRAS}]"
diff --git a/scripts/ci/cuda/ci_install_dependency.sh b/scripts/ci/cuda/ci_install_dependency.sh
index d941da6d2..3c78b1be9 100755
--- a/scripts/ci/cuda/ci_install_dependency.sh
+++ b/scripts/ci/cuda/ci_install_dependency.sh
@@ -304,28 +304,28 @@ install_sglang() {
 }
 
 install_sglang_kernel() {
-    SGL_KERNEL_VERSION_FROM_KERNEL=$(grep -Po '(?<=^version = ")[^"]*' sgl-kernel/pyproject.toml)
+    SGL_KERNEL_VERSION_FROM_KERNEL=$(grep -Po '(?<=^version = ")[^"]*' python/sglang/kernels/aot/pyproject.toml)
     SGL_KERNEL_VERSION_FROM_SRT=$(grep -Po -m1 '(?<=sglang-kernel==)[0-9A-Za-z\.\-]+' python/pyproject.toml)
     echo "SGL_KERNEL_VERSION_FROM_KERNEL=${SGL_KERNEL_VERSION_FROM_KERNEL} SGL_KERNEL_VERSION_FROM_SRT=${SGL_KERNEL_VERSION_FROM_SRT}"
 
-    if [ "${CUSTOM_BUILD_SGL_KERNEL:-}" = "true" ] && [ -d "sgl-kernel/dist" ]; then
-        ls -alh sgl-kernel/dist
+    if [ "${CUSTOM_BUILD_SGL_KERNEL:-}" = "true" ] && [ -d "python/sglang/kernels/aot/dist" ]; then
+        ls -alh python/sglang/kernels/aot/dist
         if [ "$ARCH" = "aarch64" ] || [ "$ARCH" = "arm64" ]; then
             WHEEL_ARCH="aarch64"
         else
             WHEEL_ARCH="x86_64"
         fi
-        KERNEL_WHL=$(ls sgl-kernel/dist/sglang_kernel-${SGL_KERNEL_VERSION_FROM_KERNEL}+${CU_VERSION}-cp310-abi3-manylinux2014_${WHEEL_ARCH}.whl 2>/dev/null | head -1 || true)
+        KERNEL_WHL=$(ls python/sglang/kernels/aot/dist/sglang_kernel-${SGL_KERNEL_VERSION_FROM_KERNEL}+${CU_VERSION}-cp310-abi3-manylinux2014_${WHEEL_ARCH}.whl 2>/dev/null | head -1 || true)
         if [ -z "$KERNEL_WHL" ]; then
-            echo "ERROR: No matching sgl-kernel wheel found in sgl-kernel/dist/ for version ${SGL_KERNEL_VERSION_FROM_KERNEL} arch ${WHEEL_ARCH} cuda ${CU_VERSION}"
-            ls -alh sgl-kernel/dist/
+            echo "ERROR: No matching sgl-kernel wheel found in python/sglang/kernels/aot/dist/ for version ${SGL_KERNEL_VERSION_FROM_KERNEL} arch ${WHEEL_ARCH} cuda ${CU_VERSION}"
+            ls -alh python/sglang/kernels/aot/dist/
             exit 1
         fi
         echo "Installing sgl-kernel wheel: $KERNEL_WHL"
         $PIP_CMD install "$KERNEL_WHL" --force-reinstall $PIP_INSTALL_SUFFIX
     else
-        if [ "${CUSTOM_BUILD_SGL_KERNEL:-}" = "true" ] && [ ! -d "sgl-kernel/dist" ]; then
-            echo "ERROR: CUSTOM_BUILD_SGL_KERNEL=true but sgl-kernel/dist not found."
+        if [ "${CUSTOM_BUILD_SGL_KERNEL:-}" = "true" ] && [ ! -d "python/sglang/kernels/aot/dist" ]; then
+            echo "ERROR: CUSTOM_BUILD_SGL_KERNEL=true but python/sglang/kernels/aot/dist not found."
             echo "This usually happens when rerunning a stage without the sgl-kernel-build-wheels job."
             echo "Please re-run the full workflow using /tag-and-rerun-ci to rebuild the kernel."
             exit 1
diff --git a/scripts/ci/musa/musa_install_dependency.sh b/scripts/ci/musa/musa_install_dependency.sh
index 33f2b1717..0dfe942bf 100755
--- a/scripts/ci/musa/musa_install_dependency.sh
+++ b/scripts/ci/musa/musa_install_dependency.sh
@@ -84,7 +84,7 @@ else
 
     cd "${REPO_ROOT}" && ${PIP_INSTALL} -v -e "python[dev_musa]" --user
 
-    cd "${REPO_ROOT}/sgl-kernel"
+    cd "${REPO_ROOT}/python/sglang/kernels/aot"
     rm -f pyproject.toml && mv pyproject_musa.toml pyproject.toml && MTGPU_TARGET=mp_31 python3 setup_musa.py install --user
     echo "$HOME/.local/bin" >> "$GITHUB_PATH"
 fi
diff --git a/scripts/ci/musa/rename_wheels_musa.sh b/scripts/ci/musa/rename_wheels_musa.sh
index f3816548a..25ada3b2f 100755
--- a/scripts/ci/musa/rename_wheels_musa.sh
+++ b/scripts/ci/musa/rename_wheels_musa.sh
@@ -10,7 +10,7 @@
 # Usage:
 #   rename_wheels_musa.sh <musa_suffix> [wheel_dir]
 # Example:
-#   rename_wheels_musa.sh 43 sgl-kernel/dist
+#   rename_wheels_musa.sh 43 python/sglang/kernels/aot/dist
 set -euxo pipefail
 
 if [[ $# -lt 1 || $# -gt 2 ]]; then
diff --git a/scripts/ci/npu/npu_log_print.sh b/scripts/ci/npu/npu_log_print.sh
index 92ba4fe3e..e0c4a0ce4 100755
--- a/scripts/ci/npu/npu_log_print.sh
+++ b/scripts/ci/npu/npu_log_print.sh
@@ -8,7 +8,7 @@ get_version() {
     [ -f "$1" ] && python3 -c 'import re, sys; print(sys.argv[2] + " version: v" + re.search(r"__version__\s*=\s*[\"'"'"'](.*?)[\"'"'"']", open(sys.argv[1]).read()).group(1))' "$1" "$2" 2>/dev/null || echo "$2 version: unknown"
 }
 get_version "./python/sglang/version.py" "sglang"
-get_version "./sgl-kernel/python/sgl_kernel/version.py" "sgl_kernel"
+get_version "./python/sglang/kernels/aot/python/sgl_kernel/version.py" "sgl_kernel"
 SGLANG_URL="https://github.com/sgl-project/sglang.git"
 SGL_KERNEL_URL="https://github.com/sgl-project/sgl-kernel-npu.git"
 SGLANG_BRANCH="main"
diff --git a/scripts/ci/utils/slash_command_handler.py b/scripts/ci/utils/slash_command_handler.py
index c3d4af1e0..1be998bf4 100644
--- a/scripts/ci/utils/slash_command_handler.py
+++ b/scripts/ci/utils/slash_command_handler.py
@@ -289,14 +289,14 @@ def load_permissions(user_login):
 
 def has_sgl_kernel_changes(pr):
     """
-    Check if the PR has changes to the sgl-kernel directory.
+    Check if the PR has changes to the AOT kernel directory.
     This is used to determine if we need a full workflow rerun
     (to rebuild the kernel) vs just rerunning failed jobs.
     """
     try:
         files = pr.get_files()
         for f in files:
-            if f.filename.startswith("sgl-kernel/"):
+            if f.filename.startswith("python/sglang/kernels/aot/"):
                 return True
         return False
     except Exception as e:
diff --git a/scripts/ci_monitor/ci_auto_bisect.py b/scripts/ci_monitor/ci_auto_bisect.py
index 2600d343f..2fe61d3b7 100755
--- a/scripts/ci_monitor/ci_auto_bisect.py
+++ b/scripts/ci_monitor/ci_auto_bisect.py
@@ -576,7 +576,7 @@ def _infer_related_paths(test_file: str) -> List[str]:
         "specul": ["python/sglang/srt/speculative/"],
         "vision": ["python/sglang/srt/models/"],
         "embed": ["python/sglang/srt/layers/"],
-        "kernel": ["sgl-kernel/", "python/sglang/srt/layers/"],
+        "kernel": ["python/sglang/kernels/aot/", "python/sglang/srt/layers/"],
         "bench": ["benchmark/"],
         "constrained": ["python/sglang/srt/constrained/"],
     }
diff --git a/scripts/release/README.md b/scripts/release/README.md
index 52d203bd6..6ecf3ef47 100644
--- a/scripts/release/README.md
+++ b/scripts/release/README.md
@@ -109,11 +109,11 @@ python scripts/release/bump_kernel_version.py 0.4.0
 ```
 
 **Files updated:**
-- `sgl-kernel/pyproject.toml`
-- `sgl-kernel/pyproject_cpu.toml`
-- `sgl-kernel/pyproject_rocm.toml`
-- `sgl-kernel/pyproject_musa.toml`
-- `sgl-kernel/python/sgl_kernel/version.py`
+- `python/sglang/kernels/aot/pyproject.toml`
+- `python/sglang/kernels/aot/pyproject_cpu.toml`
+- `python/sglang/kernels/aot/pyproject_rocm.toml`
+- `python/sglang/kernels/aot/pyproject_musa.toml`
+- `python/sglang/kernels/aot/python/sgl_kernel/version.py`
 
 ## Manual Testing Instructions
 
@@ -155,8 +155,8 @@ python scripts/release/bump_kernel_version.py 0.4.0
 
 3. **Check specific files contain the new version:**
    ```bash
-   grep -r "0.4.0" sgl-kernel/python/sgl_kernel/version.py
-   grep -r "0.4.0" sgl-kernel/pyproject.toml
+   grep -r "0.4.0" python/sglang/kernels/aot/python/sgl_kernel/version.py
+   grep -r "0.4.0" python/sglang/kernels/aot/pyproject.toml
    ```
 
 4. **Reset changes (if testing):**
diff --git a/scripts/release/bump_kernel_version.py b/scripts/release/bump_kernel_version.py
index 2a7f89fc5..f6c59b1d8 100755
--- a/scripts/release/bump_kernel_version.py
+++ b/scripts/release/bump_kernel_version.py
@@ -16,14 +16,14 @@ def main():
     )
     args = parser.parse_args()
 
-    version_file = Path("sgl-kernel/python/sgl_kernel/version.py")
+    version_file = Path("python/sglang/kernels/aot/python/sgl_kernel/version.py")
 
     files_to_update = [
-        Path("sgl-kernel/pyproject.toml"),
-        Path("sgl-kernel/pyproject_cpu.toml"),
-        Path("sgl-kernel/pyproject_rocm.toml"),
-        Path("sgl-kernel/pyproject_musa.toml"),
-        Path("sgl-kernel/python/sgl_kernel/version.py"),
+        Path("python/sglang/kernels/aot/pyproject.toml"),
+        Path("python/sglang/kernels/aot/pyproject_cpu.toml"),
+        Path("python/sglang/kernels/aot/pyproject_rocm.toml"),
+        Path("python/sglang/kernels/aot/pyproject_musa.toml"),
+        Path("python/sglang/kernels/aot/python/sgl_kernel/version.py"),
     ]
 
     bump_version(args.new_version, version_file, files_to_update)
diff --git a/scripts/release/bump_kernel_version_to_sglang.py b/scripts/release/bump_kernel_version_to_sglang.py
index c271835dc..9250f2ff8 100755
--- a/scripts/release/bump_kernel_version_to_sglang.py
+++ b/scripts/release/bump_kernel_version_to_sglang.py
@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 """
-Bump sglang-kernel version in SGLang files to match the version in sgl-kernel/pyproject.toml.
+Bump SGLang's kernel dependencies to match the AOT source pyproject.
 Updates:
   - python/pyproject.toml
   - python/sglang/srt/entrypoints/engine.py
@@ -18,8 +18,8 @@ except ImportError:
 
 
 def get_kernel_version_from_source() -> str:
-    """Extract version from sgl-kernel/pyproject.toml"""
-    pyproject_path = Path("sgl-kernel/pyproject.toml")
+    """Extract the version from the AOT source pyproject."""
+    pyproject_path = Path("python/sglang/kernels/aot/pyproject.toml")
 
     if not pyproject_path.exists():
         print(f"Error: {pyproject_path} not found")
@@ -30,7 +30,9 @@ def get_kernel_version_from_source() -> str:
 
     version = data.get("project", {}).get("version")
     if not version:
-        print("Error: Could not find version in sgl-kernel/pyproject.toml")
+        print(
+            "Error: Could not find version in python/sglang/kernels/aot/pyproject.toml"
+        )
         sys.exit(1)
 
     return version
diff --git a/scripts/release/check_kernel_version_to_sglang.py b/scripts/release/check_kernel_version_to_sglang.py
index ec0aeae7b..97a185d97 100755
--- a/scripts/release/check_kernel_version_to_sglang.py
+++ b/scripts/release/check_kernel_version_to_sglang.py
@@ -1,7 +1,9 @@
 #!/usr/bin/env python3
 """
-Check if sglang-kernel version from sgl-kernel/pyproject.toml matches the versions
-used in SGLang files (python/pyproject.toml, engine.py, and Dockerfile).
+Check whether SGLang's kernel dependencies match the AOT source pyproject.
+
+The dependent versions are read from python/pyproject.toml, engine.py, and
+Dockerfile.
 Sets GitHub Actions output variables to indicate if sync is needed.
 """
 
@@ -17,8 +19,8 @@ except ImportError:
 
 
 def get_kernel_version_from_source() -> str:
-    """Extract version from sgl-kernel/pyproject.toml (line 11)"""
-    pyproject_path = Path("sgl-kernel/pyproject.toml")
+    """Extract the version from the AOT source pyproject."""
+    pyproject_path = Path("python/sglang/kernels/aot/pyproject.toml")
 
     if not pyproject_path.exists():
         print(f"Error: {pyproject_path} not found")
@@ -29,7 +31,9 @@ def get_kernel_version_from_source() -> str:
 
     version = data.get("project", {}).get("version")
     if not version:
-        print("Error: Could not find version in sgl-kernel/pyproject.toml")
+        print(
+            "Error: Could not find version in python/sglang/kernels/aot/pyproject.toml"
+        )
         sys.exit(1)
 
     return version
@@ -101,7 +105,9 @@ def main():
     engine_version = get_kernel_version_from_engine()
     dockerfile_version = get_kernel_version_from_dockerfile()
 
-    print(f"Kernel version in sgl-kernel/pyproject.toml: {kernel_version}")
+    print(
+        f"Kernel version in python/sglang/kernels/aot/pyproject.toml: {kernel_version}"
+    )
     print(
         f"SGLang kernel dependency version in python/pyproject.toml: {pyproject_version}"
     )
diff --git a/scripts/release/commit_and_pr_kernel_to_sglang.sh b/scripts/release/commit_and_pr_kernel_to_sglang.sh
index 10579c85f..f63631910 100755
--- a/scripts/release/commit_and_pr_kernel_to_sglang.sh
+++ b/scripts/release/commit_and_pr_kernel_to_sglang.sh
@@ -28,7 +28,7 @@ git add -A
 git commit -m "chore: bump sglang-kernel version to ${KERNEL_VERSION} in SGLang
 
 This commit updates the sglang-kernel version across SGLang files to match
-the version defined in sgl-kernel/pyproject.toml.
+the version defined in python/sglang/kernels/aot/pyproject.toml.
 
 Files updated:
 ${COMMIT_FILES}
@@ -45,7 +45,7 @@ PR_URL=$(gh pr create \
   --title "chore: bump sglang-kernel version to ${KERNEL_VERSION}" \
   --body "## Summary
 
-This PR bumps the \`sglang-kernel\` version to \`${KERNEL_VERSION}\` across SGLang files to match the version defined in \`sgl-kernel/pyproject.toml\`.
+This PR bumps the \`sglang-kernel\` version to \`${KERNEL_VERSION}\` across SGLang files to match the version defined in \`python/sglang/kernels/aot/pyproject.toml\`.
 
 **Kernel Version:** \`${KERNEL_VERSION}\`
 
@@ -54,7 +54,7 @@ ${FILES_LIST}
 
 ## Context
 
-The kernel version in \`sgl-kernel/pyproject.toml\` has been updated. This PR ensures that all SGLang files referencing the \`sglang-kernel\` dependency are updated accordingly:
+The kernel version in \`python/sglang/kernels/aot/pyproject.toml\` has been updated. This PR ensures that all SGLang files referencing the \`sglang-kernel\` dependency are updated accordingly:
 - \`python/pyproject.toml\` - dependency specification
 - \`python/sglang/srt/entrypoints/engine.py\` - version check
 - \`docker/Dockerfile\` - Docker build argument
diff --git a/scripts/update_kernel_whl_index.py b/scripts/update_kernel_whl_index.py
index 093a7d699..05e20661e 100644
--- a/scripts/update_kernel_whl_index.py
+++ b/scripts/update_kernel_whl_index.py
@@ -34,7 +34,7 @@ def update_wheel_index(cuda_version=DEFAULT_CUDA_VERSION, rocm_version=None):
     index_dir.mkdir(exist_ok=True, parents=True)
     base_url = "https://github.com/sgl-project/whl/releases/download"
 
-    for path in sorted(pathlib.Path("sgl-kernel/dist").glob("*.whl")):
+    for path in sorted(pathlib.Path("python/sglang/kernels/aot/dist").glob("*.whl")):
         # Skip the wheel if mismatches the passed in cuda_version
         if not check_wheel_cuda_version(path.name, cuda_version):
             continue
@@ -53,7 +53,7 @@ def _update_non_cuda_wheel_index(backend, version):
     index_dir.mkdir(exist_ok=True, parents=True)
     base_url = "https://github.com/sgl-project/whl/releases/download"
 
-    for path in sorted(pathlib.Path("sgl-kernel/dist").glob("*.whl")):
+    for path in sorted(pathlib.Path("python/sglang/kernels/aot/dist").glob("*.whl")):
         # Skip the wheel if not for this backend
         if re.search(f"{backend}", path.name) is None:
             continue
diff --git a/test/registered/cpu/test_spec_kernels.py b/test/registered/cpu/test_spec_kernels.py
index 73d7fe076..95ee1efec 100644
--- a/test/registered/cpu/test_spec_kernels.py
+++ b/test/registered/cpu/test_spec_kernels.py
@@ -294,7 +294,7 @@ class TestVerifyTreeGreedy(CustomTestCase):
 
     def test_verify_tree_greedy_upstream_golden(self):
         # Golden fixture ported from the CUDA kernel UT
-        # sgl-kernel/tests/speculative/test_eagle_utils.py::test_verify_tree_greedy
+        # python/sglang/kernels/aot/tests/speculative/test_eagle_utils.py::test_verify_tree_greedy
         # (device swapped to CPU); expected outputs are the CUDA kernel's.
         candidates = torch.tensor(
             [
@@ -767,7 +767,7 @@ class TestReconstructIndicesFromTreeMask(CustomTestCase):
         bs, draft_token_num = 2, 4
         seq_lens = torch.tensor([12, 5], dtype=torch.int64)
         # Request 0: root(0) -> {1, 2}, 2 -> 3 (golden case from
-        # sgl-kernel/tests/speculative/test_ngram_utils.py).
+        # python/sglang/kernels/aot/tests/speculative/test_ngram_utils.py).
         # Request 1: plain chain 0 -> 1 -> 2 -> 3.
         tree_mask = torch.tensor(
             # fmt: off
diff --git a/test/registered/kernels/ops/attention/test_hadamard_jit.py b/test/registered/kernels/ops/attention/test_hadamard_jit.py
index 3c53bd29f..cefd5f68d 100644
--- a/test/registered/kernels/ops/attention/test_hadamard_jit.py
+++ b/test/registered/kernels/ops/attention/test_hadamard_jit.py
@@ -221,13 +221,13 @@ def hadamard_transform_mn_ref(x, multiple, scale=1.0):
 @pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.bfloat16])
 @pytest.mark.parametrize(
     "dim",
-    # Power-of-2 dims from sgl-kernel/tests/test_hadamard.py (old AOT test)
+    # Power-of-2 dims from python/sglang/kernels/aot/tests/test_hadamard.py (old AOT test)
     [1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384, 32768],
 )
 def test_hadamard_transform(dim, dtype):
     device = "cuda"
 
-    # Tolerances from sgl-kernel/tests/test_hadamard.py (old AOT test)
+    # Tolerances from python/sglang/kernels/aot/tests/test_hadamard.py (old AOT test)
     if dtype == torch.float32:
         rtol, atol = 3e-4, 3e-3
     elif dtype == torch.bfloat16:
@@ -252,7 +252,7 @@ def test_hadamard_transform(dim, dtype):
 @pytest.mark.parametrize(
     "dim",
     # Non-power-of-2 dims to test the padding path
-    # (137 from sgl-kernel/tests/test_hadamard.py, 500/1000 added for coverage)
+    # (137 from python/sglang/kernels/aot/tests/test_hadamard.py, 500/1000 added for coverage)
     [137, 500, 1000],
 )
 def test_hadamard_transform_non_power_of_two(dim, dtype):
diff --git a/test/registered/kernels/ops/moe/test_renorm.py b/test/registered/kernels/ops/moe/test_renorm.py
index 49142cfd4..7aafcf979 100644
--- a/test/registered/kernels/ops/moe/test_renorm.py
+++ b/test/registered/kernels/ops/moe/test_renorm.py
@@ -1,5 +1,5 @@
 # Adapted from https://github.com/flashinfer-ai/flashinfer/blob/main/tests/test_sampling.py
-# and /sgl-workspace/sglang/sgl-kernel/tests/test_sampling.py
+# and /sgl-workspace/sglang/python/sglang/kernels/aot/tests/test_sampling.py
 
 import sys