项目二+项目三

project2-report.md

@@ -0,0 +1,130 @@
+# 项目二报告：基于CUDA的推理引擎实现
+
+## 一、 概述
+
+本项目旨在为大语言模型推理框架 `LLAiSYS` 构建底层的 CUDA 算子库。大语言模型（如 Qwen2 系列）的自回归推理过程高度依赖 GPU 的并行计算能力与显存吞吐率。为满足框架在 NVIDIA GPU 上的运行需求，本项目基于 CUDA C++ 实现了模型推理所需的全部核心算子，涵盖 `Add`、`SwiGLU`、`RMSNorm`、`Linear`、`Embedding`、`RoPE`、`Argmax` 以及 `Self-Attention`。
+
+该算子库原生支持 `FP32`、`FP16` 及 `BF16` 数据类型。所有算子均通过了与 PyTorch 原生实现的精度对比测试，并在端到端推理验证中实现了输出 Token 序列的 100% 精度对齐，为上层推理服务提供了可靠的算力基础。
+
+## 二、 运行环境
+
+- **硬件平台**：NVIDIA GPU（测试环境基于 A100 Tensor Core GPU x8）
+- **操作系统**：Linux / Windows 跨平台支持
+- **核心语言**：C++ 17, CUDA C++, Python 3.10+
+- **构建系统**：Xmake
+- **依赖库**：CUDA Toolkit, cuBLAS (NVIDIA Basic Linear Algebra Subprograms)
+- **验证基准**：PyTorch 2.x
+
+## 三、 核心架构与具体实现
+
+### 3.1 算子构建与链接架构
+
+在算子库构建初期，C++ 静态库之间的循环依赖导致了 CUDA 宏注入失败与符号丢失（Undefined Symbol）问题。本项目通过重构 `xmake.lua`，将算子的编译与链接权限上移至动态链接库（`libllaisys.so` / `.dll`），实现了多级依赖环境下的 CUDA 文件安全编译。
+
+**核心构建配置示例 (`xmake.lua`)：**
+
+```lua
+target("llaisys")
+    set_kind("shared")
+    add_deps("llaisys-utils", "llaisys-device", "llaisys-core", "llaisys-tensor", "llaisys-ops", "llaisys-models") 
+
+    if has_config("nv-gpu") then
+        add_rules("cuda")
+        -- 将算子的 CUDA 实现统一交由拥有一切依赖的上层动态库编译
+        add_files("src/ops/*/nvidia/*.cpp", "src/ops/*/nvidia/*.cu")
+    end
+target_end()
+```
+
+### 3.2 核心算子实现细节
+
+#### 1. 基础并行计算 (Add & SwiGLU)
+
+对于 Element-wise（逐元素）操作，采用网格跨步循环（Grid-Stride Loop）以适配任意长度的 Tensor。在处理 `FP16` 和 `BF16` 类型时，通过寄存器级别的类型转换（如 `__half2float`），将数据提升至单精度进行非线性计算，以保证数值稳定性。
+
+#### 2. 并行规约算子 (RMSNorm & Argmax)
+
+规约（Reduction）操作是典型的显存带宽瓶颈。
+
+- **RMSNorm**：采用 Block 级别的并行计算。为每个 Token 分配一个 Thread Block，利用 `__shared__` 内存进行块内规约求平方和，并使用 CUDA 硬件指令 `rsqrtf` 计算均方根倒数。
+- **Argmax**：为应对输出层巨大的词表维度，通过维护线程局部极值（`local_max`）与局部索引（`local_idx`），随后在单 Block 内通过共享内存规约出全局极值。
+
+#### 3. 矩阵乘法 (Linear)
+
+大语言模型的全连接层运算由 `cuBLAS` 库接管，以充分利用 GPU 的 Tensor Cores。由于 C/C++ 采用行优先（Row-Major）存储，而 cuBLAS 基于列优先（Column-Major），系统利用转置等价性公式 $(AB^T)^T = BA^T$ 设置 `CUBLAS_OP_T` 参数，实现了零拷贝的矩阵乘法，随后通过轻量级 Kernel 注入偏置项（Bias）。
+
+#### 4. 显存寻址与相对位置编码 (Embedding & RoPE)
+
+在索引寻址类算子中，系统严格规范了数据类型对齐与内存步长：
+
+- **数据类型对齐**：确保 `index` / `pos_ids` 强制使用 `int64_t` 指针进行解引用，避免因 Python 端与 C++ 端类型错位导致的内存越界。
+- **RoPE 内存步长**：严格匹配标准 PyTorch 的 RoPE 语义，将特征维度切分为前后两半（`half_dim`），对间隔 `half_dim` 的元素对执行复数旋转。
+
+**RoPE 核心寻址逻辑片段：**
+
+```c++
+// 计算内存偏移，前一半与后一半组合为一对复数
+size_t idx_a = seq_idx * (nhead * head_dim) + head_idx * head_dim + pair_idx;
+size_t idx_b = idx_a + half_dim;
+
+float x0 = in[idx_a];
+float x1 = in[idx_b];
+out[idx_a] = x0 * cos_m - x1 * sin_m;
+out[idx_b] = x1 * cos_m + x0 * sin_m;
+```
+
+#### 5. 分组查询自注意力 (Self-Attention with GQA)
+
+该算子采用分块并行（Block-level Parallelism）设计，Grid 维度设定为 `[seqlen, nhead]`，使每个 Block 独立处理单个 Query 向量。
+
+- **动态共享内存**：在 Block 内部申请长度为 `total_len` 的动态共享内存 `extern __shared__ float scores[]`。
+- **Softmax 融合与 Causal Mask**：计算点积时，通过索引比对将未来位置的分数置为负无穷（`-1e20f`）。点积完成后，直接在共享内存中就地执行 Softmax 操作并与 Value 矩阵进行加权求和，避免了中间结果落入全局显存。
+
+## 四、 构建与测试
+
+### 4.1 项目构建说明
+
+本项目依赖 `xmake` 工具进行工程管理与编译。构建全量带有 NVIDIA GPU 支持的动态链接库与 Python 包包，执行以下标准流程：
+
+Bash
+
+```
+# 清理构建缓存并重新配置 GPU 编译选项
+xmake clean -a
+xmake f -c --nv-gpu=y
+# 编译并生成共享库
+xmake -r install
+# 将生成的 C++ 库注册至 Python 环境
+pip install ./python/
+```
+
+### 4.2 算子单元测试
+
+框架针对各算子实现了独立的 Python 测试脚本。测试脚本基于 `ctypes` 调用生成的 `libllaisys.so` 接口，并采用 PyTorch 同等运算作为对照组，利用 `torch.allclose` 验证 `atol` 与 `rtol`。
+
+```Bash
+python test/ops/add.py --device nvidia
+python test/ops/swiglu.py --device nvidia
+python test/ops/rms_norm.py --device nvidia
+python test/ops/rope.py --device nvidia
+python test/ops/self_attention.py --device nvidia
+```
+
+所有算子均能稳定通过全精度（F32）、半精度（F16/BF16）的边界测试与精度校验。
+
+### 4.3 端到端推理验证
+
+在单算子验证通过的基础上，对 Hugging Face 开源模型 `deepseek-ai/DeepSeek-R1-Distill-Qwen-1.5B` 进行自回归生成测试，对比原生 PyTorch 推理结果与 LLAiSYS 推理结果。
+
+**验证指令与结果示例：**
+
+```Bash
+python test/test_infer.py --device nvidia
+```
+
+- **输出比对**：LLAiSYS 输出的 Token ID 序列与 PyTorch 生成的 Token 序列一致。
+- **功能验证**：模型能够成功载入权重配置，正确处理 KV Cache 状态，并连续生成逻辑连贯的文本。
+
+## 五、 结论
+
+本项目成功在 LLAiSYS 框架中构建了底层 CUDA 算子生态。通过解决多级依赖构建、张量内存排布、类型边界控制等关键工程问题，实现了大语言模型所需全套核心算子的高效 GPU 并行计算。

project3-report.md

@@ -0,0 +1,131 @@
+# 项目三报告：基于 FastAPI 的大模型推理 API 服务与 Web 界面集成
+
+## 一、 概述
+
+本项目是 `LLAiSYS` 大语言模型推理框架的顶层工程应用。在项目一和项目二实现了底层 Tensor 内存管理与高效 CUDA 推理算子的基础上，本项目旨在打破底层 C++/CUDA 代码与终端用户之间的交互壁垒，将其构建为一个现代化的 AI Web 服务。
+
+本项目基于 Python 的高性能异步 Web 框架 `FastAPI`，设计并实现了一套完全兼容 OpenAI 官方标准定义（`/v1/chat/completions`）的 RESTful API 接口。同时，利用 Server-Sent Events (SSE) 技术实现了模型推理过程的实时流式（Streaming）输出，并成功将我们的本地推理引擎与业界主流的开源前端 UI `ChatGPT-Next-Web`（NextChat）无缝集成，最终交付了一个端到端的完整大语言模型对话系统。
+
+## 二、 运行环境
+
+- **硬件平台**：NVIDIA GPU / CPU
+- **操作系统**：Linux / Windows 跨平台支持
+- **核心语言**：Python 3.10+, TypeScript (前端)
+- **后端依赖库**：`fastapi`, `uvicorn`, `pydantic`, `sse-starlette`, `transformers`, `huggingface_hub`
+- **前端系统**：`ChatGPT-Next-Web` (基于 Next.js 与 React 构建)
+- **测试模型**：`deepseek-ai/DeepSeek-R1-Distill-Qwen-1.5B` (Qwen2 架构)
+
+## 三、 核心架构与具体实现
+
+### 3.1 OpenAI 兼容协议设计与数据模型
+
+为了使我们的推理框架能够直接被市面上成熟的第三方 AI 客户端（如 NextChat、LobeChat 等）调用，服务端必须严格遵守 OpenAI 的接口规范。本项目使用 `pydantic` 构建了严谨的数据校验模型。
+
+**核心数据结构实现：**
+
+```Python
+class ChatMessage(BaseModel):
+    role: str
+    content: str
+
+class ChatCompletionRequest(BaseModel):
+    model: str = "qwen2"
+    messages: List[ChatMessage]
+    stream: Optional[bool] = False
+    max_tokens: Optional[int] = 512
+    temperature: Optional[float] = 0.8
+    top_p: Optional[float] = 0.8
+    top_k: Optional[int] = 50
+```
+
+### 3.2 Prompt 模板化与 Tokenizer 接入
+
+大语言模型的对话能力高度依赖于特定的特殊占位符（如 `<|im_start|>`、`<|im_end|>` 等）。本项目通过引入 Hugging Face 的 `transformers.AutoTokenizer`，并调用其 `apply_chat_template` 机制，将前端传入的 JSON 格式历史对话上下文，自动拼装成符合 Qwen2 模型底层训练格式的单行 Prompt 字符串，随后进行 Encode 转换为 `input_ids` 序列，送入底层的 C++ `llaisys.models.Qwen2` 引擎进行推理。
+
+### 3.3 Server-Sent Events (SSE) 流式输出机制
+
+大模型自回归生成的特性决定了如果采用同步阻塞返回，用户将面临巨大的等待延迟。本项目基于 FastAPI 的 `StreamingResponse` 实现了实时流式传输。
+
+**流式生成核心逻辑：**
+
+```Python
+async def generate_stream():
+    # 执行模型自回归生成新 Token
+    for token_id in new_tokens:
+        if token_id == tokenizer.eos_token_id:
+            break
+        # 解码单步 Token 为文本
+        word = tokenizer.decode([token_id], skip_special_tokens=True)
+        # 组装符合 OpenAI 规范的 SSE 数据块
+        chunk = {
+            "id": f"chatcmpl-{int(time.time())}",
+            "object": "chat.completion.chunk",
+            "choices": [{"delta": {"content": word}}]
+        }
+        yield f"data: {json.dumps(chunk)}\n\n"
+        time.sleep(0.02) # 模拟平滑输出的打字机效果
+
+    yield "data: [DONE]\n\n"
+```
+
+通过上述生成器，底层的每一次 C++ `llaisysQwen2ModelInfer` 调用结果都能被瞬间推送到前端，极大地提升了系统的人机交互体验。
+
+### 3.4 跨域资源共享 (CORS) 与网络安全配置
+
+由于前端 Web 页面（如跑在 `localhost:3000` 的 NextChat）与 FastAPI 后端服务（运行于 `0.0.0.0:8199`）通常处于不同的端口或域名下，浏览器会触发跨域安全拦截（发送 `OPTIONS` 预检请求）。本项目通过配置 FastAPI 的 `CORSMiddleware` 中间件，全面放行了前端发起的跨域请求，彻底解决了 `405 Method Not Allowed` 的网络拦截问题。
+
+## 四、 构建与测试
+
+### 4.1 启动推理后端服务
+
+本项目使用 `argparse` 暴露了启动配置参数，支持自动从 Hugging Face Hub 下载并缓存模型权重，使用 `uvicorn` 承载 ASGI HTTP 服务。
+
+```Bash
+python python/server.py --device nvidia --port 8199
+```
+
+启动成功后，终端将输出 `[READY] Server starting on http://0.0.0.0:8199`。
+
+**终端图像：**
+
+![image-20260317000545784](assets/image-20260317000545784.png)
+
+### 4.2 接入 ChatGPT-Next-Web 前端进行验证
+
+本项目选择业界主流的开源大模型前端 `ChatGPT-Next-Web`（又称 NextChat）作为可视化交互界面。通过该 UI 验证了推理后端对 OpenAI 标准协议的兼容性及流式传输的稳定性。具体操作流程如下：
+
+#### 1. 前端环境部署
+
+- **访问方式**：可直接访问 NextChat 的 Web 托管版本，例如：https://app.nextchat.club/#/chat，或通过 Docker 分布在本地 `3000` 端口。
+- **通信准备**：确保前端所在的浏览器环境能够访问到 `server.py` 运行的后端 IP 地址及端口（如 `http://127.0.0.1:8199`）。
+
+#### 2. 自定义接口配置 (Settings)
+
+进入 NextChat 左下角的“设置”面板，进行以下关键参数的绑定：
+
+- **模型服务商 (Model Provider)**：选择 `OpenAI`。
+- **自定义接口地址 (Endpoint URL)**：填写我们的 FastAPI 服务端地址 `http://127.0.0.1:8199`。
+  - *注：NextChat 会自动在末尾拼接 `/v1/chat/completions` 路径。*
+- **API Key**：由于本地测试未开启鉴权，此处可随意填写（如 `sk-llaisys`），以绕过前端的非空校验。
+- **自定义模型 (Custom Models)**：在自定义模型列表中输入 `qwen2` 并添加。
+
+#### 3. 核心交互验证
+
+- **模型切换**：在聊天窗口顶部下拉菜单中选中刚才添加的 `qwen2` 模型。
+- **流式生成测试**：在输入框发送长文本问题（如“请写一段 200 字左右关于人工智能的介绍”）。
+- **响应观察**：
+  - **SSE 验证**：观察文字是否以“打字机”效果逐个跳出。这证明了后端的 `StreamingResponse` 正在实时推送 Token，而非等待生成结束后一次性返回。
+  - **标题自动总结**：NextChat 会在对话开始后自动发送一个 `stream: False` 的后台请求。验证左侧历史记录栏是否成功根据模型回复生成了简短标题，这证明了后端对非流式 JSON 响应格式的正确处理。
+  - **CORS 预检**：通过浏览器开发者工具（F12）观察，确认浏览器发出的 `OPTIONS` 预检请求已被 FastAPI 成功拦截并允许跨域，从而保证了 `POST` 请求的顺利下发。
+
+#### 4. 交互原理示意
+
+`用户输入` $\rightarrow$ `NextChat UI (JSON 封装)` $\rightarrow$ `HTTP POST 请求` $\rightarrow$ `FastAPI 后端 (路由解析)` $\rightarrow$ `LLAiSYS C++ 引擎` $\rightarrow$ `GPU 并行计算` $\rightarrow$ `SSE 流式写回` $\rightarrow$ `前端 Markdown 渲染`。
+
+**演示图像：**
+
+![image-20260317000612224](assets/image-20260317000612224.png)
+
+## 五、 结论
+
+本项目成功为 `LLAiSYS` 框架构建了应用层的服务端基础设施。通过实现标准化的 OpenAI API 协议、跨域中间件以及 SSE 流式传输机制，不仅使底层的 C++ 算子引擎具备了作为云端微服务独立运行的能力，还实现了与业界主流 Web UI 的零成本集成。至此，本系统已具备了从底层内存分配到前端可视化交互的完整大模型基础设施能力。

python/llaisys/libllaisys/__init__.py

@@ -12,7 +12,7 @@
 from .tensor import llaisysTensor_t
 from .tensor import load_tensor
 from .ops import load_ops
-
+from .models import load_models
 
 def load_shared_library():
     lib_dir = Path(__file__).parent
@@ -38,7 +38,7 @@ def load_shared_library():
 load_runtime(LIB_LLAISYS)
 load_tensor(LIB_LLAISYS)
 load_ops(LIB_LLAISYS)
-
+load_models(LIB_LLAISYS)
 
 __all__ = [
     "LIB_LLAISYS",

python/llaisys/libllaisys/models.py

@@ -0,0 +1,57 @@
+import ctypes
+from ctypes import POINTER, c_size_t, c_float, c_int64, c_int, Structure
+from .llaisys_types import llaisysDataType_t, llaisysDeviceType_t
+from .tensor import llaisysTensor_t
+
+class LlaisysQwen2Meta(Structure):
+    _fields_ = [
+        ("dtype", llaisysDataType_t),
+        ("nlayer", c_size_t),
+        ("hs", c_size_t),
+        ("nh", c_size_t),
+        ("nkvh", c_size_t),
+        ("dh", c_size_t),
+        ("di", c_size_t),
+        ("maxseq", c_size_t),
+        ("voc", c_size_t),
+        ("epsilon", c_float),
+        ("theta", c_float),
+        ("end_token", c_int64),
+    ]
+
+class LlaisysQwen2Weights(Structure):
+    _fields_ = [
+        ("in_embed", llaisysTensor_t),
+        ("out_embed", llaisysTensor_t),
+        ("out_norm_w", llaisysTensor_t),
+        ("attn_norm_w", POINTER(llaisysTensor_t)),
+        ("attn_q_w", POINTER(llaisysTensor_t)),
+        ("attn_q_b", POINTER(llaisysTensor_t)),
+        ("attn_k_w", POINTER(llaisysTensor_t)),
+        ("attn_k_b", POINTER(llaisysTensor_t)),
+        ("attn_v_w", POINTER(llaisysTensor_t)),
+        ("attn_v_b", POINTER(llaisysTensor_t)),
+        ("attn_o_w", POINTER(llaisysTensor_t)),
+        ("mlp_norm_w", POINTER(llaisysTensor_t)),
+        ("mlp_gate_w", POINTER(llaisysTensor_t)),
+        ("mlp_up_w", POINTER(llaisysTensor_t)),
+        ("mlp_down_w", POINTER(llaisysTensor_t)),
+    ]
+
+class LlaisysQwen2Model(Structure):
+    pass
+
+llaisysQwen2Model_t = POINTER(LlaisysQwen2Model)
+
+def load_models(lib):
+    lib.llaisysQwen2ModelCreate.argtypes = [POINTER(LlaisysQwen2Meta), llaisysDeviceType_t, POINTER(c_int), c_int]
+    lib.llaisysQwen2ModelCreate.restype = llaisysQwen2Model_t
+
+    lib.llaisysQwen2ModelDestroy.argtypes = [llaisysQwen2Model_t]
+    lib.llaisysQwen2ModelDestroy.restype = None
+
+    lib.llaisysQwen2ModelWeights.argtypes = [llaisysQwen2Model_t]
+    lib.llaisysQwen2ModelWeights.restype = POINTER(LlaisysQwen2Weights)
+
+    lib.llaisysQwen2ModelInfer.argtypes = [llaisysQwen2Model_t, POINTER(c_int64), c_size_t]
+    lib.llaisysQwen2ModelInfer.restype = c_int64