【并行计算】复习
教材为《并行程序设计导论》
1.Stencil 优化



-
#pragma omp:OpenMP 编译指示的标准前缀,告知编译器处理后面的并行指令。如果编译器未开启 OpenMP 支持(如 GCC 未加-fopenmp参数),该行会被直接忽略,代码将退化为串行执行。 -
parallel for:这是一个组合指令。-
parallel:创建一个线程组(Team of threads)。 -
for:将紧跟其后的for循环的迭代任务,划分并分配给线程组中的各个线程。
-
-
collapse(2):这是该指令中最关键的子句。它指示编译器将紧随其后的两层完美嵌套循环(即变量为i和j的循环)合并(平铺)成一个单一的循环结构,然后再进行并行任务分配。
2.前缀和




这里应该是有笔误,在加上sum_of_preceding的时候,loc_i的范围也应该是<loc_n





omm_sz (Communicator Size):
-
定义:指当前通信域(Communicator)里总共有多少个进程。
mask (掩码):
-
定义:这是在二进制运算中用来“定位”或“筛选”特定位的工具。在倍增算法里,它是一个从 1 开始、每次乘 2 的变量(1, 2, 4, 8...)。
-
作用:它决定了每一轮通信的距离。
-
mask = 1(二进制001):翻转第 0 位,跨度为 1。 -
mask = 2(二进制010):翻转第 1 位,跨度为 2。 -
mask = 4(二进制100):翻转第 2 位,跨度为 4。
-
MPI用法
MPI 在实际工程中最常见的用法和核心函数分类:
1. 环境与通信域管理 (Environment Management)
所有 MPI 程序的骨架。它负责初始化并行环境,并让每个进程知道自己的身份。
-
MPI_Init/MPI_Finalize:必须在并行代码的开始和结束时调用,用于启动和清理 MPI 环境。 -
MPI_Comm_rank:获取当前进程的 ID(即上一题代码中的my_rank)。通常用于条件分支(如if (rank == 0)),让不同进程执行不同的逻辑。 -
MPI_Comm_size:获取当前通信域内的总进程数(即上一题代码中的n)。
2. 点对点通信 (Point-to-Point Communication)
这是基础的通信方式,仅涉及两个明确的进程:一个发送,一个接收。
-
阻塞式通信 (
MPI_Send/MPI_Recv):上一题使用的方式。发送方会阻塞直到数据安全交出,接收方会阻塞直到数据完全到达。容易导致死锁,但逻辑清晰。 -
非阻塞式通信 (
MPI_Isend/MPI_Irecv):函数调用后立即返回,允许进程在后台传输数据的同时继续执行其他计算(通信与计算重叠,Overlapping)。配合MPI_Wait使用,是优化 HPC 性能的常用手段。
3. 集合通信 (Collective Communication)
这是 MPI 的核心威力所在。很多时候,我们需要所有进程共同参与一项数据交互(例如同步、汇总)。集合通信在底层由 MPI 库进行了深度优化(如使用树状拓扑、倍增法等),效率远高于程序员自己写循环进行点对点收发。
常见的集合通信包括:
-
广播 (
MPI_Bcast):将一个进程(通常是根进程,如进程 0)的数据复制给所有其他进程。常用于分发全局配置参数。 -
分发 (
MPI_Scatter):将根进程中的一个大数组切分成若干等份,分别发送给各个进程。 -
收集 (
MPI_Gather):分发的逆操作。将各个进程计算得到的小块数据,按顺序拼接回根进程的大数组中。 -
归约 (
MPI_Reduce):将所有进程提供的数据利用特定的操作符(如求和MPI_SUM、求最大值MPI_MAX)合并成一个单一结果,存放在根进程中。 -
前缀扫描 (
MPI_Scan):这正是上一题的核心! 实际上,在标准的 MPI 编程中,你根本不需要手写复杂的 $k$ 阶段倍增算法。你只需直接调用MPI_Scan,MPI 底层就会自动为你执行最高效的并行前缀和计算。
3.处理器峰值性能计算
cuda reduce G80 GPU
根据位宽和主频
如果是CPU,计算峰值计算能力 CPU/内存(明白公式)
DDR(Double Data Rate,双倍数据速率)


4.缓存映射机制
(Cache Mapping Policies)。
全相联、直接映射、组相连,能把映射方式写出来(8个内存,2个内存行,p13,14)
它旨在解决一个最基本的物理矛盾:主存容量极大(例如图中的 16 个块),而 Cache 容量极小(例如图中的 4 个位置)。当我们要把主存中的数据加载到 Cache 中时,应该把它放在这 4 个位置中的哪一个?




n-路组相联

5.奇偶排序
问题和解决(接谁)额外开销




什么是隐式barrier?

语法
1. #pragma omp parallel (区域构造器)
-
本质功能: 真正向操作系统申请创建(或唤醒)线程组的指令。
-
作用范围: 紧随其后的大括号
{ ... }构成的代码块(Parallel Region)。 -
执行逻辑: 当主线程遇到这条语句时,会派生出多个线程。每个线程都会把花括号内的所有代码完整地执行一遍。 如果里面有一个
printf("Hello\n");,有 4 个线程就会打印 4 次。 -
注意: 它不包含任务分配的功能。如果里面有一个普通的
for循环,4 个线程都会把这个for循环从头到尾跑一遍(冗余计算)。
2. #pragma omp for (工作分担指令 / Work-sharing)
-
本质功能: 任务切分与分发。
-
前置条件: 它不能独立存在,必须被包含在一个已经生效的
omp parallel区域内部。它本身绝对不会创建任何新线程。 -
执行逻辑: 它作用于紧随其后的
for循环。它告诉当前已经存在的线程组:“不要每个人都把整个循环跑一遍,大家把这 $n$ 次迭代平分了”。 -
Barrier 行为: 默认情况下,
omp for循环结束的边界存在隐式的 Barrier。如果业务逻辑不需要同步,可以添加nowait子句(即#pragma omp for nowait)来显式消除这个 Barrier。
3. #pragma omp parallel for (复合指令)
-
本质功能: 语法糖。等价于
#pragma omp parallel紧嵌套一个#pragma omp for。 -
执行逻辑: 同时完成“创建线程组”和“分配紧随其后的
for循环任务”。循环结束后,直接回收/休眠线程组。 -
适用场景: 适用于并行逻辑非常简单,只有单一的、无需与其他并行任务共享生命周期的循环。第一种优化失败的代码正是滥用了这个复合指令。
附加解析:数据环境子句 (Data-sharing Clauses)
幻灯片代码的尾部还带有长长的参数,它们用于控制变量在多线程环境下的作用域:
-
default(none):一种极佳的工程习惯。它强制要求程序员必须显式声明后续所有变量的作用域(是 shared 还是 private),否则编译器直接报错。这能从根源上杜绝许多由于默认变量作用域模糊导致的数据竞争 Bug。 -
shared(a, n):声明变量a(数组首地址)和n(数组长度)是共享的。所有线程看到和操作的都是内存中的同一份数据。 -
private(i, tmp, phase):声明这些变量是私有的。系统会在每个线程的独立栈空间(Stack)里为这些变量创建副本。-
i必须私有,否则各线程的循环索引会互相覆盖。 -
tmp必须私有,因为它是交换两个元素时的中间变量,若共享会导致严重读写冲突。 -
phase必须私有,以确保各线程记录自己当前所处的计算阶段。
-


6.加速比、效率、阿姆达尔定律
加速比,效率(描述,具体的数字进行计算)
线性加速比和超线性加速比(两种方法,并行DFS,超线性加速比的两种场景)

基础评估指标

阿姆达尔定律与理论极限
(Amdahl's Law)

程序的极限加速比,完全取决于代码中无法被并行化的那部分(串行部分)的比例
两个并行化对比的案例:(P为核心数)




超线性加速比

按照常理,加速比 S 的上限是处理器数量 P(即 S <P)。但在极少数特定场景下,程序的加速比会大于核心数(S > P),这就是超线性加速比。最后一张幻灯片给出了导致该现象的两个经典原因:
-
硬件原因:Cache 容量累加效应
单核处理器的 L1/L2 Cache 容量非常有限。当处理大规模数据时,数据无法全部装入 Cache,导致 CPU 必须频繁读取速度极慢的主存(发生 Cache Miss)。
当使用多核并行时,总体可用的 Cache 容量是随核心数线性增加的。如果切分后的子任务数据量恰好能完全塞进各自核心的 Cache 中,主存访问的性能瓶颈被瞬间消除。这种“访存降维打击”带来的时间节省,叠加多核本身的计算并行,就会产生 S > P 的效果。
-
算法原因:非确定性算法(如并行 DFS)
在进行深度优先搜索(DFS)寻找某个目标解时,串行程序可能非常倒霉,它按照固定顺序遍历,把目标解所在的右侧分支留到了最后,导致遍历了整个庞大的搜索树。
而并行 DFS 会让多个核心同时探索不同的分支。只要其中某一个核心“运气好”,迅速在它的分支里找到了目标,整个程序就可以立刻向所有核心发送中止信号,提前结束运行。并行化在这里不仅分担了计算量,更是实质性地削减了总计算量。
7.宏_OPENMP
int类型的十进制数是什么,意义是什么?


1.void Usage(char prog_name[]);,这里是个函数声明,虽然没有用到
2.#ifdef _OPENMP 类似python中的if...,不过发生在预处理阶段,若没有,那么中间代码去掉
# include <omp.h>
#endif 试探是否支持并行
3.#include <stdio.h> #include <stdlib.h>
这里是因为要用printf()这个函数,在库stdio.h中
8.缓存伪共享
矩阵向量乘,代码中有问题,写出优化或是解决方法cache
(缓存行:现代处理器中,缓存是以缓存行的形式进行存储的。每个缓存行通常是64字节大小。当一个线程访问某个变量时,该变量所在的整个缓存行都会被加载到处理器的缓存中)
线程之间没有共享任何变量(但是共享了同一个缓存行),但是它们访问主存的行为看起来好像它们共享了一个变量,这种情况称为伪共享

在原始代码中,多个线程通过 #pragma omp parallel for 并行化了外层循环。虽然 OpenMP 会处理数据划分,但如果多个线程对共享数组 y 的同一内存区域进行频繁的读写更新,会因为 False Sharing(伪共享) 或竞争条件导致 Cache 命中率下降,甚至影响并行性能。
题目要求采用“私有存储”策略,本质上是减少对共享内存的直接竞争,提高局部性。

(最开始的时候,我们把A看作矩阵,后来我们把A看作拉长的)
使用了 sum 变量作为局部累加器。这是一个非常经典的优化手段,原因如下:
-
减少写回次数:在
for (j = 0; j < n; j++)循环中,sum始终保存在寄存器或线程私有的 Cache 中,避免了每次执行y[i] += ...时对共享数组y[i]进行不必要的内存访问。 -
私有化:通过
private(sum),每个线程拥有独立的sum副本,彻底消除了计算过程中的数据竞争。 -
最终同步:计算完成后,执行一次
y[i] = sum,将最终结果写入共享内存。
9.缓存一致性
什么是cache一致性?什么能解决
Cache coherence
在现代计算机中,程序员无法直接控制各个核心的 Cache 何时更新
如图所示,Core 0 和 Core 1 都有独立的 Cache。当 x 等变量被同时存放在多个 Cache 中时,如果 Core 0 修改了 x,Core 1 里的 x 副本就变成了“脏数据”(过时了)。如果缺乏机制来强制同步,两个核心就会对同一个变量产生不同的认知,导致程序行为不可预测。


非一致性带来的后果
-
幻灯片 2 (时间线示例):展示了如果不加控制会发生什么:
-
Time 0: 两者都读取了
x=2。 -
Time 1:
Core 0将x更新为7。但此时Core 1的 Cache 里依然存着旧的2。 -
Time 2:
Core 1使用自己 Cache 里的旧值计算z1 = 4 * x。此时z1得到的是8而不是预期的28(即4 * 7)。
-
-
结论:这就是典型的 Cache 不一致 问题。程序员写的逻辑是连续的,但在硬件层面,由于缺乏同步,计算结果出错了。
-

“Statements not involving x” 指的是那些不读写共享变量 x 的程序语句。在并行计算的分析中,这通常表示核心在处理自己的私有数据(如 y0, y1, z1)或者进行其他不依赖全局状态的计算。
两个解决方法
1.Snooping(侦听机制)
这是一种基于“广播”的解决方案,适用于总线(Bus)架构,不同核心共享总线:
-
侦听(Snoop):每个核心的 Cache 控制器都会时刻“盯着”总线上的信号。
-
广播:当
Core 0修改x时,它会向总线广播一条“我修改了x”的消息。 -
联动:
Core 1听到广播后,立刻将自己 Cache 中存储的x标记为“无效”(Invalid)。下次Core 1再想用x时,它被迫去内存或其它 Cache 重新读取最新值。
2.Directory-based(基于目录的机制)
Snooping 在核心数量很多(比如几十个核)时,总线带宽会成为瓶颈(广播消息太多)。因此有了基于目录的方法:
-
状态跟踪:内存系统维护一个“目录”结构,记录哪些 Cache 缓存了哪些数据行。
-
精准控制:当变量更新时,系统不需要广播给所有人,而是只去查询“目录”,向那些“确实持有该变量”的核心发出“失效”指令。这种方式扩展性更好,适用于超大规模多核处理器。

10.集合通信与点对点通信
集合通信:
1.同一个通信子(Communicator)内的所有进程,必须在逻辑流中都执行到同一个集合函数。
(如果你在进程 A 中写了 MPI_Reduce,但在进程 B 中为了某种逻辑判断将其换成了 MPI_Recv,程序会立刻陷入死锁(Deadlock)或崩溃。因为 MPI 的底层设计要求集合通信在全局范围内保持同步状态,不能混用不同机制。)
2.每个进程传递给MPI集体通信的参数必须具有“兼容性”。集合通信函数对参数的一致性要求极高。所有参与集合操作的进程,必须对通信的全局属性达成共识。
(如果 MPI_Reduce 需要指定一个“目的地进程”(dest_process),所有参与的进程必须传递完全相同的 dest_process 参数(比如都传 0)。如果进程 A 传 0,进程 B 传 1,系统会因无法确定最终结果应送往何处而导致不可预测的行为。)
3.为了满足 MPI 函数调用的完整性,即便某些参数只对特定进程有用,所有进程也必须提供该参数位
(以 output_data_p(结果存放位置)为例,它通常只在 root 进程(接收结果的进程)中有意义。但在代码语法上,其他非 root 进程不能空缺该参数,必须显式填入一个值(通常传入 NULL)。MPI 需要这种参数结构在全员函数调用中保持一致。)
4.点对点通信是根据标签和通信设备进行匹配的;而集合通信则不使用标签,仅依据通信设备及其调用顺序进行匹配。
集合通信与点到点通讯的异同点(课本 p69)

11.数据划分
划分方式:块划分,循环划分等
4个进程,14个数,正确的划分出来
块划分:将连续的组件块分配给每个进程。
循环划分:以轮询方式分配组件。
块-循环划分:采用组件块的循环分布方式。

1. 块划分(Block Partitioning)
-
逻辑:将向量“切分”成连续的几大块,每块分给一个进程。
-
特点:
-
空间局部性极好:每个进程处理的数据在内存中是完全连续的,非常利于 Cache 命中。
-
通信量小:因为数据连续,如果有邻居通信需求,通常只需要处理边界。
-
-
缺点:负载不均衡(Load Imbalance)。如果计算量与数值大小相关(比如某些位置的计算比其他位置复杂),分配到“重任务”数据的进程会成为整个程序的瓶颈(木桶效应)。
2. 循环划分(Cyclic Partitioning)
-
逻辑:像发牌一样,按照顺序一个一个分给进程(0给P0,1给P1,2给P2,3又给P0……)。
-
特点:
-
负载均衡能力极强:如果向量中包含了某些“计算难度很高”的异常点,这种划分方式可以将这些负载分散到所有进程中,避免某个进程过载。
-
-
缺点:数据散乱。进程处理的数据在内存中是不连续的,会导致严重的 Cache Miss,降低单个进程的计算速度。
3. 块-循环划分(Block-Cyclic Partitioning)
-
逻辑:这是前两者的折中。它先将数据切成小块(例如图中
Blocksize = 2),然后再像“发牌”一样把这些小块轮流分给进程。 -
特点:
-
平衡了局部性与负载均衡:既保留了一定程度的连续性(小块内部是连续的),又通过轮流分配实现了负载均衡。
-
-
应用场景:这在高性能计算(HPC)中最常用,尤其是在大规模矩阵运算库(如 ScaLAPACK)中,这种方式能有效应对复杂的并行计算拓扑。
12.流水线
指令级并行:流水线结构——功能单元按阶段排列;多发指令——可同时执行多条指令。



a.
-
取操作数 (Fetch operands)
-
比较阶码 (Compare exponents)
-
对阶 (Shift one of the operands)
-
尾数相加 (Add)
-
规格化结果 (Normalize result)
-
舍入结果 (Round result)
-
存结果 (Store result)
9ns
b.在非流水线(Unpipelined)架构中,每个计算任务必须等待上一个任务完全结束后才能进入处理器。总耗时为单次耗时乘以任务总数 N

c.


表格理解:
这是另一个所有指令花费时长都为1的问题


13.缓存
ache 两种循环方式的效率
内外循环的调用次序不同
哪种效率高?解释(第1种 p15)
一组可在比其他某些内存位置更短时间内访问的存储地址。
CPU缓存通常位于同一芯片上,或其访问速度远快于普通内存。
Cache Line(缓存行)机制:当 CPU 想读取内存中的某个数据(比如 A[0][0])时,它绝不会只拿那一个字节。硬件会想:“既然你用了这个数据,你接下来大概率会用它旁边的数据。” 因此,Cache 每次都会从内存中搬运一整块连续的数据(比如图中的 4 个元素)进来,这就叫一个 Cache Line。

在 C 语言中,二维数组是按行优先(Row-Major)在内存中连续排列的。 A[0][0] 旁边紧挨着的是 A[0][1],而不是 A[1][0]。

14.Flynn 分类法

费林分类法(Flynn's Taxonomy)是计算机体系结构中最经典的宏观分类标准。它仅仅根据系统在同一时间内处理的指令流(Instruction Stream)和数据流(Data Stream)的数量,将计算机划分为四个象限。
以下是对这四种分类的严谨剖析:
1. SISD(单指令流单数据流)
-
运行机制:系统在任何一个时刻,只能从内存中读取一条指令,并对单一的数据单元执行该指令。
-
本质:这是最纯粹的经典冯·诺依曼架构。
-
特性:正如上一题所讨论的,即便现代处理器在底层使用了流水线、超标量(多发射)等技术来压榨硬件性能,只要它在软件逻辑层面依然是单线程顺序执行的,它就仍被归类为 SISD。
-
实例:早期的单核单线程处理器(如 Intel 8086、早期 ARM 芯片)。
2. SIMD(单指令流多数据流)
-
运行机制:系统读取一条指令,但该指令会被广播给多个运算单元,同时作用于多个不同的数据。
-
本质:这是一种数据级并行(Data-Level Parallelism, DLP)。
-
特性:在处理同构的大规模数据(如图像像素、矩阵计算)时效率极高,因为省去了大量重复取指令的开销。
-
实例:
-
现代 CPU 中的向量指令集(如 Intel 的 SSE, AVX)。
-
GPU(图形处理器)的核心工作模式。当你编写 CUDA 程序时,大量线程执行同一套内核(Kernel)代码但处理不同的数据块,就是典型的 SIMD(在 NVIDIA 术语中称为 SIMT)。
-
3. MISD(多指令流单数据流)
-
运行机制:多个不同的处理器(或核心)接收同一个数据流,但各自对该数据执行不同的指令。
-
本质:在通用计算领域几乎没有实际价值,因此你的幻灯片上标注了“not covered”。
-
实例:主要见于极少数特种场景。
-
高可靠性容错系统:例如航天飞机的飞行控制系统,多个处理器对同一个传感器数据进行不同的计算验证,以防止硬件故障导致错误。
-
某些特定的流式密码破译或信号过滤硬件。
-
4. MIMD(多指令流多数据流)
-
运行机制:系统包含多个独立工作的处理器,每个处理器都可以独立地提取自己的指令流,并处理各自独立的数据流。
-
本质:这代表了任务级并行(Task-Level Parallelism)或线程级并行(Thread-Level Parallelism)。
-
特性:这是当今高性能并行计算(HPC)的绝对主力。它具有极高的灵活性,处理器之间可以执行完全无关的程序。
-
实例:
-
现代的多核处理器(如八核酷睿、AMD Ryzen)。
-
分布式的超级计算机集群。
-
-
进阶细分:MIMD 在实际应用中通常被细分为两类:
-
共享内存 MIMD:所有处理器访问同一块物理内存。我们之前讨论的一致性协议(Cache Coherence)和 OpenMP 就是在这个背景下工作的。
-
分布式内存 MIMD:每个处理器有自己独立的本地内存,互相之间通过网络通信。你之前看到的 MPI(如
MPI_Send、集合通信)就是为这种架构设计的。
-
Flynn 分类法的核心评判标准只有两个:指令流(Instruction Stream)和数据流(Data Stream)的数量。只要系统在逻辑上依然是按照单一的顺序序列读取指令并处理数据,它就依然属于 SISD(单指令单数据流)。

不变,改变,改变。说出原因
指令级并行:
流水线 - 功能单元按阶段排列。
多发射 - 可以同时启动多条指令。
①加人缓存和虚内存不会改变系统的SISD类型。因为缓存和虚内存的加入没有改变一次可以执行的指令数或者一次可以操作的数据量,因此系统仍为SISD类型。
②加入流水线会使得系统变为SIMD类型。因为加入流水线相当于将一个指令应用于多个数据项,也就是SIMD。
③多发射和硬件多线程 会使得系统变为MIMD类型。
因为多发射相当于一个指令流在多个执行单元(如:ALU)上同时启动一部分,分裂为多个指令流,多个指令流用于不同的数据项,就是MIMD。
硬件多线程是指当前任务被阻塞时,系统尝试切换到其他线程继续工作,这允许了多线程的存在,多个指令流用于不同的数据项,也是MIMD。

15.集合通信的匹配
函数调用的次序有关,实际的b和d的值
1号进程写错,原本正确结果?实际结果?
集合通信(如 MPI_Reduce)的匹配,绝对不看变量名,只看函数调用的先后顺序!

原本正确的结果:b=3,d=6
实际结果:b=4,b=5


16.矩阵向量乘法-无法整除
(4.1)m和n不能被线程整除的情况,如何处理余数


【能求出这一个线程处理的任务块的起始索引和结束索引】

代码最后的红圈 -1 是因为 C/C++ 数组是 0 索引的闭区间。如果从 my_first_i 开始,总共分配了 my_n_count 个任务,那么最后一个任务的索引理应是起始索引加上任务数再减一。
我们可以代入一组具体数据:假设任务总数 n = 14,进程数 p = 4。
-
quotient= 14 / 4 = 3 -
remainder= 14 % 4 = 2 (意味着前 2 个进程要多拿一个任务)
| 进程 (my_rank) | 是否满足条件 (< 2) | 任务数 (my_n_count) | 起始索引 (my_first_i) | 结束索引 (my_last_i) |
| 0 | 是 | 3 + 1 = 4 | ||
| 1 | 是 | 3 + 1 = 4 | ||
| 2 | 否 | 3 | ||
| 3 | 否 | 3 |
17.矩阵向量乘法-Ptheads并行化


ptheads并行化后:

-
获取身份:
long my_rank = (long) rank;每个线程要知道自己的编号(比如 0, 1, 2, 3)。 -
计算份额:
int local_m = m / thread_count;-
注意:这段代码采用的是最简单的理想模型,即假设总行数
m能被线程数thread_count完美整除。这就是为什么这里没有看到我们上一节讲的 quotient 和 remainder 逻辑。
-
-
划定边界(核心!):
-
my_first_row = my_rank * local_m; -
my_last_row = (my_rank + 1) * local_m - 1; -
如果 $m=12$,4 个线程,那么 1 号线程(
my_rank=1)会算出自己的起始行是 3,结束行是 5。这正是块划分的完美体现。
-
-
执行计算:
-
仔细看代码的下半部分,双层
for循环内部的计算逻辑与串行代码一模一样! -
唯一的区别是:外层循环不再是
for (i = 0; i < m; ...),而是被替换成了for (i = my_first_row; i <= my_last_row; ...)。
-
18.CUDA基础
i的位置如何计算(找错)
同步问题

转为cuda形式,需要计算i的坐标:int i = blockIdx.x * blockDim.x + threadIdx.x;



可能的陷阱:

块同步与死锁

CUDA 中极其重要的块内屏障指令:__syncthreads(),在读取其他线程写的值时,保证正确
块内的所有线程必须执行相同的 __syncthreads(),否则 GPU 将会死锁

补充
threadIdx
blockIdx
blockDim
gridDim
threadIdx、blockIdx、blockDim 和 gridDim 是用于定位线程在网格(Grid)和线程块(Block)中物理位置的内置变量。
1.一维网格,一维线程块(1D Grid, 1D Block)
int idx = threadIdx.x + blockIdx.x * blockDim.x;
2.二维网格,二维线程块 (2D Grid, 2D Block)
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
// 全局网格的宽度
int width = gridDim.x * blockDim.x;
// 映射到一维数组
int idx = x + y * width;
3.三维网格,三维线程块 (3D Grid, 3D Block)
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int z = threadIdx.z + blockIdx.z * blockDim.z;
int width = gridDim.x * blockDim.x;
int height = gridDim.y * blockDim.y;
int idx = x + y * width + z * (width * height);
19.CUDA优化
理解前3个kernel,指出存在的问题和改正,核心代码的修正
v1问题:
1.除法和区域取余,会编译成20个加法指令,非常慢(特殊情况下可以用位操作)
2.32个线程组成的warps一起使用效率才更好
如何解决上述问题?(去掉non-divergent)
v2:所有线程都进行操作;且没有取余操作;cuda奇数块不冲突,偶数块冲突(Bank conflicts),偶数块有线程
针对这个偶数块冲突问题,如何优化?
v3:(这里不用除法,而是用位操作)
重点想v1-v2改变线程的对应模式(divergent->non-divergent)
不过现在只有一半线程被使用


问题:1.线程束分化(Warp Divergence)使得效率低 2.求余运算非常慢
GPU 调度线程的最小单位是线程束(Warp),通常包含 32 个线程。这 32 个线程在硬件上是绑定在一起的,必须执行相同的指令。
如果在代码里写了 if-else 分支,导致这 32 个线程里有一部分想走 if,另一部分想走 else,GPU 无法让它们分道扬镳。它只能让整组线程先把 if 走一遍(此时不需要走 if 的线程闲置),再把 else 走一遍(此时不需要走 else 的线程闲置)。这会极大地拉低执行效率。
-
在 $s=1$ 时的表现:条件是
tid % 2 == 0。这意味着线程 0, 2, 4, 6 是活跃的,而 1, 3, 5, 7 是闲置的。在一个 Warp 内,活跃线程和闲置线程像斑马线一样交替排列。这就引发了极其严重的 Warp Divergence,并行效率直接砍半。



新的问题:由于线程访问共享内存的步长(Stride)是 2 的幂次递增的,这导致同一线程束(Warp)内的多个线程会同时访问映射到同一个 Shared Memory Bank 的不同地址。这种多路 Bank 冲突(Bank Conflicts)迫使硬件将原本可以并行的访存请求串行化(Serialization)处理,从而导致访存带宽大幅下降,成为算法的性能瓶颈。”


将循环倒置,使得最内层的访问地址直接与原生连续的 tid 绑定。它顺应了硬件的取模机制,使得任何时刻一个 Warp 内的访存请求都是连续的,从而实现了 100% 的无冲突满带宽访问

20.主观题
与实验相关(优化思路,写出来,编程模型的特点,在实践中的特点,实验中的优化思路的优化方法等,收获)
对并行编程模式的理解,实验的优化思路,用到的方法
一、 宏观认知:并行编程模式的理解与特点
【问题】1. 我使用了哪些并行编程模型?它们在理论上的核心特点是什么? 【答案】本实验使用了 OpenMP 并行编程模型 。其核心特点是基于共享内存架构,通过编译制导指令(如 #pragma omp)实现多线程并发,编程相对简单且易于在现有的串行代码上进行扩展 。
【问题】2. 这些模型在实际工程实践中,表现出了哪些与理论不同的特点或痛点? 【答案】在处理生物序列高维稀疏特征比对时,单纯增加线程并不能带来理想的线性加速 。具体的工程痛点包括:动态内存分配(如 C++ 中 vector 的频繁扩容)在多并发下会引发内核态的锁争用,且内存碎片化会导致严重的 Cache 缺失惩罚 ;底层的控制流分支跳转会打断 CPU 的预取与流水线 ;此外,由于不同序列长度差异巨大导致的数据倾斜,外加物理多核满载时的降频机制,多核并行极易出现严重的负载失衡与停机空转 。
二、 核心实践:实验中的优化思路与具体方法
【问题】3. 面对实验任务,我的整体优化思路(优化方向)是什么? 【答案】为了解决上述瓶颈,整体实施了自顶向下的五个步进式深度优化策略 :通过数学推导进行算法级剪枝以控制空间复杂度;通过物理连续性重构打破访存“内存墙”;通过架构级解耦实现无分支化;配合数据结构梳理去激发编译器的现代硬件向量化能力;最后利用动态启发式调度抹平计算负载的不均衡 。
【问题】4. 针对上述思路,我具体落实了哪些优化方法?
【答案】
算法重构:
传统的正向索引记录的是“每个序列包含哪些 k-mer 特征” ;而倒排索引记录的是“每个 k-mer 特征被哪些序列所包含”,从而用于极速检索具有相同特征的候选序列 。
算法级剪枝优化:利用加权 Jaccard 相似度定义的数学性质推导候选序列的物理长度上限,结合 O(log N) 的 std::upper_bound 二分检索确定截断下标,将 O(N) 的无效扫描降维至极小的局部区间 。
访存级内存重构:将动态嵌套容器 std::vector<std::vector<Posting>> 彻底重构为全静态的 CSR 扁平化连续物理内存架构,利用前缀和数组与连续分配规避随机跳跃,实现硬件预取器的完美 Cache Line 对齐 。
架构级去分支化:剥离生物序列中占比极高的频次为 1 的孤立特征,将内层循环精简为 shared_mins[*ptr]++ 形式的无分支寻址自增指令,保持 CPU 流水线不中断 。
激发硬件级向量化:抛弃繁杂的手写底层位宽优化,恢复最贴合 x86 体系结构的 32 位整型,利用 #pragma GCC ivdep 原语消除循环依赖,让编译器自动生成 AVX2 / AVX-512 的 SIMD 掩码测试指令加速数组清理 。
多核级动态调度:抛弃 OpenMP 默认的静态分配(schedule(static)),在并行主体循环中引入 #pragma omp for schedule(guided, 8) 的动态启发式负载均衡,让未被降频的核心自适应接管更多任务块 。
三、 总结升华:实验收获与反思
【问题】5. 实际测试结果如何验证了我的优化方法? 【答案】最终程序在 14核20线程的硬件环境下,16个线程将运行时间从单核的 14.69 秒缩短至 1.92 秒,加速比达到 7.65 。消融实验有力地证明了各模块的有效性:算法级边界剪枝的作用最具决定性,剥离该模块会导致耗时剧增约 75% ;同时,动态调度也成功证明了其在异构多核平台(如 P-Core 与 E-Core 混合架构)面对不均匀数据分布时,表现明显优于传统静态均分 。
【问题】6. 通过本次实验,我获得了什么核心工程经验?(收获)
【答案】
性能瓶颈的转移认知:并行程序的加速绝不仅仅等于“增加线程数”,很多时候木桶效应的短板不在纯粹的算力上,而是隐藏在访存模式的随机性、内存连续性以及多核任务的分配倾斜中 。
对“手工微操”的克制与反思:消融实验揭示了一个反常识规律,某些看似极客的底层干预(如手动降维 16 位或设计 64 位跳跃重组)反而会干扰现代编译器(如 GCC)建立依赖图 。保持数据结构清爽和循环逻辑干净,充分信任并依靠编译器去做 SIMD 向量化,才是当前算力平台下获得极致性能的更优解 。
课程收获
一、 课程核心知识点梳理
这门课程通常采用自底向上与实践相结合的脉络,核心模块主要包括以下几个方面:
1. 硬件与架构基础 (Hardware & Architecture)
-
冯·诺依曼架构的演进:从单核到多核处理器,缓存(Cache)层级结构对性能的影响。
-
内存架构:深入理解共享内存(Shared-Memory)与分布式内存(Distributed-Memory)系统的本质区别。
-
互连网络:集群内部节点是如何通过总线、交叉开关或超立方体网络进行通信的。
2. 并行编程模型与核心API (Programming Models)
-
共享内存编程 (Pthreads & OpenMP):
-
Pthreads:底层的线程管理,涉及互斥锁(Mutex)、信号量(Semaphore)、条件变量以及读写锁的设计。
-
OpenMP:基于编译指示(Pragmas)的高级并行,包括并行 for 循环、规约操作(Reduction)以及不同调度策略的使用。
-
-
分布式内存编程 (MPI):
-
消息传递接口(Message Passing Interface)的基础。
-
点对点通信与集合通信(如广播、归约、发散与收集)。
-
MPI 数据类型的派生与自定义通信域。
-
-
异构计算与 GPU 编程 (CUDA) (注:现代课程通常会加入此部分):
-
GPU 线程层次结构(Grid, Block, Thread)。
-
显存架构与内存合并(Memory Coalescing),加速大规模矩阵操作。
-
3. 并行算法设计与模式 (Algorithms & Patterns)
-
经典算法的并行化:并行排序(如奇偶转置排序)、前缀和(Scan/Prefix Sum)、矩阵与向量乘法、N体问题(N-Body Solvers)。
-
负载均衡:静态与动态任务分配策略。
4. 性能评估与优化 (Performance & Optimization)
-
核心指标:加速比(Speedup)、效率(Efficiency)以及阿姆达尔定律(Amdahl's Law)的理论极限。
-
常见陷阱:死锁(Deadlock)、竞争条件(Race Condition)、伪共享(False Sharing)及其在缓存一致性协议下的影响。
二、 课程收获
学习这门课程带来的不仅是掌握几个编程库,更重要的是底层思维的升华:
-
从串行到并发的思维跃迁:习惯于拆解问题,将单一的计算流转化为多维的并发任务,并能精准识别出算法中的数据依赖与通信瓶颈。
-
打破“软件与硬件隔离”的错觉:深刻理解代码在真实物理硬件上的运行机制。你会意识到,有时候调整一下数组的遍历顺序(利用 Cache 命中率)或者避免伪共享,带来的性能提升甚至远超算法时间复杂度常数级的优化。
-
突破前沿技术的算力瓶颈:在实际工程和科研中,这门课的价值不可估量。例如,在处理海量工业级数据、训练复杂的推荐系统模型(如 DeepFM 或多任务 MMoE 结构)时,分布式计算是打破算力瓶颈的唯一途径。同样,在计算机视觉领域(比如基于图神经网络的时空预测,或视频的隐式表示等研究),底层的大规模矩阵运算极其依赖 CUDA 等异构计算手段的优化。掌握并行程序设计,意味着你能真正将这些前沿算法落地,并将其性能发挥到极致。
注
(注:部分答案参考山东大学软件学院2023-2024第二学期多核平台上的并行计算期末复习总结_山东大学多核期末-CSDN博客,如缓存一致性部分和Flynn分类法部分,侵删)
更多推荐




所有评论(0)