diff --git a/csrc/CMakeLists.txt b/csrc/CMakeLists.txt
index d17ea61..eecf091 100644
--- a/csrc/CMakeLists.txt
+++ b/csrc/CMakeLists.txt
@@ -90,6 +90,7 @@ foreach(i RANGE ${_kernel_last})
     target_compile_definitions(${name} PRIVATE TORCH_EXTENSION_NAME=${name})
 
     target_include_directories(${name} PRIVATE
+        "${CMAKE_CURRENT_SOURCE_DIR}/kernels"
         "${TORCH_HOME}/include"
         "${TORCH_HOME}/include/torch/csrc/api/include"
         "${PYTHON_INCLUDE_DIR}")
diff --git a/csrc/kernels/attention/decode_split_kv.cuh b/csrc/kernels/attention/decode_split_kv.cuh
index 56c7f19..704a279 100644
--- a/csrc/kernels/attention/decode_split_kv.cuh
+++ b/csrc/kernels/attention/decode_split_kv.cuh
@@ -2,8 +2,8 @@
 #include <cuda_bf16.h>
 #include <float.h>
 #include "common.h"
+#include "common/reduce.cuh"
 #include "layout_policies.cuh"
-#include "../common/reduce.cuh"
 
 namespace astrai {
 namespace attention {
diff --git a/csrc/kernels/attention/mma_utils.cuh b/csrc/kernels/attention/mma_utils.cuh
index 7cceb33..092aded 100644
--- a/csrc/kernels/attention/mma_utils.cuh
+++ b/csrc/kernels/attention/mma_utils.cuh
@@ -3,8 +3,8 @@
 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
 
-#include "../common/cp_async.cuh"
-#include "../common/mma.cuh"
+#include "common/cp_async.cuh"
+#include "common/mma.cuh"
 
 // Predicated cp.async (4-operand form) requires CUDA 11.2+.
 // bf16 mma.sync requires sm_80+ (guarded at build time by ASTRAI_NO_MMA).
diff --git a/csrc/kernels/attention/prefill_split_q.cuh b/csrc/kernels/attention/prefill_split_q.cuh
index 2717a38..0e0587e 100644
--- a/csrc/kernels/attention/prefill_split_q.cuh
+++ b/csrc/kernels/attention/prefill_split_q.cuh
@@ -2,8 +2,8 @@
 #include <cfloat>
 #include <cuda_bf16.h>
 #include "common.h"
+#include "common/reduce.cuh"
 #include "layout_policies.cuh"
-#include "../common/reduce.cuh"
 
 namespace astrai {
 namespace attention {
diff --git a/csrc/kernels/fp8/gemm.cuh b/csrc/kernels/fp8/gemm.cuh
index bf96b79..b5456e3 100644
--- a/csrc/kernels/fp8/gemm.cuh
+++ b/csrc/kernels/fp8/gemm.cuh
@@ -11,8 +11,8 @@
 #include <cuda_runtime.h>
 #include <type_traits>
 
-#include "../common/cp_async.cuh"
 #include "common.h"
+#include "common/cp_async.cuh"
 #include "gemm/epilogue.cuh"
 #include "gemm/load.cuh"
 #include "gemm/mainloop.cuh"
diff --git a/csrc/kernels/fp8/gemm/epilogue.cuh b/csrc/kernels/fp8/gemm/epilogue.cuh
index c9456c4..3d0caf7 100644
--- a/csrc/kernels/fp8/gemm/epilogue.cuh
+++ b/csrc/kernels/fp8/gemm/epilogue.cuh
@@ -2,7 +2,7 @@
 // Collective epilogue: fused bias, the bf16 scatter of the fp32 accumulators
 // through the reclaimed operand shared memory, and the coalesced copy-out.
 
-#include "../common.h"
+#include "fp8/common.h"
 #include "policy.cuh"
 
 namespace astrai {
diff --git a/csrc/kernels/fp8/gemm/load.cuh b/csrc/kernels/fp8/gemm/load.cuh
index 7d28e0e..f5d26b3 100644
--- a/csrc/kernels/fp8/gemm/load.cuh
+++ b/csrc/kernels/fp8/gemm/load.cuh
@@ -5,8 +5,8 @@
 // The staging invariants and the swizzle derivation live in
 // docs/developer/cuda_kernels.md.
 
-#include "../../common/cp_async.cuh"
-#include "../common.h"
+#include "common/cp_async.cuh"
+#include "fp8/common.h"
 #include "policy.cuh"
 
 namespace astrai {
diff --git a/csrc/kernels/fp8/gemm/mainloop.cuh b/csrc/kernels/fp8/gemm/mainloop.cuh
index 730c328..e8e5d88 100644
--- a/csrc/kernels/fp8/gemm/mainloop.cuh
+++ b/csrc/kernels/fp8/gemm/mainloop.cuh
@@ -7,8 +7,8 @@
 
 #include <type_traits>
 
-#include "../../common/mma.cuh"
-#include "../common.h"
+#include "common/mma.cuh"
+#include "fp8/common.h"
 #include "load.cuh"
 #include "policy.cuh"
 
diff --git a/csrc/kernels/fp8/gemm/policy.cuh b/csrc/kernels/fp8/gemm/policy.cuh
index ca3b53f..e7e074b 100644
--- a/csrc/kernels/fp8/gemm/policy.cuh
+++ b/csrc/kernels/fp8/gemm/policy.cuh
@@ -5,7 +5,7 @@
 
 #include <type_traits>
 
-#include "../common.h"
+#include "fp8/common.h"
 
 namespace astrai {
 namespace fp8 {
diff --git a/csrc/kernels/fp8/ops.cu b/csrc/kernels/fp8/ops.cu
index 14bfc2e..0dd7a60 100644
--- a/csrc/kernels/fp8/ops.cu
+++ b/csrc/kernels/fp8/ops.cu
@@ -8,7 +8,7 @@
 #include <mutex>
 #include <unordered_map>
 
-#include "../common/device.cuh"
+#include "common/device.cuh"
 #include "gemm.cuh"
 #include "quantize.cuh"
 
diff --git a/csrc/kernels/fp8/quantize.cuh b/csrc/kernels/fp8/quantize.cuh
index 97f2bbd..8522fea 100644
--- a/csrc/kernels/fp8/quantize.cuh
+++ b/csrc/kernels/fp8/quantize.cuh
@@ -10,7 +10,7 @@
 #include <cstdint>
 
 #include "common.h"
-#include "../common/reduce.cuh"
+#include "common/reduce.cuh"
 
 namespace astrai {
 namespace fp8 {
diff --git a/csrc/tests/attn_paged_test.cu b/csrc/tests/attn_paged_test.cu
index 632b7ea..04c5cc2 100644
--- a/csrc/tests/attn_paged_test.cu
+++ b/csrc/tests/attn_paged_test.cu
@@ -1,5 +1,5 @@
 // Compile:
-//   nvcc -I csrc -arch=sm_89 -O3 --use_fast_math --ptxas-options=-O3 \
+//   nvcc -I csrc/kernels -arch=sm_89 -O3 --use_fast_math --ptxas-options=-O3 \
 //        --extra-device-vectorization -Xcompiler -fopenmp \
 //        csrc/tests/attn_paged_test.cu \
 //        -o /tmp/test_paged && /tmp/test_paged
@@ -7,7 +7,7 @@
 #include <cstring>
 #include <vector>
 #include "test_utils.cuh"
-#include "../kernels/attention/dispatchers.cuh"
+#include "attention/dispatchers.cuh"
 
 using namespace astrai::attention;
 
diff --git a/csrc/tests/attn_test.cu b/csrc/tests/attn_test.cu
index 9ef5e3b..4c972ea 100644
--- a/csrc/tests/attn_test.cu
+++ b/csrc/tests/attn_test.cu
@@ -1,13 +1,13 @@
 /*
 Pure-C test — uses shared dispatcher.  Combines the decode (split-KV) and
 prefill (split-Q) correctness checks + benchmarks into one binary.
-nvcc -I csrc -arch=sm_89 -O3 \
+nvcc -I csrc/kernels -arch=sm_89 -O3 \
     --use_fast_math --ptxas-options=-O3 --extra-device-vectorization \
     -Xcompiler -fopenmp csrc/tests/attn_test.cu -o test && ./test
 */
 
 #include "test_utils.cuh"
-#include "../kernels/attention/dispatchers.cuh"
+#include "attention/dispatchers.cuh"
 
 using namespace astrai::attention;
 
diff --git a/csrc/tests/fp8_test.cu b/csrc/tests/fp8_test.cu
index a994a0c..ca97f2c 100644
--- a/csrc/tests/fp8_test.cu
+++ b/csrc/tests/fp8_test.cu
@@ -6,7 +6,7 @@ Part 1 exercises one bf16 -> fp8 -> mma.sync m16n8k32 instruction pair
 Part 2 checks launch_fp8_gemm across all four operand layouts, both K
 tiles, and ragged shapes against an fp32 CPU reference.
 
-nvcc -I csrc -arch=sm_89 -std=c++17 -O3 csrc/tests/fp8_test.cu -o /tmp/fp8_test \
+nvcc -I csrc/kernels -arch=sm_89 -std=c++17 -O3 csrc/tests/fp8_test.cu -o /tmp/fp8_test \
     && /tmp/fp8_test
 */
 
@@ -21,8 +21,8 @@ nvcc -I csrc -arch=sm_89 -std=c++17 -O3 csrc/tests/fp8_test.cu -o /tmp/fp8_test
 #include <type_traits>
 #include <vector>
 
-#include "../kernels/common/mma.cuh"
-#include "../kernels/fp8/gemm.cuh"
+#include "common/mma.cuh"
+#include "fp8/gemm.cuh"
 
 using namespace astrai::fp8;
 
diff --git a/docs/developer/cuda_kernels.md b/docs/developer/cuda_kernels.md
index c8b600d..eca745f 100644
--- a/docs/developer/cuda_kernels.md
+++ b/docs/developer/cuda_kernels.md
@@ -407,7 +407,7 @@ blocks.
 Each `csrc/tests/*.cu` file has the `nvcc` compile command in its header comment. Example:
 
 ```bash
-nvcc -I csrc -arch=sm_89 -O3 --use_fast_math \
+nvcc -I csrc/kernels -arch=sm_89 -O3 --use_fast_math \
      --ptxas-options=-O3,-v --extra-device-vectorization \
      -Xcompiler -fopenmp csrc/tests/attn_test.cu -o /tmp/test && /tmp/test
 ```
@@ -423,7 +423,7 @@ Hardware: NVIDIA L20 (sm_89, 46 GB), CUDA 12.8, driver 570.86.
 
 Reproduce (decode + prefill in `attn_test.cu`, paged in `attn_paged_test.cu`):
 ```bash
-nvcc -I csrc -arch=sm_89 -O3 --use_fast_math \
+nvcc -I csrc/kernels -arch=sm_89 -O3 --use_fast_math \
      --ptxas-options=-O3,-v --extra-device-vectorization \
      -Xcompiler -fopenmp csrc/tests/attn_test.cu -o /tmp/test && /tmp/test
 ```