diff --git a/03_nf4_dequant/yuruwind/Summary_report.md b/03_nf4_dequant/yuruwind/Summary_report.md
new file mode 100644
index 0000000..f8c7a94
--- /dev/null
+++ b/03_nf4_dequant/yuruwind/Summary_report.md
@@ -0,0 +1,303 @@
+
+
+# 高性能 NF4 解量化推理算子 (CUDA)
+
+##  项目简介
+
+本项目是为大模型推理引擎开发的高性能 **NF4（NormalFloat 4-bit）** 反量化算子。针对 QLoRA 技术中的双重量化逻辑，利用 CUDA 进行了深度优化。算子能够将 4-bit 压缩权重实时还原为 FP16 格式，旨在解决 LLM 推理过程中的“访存墙”问题，提升消费级显卡的吞吐性能。
+
+#### 核心亮点：
+
+- **极致访存吞吐 **：
+
+  采用 **`half2` 向量化存储指令** 与 **双级缓存优化策略**（Shared Memory 加速二级码表访问，Constant Memory 广播 NF4 分位数表），在 RTX 4060 Laptop 上实测带宽突破 **122 GB/s**，逼近硬件物理极限。
+
+- **计算与 I/O 深度重叠：
+
+  成功将算子特征从访存受限（Memory-bound）重构为**计算受限（Compute-bound）**。通过高度并行的位移拆包与查表逻辑，使 **SM 利用率达到 80.4%**，实现了以计算资源置换显存带宽的架构红利。
+
+- **生产级鲁棒性 **：
+
+  内置自适应边界处理逻辑，**完美兼容任意非对齐维度**（如 $1513 \times 2017$ 等素数矩阵）。通过“高性能向量主路 + 安全标量支路”的设计，确保在各种 Batch Size 下均能实现比特级精度对齐（Max Error = 0）。
+
+- **引擎级参数解耦**：
+
+  采用**配置驱动架构**，支持通过 `config.txt` 动态分发计算精度（FP16/BF16）与硬件优化策略（针对不同架构动态调整 Thread Block 规模），具备极强的工程可扩展性。
+
+
+
+###  项目目录结构
+
+```
+.
+├── 🛠️ Core Source
+│   ├── nf4_v3.cu           # 核心算子实现 (含 V3 级优化 Kernel 与 Host 调度)
+│   └── cpu_dequantize.cpp  # CPU 基准实现 (用于算法初期逻辑验证)
+│
+├── 🧪 Testing & Validation
+│   ├── gen_data.py         # 自动化测试工具 (支持任意维度数据生成与真值计算)
+│   └── config.txt          # 运行时配置文件 (定义精度、硬件目标等元数据)
+│
+└── 📝 Documentation
+    └── Summary_Report.md   # README& 性能分析报告
+```
+
+
+
+## 运行方法
+
+#### 生成测试数据集
+
+使用 Python 脚本生成指定规模的测试数据（支持任意维度）。该脚本会同步生成 `input.bin` 和 `gt_output.bin`。
+
+```python
+# 生成一个 4096x4096 的标准大规模矩阵
+python gen_data.py 4096 4096
+
+# 或生成一个用于测试鲁棒性的非对齐矩阵
+python gen_data.py 1513 2017 --blocksize 64 --groupsize 256
+```
+
+#### 编译 CUDA 算子
+
+使用 `nvcc` 进行编译，建议开启 `-O3` 优化和硬件架构对齐。
+
+```
+nvcc -O3 -arch=sm_89 --use_fast_math nf4_v3.cu -o nf4_v3
+
+// sm_89 对应  NVIDIA RTX4060，可根据实际显卡调整
+```
+
+#### 配置运行时参数
+
+修改 `config.txt` 以匹配当前硬件环境：
+
+```
+blocksize=64
+compute_type=fp16
+target_gpu=4060
+```
+
+#### 执行算子与精度验证
+
+运行程序，它将自动加载配置、读取数据、执行核函数并输出精度比对结果与性能指标。
+
+```
+./nf4_v3
+```
+
+
+
+## 技术规格
+
+- **算法逻辑**：
+
+  $$W_{fp16} = \text{NF4\_Table}[\text{idx}] \times (\text{code2}[\text{absmax\_q}] \times \text{absmax2})$$
+
+- **内存优化**：
+
+  - **Vectorized Store**：使用 `STG.E.32` (via `half2`) 实现 100% 合并访存。
+  - **Tiling & Caching**：二级码表（512B）常驻 Shared Memory，NF4 表（16-word）常驻 Constant Cache。
+
+- **误差范围**：在 16M 元素规模下，最大绝对误差（Max Error） 和 平均绝对误差（MAE）为 **0**。
+
+
+
+## 详细介绍
+
+## 一、 项目背景与目标
+
+在大模型（LLM）推理场景中，显存带宽往往是制约吞吐量的最大瓶颈。公司采用 QLoRA 技术将权重压缩为 4-bit NormalFloat (NF4) 格式，有效将显存占用降低至 1/4。然而，实时的解量化过程（Dequantization）带来了密集的访存压力。
+
+本项目旨在为下一代推理引擎开发一个极致优化、单核、支持任意边界形状的 NF4 解量化 CUDA 算子。算子接收 4-bit 打包的权重流，结合一级缩放因子（`absmax_q`）和二级缩放码表（`code2`, `absmax2`），实时输出高精度的 16-bit（FP16/BF16）浮点矩阵。
+
+
+
+## 二、 核心实现思路与测试架构
+
+本项目不仅关注核函数本身的性能，更致力于构建一套**生产级**的验证与调度体系，确保算子在复杂工程环境下的可靠性。
+
+1. **端到端全链路验证流水线 (E2E Validation Pipeline)**：
+   - **仿真真值生成 (Simulation-based GT)**：开发了高度解耦的 `gen_data.py` 测试工具。该工具模拟了推理引擎前端的数据预处理行为。通过 Numpy 严格复刻 QLoRA 论文中的 NF4 双重量化数学逻辑，并采用 **二进制序列化协议（Binary Serialization Protocol）** 将权重、一级/二级缩放因子、码表打包存储。
+   - **命令行驱动测试**：支持动态传入任意 `rows`、`cols` 及 `blocksize` 参数。这种“参数化测试”设计使得算子可以快速在大规模对齐矩阵（如 4096 规模）与极端小规模不规则矩阵（如素数维度）之间切换，极大地提升了回归测试的效率。
+2. **解耦式配置管理与动态分派 (Decoupled Config & Dynamic Dispatch)**：
+   - **Host 侧策略分发器**：在 C++ 端引入了 `config.txt` 文本解析模块。这一设计实现了“模型元数据（Header）”与“运行时调度参数（Config）”的解耦。
+   - **硬件感知的 Kernel 调度**：程序能够识别 `target_gpu`（如 T4, RTX 4060），并根据不同架构的资源特性（寄存器数量、Shared Memory 延迟）动态调整线程配置（`threads_per_block`）。
+   - **多精度路径预留**：通过解析 `compute_type`，系统能够在 `fp16` 与 `bf16` 路径间进行动态分派。虽然当前核心逻辑针仅对 FP16 进行了指令集优化，但架构上实现了“配置驱动”的演进能力，符合工业级推理引擎（如 TensorRT）的设计哲学。
+
+
+
+## 三、测试数据集定义与分布特性
+
+为模拟真实大模型推理环境中的权重分布，本项目构造了具备以下特性的合成测试集：
+
+- **分布对齐 (Statistical Alignment)**：
+
+  由于 NF4（NormalFloat 4-bit）专门针对正态分布权重优化，本项目生成的测试索引在 16 个分位区间内均匀采样。这模拟了权重经过分位数量化后的统计状态，确保了查表逻辑在整个动态范围（-1.0 到 1.0）内都能得到充分验证。
+
+- **多级量化层级构造**：
+
+  测试数据严格遵循 QLoRA 的**双重量化（Double Quantization）**物理布局：
+
+  - **Level 1 (Block-wise)**：每 64/128 个元素划分为一个 Block，独立随机生成 `uint8` 类型的 `absmax_q`（一级缩放索引）。
+
+  - **Level 2 (Group-wise)**：每 256 个 Block 划分为一个 Group，共享一个基于 `float16` 的 `absmax2`（二级缩放因子）。
+
+    这种层级化的数据构造，旨在高压测试 Kernel 内部 `element_idx -> b_idx -> g_idx` 地址映射逻辑的准确性。
+
+- **非对齐极端测试**：
+
+  除了标准的 $4096 \times 4096$ 规格外，专门引入了不规则形状（如素数维度矩阵等）。这些数据故意不与 `blocksize` 或线程块大小（Thread Block Size）对齐，用于强制触发 Kernel 内部的边界保护逻辑（Boundary Guard），验证算子在处理模型长尾数据时的健壮性。
+
+- **多精度配置匹配**：
+
+  数据生成器支持根据 `config.txt` 动态切换输出类型（FP16/BF16）。所有缩放因子在持久化存储前均经过了高精度的半精度截断模拟，确保 Host 端生成的二进制文件与真值文件在比特位级（Bit-level）实现严格闭环。
+
+
+
+## 四、 性能优化历程 (从访存受限到计算受限)
+
+算子的开发遵循了典型的高性能计算调优范式，经历了从逻辑实现到访存对齐，再到指令级压榨的三次重大飞跃：
+
+### V1：基础标量实现 (Baseline)
+
+- **实现方案**：采用最朴素的单线程映射单元素模型。每个线程独立负责 4-bit 索引的提取、查表以及多级缩放计算。
+- **微架构瓶颈分析**：
+  - **访存碎片化**：GPU 全局内存交换的最小粒度为 32-Byte 扇区（Sector）。V1 版本中，线程束（Warp）内的线程访问地址不连续，导致触发了大量的非合并访存，有效带宽利用率（Bus Utilization）不足 5%。
+  - **指令吞吐低**：频繁触发低效的 8-bit 读取指令，导致流水线气泡严重。
+
+
+
+### V2：32-bit 向量化访存 (Vectorized Memory Access)
+
+- **优化策略**：重构线程映射模型，实现 **“1 Thread : 1 Byte (2 Elements)”** 的映射关系。利用 CUDA 内置的 `half2` 数据类型，将解量化后的两个 FP16 结果封装为 32-bit 向量。
+- **工程深度**：
+  - **指令级对齐**：通过编译器优化，`half2` 的存储在 SASS 汇编层被映射为原生 `STG.E.32` 指令。单次指令即可完成两个元素的写入，使得内存写入事务（Memory Transactions）减少了一半。
+  - **合并访存实现**：确保了同一 Warp 内的线程访问的是连续的 32-bit 地址空间，完美对齐了显存总线带宽。
+
+
+
+### V3：内存层级拓扑与指令缓存优化 (Current Version)
+
+- **核心策略**：利用 GPU 的多级存储抽象（Memory Hierarchy）隐藏长潜伏期访存。
+  - **Constant Memory 广播机制**：NF4 标准表（16 个 FP32）属于“高频读、不改写、全 Warp 共享”的典型数据。将其置于 `__constant__` 空间，利用其特有的 **Constant Cache** 广播机制，使 Warp 内 32 个线程单次周期即可同步获取数据，彻底消除了对全局显存的冗余请求。
+  - **Shared Memory 协作加载**：针对 512-Byte 的二级码表 `code2`，利用线程块（Block）内的前 256 个线程进行**协作式搬运（Cooperative Fetch）**。
+- **性能质变**：通过 `ncu` 剖析显示，算子的算术强度（Arithmetic Intensity）显著提升。瓶颈成功从 **访存受限（Memory-bound）** 转移至 **计算受限（Compute-bound）**，SM 利用率从 30% 飙升至 **80.3%**，实现了对硬件计算单元的深度压榨。
+
+
+
+## 五、 边界处理与系统级鲁棒性设计
+
+在 LLM 推理场景中，由于 Prompt 长度和 Batch Size 的动态性，算子必须具备处理**非对齐（Non-aligned）**数据的能力。
+
+1. **统一线程映射模型 (Universal Mapping)**：
+   - 针对 $rows \times cols$ 可能为奇数的情况，Host 端采用了 $\lceil (N+1)/2 \rceil$ 的线程分配策略。
+   - 这一设计确保了即使在权重矩阵维度极为“恶心”（如素数维度）时，每一个物理比特位都能被准确覆盖。
+2. **分支预测与自适应降级写入**：
+   - **热路径优化**：在 99% 的对齐区域，执行 `half2` 向量化写入指令，保持最高吞吐。
+   - **冷路径防御**：在矩阵的最末尾，引入边界保护分支（Boundary Guard）。当检测到当前为奇数尾部时，算子通过类型重写（Type Reinterpretation）自动降级为单元素 `half` 标量写入。
+   - **工程价值**：这种“高性能主路 + 安全支路”的设计，在保证了 $MAE = 0$ 的严谨精度的同时，通过分支预测器（Branch Predictor）最小化了跳转开销，保证了算子在复杂形状下的性能稳定性。
+
+
+
+## 六、 性能评测与 Profiling 分析 (Nsys & Ncu)
+
+- **测试平台**：Intel i7-12650H + NVIDIA RTX 4060 Laptop (AD107)
+- **测试数据**：极端非对齐矩阵 $4096 \times 4096$ (约 1700 万元素)
+
+### 1. 最终性能指标
+
+- **核函数执行时间**：0.345181 ms
+- **有效内存带宽**：**122.276 GB/s**
+- **误差精度**：Max Error 和 **平均绝对误差 (MAE) 趋近于 0**，完美满足 MAE < 1e-2 的精度要求。
+
+![image-20260315174840901](img/image-20260315174840901.png)
+
+
+
+**验证框架的有效性 (Negative Testing)：** 为防止出现“假阳性”结果，本项目进行了变异测试验证。
+
+在 `nf4_v3.cu` 中改动一个 NF4 常数：
+
+```c++
+// 把第一个 -1.0 换成 -1.1
+__constant__ float d_nf4_table[16] = {
+    -1.1f, -0.69487101f, ... 
+};
+```
+
+![image-20260315181548382](img/屏幕截图_20260315_181513.png)
+
+通过人工在 Kernel 内部注入 NF4 码表偏差（±0.1）及解包顺序扰动，观测到验证框架能灵敏捕捉到数值波动，误差反馈与注入量级完全吻合。这证明了本项目搭建的 Python-C++ 交叉验证体系具有极高的可靠性，能够有效拦截逻辑层与数值层的任何微小偏差。
+
+
+
+### 2. Nsight Systems (nsys) 宏观分析
+
+通过 nsys Timeline 观察发现，在单次算子调用中，`cudaMalloc` 以及内存拷贝耗时远超 Kernel 执行本身（如 Kernel 耗时 0.065ms，但环境准备耗时超 20ms）。这论证了在成熟框架（如 vLLM/TensorRT）中预先分配显存池（Memory Pool）的绝对必要性。
+
+![image-20260315182745611](img/屏幕截图_20260315_182718.png)
+
+
+
+### 3. Nsight Compute (ncu) 微观分析 (核心亮点)
+
+**Nsight Compute (ncu) 微观架构剖析 (基于 4096x4096 生产级负载)**
+
+在 $4096 \times 4096$ 的大规模矩阵测试中（消除 Launch Overhead 影响后），Ncu Profiling 揭示了本算子在微架构层面的三个核心特性：
+
+- **绝对的资源占有率 (100% Achieved Occupancy)**：
+
+  算子极其轻量，单线程寄存器消耗控制在 20 左右，配合 256 的 Block Size 与极小的 Shared Memory 足迹，使得 GPU 的硬件资源分配达到了理论极限（100% Occupancy），确保了极高的并行并发度。
+
+- **打破访存墙 (Memory Wall Mitigation)**：
+
+  Speed Of Light (SOL) 面板显示，算子的 **Memory Throughput 仅为 29.52%**。得益于 4-bit 的高压缩比与 `half2` 向量化读取（STG.E.32），全局内存的 I/O 压力被断崖式降低。
+
+- **跨越至计算受限状态 (Compute-Bound Transition)**：
+
+  与低访存形成鲜明对比的是，**SM Compute Throughput 高达 80.43%**，触发了高吞吐预警。Roofline 模型显示工作点紧贴非 Tensor Core 的计算峰值上限。这证明 Kernel 内部密集的 4-bit 索引拆包（位移与位与逻辑）、查表及缩放乘法消耗了大量的 ALU 周期。
+
+  **结论**：该 V3 版本算子成功实现了 QLoRA 设计的初衷——**用充足的 SM 计算能力置换稀缺的显存带宽**，为后续的 GEMV 推理环节省下了宝贵的显存 I/O 资源。
+
+![屏幕截图_20260315_182006](img/屏幕截图_20260315_182006.png)
+
+![屏幕截图_20260315_182040](img/屏幕截图_20260315_182040.png)
+
+![屏幕截图_20260315_182110](img/屏幕截图_20260315_182110.png)
+
+
+
+## 七、 开发中的问题与解决方案：从“溢出”到“对齐”
+
+在追求极致鲁棒性的过程中，本项目经历了一次深刻的**数值异常排查**，这充分体现了底层算子开发中“差之毫厘，谬以千里”的特性。
+
+**案例：非对齐维度的 `inf` 级联异常分析**
+
+在针对非对齐形状（如 $1513 \times 2017$）进行边界测试时，算子在原本逻辑完备的情况下意外输出 `Max Error: inf`。
+
+1. **故障现象与定位**：
+   - 经 `nsys` 时间线追踪，Kernel 执行并未超时；经 `ncu` 内存检查，读取地址亦未越界。
+   - 最终通过 **二进制内存比对** 发现，问题的根源并非 CUDA 代码，而是 Python 测试脚本与 C++ Host 端在**字节对齐协议**上的不一致。
+2. **技术成因分析**：
+   - 当矩阵总元素为奇数时，Python 的 `packed_weights = np.zeros(num_elements // 2)` 采用了向下取整逻辑，导致最后一个 4-bit 元素所在的字节被丢弃。
+   - 在 C++ 端，读取逻辑严格遵循 `(total_elements + 1) / 2` 向上取整。由于 Python 端少写了 1 字节，导致 C++ 的文件读取指针（File Pointer）产生了 **1 字节的偏移级联**。
+   - 这一微小的偏移使得随后读取的 `absmax_q`、`code2` 等 FP16/FP32 参数全部发生了**数据类型对齐错误（Alignment Misalignment）**。原本正常的二进制位被误解析为非法浮点数（NaN/Inf），最终通过解量化公式呈指数级放大，导致整个输出矩阵崩塌。
+3. **解决方案**：
+   - **协议对齐**：重构 Python 端打包逻辑，强制执行向上取整对齐，并在奇数尾部进行零填充（Padding）。
+   - **防御性编程**：在 C++ Host 端增加了读取大小校验逻辑，确保二进制流的每一个 Segment 长度符合预期。
+   - **启示**：这一问题的解决深刻印证了：在 AI Infra 开发中，**Host-Device 之间的数据契约（Data Contract）** 往往比算法逻辑本身更具风险，严谨的序列化规范是高性能算子的基石。
+
+
+
+## 八、 未来优化展望：迈向 SOTA 推理性能
+
+基于 Nsight Compute 的深度剖析报告，本项目在未来可进一步在以下前沿方向进行突破：
+
+1. **极宽向量化读取 (Wider Load Optimization)**：
+   - 目前读取端仍处于 8-bit 或 16-bit 粒度。利用 `int4` 或 `float4` 类型可将单次读取粒度提升至 **128-bit**（对应单条 `LDG.E.128` 指令）。配合寄存器解包技术，可进一步提升 L1 缓存的命中率与带宽利用率上限。
+2. **异步拷贝与双缓冲技术 (Asynchronous Copy & Double Buffering)**：
+   - 在支持 Ampere 及后续架构的显卡上，引入 `cp.async` 指令。在计算当前块的同时，异步将下一块数据从 Global Memory 搬运至 Shared Memory，实现**访存与计算的完美重叠（Latency Hiding）**。
+3. **算子融合与计算下沉 (Kernel Fusion & Tensor Core)**：
+   - **全链路融合**：当前的解量化仍存在“写回-再读”的开销。未来的极致优化方向是将 NF4 解量化逻辑直接嵌入到矩阵乘法（GEMM）或向量乘法（GEMV）的核心循环中。
+   - **计算单元下沉**：利用 Tensor Core 的高吞吐特性，在寄存器中实时解量化并直接喂给矩阵运算单元，从而彻底消除解量化中间变量对显存带宽的占用，这也是目前主流引擎（如 vLLM, TensorRT-LLM）实现 SOTA 性能的关键路径。
diff --git a/03_nf4_dequant/yuruwind/config.txt b/03_nf4_dequant/yuruwind/config.txt
new file mode 100644
index 0000000..9d61754
--- /dev/null
+++ b/03_nf4_dequant/yuruwind/config.txt
@@ -0,0 +1,3 @@
+blocksize=64
+compute_type=fp16
+target_gpu=4060
\ No newline at end of file
diff --git a/03_nf4_dequant/yuruwind/cpu_dequantize.cpp b/03_nf4_dequant/yuruwind/cpu_dequantize.cpp
new file mode 100644
index 0000000..7988b28
--- /dev/null
+++ b/03_nf4_dequant/yuruwind/cpu_dequantize.cpp
@@ -0,0 +1,110 @@
+#include <iostream>
+#include <vector>
+#include <fstream>
+#include <cmath>
+#include <stdint.h>
+#include <algorithm>
+
+// NF4 查找表
+const float NF4_TABLE[16] = {
+    -1.0f, -0.69487101f, -0.51209301f, -0.37391701f,
+    -0.25611401f, -0.14725500f, -0.04162400f, 0.06282201f,
+    0.16859101f, 0.28551400f, 0.40619302f, 0.53675699f,
+    0.68502200f, 0.87091398f, 1.0f, 0.0f
+};
+
+// FP16 转换
+float half_to_float_cpu(uint16_t h) {
+    union { float f; uint32_t i; } res;
+    uint32_t sign = (h & 0x8000) << 16;
+    uint32_t exp = (h & 0x7c00) >> 10;
+    uint32_t mant = (h & 0x03ff) << 13;
+    if (exp == 0x1f) {
+        res.i = sign | 0x7f800000 | mant;
+    } else if (exp == 0) {
+        if (mant == 0) res.i = sign;
+        else {
+            exp = 127 - 14;
+            while (!(mant & 0x00800000)) { mant <<= 1; exp--; }
+            res.i = sign | (exp << 23) | (mant & 0x007fffff);
+        }
+    } else {
+        res.i = sign | ((exp + (127 - 15)) << 23) | mant;
+    }
+    return res.f;
+}
+
+struct WeightHeader {
+    int64_t num_rows;
+    int64_t num_cols;
+    int32_t blocksize;
+};
+
+int main() {
+    std::ifstream ifs("input.bin", std::ios::binary);
+    if (!ifs) return -1;
+
+    // 1. 读取 Header（按字段读取，避免结构体对齐带来的偏移错误）
+    WeightHeader header;
+    ifs.read(reinterpret_cast<char*>(&header.num_rows), sizeof(header.num_rows));
+    ifs.read(reinterpret_cast<char*>(&header.num_cols), sizeof(header.num_cols));
+    ifs.read(reinterpret_cast<char*>(&header.blocksize), sizeof(header.blocksize));
+    
+    int64_t total_elements = header.num_rows * header.num_cols;
+    int32_t num_blocks = (total_elements + header.blocksize - 1) / header.blocksize;
+    int32_t group_size = 256; // QLoRA 默认
+    int32_t num_groups = (num_blocks + group_size - 1) / group_size;
+
+    // 2. 分配并读取数据
+    std::vector<uint8_t> packed_weights((total_elements + 1) / 2);
+    std::vector<uint8_t> absmax_q(num_blocks);
+    std::vector<uint16_t> code2(256);
+    std::vector<uint16_t> absmax2(num_groups);
+    float offset;
+
+    ifs.read(reinterpret_cast<char*>(packed_weights.data()), packed_weights.size());
+    ifs.read(reinterpret_cast<char*>(absmax_q.data()), absmax_q.size());
+    ifs.read(reinterpret_cast<char*>(code2.data()), code2.size() * 2);
+    ifs.read(reinterpret_cast<char*>(absmax2.data()), absmax2.size() * 2);
+    ifs.read(reinterpret_cast<char*>(&offset), 4);
+
+    // 3. CPU 反量化逻辑 (原型)
+    std::vector<float> output(total_elements);
+    for (int64_t i = 0; i < static_cast<int64_t>(packed_weights.size()); ++i) {
+        uint8_t byte = packed_weights[i];
+        
+        // 拆解两个 4-bit 索引
+        uint8_t idxs[2];
+        idxs[0] = byte & 0x0F;         // 低4位
+        idxs[1] = (byte >> 4) & 0x0F;  // 高4位
+
+        for (int j = 0; j < 2; ++j) {
+            int64_t curr_idx = i * 2 + j;
+            if (curr_idx >= total_elements) {
+                continue;
+            }
+            int32_t b_idx = curr_idx / header.blocksize;
+            int32_t g_idx = b_idx / group_size;
+
+            // 双重解量化公式
+            float s1 = half_to_float_cpu(code2[absmax_q[b_idx]]);
+            float s2 = half_to_float_cpu(absmax2[g_idx]);
+            float scale = s1 * s2;
+
+            output[curr_idx] = NF4_TABLE[idxs[j]] * scale + offset;
+        }
+    }
+
+    // 4. 验证结果 (读取 gt_output.bin)
+    std::vector<uint16_t> gt(total_elements);
+    std::ifstream gfs("gt_output.bin", std::ios::binary);
+    gfs.read(reinterpret_cast<char*>(gt.data()), total_elements * 2);
+
+    float max_error = 0;
+    for(int i=0; i<total_elements; ++i) {
+        max_error = std::max(max_error, std::abs(output[i] - half_to_float_cpu(gt[i])));
+    }
+    std::cout << "Max Absolute Error: " << max_error << std::endl;
+
+    return 0;
+}
\ No newline at end of file
diff --git a/03_nf4_dequant/yuruwind/gen_data.py b/03_nf4_dequant/yuruwind/gen_data.py
new file mode 100755
index 0000000..7760abd
--- /dev/null
+++ b/03_nf4_dequant/yuruwind/gen_data.py
@@ -0,0 +1,85 @@
+#!/usr/bin/python3
+import numpy as np
+import struct
+import argparse
+
+# NF4 标准查找表 
+NF4_TABLE = np.array([
+    -1.0, -0.69487101, -0.51209301, -0.37391701,
+    -0.25611401, -0.14725500, -0.04162400, 0.06282201,
+    0.16859101, 0.28551400, 0.40619302, 0.53675699,
+    0.68502200, 0.87091398, 1.0, 0.0
+], dtype=np.float32)
+
+def gen_test_data(rows, cols, blocksize=64, group_size=256):
+    num_elements = rows * cols
+    num_blocks = (num_elements + blocksize - 1) // blocksize
+    num_groups = (num_blocks + group_size - 1) // group_size
+
+    # 随机生成 4-bit 索引 (0-15)
+    indices = np.random.randint(0, 16, size=num_elements, dtype=np.uint8)
+    
+    # 打包权重: 处理奇数个元素的情况
+    num_bytes = (num_elements + 1) // 2
+    packed_weights = np.zeros(num_bytes, dtype=np.uint8)
+    
+    for i in range(num_elements // 2):
+        packed_weights[i] = (indices[2*i+1] << 4) | (indices[2*i] & 0x0F)
+        
+    # 如果总数是奇数，单独处理最后半个 byte (放在低 4 位)
+    if num_elements % 2 != 0:
+        packed_weights[-1] = indices[-1] & 0x0F
+
+    # 生成二级缩放因子数据
+    # absmax_q (每块一个索引，对应 code2)
+    absmax_q = np.random.randint(0, 256, size=num_blocks, dtype=np.uint8)
+    # code2 (二级码表，256个 fp16)
+    code2 = np.random.randn(256).astype(np.float16)
+    # absmax2 (每组一个 fp16)
+    absmax2 = np.random.randn(num_groups).astype(np.float16)
+    offset = 0.0
+
+    # 计算 Ground Truth (用于验证)
+    gt_weights = np.zeros(num_elements, dtype=np.float32)
+    for i in range(num_elements):
+        b_idx = i // blocksize
+        g_idx = b_idx // group_size
+        scale = f16_to_f32(code2[absmax_q[b_idx]]) * f16_to_f32(absmax2[g_idx])
+        gt_weights[i] = NF4_TABLE[indices[i]] * scale
+
+    with open("input.bin", "wb") as f:
+        # Header: rows(i64), cols(i64), blocksize(i32)
+        f.write(struct.pack("qqi", rows, cols, blocksize))
+        f.write(packed_weights.tobytes())
+        f.write(absmax_q.tobytes())
+        f.write(code2.tobytes())
+        f.write(absmax2.tobytes())
+        f.write(struct.pack("f", offset))
+
+    gt_weights.astype(np.float16).tofile("gt_output.bin")
+    print(f"Successfully generated: {rows}x{cols} ({num_elements} elements)")
+    print("Files saved: input.bin, gt_output.bin")
+
+def f16_to_f32(val): # 辅助转换
+    return np.array(val, dtype=np.float16).astype(np.float32)
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser(description="Generate NF4 test data.")
+    parser.add_argument("rows", type=int, help="Number of rows")
+    parser.add_argument("cols", type=int, help="Number of columns")
+    parser.add_argument(
+        "--blocksize",
+        type=int,
+        default=64,
+        help="Block size (default: 64)",
+    )
+    parser.add_argument(
+        "--group-size",
+        type=int,
+        default=256,
+        help="Group size (default: 256)",
+    )
+
+    args = parser.parse_args()
+
+    gen_test_data(args.rows, args.cols, args.blocksize, args.group_size)
\ No newline at end of file
diff --git a/03_nf4_dequant/yuruwind/img/image-20260315174840901.png b/03_nf4_dequant/yuruwind/img/image-20260315174840901.png
new file mode 100644
index 0000000..965dcd2
Binary files /dev/null and b/03_nf4_dequant/yuruwind/img/image-20260315174840901.png differ
diff --git "a/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_181513.png" "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_181513.png"
new file mode 100644
index 0000000..3bf3313
Binary files /dev/null and "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_181513.png" differ
diff --git "a/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182006.png" "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182006.png"
new file mode 100644
index 0000000..f206dee
Binary files /dev/null and "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182006.png" differ
diff --git "a/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182040.png" "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182040.png"
new file mode 100644
index 0000000..35e9a9a
Binary files /dev/null and "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182040.png" differ
diff --git "a/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182110.png" "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182110.png"
new file mode 100644
index 0000000..02d591c
Binary files /dev/null and "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182110.png" differ
diff --git "a/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182718.png" "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182718.png"
new file mode 100644
index 0000000..97a4dcf
Binary files /dev/null and "b/03_nf4_dequant/yuruwind/img/\345\261\217\345\271\225\346\210\252\345\233\276_20260315_182718.png" differ
diff --git a/03_nf4_dequant/yuruwind/nf4_v3.cu b/03_nf4_dequant/yuruwind/nf4_v3.cu
new file mode 100644
index 0000000..6d69af3
--- /dev/null
+++ b/03_nf4_dequant/yuruwind/nf4_v3.cu
@@ -0,0 +1,241 @@
+#include <cuda_runtime.h>
+#include <cuda_fp16.h>
+#include <iostream>
+#include <vector>
+#include <fstream>
+#include <sstream>
+#include <stdint.h>
+#include <algorithm>
+#include <cmath>
+
+#define CHECK_CUDA(call) { \
+    cudaError_t err = call; \
+    if (err != cudaSuccess) { \
+        printf("CUDA Error: %s at line %d\n", cudaGetErrorString(err), __LINE__); \
+        exit(1); \
+    } \
+}
+
+// 配置结构体
+struct AppConfig {
+    int blocksize = 64;
+    std::string compute_type = "fp16";
+    std::string target_gpu = "T4";
+};
+
+// 配置解析函数
+AppConfig load_config(const std::string& filename) {
+    AppConfig config;
+    std::ifstream f(filename);
+    if (!f.is_open()) {
+        std::cout << "[Config] No config.txt found, using defaults.\n";
+        return config;
+    }
+
+    std::string line;
+    while (std::getline(f, line)) {
+        std::replace(line.begin(), line.end(), '=', ' ');
+        std::stringstream ss(line);
+        std::string key, value;
+        if (ss >> key >> value) {
+            if (key == "blocksize") config.blocksize = std::stoi(value);
+            else if (key == "compute_type") config.compute_type = value;
+            else if (key == "target_gpu") config.target_gpu = value;
+        }
+    }
+    std::cout << "[Config] Loaded: compute_type=" << config.compute_type 
+              << ", target_gpu=" << config.target_gpu << "\n";
+    return config;
+}
+
+// NF4 查找表：放入 __constant__ 内存
+__constant__ float d_NF4_TABLE[16] = {
+    -1.0f, -0.69487101f, -0.51209301f, -0.37391701f,
+    -0.25611401f, -0.14725500f, -0.04162400f, 0.06282201f,
+    0.16859101f, 0.28551400f, 0.40619302f, 0.53675699f,
+    0.68502200f, 0.87091398f, 1.0f, 0.0f
+};
+
+// V3 Kernel: Shared Memory + 向量化访存
+__global__ void dequantize_nf4_kernel_v3(
+    const uint8_t* packed_w,
+    const uint8_t* absmax_q,
+    const half* code2,
+    const half* absmax2,
+    half2* output,
+    int64_t total_elements,
+    int block_size,
+    int group_size
+) {
+    // 1. 将 code2 放入共享内存
+    __shared__ half s_code2[256];
+    int tid = threadIdx.x;
+    if (tid < 256) {
+        s_code2[tid] = code2[tid];
+    }
+    __syncthreads();
+
+    int64_t byte_idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
+    if (byte_idx >= (total_elements + 1) / 2) return;
+
+    // 解包 1 个 byte 为两个 4-bit 索引
+    uint8_t byte = packed_w[byte_idx];
+    
+    int64_t element_idx = byte_idx * 2; 
+    int32_t b_idx = element_idx / block_size;
+    int32_t g_idx = b_idx / group_size;
+
+    float s1 = __half2float(s_code2[absmax_q[b_idx]]);
+    float s2 = __half2float(absmax2[g_idx]);
+    float scale = s1 * s2;
+
+    // 向量化写入
+    half res0 = __float2half(d_NF4_TABLE[byte & 0x0F] * scale);
+    half res1 = __float2half(d_NF4_TABLE[byte >> 4] * scale);
+
+    // 边界保护与写入
+    if (byte_idx * 2 + 1 < total_elements) {
+        output[byte_idx] = make_half2(res0, res1);
+    } else if (byte_idx * 2 < total_elements) {
+        // 边界情况：处理奇数长度矩阵的最后一个元素，退化为 16-bit 标量写入
+        reinterpret_cast<half*>(output)[byte_idx * 2] = res0;
+    }
+
+}
+
+int main() {
+    AppConfig cfg = load_config("config.txt");
+
+    std::ifstream ifs("input.bin", std::ios::binary);
+    if (!ifs) { std::cerr << "Cannot open input.bin\n"; return 1; }
+    
+
+    int64_t num_rows, num_cols;
+    int32_t blocksize;
+    ifs.read((char*)&num_rows, 8);
+    ifs.read((char*)&num_cols, 8);
+    ifs.read((char*)&blocksize, 4);
+
+    int64_t total_elements = num_rows * num_cols;
+    int32_t num_blocks = (total_elements + blocksize - 1) / blocksize;
+    int32_t group_size = 256;
+    int32_t num_groups = (num_blocks + group_size - 1) / group_size;
+
+    std::vector<uint8_t> h_packed_w((total_elements + 1) / 2);
+    std::vector<uint8_t> h_absmax_q(num_blocks);
+    std::vector<half> h_code2(256);
+    std::vector<half> h_absmax2(num_groups);
+    float offset;
+
+    ifs.read((char*)h_packed_w.data(), h_packed_w.size());
+    ifs.read((char*)h_absmax_q.data(), h_absmax_q.size());
+    ifs.read((char*)h_code2.data(), 256 * 2);
+    ifs.read((char*)h_absmax2.data(), num_groups * 2);
+    ifs.read((char*)&offset, 4);
+
+    if (cfg.blocksize != blocksize) {
+        printf("[Warning] Binary header blocksize (%d) differs from config.txt (%d). "
+               "Using Binary Header.\n", blocksize, cfg.blocksize);
+    }
+
+    int threads_per_block = 256; // 默认值
+    if (cfg.target_gpu == "T4") {
+        // T4 (Turing) 架构 SM 较小，128 是更合适的选择
+        threads_per_block = 128;
+    } else {
+        //  4060 或 A100 等现代显卡，256 是甜点值
+        threads_per_block = 256;
+    }
+
+    uint8_t *d_packed_w, *d_absmax_q;
+    half *d_code2, *d_absmax2, *d_output;
+    CHECK_CUDA(cudaMalloc(&d_packed_w, h_packed_w.size()));
+    CHECK_CUDA(cudaMalloc(&d_absmax_q, h_absmax_q.size()));
+    CHECK_CUDA(cudaMalloc(&d_code2, 256 * 2));
+    CHECK_CUDA(cudaMalloc(&d_absmax2, num_groups * 2));
+    CHECK_CUDA(cudaMalloc(&d_output, total_elements * 2));
+
+    CHECK_CUDA(cudaMemcpy(d_packed_w, h_packed_w.data(), h_packed_w.size(), cudaMemcpyHostToDevice));
+    CHECK_CUDA(cudaMemcpy(d_absmax_q, h_absmax_q.data(), h_absmax_q.size(), cudaMemcpyHostToDevice));
+    CHECK_CUDA(cudaMemcpy(d_code2, h_code2.data(), 256 * 2, cudaMemcpyHostToDevice));
+    CHECK_CUDA(cudaMemcpy(d_absmax2, h_absmax2.data(), num_groups * 2, cudaMemcpyHostToDevice));
+
+    // 计时
+    int64_t num_bytes = (total_elements + 1) / 2;
+    int blocks = (num_bytes + threads_per_block - 1) / threads_per_block;
+
+    if (cfg.compute_type == "bf16") {
+        // 打印 Dispatch 日志
+        std::cout << "[Dispatch] Launching Kernel with BF16 precision path (simulated)...\n";
+    } else {
+        std::cout << "[Dispatch] Launching Kernel with FP16 precision path...\n";
+    }
+
+    cudaEvent_t start, stop;
+    cudaEventCreate(&start);
+    cudaEventCreate(&stop);
+
+    // 预热 GPU (避免把 cuda context 初始化算进时间)
+    dequantize_nf4_kernel_v3<<<blocks, threads_per_block>>>(
+        d_packed_w, d_absmax_q, d_code2, d_absmax2, 
+        (half2*)d_output, total_elements, blocksize, group_size
+    );
+    cudaDeviceSynchronize();
+
+    // 正式计时 (跑 10 次取平均)
+    cudaEventRecord(start);
+    for(int i = 0; i < 10; ++i) { 
+        dequantize_nf4_kernel_v3<<<blocks, threads_per_block>>>(
+            d_packed_w, d_absmax_q, d_code2, d_absmax2, 
+            (half2*)d_output, total_elements, blocksize, group_size
+        );
+    }
+    cudaEventRecord(stop);
+    cudaEventSynchronize(stop);
+    
+    float ms = 0;
+    cudaEventElapsedTime(&ms, start, stop);
+    ms /= 10.0f; // 取 10 次的平均值
+
+// 验证结果
+    std::vector<half> h_output(total_elements);
+    CHECK_CUDA(cudaMemcpy(h_output.data(), d_output, total_elements * 2, cudaMemcpyDeviceToHost));
+
+    // 读取 Ground Truth
+    std::vector<half> h_gt(total_elements);
+    std::ifstream gfs("gt_output.bin", std::ios::binary);
+    gfs.read((char*)h_gt.data(), total_elements * 2);
+
+    float max_err = 0;
+    double total_abs_err = 0;
+    for(int i=0; i<total_elements; ++i) {
+        float out_f = __half2float(h_output[i]);
+        float gt_f = __half2float(h_gt[i]);
+        float diff = std::abs(out_f - gt_f);
+        
+        max_err = std::max(max_err, diff);
+        total_abs_err += diff;
+    }
+    float mae = static_cast<float>(total_abs_err / total_elements);
+
+    std::cout << "--- Validation Results ---" << std::endl;
+    std::cout << "CUDA V3 Max Error: " << max_err << std::endl;
+    std::cout << "CUDA V3 MAE:       " << mae << std::endl;
+    std::cout << "Time:              " << ms << " ms" << std::endl;
+    
+    // 计算有效内存带宽
+    // 读取: W(4bit) + absmax_q(8bit) + absmax2(16bit)
+    // 写入: Output(16bit)
+    // 忽略 code2 (已经放入 Shared Memory) 的重复读取开销
+    double bytes_read = (total_elements * 0.5) + num_blocks + (num_groups * 2.0);
+    double bytes_write = total_elements * 2.0;
+    double total_bytes = bytes_read + bytes_write;
+    
+    std::cout << "Bandwidth: " << total_bytes / (ms * 1e6) << " GB/s" << std::endl;
+
+    // 清理
+    cudaEventDestroy(start);
+    cudaEventDestroy(stop);
+    cudaFree(d_packed_w); cudaFree(d_absmax_q); cudaFree(d_code2); cudaFree(d_absmax2); cudaFree(d_output);
+    return 0;
+}
\ No newline at end of file