📅 创建时间：2026-06-03 🏷️ 标签：#调度 #Scheduling #依赖分析 #延迟隐藏 #SoftwarePipelining #指令级并行 📚 前置知识：[[/04-ai/01-llm-engineering/07-llm-evolution]]（LLM 发展脉络） 📚 相关知识：[[11-dynamic-shape]]（动态 Shape） [[13-codegen-architecture]]（代码生成架构） [[08-operator-fusion]]（算子融合）

硬件约束下的操作调度 / Operation Scheduling Under Hardware Constraints

┌──────────────────────────────────────────────────────────────────────────────┐ │ 情境描述 │ ├──────────────────────────────────────────────────────────────────────────────┤ │ 你有两个 GPU 操作：矩阵乘法和数据加载。直觉上应该先加载数据再计算，但编译器 │ │ 选择了先计算——然后等待数据。你开始怀疑编译器是不是 bug 了。实际上这是调度器 │ │ 的问题：编译器需要决定操作在哪个时间点、哪个硬件单元上执行。 │ └──────────────────────────────────────────────────────────────────────────────┘

第1节：什么是调度——硬件约束下的操作排序

1.1 调度问题的本质

调度（Scheduling） 是编译器后端最核心的优化之一。它的目标是：在给定的硬件资源约束下，决定每个操作的执行时间和执行位置，使得整体执行时间最短或吞吐量最高。

调度问题可以形式化为：

给定：
  - 一组操作 (operations) O = {o₁, o₂, ..., oₙ}
  - 操作之间的依赖关系 (dependencies)
  - 硬件资源约束 (resource constraints)

求解：
  - 每个操作 oᵢ 的开始时间 start(oᵢ)
  - 每个操作 oᵢ 的执行单元 assignment(oᵢ)
  
目标：
  minimize max(start(oᵢ) + duration(oᵢ))  // 最小化总执行时间
  或
  maximize throughput  // 最大化吞吐量

1.2 调度的层次

调度发生在多个层次：

层次	粒度	调度器	典型优化
指令级调度	单条指令	编译器/硬件	隐藏 latency，暴露 ILP
基本块调度	单个基本块	编译器	减少 stall
超全局调度	跨基本块/函数	编译器	Software Pipelining
CTA/Grid 调度	CUDA Thread Block	CUDA Runtime	块间并行
Kernel 调度	整个 Kernel	Runtime/Scheduler	计算与通信重叠

1.3 为什么调度如此重要

错误的调度会导致严重的性能问题：

python

# 错误调度的后果示意（伪代码）
def bad_schedule():
    # 调度器选择了这种顺序：
    
    # Step 1: 启动计算（需要数据 A）
    compute_start = time.time()
    result = gpu.matmul(X, W)  # 需要 X 就绪
    # 问题：X 还没有从内存加载！
    # GPU 会在这里 stall，等待 X 从 HBM 传输完成
    
    # Step 2: 加载数据（太晚了！）
    X = load_from_hbm(data_a)   # 此时计算已经在等了
    
    return result

def good_schedule():
    # 正确的调度顺序：
    
    # Step 1: 先异步加载数据
    X_future = async_copy_to_gpu(data_a)  # 异步发起传输
    
    # Step 2: 在数据传输的同时，做其他不依赖 X 的计算
    result_b = compute_something_else()
    
    # Step 3: 等数据就绪后再计算
    X = X_future.wait()  # 此时数据应该已经传输完成
    result_a = gpu.matmul(X, W)
    
    return result_a

第2节：依赖分析——调度的前提

2.1 数据依赖（Data Dependency）

数据依赖描述了操作之间的数据流动关系：

// 依赖分析示例
a = load(x)        // Load1: 加载 x 到 a
b = a + 1          // Add1: 依赖 Load1 的结果
store(y, b)        // Store1: 依赖 Add1 的结果

// 依赖类型：
// Load1 → Add1: RAW (Read After Write) - 真依赖
// Add1 → Store1: RAW (Read After Write) - 真依赖

三种基本数据依赖类型：

依赖类型	英文	含义	示例
RAW	Read After Write	写后读（真依赖）	`b = a + 1` 依赖 `a = load(x)`
WAR	Write After Read	读后写（反依赖）	`a = 1; b = a + 1` 交换后语义变
WAW	Write After Write	写后写（输出依赖）	两次写同一变量

2.2 控制依赖（Control Dependency）

控制依赖决定了操作是否需要执行：

// 控制依赖示例
if (condition) {
    a = load(x);  // 只有 condition=true 才执行
}
b = a + 1;        // 依赖 if 语句的执行结果

// 依赖图表示：
// [condition] ──控制依赖──→ [a = load(x)]
// [if-block]  ──控制依赖──→ [b = a + 1]

2.3 内存依赖（Memory Dependency）

内存依赖是最复杂的，因为需要分析指针别名：

// 内存依赖分析挑战
void compute(float* a, float* b, int n) {
    for (int i = 0; i < n; i++) {
        a[i] = b[i] * 2;      // 写 a[i]
        float t = a[i-1] + 1;  // 读 a[i-1]，可能与上一行冲突
    }
}

// 编译器需要证明：
// 1. a[i] 和 a[i-1] 不是同一个地址（i != i-1）
// 2. 没有指针别名使得 a == b

2.4 依赖图的构建

python

# 依赖图数据结构
class DependencyGraph:
    def __init__(self):
        self.nodes = []      # 操作节点
        self.edges = []      # 依赖边 (from, to, type, latency)
        self.adj_list = {}   # 邻接表表示
    
    def add_edge(self, src, dst, dep_type, latency=1):
        """添加依赖边
        src: 源操作
        dst: 目标操作
        dep_type: 'RAW', 'WAR', 'WAW', 'CTRL'
        latency: 源操作到目标操作的数据延迟（cycles）
        """
        self.edges.append((src, dst, dep_type, latency))
        if src not in self.adj_list:
            self.adj_list[src] = []
        self.adj_list[src].append((dst, latency))
    
    def topological_sort(self):
        """返回合法的调度顺序"""
        # 使用 Kahn 算法
        in_degree = {n: 0 for n in self.nodes}
        for src in self.adj_list:
            for dst, _ in self.adj_list[src]:
                in_degree[dst] += 1
        
        queue = [n for n in self.nodes if in_degree[n] == 0]
        result = []
        
        while queue:
            node = queue.pop(0)
            result.append(node)
            for dst, _ in self.adj_list.get(node, []):
                in_degree[dst] -= 1
                if in_degree[dst] == 0:
                    queue.append(dst)
        
        return result

第3节：并行性暴露——调度的目标

3.1 指令级并行 (ILP)

ILP 指的是同一条指令流中，多条指令可以同时执行（流水线并行）：

时间 →  T0   T1   T2   T3   T4   T5   T6   T7
指令1: [Load] [Mul ] [Add] [Store]
指令2:       [Load] [Mul ] [Add] [Store]     ← 与指令1并行
指令3:             [Load] [Mul ] [Add] [Store]  ← 三发射

ILP = 3（同时执行3条指令）

3.2 数据级并行 (DLP)

DLP 指的是同一条指令操作多个数据（SIMD/SIMT）：

时间 →  T0   T1   T2   
向量加: [a0+a1] [a2+a3] [a4+a5] ...  ← 一次处理多个数据

SIMD 宽度 = 4（同时处理4个元素）

3.3 任务级并行 (TLP)

TLP 指的是多个独立任务可以并行执行：

核0: [Task A]                   
核1: [Task B]  ← 完全独立，可并行
核2: [Task C]  ← 完全独立，可并行
核3: [Task D]  ← 完全独立，可并行

3.4 并行性暴露策略

// 暴露 ILP 的关键：打破依赖链
// 原始代码（依赖链长）：
for (int i = 2; i < N; i++) {
    a[i] = (a[i-1] + a[i-2]) * 0.5;  // 每次迭代依赖前两次的结果
}

// 循环展开暴露 ILP（2路展开）：
for (int i = 2; i < N; i += 2) {
    a[i]   = (a[i-1]   + a[i-2])   * 0.5;
    a[i+1] = (a[i]     + a[i-1])   * 0.5;  // 新展开的迭代可以利用前一个展开的结果
}

第4节：延迟隐藏——计算与内存访问重叠

4.1 延迟隐藏的核心思想

GPU 和 CPU 的内存访问延迟很高（数百个时钟周期），但计算单元可以在等待期间执行其他操作：

延迟隐藏成功示例：
时间 →  T0   T1   T2   T3   T4   T5   T6   T7   T8
Warp0: [MemRd] [MemRd] [MemRd] [MemRd] [MemRd] [Compute]
       ↑ 每个 warp 等待内存时，硬件切换到其他 warp

Warp1:       [Compute] [Compute] [Compute] [Compute] [MemRd]
Warp2: [Compute] [Compute] [Compute] [MemRd] [MemRd] [MemRd]

总延迟被有效隐藏！

4.2 Double Buffering——预加载下一批次

Double Buffering 是深度学习中最常用的延迟隐藏技术：计算当前批次的同时，预加载下一批次的数据。

python

import numpy as np

def double_buffered_gemm_cpu(BATCH_SIZE, M, N, K, iterations):
    """
    Double Buffering 示例：
    - Buffer A[2] 交替用于"当前批次"和"预加载下一批次"
    - Buffer B[2] 同理
    - 计算和内存传输并行执行
    """
    
    # 分配双缓冲
    # buf_a[0], buf_a[1] 交替作为 compute_buf 和 load_buf
    # buf_b[0], buf_b[1] 同理
    buf_a = [np.zeros((BATCH_SIZE, M, K), dtype=np.float32) for _ in range(2)]
    buf_b = [np.zeros((BATCH_SIZE, K, N), dtype=np.float32) for _ in range(2)]
    buf_c = [np.zeros((BATCH_SIZE, M, N), dtype=np.float32) for _ in range(2)]
    
    # 初始化状态
    # cur_buf: 当前正在使用的 buffer 索引
    # next_buf: 预加载下一批次的 buffer 索引
    cur_buf = 0
    next_buf = 1
    
    total_time = 0.0
    
    for iter_id in range(iterations):
        batch_id = iter_id % NUM_BATCHES
        
        # ========== 阶段1: 异步预加载下一批次 ==========
        # 这个操作在计算当前批次时并行执行
        next_batch_id = (batch_id + 1) % NUM_BATCHES
        
        # 模拟异步内存拷贝（GPU 上使用 cudaMemcpyAsync）
        # CPU 端使用线程池并行加载
        def async_load(buffer_idx, batch):
            buf_a[buffer_idx][:] = load_batch_A(batch)
            buf_b[buffer_idx][:] = load_batch_B(batch)
        
        # 启动异步加载线程（不阻塞）
        load_thread = Thread(target=async_load, args=(next_buf, next_batch_id))
        load_thread.start()
        
        # ========== 阶段2: 计算当前批次（与预加载并行） ==========
        compute_start = time.time()
        
        # 执行 GEMM: C = A @ B
        # 使用当前 buf 进行计算
        buf_c[cur_buf] = np.matmul(buf_a[cur_buf], buf_b[cur_buf])
        
        compute_time = time.time() - compute_start
        
        # 等待预加载完成（如果计算比加载快，需要等）
        load_thread.join()
        
        # ========== 阶段3: 交换缓冲区 ==========
        cur_buf, next_buf = next_buf, cur_buf
        
        total_time += compute_time
        print(f"Iter {iter_id}: Compute={compute_time*1000:.2f}ms, "
              f"Effective throughput={BATCH_SIZE*M*N*2/1e9/compute_time:.2f} GFLOPS")
    
    return total_time / iterations


def single_buffered_gemm_cpu(BATCH_SIZE, M, N, K, iterations):
    """
    对比：单缓冲（没有 Double Buffering）
    每个批次必须：先加载，再计算，总时间 = 加载时间 + 计算时间
    """
    buf_a = np.zeros((BATCH_SIZE, M, K), dtype=np.float32)
    buf_b = np.zeros((BATCH_SIZE, K, N), dtype=np.float32)
    buf_c = np.zeros((BATCH_SIZE, M, N), dtype=np.float32)
    
    total_time = 0.0
    
    for iter_id in range(iterations):
        batch_id = iter_id % NUM_BATCHES
        
        # 加载数据（阻塞）
        buf_a[:] = load_batch_A(batch_id)
        buf_b[:] = load_batch_B(batch_id)
        
        # 计算
        compute_start = time.time()
        buf_c = np.matmul(buf_a, buf_b)
        compute_time = time.time() - compute_start
        
        total_time += compute_time
    
    return total_time / iterations


# CUDA Double Buffering 伪代码
def double_buffered_gemm_cuda():
    """
    CUDA Double Buffering 核心模式
    """
    # 分配两个 A 缓冲区和两个 B 缓冲区
    d_A0 = cuda.alloc((M, K), dtype=np.float32)
    d_A1 = cuda.alloc((M, K), dtype=np.float32)
    d_B0 = cuda.alloc((K, N), dtype=np.float32)
    d_B1 = cuda.alloc((K, N), dtype=np.float32)
    d_C  = cuda.alloc((M, N), dtype=np.float32)
    
    cur_buf = 0  # 当前 buffer
    next_buf = 1  # 下一批次 buffer
    
    for batch_id in range(NUM_BATCHES):
        next_batch = (batch_id + 1) % NUM_BATCHES
        
        # 步骤1: 异步预加载下一批次到 next_buf
        # cudaMemcpyAsync 是非阻塞的，传输与计算并行
        cuda.memcpy_async(d_A[next_buf], h_A[next_batch], stream=compute_stream)
        cuda.memcpy_async(d_B[next_buf], h_B[next_batch], stream=compute_stream)
        
        # 步骤2: 计算当前批次（使用 cur_buf）
        # 这个 kernel 执行时，next_batch 的数据正在异步传输
        gemm_kernel(d_C, d_A[cur_buf], d_B[cur_buf], M, N, K)
        
        # 步骤3: 等待预加载完成
        cuda.stream_synchronize(compute_stream)
        
        # 步骤4: 交换缓冲区
        cur_buf, next_buf = next_buf, cur_buf

4.3 Async CUDA——异步内存操作

CUDA 提供了丰富的异步内存操作 API：

cuda

// CUDA 异步操作示例
__global__ void async_compute_kernel(float* C, float* A, float* B, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] * B[i];
    }
}

void launch_async_gemm() {
    cudaStream_t compute_stream, copy_stream;
    cudaStreamCreate(&compute_stream);
    cudaStreamCreate(&copy_stream);
    
    // 创建事件用于同步
    cudaEvent_t copy_done, compute_done;
    cudaEventCreate(&copy_done);
    cudaEventCreate(&compute_done);
    
    float *d_A, *d_B, *d_C;
    float *h_A, *h_B;
    
    // 分配设备内存
    cudaMalloc(&d_A, N * sizeof(float));
    cudaMalloc(&d_B, N * sizeof(float));
    cudaMalloc(&d_C, N * sizeof(float));
    
    // 分配页锁定主机内存（支持异步传输）
    cudaMallocHost(&h_A, N * sizeof(float));
    cudaMallocHost(&h_B, N * sizeof(float));
    
    // 主循环
    for (int iter = 0; iter < NUM_ITERS; iter++) {
        // 1. 发起异步内存拷贝（copy_stream）
        // 数据从主机传输到设备
        cudaMemcpyAsync(d_A, h_A, N * sizeof(float), 
                        cudaMemcpyHostToDevice, copy_stream);
        cudaMemcpyAsync(d_B, h_B, N * sizeof(float),
                        cudaMemcpyHostToDevice, copy_stream);
        
        // 2. 标记拷贝完成
        cudaEventRecord(copy_done, copy_stream);
        
        // 3. 让计算流等待拷贝完成
        cudaStreamWaitEvent(compute_stream, copy_done, 0);
        
        // 4. 发起计算（compute_stream）
        // 此时 d_A, d_B 已经可用
        int blocks = (N + 255) / 256;
        async_compute_kernel<<<blocks, 256, 0, compute_stream>>>(
            d_C, d_A, d_B, N);
        
        // 5. 标记计算完成
        cudaEventRecord(compute_done, compute_stream);
        
        // 6. 主线程可以在这里做其他事情
        // ...
        
        // 7. 最终同步计算流
        cudaStreamSynchronize(compute_stream);
    }
}

第5节：Software Pipelining——循环迭代重叠

5.1 Software Pipelining 原理

Software Pipelining 将循环的多次迭代错开执行，让不同迭代的指令在硬件上并行：

原始循环（无流水线）：
迭代1: [Load] [Compute] [Store] ────────────────────→ 3 cycles
迭代2:         [Load] [Compute] [Store] ────────────→ 3 cycles  
迭代3:                 [Load] [Compute] [Store] → 3 cycles
总时间: 9 cycles

Software Pipelining（流水线）：
时间:  T0   T1   T2   T3   T4   T5   T6
Iter1: [Load] [Compute] [Store]
Iter2:       [Load] [Compute] [Store]
Iter3:             [Load] [Compute] [Store]
      ──────────────────────────────────────→ Initiation Interval (II) = 1
总时间: 5 cycles（稳态后每个迭代只需1个cycle）

加速比: 9/5 = 1.8x

5.2 依赖环与 Initiation Interval

Software Pipelining 的关键参数是 II (Initiation Interval)：连续两次迭代开始的间隔周期数。

python

def analyze_pipelining_feasibility():
    """
    分析软件流水线的可行性
    II 必须大于等于所有依赖环的最小距离
    """
    
    # 假设有以下依赖关系（每个操作的延迟）
    # Load → Compute: 5 cycles
    # Compute → Store: 3 cycles
    
    # 依赖环分析：
    # 一个完整的迭代必须执行完 Load → Compute → Store
    # 最小总延迟 = 5 + 3 = 8 cycles
    
    # 如果硬件支持 3 路并行发射：
    # II = max(依赖延迟) / 并行度 = 8 / 3 = 2.67 → 向上取整 = 3
    
    print("最小可能的 II = 3")
    print("这意味着每3个cycle可以开始一个新迭代")
    print("理论最大 throughput = 1/3 iterations per cycle")
    

def loop_carried_dependency():
    """
    循环携带依赖（Loop-Carried Dependency）
    这会限制 Software Pipelining 的效果
    """
    
    # 循环携带依赖示例
    # 这个循环不能被流水线化，因为每次迭代依赖上一次迭代的结果
    for i in range(1, N):
        a[i] = (a[i-1] + a[i]) * 0.5  # 依赖 a[i-1]
    
    # 依赖距离 = 1（相邻迭代之间）
    # 这意味着下一个迭代必须等待上一个迭代完成
    # II 至少为 1（必须串行）
    
    # 没有循环携带依赖的示例
    # 可以完全并行化
    for i in range(N):
        b[i] = a[i] * 2  # 每个迭代独立

5.3 循环展开 + 调度

循环展开是 Software Pipelining 的常见辅助手段：

// 原始循环
for (int i = 0; i < N; i++) {
    c[i] = a[i] + b[i];
}

// 2路展开
for (int i = 0; i < N; i += 2) {
    c[i]   = a[i]   + b[i];     // 迭代i
    c[i+1] = a[i+1] + b[i+1];   // 迭代i+1，理论上可并行
}

// Software Pipelining 调度后（假设 Load 3 cycles，Add 2 cycles）：
// T0: Load a[0], a[1]
// T1: Load b[0], b[1] (与 T0 的 Load 并行)
// T2: Add c[0], c[1] (与 T0/T1 并行)
// T3: Store c[0], c[1]

第6节：调度策略——算法实现

6.1 List Scheduling（列表调度）

List Scheduling 是最常用的启发式调度算法：

python

def list_scheduling(graph, num_resources, resource_types):
    """
    List Scheduling 算法
    
    思想：
    1. 按优先级排序所有操作（优先级高的先调度）
    2. 从优先级最高的开始，贪心地分配时间槽
    3. 重复直到所有操作都被调度
    
    优先级计算策略：
    - 高度数优先（操作的后继越多越早调度）
    - 关键路径长度优先（距离终点越近越早调度）
    - 最长路径优先（上游依赖越长的越早调度）
    """
    
    # 初始化
    scheduled = set()           # 已调度的操作
    ready_queue = PriorityQueue()  # 就绪队列（按优先级排序）
    
    # 初始就绪操作（无依赖的操作）
    for op in graph.roots():
        priority = compute_priority(op, graph)
        ready_queue.push(op, priority)
    
    schedule = {}  # 操作 → 开始时间
    
    while not ready_queue.empty():
        # 取出优先级最高的就绪操作
        op = ready_queue.pop()
        
        # 找到最早的可行时间槽
        earliest_time = 0
        
        for pred, latency in graph.pred_edges(op):
            if pred in schedule:
                pred_end = schedule[pred] + latency
                earliest_time = max(earliest_time, pred_end)
        
        # 检查资源约束
        time_slot = earliest_time
        while not resource_available(op, time_slot, resource_types):
            time_slot += 1
        
        # 调度
        schedule[op] = time_slot
        scheduled.add(op)
        
        # 更新后继操作
        for succ in graph.succ(op):
            if all(pred in scheduled for pred in graph.preds(succ)):
                priority = compute_priority(succ, graph)
                ready_queue.push(succ, priority)
    
    return schedule


def compute_priority(op, graph):
    """
    计算操作的调度优先级
    使用关键路径长度作为优先级
    """
    # 递归计算从该操作到终点的最长路径
    max_succ_dist = 0
    for succ in graph.succ(op):
        dist = graph.edge_weight(op, succ) + succ.critical_path_length
        max_succ_dist = max(max_succ_dist, dist)
    
    return max_succ_dist  # 关键路径长的优先级高

6.2 ASAP / ALAP 调度

python

def asap_schedule(graph):
    """
    ASAP (As Soon As Possible) 调度
    所有操作尽可能早地调度
    """
    schedule = {}
    
    # 按拓扑序调度
    for op in graph.topological_order():
        earliest = 0
        for pred in graph.preds(op):
            earliest = max(earliest, schedule[pred] + pred.duration)
        schedule[op] = earliest
    
    return schedule


def alap_schedule(graph, makespan):
    """
    ALAP (As Late As Possible) 调度
    所有操作尽可能晚地调度，但仍满足依赖和目标 makespan
    """
    schedule = {}
    
    # 逆拓扑序调度
    for op in reversed(graph.topological_order()):
        latest = makespan - op.duration
        for succ in graph.succs(op):
            latest = min(latest, schedule[succ] - op.duration - graph.edge_weight(op, succ))
        schedule[op] = latest
    
    return schedule


def compute_mobility(schedule_asap, schedule_alap):
    """
    计算每个操作的调度机动性（mobility）
    mobility = ALAP时间 - ASAP时间
    """
    mobility = {}
    for op in schedule_asap:
        mobility[op] = schedule_alap[op] - schedule_asap[op]
    return mobility

6.3 调度策略对比

策略	优点	缺点	适用场景
List Scheduling	简单高效，质量好	需要好的优先级函数	大多数场景
ASAP	最小化 stall	可能产生资源冲突	需要最小化 latency
ALAP	最大化寄存器共享	增加 stall	资源受限场景
Critical Path First	保证关键路径最短	可能忽略资源平衡	延迟敏感场景
ILP Oriented	暴露最大并行	复杂，调度时间长	编译器优化
List Scheduling + Force Directed	平衡资源利用	实现复杂	嵌入式系统

第7节：调度对性能的影响

7.1 调度错误的代价

错误的调度可能导致严重的性能下降：

场景：GEMM 计算，Load 延迟 = 400 cycles

错误调度（先计算后加载）：
  T0:   [MatMul Start] ──→ stall（等待数据）──→ [数据到达 T400]
  T400: [计算开始]
  T500: [计算完成]
  总时间: 500 cycles

正确调度（先加载后计算）：
  T0:   [异步加载数据] ←────────┐
  T1:   [其他独立计算] ←──┐    │
  ...                       传输和计算并行
  T350: [独立计算完成] ←──┘    │
  T400: [数据到达] ────────────┘
  T400: [计算开始]
  T500: [计算完成]
  总时间: 500 cycles，但 GPU 利用率 100%

7.2 性能指标

python

def compute_scheduling_metrics(schedule, hardware):
    """
    计算调度相关的性能指标
    """
    
    # 1. 资源利用率
    total_resource_slots = hardware.num_units * hardware.cycle_count
    used_resource_slots = sum(op.duration for op in schedule.values())
    resource_utilization = used_resource_slots / total_resource_slots
    
    # 2. 关键路径长度
    makespan = max(schedule[op] + op.duration for op in schedule)
    
    # 3. 并行度 (Parallelism)
    # 在任意时刻，同时执行的操作数
    parallelism = []
    for cycle in range(makespan):
        active_ops = sum(1 for op, start in schedule.items() 
                        if start <= cycle < start + op.duration)
        parallelism.append(active_ops)
    
    avg_parallelism = sum(parallelism) / len(parallelism)
    peak_parallelism = max(parallelism)
    
    # 4. Load-Compute Overlap（负载隐藏效率）
    # 计算与内存访问重叠的比例
    compute_cycles = sum(op.duration for op in schedule if op.is_compute)
    memory_stall_cycles = sum(op.stall_cycles for op in schedule if op.is_memory)
    
    overlap_ratio = 1 - (memory_stall_cycles / compute_cycles)
    
    return {
        'resource_utilization': resource_utilization,
        'makespan': makespan,
        'avg_parallelism': avg_parallelism,
        'peak_parallelism': peak_parallelism,
        'overlap_ratio': overlap_ratio
    }

升华

┌─────────────────────────────────────────────────────────────────────────────────┐ │ 调度的核心原则 │ ├─────────────────────────────────────────────────────────────────────────────────┤ │ │ │ 1. 依赖分析是调度的基础：没有正确的依赖图，就没有正确的调度 │ │ │ │ 2. 并行性暴露是调度的目标：在满足依赖的前提下，最大化指令/数据/任务级并行 │ │ │ │ 3. 延迟隐藏是性能的关键：计算与内存访问的重叠决定了硬件利用率 │ │ │ │ 4. 调度空间巨大：需要启发式算法（List Scheduling）而非穷举搜索 │ │ │ └─────────────────────────────────────────────────────────────────────────────────┘

"AI 可查 vs 必须理解"清单

必须理解（不理解就等于不会）：

🔴 依赖分析的三种类型：RAW、WAR、WAW——不知道这些就无法理解为什么代码不能重排
🔴 Double Buffering 的原理——不知道就无法理解深度学习训练中的数据加载优化
🔴 Software Pipelining 与 Initiation Interval 的关系——不知道就无法理解循环展开的动机
🔴 List Scheduling 的贪心本质——不知道就无法理解为什么调度器有时选择"次优"顺序

AI 可查（知道去哪查就行）：

✅ 特定硬件的延迟数字（如 HBM 400 cycles）——不同代际差异很大
✅ 特定调度算法的论文细节——如 Force-Directed Scheduling 的数学推导
✅ 特定硬件的调度约束（如某 GPU 的最大并发 Warp 数）——需要查文档
✅ 最新的调度策略研究进展——如 ML 辅助的调度算法

学习状态：🟡 开始学习

硬件约束下的操作调度 / Operation Scheduling Under Hardware Constraints ​

第1节：什么是调度——硬件约束下的操作排序 ​

1.1 调度问题的本质 ​

1.2 调度的层次 ​

1.3 为什么调度如此重要 ​

第2节：依赖分析——调度的前提 ​

2.1 数据依赖（Data Dependency） ​

2.2 控制依赖（Control Dependency） ​

2.3 内存依赖（Memory Dependency） ​

2.4 依赖图的构建 ​

第3节：并行性暴露——调度的目标 ​

3.1 指令级并行 (ILP) ​

3.2 数据级并行 (DLP) ​

3.3 任务级并行 (TLP) ​

3.4 并行性暴露策略 ​

第4节：延迟隐藏——计算与内存访问重叠 ​

4.1 延迟隐藏的核心思想 ​

4.2 Double Buffering——预加载下一批次 ​

4.3 Async CUDA——异步内存操作 ​

第5节：Software Pipelining——循环迭代重叠 ​

5.1 Software Pipelining 原理 ​

5.2 依赖环与 Initiation Interval ​

5.3 循环展开 + 调度 ​

第6节：调度策略——算法实现 ​

6.1 List Scheduling（列表调度） ​

6.2 ASAP / ALAP 调度 ​

6.3 调度策略对比 ​

第7节：调度对性能的影响 ​

7.1 调度错误的代价 ​

7.2 性能指标 ​

升华 ​

"AI 可查 vs 必须理解"清单 ​

硬件约束下的操作调度 / Operation Scheduling Under Hardware Constraints

第1节：什么是调度——硬件约束下的操作排序

1.1 调度问题的本质

1.2 调度的层次

1.3 为什么调度如此重要

第2节：依赖分析——调度的前提

2.1 数据依赖（Data Dependency）

2.2 控制依赖（Control Dependency）

2.3 内存依赖（Memory Dependency）

2.4 依赖图的构建

第3节：并行性暴露——调度的目标

3.1 指令级并行 (ILP)

3.2 数据级并行 (DLP)

3.3 任务级并行 (TLP)

3.4 并行性暴露策略

第4节：延迟隐藏——计算与内存访问重叠

4.1 延迟隐藏的核心思想

4.2 Double Buffering——预加载下一批次

4.3 Async CUDA——异步内存操作

第5节：Software Pipelining——循环迭代重叠

5.1 Software Pipelining 原理

5.2 依赖环与 Initiation Interval

5.3 循环展开 + 调度

第6节：调度策略——算法实现

6.1 List Scheduling（列表调度）

6.2 ASAP / ALAP 调度

6.3 调度策略对比

第7节：调度对性能的影响

7.1 调度错误的代价

7.2 性能指标

升华

"AI 可查 vs 必须理解"清单