写在前面

考试时间：2025.6.6
授课教师：lwg、yzk
考试内容：与 lwg 老师在最后一节课强调的重点内容完全吻合

总之是强烈推荐这门课！去选！都去选！

CUDA Reduction

下面的代码存在什么潜在的问题？如何解决？

• %操作运算效率慢： %操作符在CUDA中被编译成20多条指令，运算速度慢，性能下降

• highiy divergent ，warp分支发散： CUDA 的 GPU采用了 SIMT架构，无法让同一 warp 的线程执行不同指令。if判断会使warp中偶数号线程运行，奇数号线程空闲，导致warp分支严重，性能下降

  SIMT的核心思想

  • 单指令流：一个线程束（Warp，通常是32个线程）在同一周期内执行相同的指令。
  • 多线程并行：每个线程可以处理不同的数据（独立寄存器、私有内存），并拥有自己的执行路径（如分支判断）。
  • 动态分支处理：当线程间存在分支（如if-else）时，SIMT会串行化所有分支路径（称为分支发散），牺牲部分效率但保持正确性。

第一次优化

• 用乘法计算 index，去掉%操作
• if 分支改写成 index< blockDim.x ，连续条件判断，使得同一个 warp 中的线程统一执行或统一不执行该语句

**但是引入了新的共享内存冲突的问题：**当同一个warp中的多个线程同时访问同一个bank的不同地址时，就会发生bank conflict，导致这些访问被序列化，降低性能（如图）

第二次优化

使用了逆序循环和基于线程ID的索引来替代之前版本中的间隔索引方式

下面的CUDA代码存在两处错误，请指出错误，给出原因并纠正

• 第四行计算i时出错，不能使用grimDim.x
  • i是全局的标号，应该由线程所在的块的标号blockIdx.x乘以块的维度blockDim.x，然后再加上局部的线程号threadIdx.x，即 int i=blockIdx.x*blockDim.x+ threadIdx.x；
• 倒数第二行调用vecAdd时传入第一个参数出错，不能使用N/256
  • vecAdd<<<grid_size，block_size>>>执行配置的第一个参数是启动块的数目，不能小于1，应为向上取整
  • 纠错后为：vecAdd<<<(N+255)/256，256>>>(a， b，c，N);

完整并行代码如下：

线程块内同步（Intra-block Synchronization）

__syncthreads() 同步一个线程块中的所有线程，如果代码中有shared memory操作，一定要写同步函数，否则会出错。

__syncthreads()的作用：

• 确保同一个线程块（Block）内的所有线程在执行到该指令时暂停，直到所有线程都到达此同步点。
• 主要解决共享内存（__shared__）和全局内存（Global Memory）的读写顺序问题：
  • 线程A写入共享内存 → 线程B读取该值 → 必须通过__syncthreads()保证A的写入在B读取前完成。

矩阵转置（Transpose）

extern __shared__ float s[];  // 动态分配的共享内存

__device__ void transpose(float* a, int lda) {
    int i = threadIdx.x, j = threadIdx.y;
    s[i + lda * j] = a[i + lda * j];  // 将全局内存数据拷贝到共享内存
    __syncthreads();                  // 同步！确保所有数据已写入共享内存
    a[i + lda * j] = s[j + lda * i];  // 从共享内存读取转置后的数据
}

为什么必须同步？

• 若省略__syncthreads()：
  • 部分线程可能过早读取s[j][i]，而此时该位置尚未被其他线程写入，导致数据错误。
• 例如：线程(0,1)可能在线程(1,0)写入s[1][0]之前读取它，得到旧值或未定义值。

矩阵向量乘

Pthreads 并行化

当讨论矩阵-向量乘法时，我们通常假设矩阵为m*n，m、n 都能够被 t 整除，t 是线程的个数。但是，如果m和n不满足能被t整除的条件，那么用什么公式来分配数据?

MPI_Reduce：process1 的调用顺序出错，则程序结束后，b和d分别为多少

假设 operator MPI_SUM

若 destination process 0：

• time1：b=a+c+a=1+2+1=4
• time2：d=c+a+c=2+1+2=5

若 destination process 1：

• time1：d=a+c+a=1+2+1=4
• time2：b=c+a+c=2+1+2=5

注意：只有目标进程的output_data _p有效，其它进程的output_data _p写NULL都无所谓

Flynn’s Taxonomy （弗林分类法）

1. 指令流(Instruction Stream): 处理器执行的指令序列
2. 数据流(Data Stream): 由指令流操作的数据序列

1. SISD (Single Instruction, Single Data)

• 描述: 单指令流单数据流
• 特点:
  • 经典的冯·诺依曼架构
  • 一次执行一条指令，操作一个数据项
  • 所有传统单核处理器都属于此类

2. SIMD (Single Instruction, Multiple Data)

• 描述: 单指令流多数据流
• 特点:
  • 一条指令同时作用于多个数据
  • 非常适合数据并行任务
  • 需要高度规则的数据处理模式

3. MISD (Multiple Instruction, Single Data)

• 描述: 多指令流单数据流
• 特点:
  • 多个处理单元对同一数据执行不同指令
  • 实际应用非常少见
  • 主要用于容错系统

4. MIMD (Multiple Instruction, Multiple Data)

• 描述: 多指令流多数据流
• 特点:
  • 多个处理器独立执行不同指令，处理不同数据
  • 最通用的并行计算形式
  • 可分为共享内存和分布式内存两种子类

在冯·诺依曼系统中加入缓存和虚拟内存改变了它作为SISD系统的类型吗？如果加人流水线呢？多发射或硬件多线程呢？

• 加人缓存和虚拟内存不会改变系统的SISD类型，因为对存储管理的优化没有改变一次可以执行的指令数或者一次可以操作的数据量
• 加入流水线会使得系统变为SIMD类型，因为加入流水线相当于将一个指令应用于多个数据项
• 多发射和硬件多线程会使得系统变为MIMD类型，
  • 多发射相当于将一个指令流在多个ALU上同时启动一部分，分裂为多个指令流，多个指令流用于不同的数据项
  • 硬件多线程是指当前任务被阻塞时，系统尝试切换到其他线程继续工作，允许了多线程的存在，将多个指令流用于不同的数据项

考虑Cache缓存命中，对矩阵是按行访问快还是按列访问快

按行访问快

• C语言以行主序来存储二维数组，多维数组在内存中也是一个超大的一维数组，每次都按高速缓存行可容纳的数量换入换出
• 按列访问在内存中是跳跃的，会让Cache命中率降低，导致更频繁的换入换出，耗费更多时间，降低运算效率

总之是按行访问时Cache命中率高（可以认为第一个循环发生4次Cache缺失，第二个循环发生16次Cache缺失）

指令级并行（Instruction Level Parallelism, ILP）

• 流水线（Pipelining）是一种将指令执行过程划分为多个独立阶段（如取指、译码、执行、访存、写回）的技术。每个阶段由专用的功能单元处理，不同指令的不同阶段可以同时进行，类似于工厂的装配线
• 多指令发射（Multiple Issue）允许同时启动（发射）多条指令到不同的执行单元

a：如图，每个浮点加法要耗费9ns，包括7个阶段（取指令、比较、移位、相加、规格化、舍入、存数）

b：1000 x 9 = 9000 ns

c：9 + 999 x 2 = 2007 ns

Partitioning options (分区)

1. 块分区 (Block Partitioning)

• 定义：将连续的components分配给每个处理进程。
• 特点：
  • 简单直接，易于实现
  • 保持了数据的局部性，相邻数据存储在同一个进程中
  • 适合数据访问模式具有空间局部性的应用
• 示例：如果有100个数据元素和4个进程，进程0获取元素1-25，进程1获取26-50，以此类推。

2. 循环分区 (Cyclic Partitioning)

• 定义：以轮询(round robin)方式分配components。
• 特点：
  • 提供了极好的负载均衡
  • 破坏了数据的局部性
  • 适合计算负载不均匀的情况
• 示例：100个元素和4个进程，进程0获取元素1,5,9,…，进程1获取元素2,6,10,…等。

3. 块循环分区 (Block-Cyclic Partitioning)

• 定义：以循环方式分配数据块。
• 特点：
  • 结合了块分区和循环分区的优点
  • 通过调整块大小可以在负载均衡和数据局部性之间取得平衡
• 示例：设块大小为10，100个元素和4个进程，进程0获取元素1-10,41-50,81-90，进程1获取11-20,51-60,91-100等。

12-component vector among 3 processes

14任务，3线程，如何进行三种划分?

集合通讯与点对点通讯的异同

• 通信子中的所有进程必须调用相同的集合通信函数。
  • 例如，若一个程序试图将某进程的MPI_Reduce调用与其他进程的MPI_Recv调用进行匹配，则属于错误行为。此类程序极可能导致死锁或崩溃。
• 每个进程传递给MPI集合通信函数的参数必须“兼容”。
  • 例如，若一个进程传入的目标进程（dest_process）参数为0，而另一个进程传入1，则MPI_Reduce调用的结果将出现错误。
• output_data_p参数仅由目标进程（dest_process）使用。
  • 但其他所有进程仍需传递一个有效的实参（即使传入NULL）。
• 点对点通信通过标签（tag）和通信子（communicator）进行匹配，而集合通信不使用标签——其匹配仅依赖于通信子及函数调用的顺序。

Cache coherence

Cache一致性问题

在多核系统中，各个核的Cache存储相同变量的副本，当一个处理器更新Cache中该变量的副本时，其他处理器应该知道该变量已更新，即其他处理器中的Cache副本也应该更新。

两种保证Cache一致性的主要方法

• 监听Cache一致性协议：

  • 工作原理：当一个核更新Cache中x的副本时，将更新消息广播在总线上，若其他核正在监听，则知道x已更新，并将自己Cache中的x的副本标记位非法。实际上，广播会通知其他核包含x的整个Cache行已经更新，而不是只有x更新。

  • 特点：

    处理器间通过互连进行广播。

    写直达Cache不需要额外的互连网络开销，因为每个核都能检测“写”；写回Cache需要额外的通信，因为对Cache的更新不会立即发送给内存。

    不可扩展，因为对于大型系统，每有更新就会广播，会导致性能的下降。

• 基于目录的Cache一致性协议：

  • 工作原理：使用“目录”数据结构，存储每个内存行的状态。当一个变量更新时，就会查询目录，将所有包含该变量的高速缓存行设置为非法。

  • 特点：

    目录需要大量额外的存储空间。

    当一个Cache变量更新时，只需要与存储这个变量的核交涉。

Cache伪共享

线程之间没有共享任何变量（但是共享了同一个缓存行），但是它们访问主存的行为看起来好像它们共享了一个变量，这种情况称为伪共享。

原始并行代码（存在伪共享）

#pragma omp parallel for num_threads(thread_count) \
default(none) private(i,j) shared(A, x, y, m, n)
for (i = 0; i < m; i++) {
    y[i] = 0.0;                     // 直接写入共享内存
    for (j = 0; j < n; j++)
        y[i] += A[i][j] * x[j];     // 频繁更新共享内存
}

问题：

• 所有线程直接读写共享数组 y。
• 内循环中频繁更新 y[i]，导致缓存行在不同线程间反复失效和同步（伪共享）。

在并行计算中，当多个线程同时修改同一缓存行（Cache Line）中的不同变量时，会触发缓存一致性协议频繁同步数据，导致性能下降。

• 在矩阵-向量乘法中，结果向量 y 的相邻元素（如 y[i] 和 y[i+1]）可能位于同一缓存行。
• 若不同线程同时修改相邻的 y 元素（即使逻辑独立），硬件会强制刷新整个缓存行，造成大量不必要的内存访问。

优化方案（消除伪共享）

#pragma omp parallel num_threads(thread_count) \
default(none) private(i, j, sum) shared(A, x, y, m, n)
{ //超过一条语句的话要用{}括起来
    #pragma omp for
    for (i = 0; i < m; i++) {
        sum = 0.0;                   // 使用线程私有变量累加
        for (j = 0; j < n; j++)
            sum += A[i*n+j] * x[j];  // 计算在私有变量中完成
        y[i] = sum;                  // 仅最终结果写入共享内存
    }
}

优化点：

1. 私有累加变量（sum）：
   • 每个线程用私有变量 sum 暂存计算结果，避免频繁写共享内存。
   • 仅在循环结束时将 sum 一次性写入 y[i]，减少共享内存访问次数。
2. 消除伪共享：
   • 线程间不再竞争同一缓存行（私有 sum 位于线程栈中，互不干扰）。
   • 最终写入 y[i] 时，相邻元素的写入由不同线程完成，但写入频率大幅降低（从内循环每次更新 → 每行一次写入）。

如果已经定义了宏_OPENMP，它是一个int类型的十进制数。编写一个程序打印它的值。这个值的意义是什么?

#include <omp.h>
#include <iostream>
using namespace std;
int main(int argc,char **argv){
cout<<" _OPENMP :"<<_OPENMP;
return 0;
}

"_OPENMP"是一个预处理器宏定义，用于标识 OpenMP 版本。_OPENMP的值是一个具有yyyymm形式的日期。OpenMP标准规定，当定义宏时，它将是已实现的OpenMP标准版本的年份和月。

加速比，效率，阿姆达尔定律

Amdahl’s Law

阿姆达尔定律指出：除非串行程序几乎全部被并行化，否则无论可用核心数有多少，其可能的加速比都将非常有限。

加速比，效率

假设，一个串行程序中，可以并行化其中的90%。进一步假设，并行化是“理想”的，也就是说，如果使用p个核，则程序可并行化部分的加速比就是p。若该程序串行版本运行时间T=20秒，则并行化后，其中的可并行部分(即90%的部分)的运行时间就是(90% x T)/p=18/p秒，不可并行化部分的运行时间为10% x T = 2 秒。

那么，程序的全部并行运行时间为：

加速比为：

因此，

若可并行部分为 99% ，则有 S=20/(19.8/P + 0.2)<100

超线性加速比（Superlinear Speedup）

超线性加速比源于并行化消除了串行程序的瓶颈，而非单纯计算能力叠加。常见原因包括：

1. 更高效利用缓存：单核程序cache有限，多核程序可以将所有数据放入cache，减少主存访问延迟。
2. 算法特性优化：并行DFS可以快速搜索目标结果。

奇偶排序(Odd-Even Sort) ：OpenMP并行优化

奇偶排序是一种基于冒泡排序的并行算法，通过交替执行两种操作：

1. 偶数阶段：比较并交换奇数索引元素与其前驱元素（a[i-1]和a[i]，其中i为奇数）
2. 奇数阶段：比较并交换奇数索引元素与其后继元素（a[i]和a[i+1]，其中i为奇数）
   需要执行n个阶段（n为数组长度）确保完全排序。

第一版实现的主要问题

for (phase = 0; phase < n; phase++) {
    if (phase % 2 == 0)
        #pragma omp parallel for num_threads(thread_count) \
           default(none) shared(a,n) private(i,tmp)  // 问题点
        for (i = 1; i < n; i += 2) { ... } // 偶数阶段
    else
        #pragma omp parallel for num_threads(thread_count) \
           default(none) shared(a,n) private(i,tmp)  // 问题点
        for (i = 1; i < n-1; i += 2) { ... } // 奇数阶段
}

性能缺陷：

1. 频繁线程创建/销毁：
   • 每个phase循环中都会通过#pragma omp parallel for创建新的线程组
   • 线程的重复初始化和销毁造成显著开销（尤其当n很大时）
2. 隐式屏障同步：
   • 每个parallel for结束时存在隐式屏障(barrier)

第二版优化方案

#pragma omp parallel num_threads(thread_count) \
   default(none) shared(a,n) private(i,tmp,phase)  // 线程组只创建一次
{
    for (phase = 0; phase < n; phase++) {
        if (phase % 2 == 0)
            #pragma omp for // 仅分配循环迭代
            for (i = 1; i < n; i += 2) { ... }
        else
            #pragma omp for // 仅分配循环迭代
            for (i = 1; i < n-1; i += 2) { ... }
    }
}

优化改进：

1. 线程复用：
   • 通过外层#pragma omp parallel一次性创建线程组
   • 所有phase循环复用同一组线程，消除创建/销毁开销
2. 按需同步：
   • #pragma omp for仅分配循环迭代，不创建新线程
   • 隐式屏障仍在每个内层循环后存在，但这是算法必需的（确保阶段间数据一致性）

Cache mappings

• 全相联缓存（Full associative）——新缓存行可放置在缓存中的任意位置。
• 直接映射缓存（Direct mapped）——每个缓存行在缓存中只有一个固定位置可存放。
• n路组相联缓存（n-way set associative）——每个缓存行可以放置在缓存中n个不同位置之一。
  • 当超过一个内存行可以映射到缓存中的多个不同位置时，我们依然需要决定应该替换哪一个缓存行。

Assignments of a 16-line main memory to a 4-line cache

• Memory Index： 0—15
• Cache Location：0—3

处理器峰值性能计算

1. GPU 内存带宽计算

• 硬件参数：
  • 384-bit memory interface: 内存总线宽度为 384 位 (bit)。这表示 GPU 一次可以从显存中读取或写入 384 位的数据。
  • 900 MHz DDR: 内存时钟频率为 900 MHz，并且是 DDR (Double Data Rate) 类型。
    • 关键点： DDR 内存在一个时钟周期内可以传输两次数据（上升沿和下降沿各一次）。因此，它的有效数据传输频率是物理时钟频率的两倍。
    • 有效数据传输频率 = 900 MHz * 2 = 1800 MHz (或 1800 MT/s - MegaTransfers per second)。
• 计算公式：
  • 理论峰值带宽 (Bytes/s) = (内存接口宽度 (bits) * 有效数据传输频率 (Hz)) / 8
• 结论： 该 G80 GPU 的理论峰值内存带宽是 86.4 GB/s。这个数字代表了在最理想情况下（无任何延迟、冲突、瓶颈），GPU 每秒钟能从显存中读取或写入的最大数据量。实际应用中通常达不到这个理论峰值。

2. CPU 峰值性能计算

• A. 双精度浮点理论峰值性能 (GFLOPS)

• 

  • 硬件参数：
    • 8核心: CPU 有 8 个物理核心。
    • 支持SIMD指令（每周期16个双精度浮点运算）
    • 主频4.0 GHz: CPU 每个核心的时钟频率为 4.0 GHz (4.0 * 10⁹ Hz)。
  • 计算公式：
    • 理论峰值 FLOPS = 核心数 * (每周期 FLOP) * 频率 (Hz)

• B. 内存理论峰值带宽 (GB/s)

• 

  • 硬件参数：
    • 内存：4通道DDR5-4800: CPU 支持 4 通道内存，使用 DDR5 标准，内存模块标称速度为 4800 MHz (指 有效数据传输频率，即 I/O 总线时钟)。
    • 每通道位宽64位: 每个内存通道的数据总线宽度是 64 位 (bit)。
  • 计算公式：
    • 理论峰值带宽 (Bytes/s) = 通道数 * (通道位宽 (bits) * 有效数据传输频率 (Hz)) / 8

区分 物理时钟频率 与 有效传输率

• 当你看到内存规格是 XXX MHz DDR 这种格式（尤其是较老的规格或某些显卡规格），XXX MHz 很可能指的是物理时钟频率，计算带宽时需要 ×2 得到有效传输率。
• 当你看到内存规格是 DDRX-YYYY 这种标准命名（如 DDR4-3200, DDR5-4800, DDR3-1600），数字 YYYY 一定指的是有效数据传输率 (YYYY MT/s)，计算带宽时 直接使用 YYYY，不需要再 ×2。

MPI编程 求解前缀和

串行实现（对应问题a）

prefix_sums[0] = vect[0];
for (i=1;i<n;i++)
   prefix_sums[i] = prefix_sums[i-1] + vect[i];

MPI 实现（对应问题 c，n=2^k）

加油，自己想（