添加readme

Author: 330-shh

.clang-format

@@ -0,0 +1,27 @@
+---
+Language:        Cpp
+BasedOnStyle:  Google
+Standard:        c++17
+ColumnLimit:     120
+IndentWidth:     4
+
+# CUDA-specific formatting
+SpaceBeforeParens: ControlStatements
+AlignConsecutiveAssignments: true
+AlignConsecutiveDeclarations: true
+DerivePointerAlignment: false
+PointerAlignment: Right
+
+# Keep blocks concise
+AllowShortBlocksOnASingleLine: true
+AllowShortFunctionsOnASingleLine: Inline
+AllowShortIfStatementsOnASingleLine: false
+AllowShortLoopsOnASingleLine: false
+
+# Includes
+SortIncludes: true
+IncludeBlocks: Regroup
+
+# CUDA macro handling
+MacroBlockBegin: "^__global__.*$"
+StatementMacros: ["__syncthreads", "CUDA_CHECK", "CUDA_CHECK_LAST"]

.github/workflows/build.yml

@@ -0,0 +1,33 @@
+name: CUDA Build Check
+
+on:
+  push:
+    branches: [ main, master ]
+  pull_request:
+    branches: [ main, master ]
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    
+    steps:
+    - uses: actions/checkout@v3
+
+    # 由于 Ubuntu 最新环境可能已经自带一定的 CUDA/C++ 环境
+    # 我们使用现成的 action 或者 ubuntu-latest 自带环境安装完整的 CUDA Toolkit
+    - name: Install CUDA Toolkit
+      uses: Jimver/cuda-toolkit@v0.2.11
+      id: cuda-toolkit
+      with:
+        cuda: '12.2.0'
+        
+    - name: Check nvcc
+      run: nvcc -V
+
+    - name: Configure CMake
+      # 这里使用我们项目的根目录 CMakeLists.txt
+      run: cmake -S . -B build
+
+    - name: Build All Projects
+      # 并发编译所有项目
+      run: cmake --build build --parallel 4

.gitignore

@@ -2,6 +2,7 @@
 build/
 docs/
 cmake-build*/
+out/
 *.o
 *.obj
 *.exe
@@ -22,3 +23,20 @@ cmake-build*/
 *.ptx
 *.fatbin
 *.gpu
+
+# Python & PyTorch C++ Extension 产物
+__pycache__/
+*.pyc
+*.pyo
+*.pyd
+build/lib.*
+build/temp.*
+*.egg-info/
+dist/
+
+# AI 助手与工作流存储
+.agents/
+.agent/
+_agents/
+_agent/
+.gemini/

01_Basics/01_vector_add/vector_add.cu

@@ -1,11 +1,5 @@
-#include <iostream>
-#include <vector>
-#include <cmath>
-#include <iomanip>
-#include "cuda_utils.cuh"
-#include "timer.cuh"
+#include <code_abbreviation.h>
 
-using namespace std;
 
 // 向量加法（GPU kernel，手写）
 __global__ void vector_add(const float* A, const float* B, float* C, const int n) {
@@ -23,13 +17,6 @@ void vector_add_cpu(const vector<float>& h_a, const vector<float>& h_b,
     }
 }
 
-// GPU 计时结果结构体（AI 生成）
-struct GpuTimingResult {
-    float h2d_ms;      // Host to Device 传输时间
-    float kernel_ms;   // Kernel 执行时间（多次平均）
-    float d2h_ms;      // Device to Host 传输时间
-    float total_ms;    // 总时间
-};
 
 // 向量加法（GPU，手写）
 GpuTimingResult vector_add_device(const vector<float>& h_a, const vector<float>& h_b,
@@ -96,7 +83,7 @@ bool verify_results(const vector<float>& h_a, const vector<float>& h_b,
     error_count = 0;
     for (int i = 0; i < n; ++i) {
         float expected = h_a[i] + h_b[i];
-        if (std::fabs(h_c[i] - expected) > 1e-5) {
+        if (fabs(h_c[i] - expected) > 1e-5) {
             error_count++;
             // 打印前 5 个错误
             if (error_count <= 5) {

01_Basics/02_matrix_mul_naive/matrix_mul_naive.cu

@@ -1,14 +1,8 @@
-#include <iostream>
-#include <vector>
-#include <cmath>
-#include <iomanip>
-#include "cuda_utils.cuh"
-#include "timer.cuh"
-
-using namespace std;
-using matrix = vector<float>;           // 行主序
-using Int = const int;
-constexpr size_t sizeF = sizeof(float);
+#include <code_abbreviation.h>
+
+using matrix = Matrix;            // 兼容本文件中的 'matrix' 别名
+using Int = CInt;                 // 兼容本文件中的 'Int' 别名
+constexpr size_t sizeF = FSIZE;   // 兼容本文件中的 'sizeF' 别名
 
 // 矩阵乘法（GPU kernel，手写）
 __global__ void matrix_mul_naive(const float* A, const float* B, float* C, Int m, Int n, Int k){
@@ -36,13 +30,6 @@ void matrix_mul_cpu(const float* A, const float* B, float* C, Int m, Int n, Int
     }
 }
 
-// GPU 计时结果结构体（AI 生成）
-struct GpuTimingResult {
-    float h2d_ms;      // Host to Device 传输时间
-    float kernel_ms;   // Kernel 执行时间（多次平均）
-    float d2h_ms;      // Device to Host 传输时间
-    float total_ms;    // 总时间
-};
 
 // 矩阵乘法（GPU，手写）
 GpuTimingResult matrix_mul_naive_device(const matrix& h_a, const matrix& h_b, matrix& h_c, Int m, Int n, Int k, Int iterations = 100) {

01_Basics/03_matrix_mul_tiled/matrix_mul_tiled.cu

@@ -1,20 +1,6 @@
-#include <iostream>
-#include <vector>
-#include <cmath>
-#include <iomanip>
-#include "cuda_utils.cuh"
-#include "timer.cuh"
-
-using namespace std;
-using CInt = const int;
-using CSize = const size_t;
-using CPFloat = const float*;
-using PFloat = float*;
-using CRMatrix = const vector<float>&;
-using RMatrix = vector<float>&;
-using Matrix = vector<float>;
+#include <code_abbreviation.h>
+
 constexpr int TILE_WIDTH = 32;
-constexpr size_t FSIZE = sizeof(float);
 
 // 矩阵乘法-使用共享内存分块优化（GPU kernel，手写）
 __global__ void matrix_mul_tiled(CPFloat a, CPFloat b, PFloat c, CInt m, CInt n, CInt k) {
@@ -78,13 +64,6 @@ void matrix_mul_cpu(CPFloat a, CPFloat b, PFloat c, CInt m, CInt n, CInt k) {
     }
 }
 
-// GPU 计时结果结构体（AI 生成）
-struct GpuTimingResult {
-    float h2d_ms;      // Host to Device 传输时间
-    float kernel_ms;   // Kernel 执行时间（多次平均）
-    float d2h_ms;      // Device to Host 传输时间
-    float total_ms;    // 总时间
-};
 
 // 矩阵乘法（GPU，手写）
 GpuTimingResult matrix_mul_tiled_device(CRMatrix a, CRMatrix b, RMatrix c, CInt m, CInt n, CInt k, CInt iterations = 100) {

01_Basics/README.md

@@ -0,0 +1,83 @@
+# 01_Basics: CUDA 基础编程与执行模型
+
+## 1. 全景导览与学习目标 (Overview & Learning Objectives)
+
+本章是 CUDA 学习的起点，旨在让开发者熟悉 CUDA 的异构编程模型、线程层级结构（Grid, Block, Thread）以及最基础的全局内存（Global Memory）访问。它涵盖了如何从 CPU 侧向 GPU 发起计算请求，以及 GPU 计算内核（Kernel）的基本编写范式。这些是最基础也是最核心的 CUDA 概念。
+
+目录下的实现逐步加深对计算与内存瓶颈的理解：
+
+- `01_vector_add/`：演示最基本的单维度网格与线程块划分，以及内存的 Allocate、Copy 与 Free 流程。
+- `02_matrix_mul_naive/`：演示二维线程网格映射，实现基础的矩阵乘法，暴露出未优化的全局内存的冗余访存问题。
+- `03_matrix_mul_tiled/`：引入基于共享内存（Shared Memory）的分块（Tiling）策略，大幅减少全局内存宽带的浪费，为后续更高级的优化打下基础。
+
+## 2. 原理推导与数学表达 (Math & Logic)
+
+向量加法是最纯粹的 Element-wise 操作：
+$$ C_i = A_i + B_i $$
+
+对于矩阵乘法 $C = A \times B$（假设矩阵大小为 $M \times K$ 和 $K \times N$），目标元素计算式为：
+$$ C_{i, j} = \sum_{k=0}^{K-1} A_{i, k} \cdot B_{k, j} $$
+在朴素实现中，这需要执行 $O(M \cdot N \cdot K)$ 次全局内存访问。而在 Tiling 优化中，我们将大矩阵分解为大小为 $B_s \times B_s$ 的小块，此时访存量缩减至原先的 $1/B_s$（理论上）。
+
+## 3. 硬核内存映射解析 (Memory & Thread Mapping)
+
+以 2D Tiling 矩阵乘法为例，使用 Shared Memory 降低全局内存访存带宽（假定 Block 大小为 $16 \times 16$）：
+
+```text
+[Global Memory] 矩阵 A 与 B
+---------------------------------------------------
+|                                                 |
+|  [Shared Mem 块 `sA` 16x16]                      |
+|  +--------------------+                         |
+|  | Thread 映射区(0,0)  | <-- Block内的(tx, ty)   |
+|  | 到(15, 15)协作加载   |                         |
+|  +--------------------+                         |
+---------------------------------------------------
+           ||
+           \/  (寄存器做内积累加)
+           
+计算局部 $C_{sub}$ 积累至对应的 Global Memory / Register
+```
+
+每个 Block 合作将数据从 Global Memory 读入 Shared Memory 后，需要调用 `__syncthreads()` 确保数据就绪，然后再执行乘累加。
+
+## 4. 关键源码逐行解剖 (Code Deep-Dive)
+
+来自 `03_matrix_mul_tiled/matrix_mul_tiled.cu` 的共享内存同步读取：
+
+```cpp
+// 声明 2D 共享内存，用于存储 A 和 B 的子块，利用高带宽和极低延迟
+__shared__ float s_A[TILE_WIDTH][TILE_WIDTH];
+__shared__ float s_B[TILE_WIDTH][TILE_WIDTH];
+
+// 协作加载：当前 Thread (tx, ty) 负责把全局内存中属于它的那个元素搬到共享内存
+s_A[ty][tx] = A[Row * k + (ph * TILE_WIDTH + tx)];
+s_B[ty][tx] = B[(ph * TILE_WIDTH + ty) * n + Col];
+
+// ⚠️ 极其关键的屏障：必须等该 Block 内部所有 Thread 完成搬运，才能开始计算
+__syncthreads();
+
+// 在共享内存上计算局部的乘积累加
+for (int j = 0; j < TILE_WIDTH; ++j) {
+    Cvalue += s_A[ty][j] * s_B[j][tx];
+}
+
+// ⚠️ 释放屏障：防止迭代过快导致下一轮的加载覆盖了当前还在计算的数据
+__syncthreads();
+```
+
+## 5. 性能基准与分析视角 (Performance & Profiling)
+
+- **基准**：对比 CPU 上的顺序循环矩阵乘法，以及简单的朴素版（Naïve）CUDA 矩阵乘法。
+- **典型分析**：使用 Tiled 分块后，全局显存带宽的使用率大幅度降低。在 NCU 中观察 `sm__throughput` 和 `l1tex__t_sectors_pipe_lsu_mem_global_op_ld` 的比值，明显改善了 Memory Workload，使得计算受限（Compute Bound）的比重增加。
+
+## 6. 编译指引与参考资料 (Compile & References)
+
+```bash
+# 通用编译指令
+nvcc -O3 -arch=sm_89 matrix_mul_tiled.cu -o run_tiled
+# NCU 性能分析指令（提取内存带宽和SM利用率）
+ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,dram__throughput.avg.pct_of_peak_sustained_elapsed ./run_tiled
+```
+
+- 参考资料: [CUDA C++ Programming Guide - Shared Memory](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory)

02_Reduction/01_reduce_sum/reduce_sum.cu

@@ -1,21 +1,5 @@
-#include <iostream>
-#include <vector>
-#include <cmath>
-#include <iomanip>
-#include "cuda_utils.cuh"
-#include "timer.cuh"
-
-using namespace std;
-using CInt = const int;
-using CSize = const size_t;
-using PFloat = float*;
-using CPFloat = const float*;
-using Matrix = vector<float>;
-using RMatrix = vector<float>&;
-using CRMatrix = const vector<float>&;
-
-constexpr size_t FSIZE = sizeof(float);
-constexpr int BLOCK_SIZE = 1024;
+#include <code_abbreviation.h>
+
 
 // 归约（GPU kernel，手写）
 __global__ void simple_reduce_sum(PFloat input, PFloat output) {
@@ -70,13 +54,6 @@ float reduce_sum_cpu(CPFloat data, CInt length) {
     return static_cast<float>(total);
 }
 
-// GPU 计时结果结构体（AI 生成）
-struct GpuTimingResult {
-    float h2d_ms;      // Host to Device 传输时间
-    float kernel_ms;   // Kernel 执行时间（多次平均）
-    float d2h_ms;      // Device to Host 传输时间
-    float total_ms;    // 总时间
-};
 
 // 验证结果（AI 生成）
 bool verify_results(float gpu_result, float cpu_result, const string& kernel_name) {

02_Reduction/02_reduce_optimized/reduce_optimized.cu

@@ -1,23 +1,5 @@
-#include <iostream>
-#include <vector>
-#include <cmath>
-#include <iomanip>
-#include "cuda_utils.cuh"
-#include "timer.cuh"
-
-using namespace std;
-using CInt = const int;
-using CFloat = const float;
-using CSize = const size_t;
-using PFloat = float*;
-using CPFloat = const float*;
-using Matrix = vector<float>;
-using RMatrix = vector<float>&;
-using CRMatrix = const vector<float>&;
-
-constexpr int BLOCK_SIZE = 1024;
-constexpr size_t FSIZE = sizeof(float);
-constexpr int COARSE_FACTOR = 4;
+#include <code_abbreviation.h>
+
 
 // 归约-任意长度（GPU kernel，手写）
 __global__ void segmented_reduce_sum(PFloat input, PFloat output, CInt length) {
@@ -143,13 +125,6 @@ bool verify_results(float gpu_result, float cpu_result, const string& kernel_nam
     }
 }
 
-// GPU 计时结果结构体（AI 生成）
-struct GpuTimingResult {
-    float h2d_ms;      // Host to Device 传输时间
-    float kernel_ms;   // Kernel 执行时间（多次平均）
-    float d2h_ms;      // Device to Host 传输时间
-    float total_ms;    // 总时间
-};
 
 // 通用归约求和 GPU 封装（GPU，手写）
 template<typename KernelFunc>

02_Reduction/03_dot_product/dot_product.cu

@@ -1,4 +1,4 @@
-#include <code_abbreviation.h>
+#include <code_abbreviation.h>
 
 // 点积-共享内存（GPU kernel，手写）
 __global__ void shared_dot_product(CPFloat a, CPFloat b, PFloat output, CInt size) {
@@ -104,13 +104,6 @@ bool verify_results(RFloat gpu_result, RFloat cpu_result, const string& kernel_n
     }
 }
 
-// GPU 计时结果结构体（AI 生成）
-struct GpuTimingResult {
-    float h2d_ms;      // Host to Device 传输时间
-    float kernel_ms;   // Kernel 执行时间（多次平均）
-    float d2h_ms;      // Device to Host 传输时间
-    float total_ms;    // 总时间
-};
 
 // 通用点积 GPU 封装（GPU，手写）
 template<typename KernelFunc>