基础概念

并发计算、并行计算与分布式计算

冯·诺依曼架构

主存：存储指令和数据的集合，每个位置有地址和内容。

CPU：

• 控制单元：决定程序中哪条指令被执行。
• 算术逻辑单元（ALU）：执行实际指令。

关键术语：

• 寄存器：CPU中非常快的存储部分。
• 程序计数器：存储下一条指令的地址。
• 总线：连接CPU和内存的硬件。

瓶颈问题：内存访问速度远低于CPU处理速度，导致CPU等待数据。

进程与线程

进程：正在执行的计算机程序的实例。

• 组件：可执行程序、内存块、资源描述符、安全信息、进程状态。

多任务：通过时间片轮转，让单个处理器系统看起来同时运行多个程序。

线程：进程内的独立执行单元。

• 优势：当一个线程阻塞时，其他线程可以继续运行。
• 操作：
  • 创建线程：称为forking。
  • 终止线程：称为joining。

缓存基础

缓存：比普通内存更快访问的存储集合。

局部性原理：

• 空间局部性：访问一个位置后，可能会访问附近的位置。
• 时间局部性：最近访问的位置可能会再次被访问。

缓存级别

L1缓存：最小且最快。

L2缓存：较大且较慢。

L3缓存：最大且最慢。

缓存命中与未命中

缓存命中：数据在缓存中找到。

缓存未命中：数据不在缓存中，需要从主存中获取。

缓存更新策略

写穿：在数据写入缓存时同时更新主存中的数据。

写回：将缓存中的数据标记为脏。当缓存行被新缓存行替换时，脏行会被写入主存。

缓存映射

全相联映射：新行可以放在缓存的任何位置。

直接映射缓存：每个缓存行有唯一的缓存位置。

$n$路组相联映射：每个缓存行可以放在$n$个不同的缓存位置之一。

虚拟内存

虚拟内存：作为二级存储的缓存，利用空间和时间局部性原理。

• 交换空间：存储不活动的部分。
• 页面：数据和指令的块，通常较大（4-16KB）。

并行

指令级并行（ILP）

指令级并行：通过多个处理器组件或功能单元同时执行指令来提高性能。

• 流水线：功能单元按阶段排列。
• 多发射：同时启动多个指令。

硬件多线程

硬件多线程：当当前执行的任务停滞时，系统继续执行其他任务。

• 细粒度多线程：每条指令后切换线程。
• 粗粒度多线程：仅在等待耗时操作完成时切换线程。
• 同时多线程（SMT）：允许多个线程使用多个功能单元。

程序、进程和线程

• 程序：存储在计算机上的代码。
• 进程：程序的运行实例，具有独立的内存地址空间。
• 线程：进程内的执行单元，共享堆但具有独立的栈。

CPU、核心和处理器

• CPU：包含一个或多个核心。
• 核心：CPU的基本计算单元，可以运行单个程序上下文或多个上下文（如果支持硬件线程）。
• 多处理器系统：在多处理器系统中，多个 CPU（每个 CPU 可能有多个核心）协同工作
• 超线程技术：将一个物理核心模拟为两个逻辑核心，允许多个线程同时运行。

多处理器和多核CPU

• 多处理器系统：包含多个CPU，允许并行工作（SMP，对称多处理）。
• 多核CPU：单个CPU芯片上包含多个核心。

线程和核心的关系

• 单个进程：可以在多个核心上运行多个线程。
• 单个线程：不能同时在多个核心上运行，但可以通过指令流水线和乱序执行实现指令级并行。

并行计算机的分类

传统分类方法

这样分类不能反映机器的系统结构特征

弗林分类法

为了解决原有分类方法的缺点，1972年Michael J.Flynn提出了一种基于数据流和指令流的并行

SISD

SISD计算机是一种传统的串行计算机。它从硬件上就不支持任何并行化计算，所有的指令都是串行执行

SIMD

SIMD计算机可以实现数据级并行，对多个不同的数据流并行执行相同的数据处理操作。

比如同时对$n$对数据（x[i]、y[i]）做加法。

实例有阵列计算机，阵列计算机的基本思想是用一个单一的控制单元提供信号驱动多个处理单元同时运行。

和SIMT的区别

SIMT用一条指令控制多个独立线程，每个线程可处理不同数据，甚至走不同逻辑分支（有的线程走if有的线程走else）。但是SIMD只能执行相同的操作，比如都执行加法，不能有不同的逻辑分支。

MIMD

MIMD中的每个处理器都有自己的指令流，也可以和其他处理器共享指令流，对自己的数据进行处理

MIMD细化分类

根据不同的CPU是如何组织和共享内存的， MIMD机器分为如下的两类：

• 共享式内存（UMA）：共享式内存系统就是处理器之间共享内存
• 分布式内存（NUMA）：分布式内存系统（消息驱动）是处理器之间不共享内存，通过消息驱动来通信。分布式内存系统可以分为如下两类：
  • 大规模并行处理器系统（MPP）：MPP系统是由成百上千台计算机组成的大规模并行计算机系统。MPP中一般每个节点可以认为是一个没有硬盘的计算机
  • 工作站机群系统（COW）：COW系统是由大量的家用计算机或者工作站通过商用网络连接在一起而构成的多计算机系统。

互联网络

互联网络（ICN）是一种可编程系统，用于在终端之间传输数据。分为：

• 广域网

• 本地区域网络

• 系统区域网络

• 片上网络

将不同计算的设备连接在一起

网络拓扑

网络拓扑结构等同于网络形状，要确定消息从一个节点到另一个节点需要传输的距离。其中：

• 直径：最大距离
• 平均距离：所有距离之和/结点对数

共享内存互连和分布式内存互连

共享内存互连

总线互连：一组并行通信线以及一些控制总线访问的硬件。随着连接到总线的设备数量增加，总线使用竞争加剧，性能下降。

交换互联：使用交换机来控制连接设备间数据路由。

分布式存储互连

直接互连：每个交换机直接连接到一个进程或内存对，交换机相互连接。

• Bisection width（二分宽度）：用于衡量一个网络或图在结构上的某种“宽度”或“连通性”。具体来说，它指的是将图的顶点集分成两个大小相等的子集时，需要移除的最少边数。
• 带宽：链路输数据的速度
• 二分带宽：指在将网络的节点集分成两个大小相等的子集时，这两个子集之间可以实现的最大的通信带宽。

间接连接：交换机不能直接连接到处理器。

并行软件

概念

SPMD（单程序多数据流）：单个可执行文件通过条件分支模拟多个不同程序的行为（如if (I’m process i)控制逻辑）

竞态条件（Race Condition）：多个线程同时修改共享变量导致不可预测结果

临界区（Critical Section）：需互斥访问的代码段，通过锁（Lock/Unlock）保护

忙等待（Busy-Waiting）：通过轮询标志位实现同步（如while (!ok_for_1);）

性能分析指标

加速比

并行时间（T_parallel）与串行时间（T_serial）的比值，理想为线性加速（T_parallel = T_serial/p）

实际加速比：需考虑开销（T_parallel = T_serial/p + T_overhead）

效率

定义：加速比与核心数之比（E = Speedup/p），衡量资源利用率

Amdahl定律

程序仅有部分可并行化（如90%），加速比上限受限于串行部分（如最高加速比5倍）

时间类型

CPU时间（计算耗时） vs. 墙钟时间（实际总耗时）

可扩展性

强可拓展性

强可扩展性：在问题规模（即任务大小）固定的情况下，通过增加计算资源（如处理器数量）来缩短计算时间。

理想情况：资源增加一倍，计算时间减半（线性加速）。

弱可拓展性

弱可拓展性：在增加计算资源的同时，问题规模也按比例增加，目标是保持每个处理器的负载不变，计算时间基本稳定。

理想情况：资源增加一倍，问题规模也增加一倍，计算时间保持不变。（效率不变）

Foster（福斯特）方法

1. 任务划分（Partitioning）：分解计算与数据为可并行任务
2. 通信分析（Communication）：定义任务间必要的数据交换
3. 任务聚合（Agglomeration）：合并任务以减少通信开销
4. 映射（Mapping）：分配任务至进程/线程，均衡负载并最小化通信

MPI

基础概念

采用SPMD（单程序多数据）模式，所有进程执行同一程序，通过条件分支实现差异化操作。

进程标识：非负整数rank（0到p-1），p为总进程数

通信器（Communicator）：进程集合，默认使用MPI_COMM_WORLD包含所有进程。

必须包含头文件mpi.h。

操作详解

编译与运行

通过以下指令进行编译：

1

mpicc -g -Wall -o mpi_hello mpi_hello.c

执行：

1

mpiexec -n <number of processes> <executable>

MPI_Init、MPI_Finalize

MPI_Init是用来进行MPI初始化的，MPI_Finalize是用来在程序结束后做清理的

1
2
3
4
5
6
7
8

# include "mpi.h"

int mian(int argc,char* argv[])
{
  MPI_Init(&argc,&argv);
  //...
  MPI_Finalize();
}

通信子

概念

通信子是一组可以互相发送消息的进程的集合。

所有进程都包含在MPI_COMM_WORLD当中。

MPI_Comm_size、MPI_ Comm_rank

MPI_Comm_size用来获取有多少个进程在这个通信子中，后者用来获得进程在通信子中的标识符。

函数原型：

1
2
3
4
5
6
7
8
9

int MPI_Comm_size (
  MPI_Comm comm
  int* comm_sz_p
);

int MPI_ Comm_rank (
  MPI_Comm comm
  int* my_rank_p
);

第一个参数是通信子，第二个参数用于获得结果，函数返回的int值代表操作有没有执行成功。

Send、Recv

Send、Recv属于点对点通信。

函数原型：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

int MPI_Send(
  void* msg_buf_p //发送缓冲区
  int msg_size //缓冲区中MPI_Datatype值的个数(msg_buf_p的字节数/MPI_Datatype字节数)
  MPI_Datatype msg_type, //MPI数据类型
  int dest, //目标进程
  int tag, //表示该通信的标识符
  MPI_Comm communicator //通信子
);

int MPI_Recv(
  void* msg_buf_p, //接受缓冲区
  int buf_size, //接收缓冲区长度
  MPI_Datatype buf_type, //数据类型
  int source, //源进程
  int tag //表示该通信的标识符
	MPI_Comm communicator, //通信子
  MPI_Status* status_P //通信是否成功
);

Gather、Reduce、Broadcast、Scatter

Gather、Reduce、Broadcast、Scatter属于集合通信，通信子中的所有进程必须调用相同函数（如MPI_Reduce、MPI_Bcast），不使用tag

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58

int MPI_Reduce(
    void* input_data_p,    /* 输入数据指针，指向当前进程的数据 */
    void* output_data_p,   /* 输出数据指针，仅在目标进程中有效 */
    int count,             /* 数据元素个数 */
    MPI_Datatype datatype, /* 数据类型 */
    MPI_Op operator,       /* 归约操作符 */
    int dest_process,      /* 目标进程排名 */
    MPI_Comm comm          /* 通信域 */
);

int MPI_Allreduce(
    void* input_data_p,    /* 输入数据指针，指向当前进程的输入缓冲区 */
    void* output_data_p,   /* 输出数据指针，指向当前进程的输出缓冲区 */
    int count,             /* 数据元素个数，表示每个进程处理的数据项数量 */
    MPI_Datatype datatype, /* 数据类型，如 MPI_INT、MPI_DOUBLE 等 */
    MPI_Op operator,       /* 归约操作符，如 MPI_SUM、MPI_PROD、MPI_MAX 等 */
    MPI_Comm comm          /* 通信域，指定参与此操作的进程组 */
);

int MPI_Bcast(
    void* data_p,      /* 指向要广播的数据的指针，输入/输出参数 */
    int count,         /* 要广播的数据项数量，输入参数 */
    MPI_Datatype datatype, /* 数据项的MPI数据类型，输入参数 */
    int source_proc,   /* 发送广播数据的源进程的rank，输入参数 */
    MPI_Comm comm      /* 通信域，指定广播操作的范围，输入参数 */
);

int MPI_Scatter(
    void *send_buf_p,      /* 发送缓冲区指针，包含要分散的数据，输入参数 */
    int send_count,        /* 发送到每个进程的数据项数量，输入参数 */
    MPI_Datatype send_type, /* 发送数据的数据类型，输入参数 */
    void *recv_buf_p,      /* 接收缓冲区指针，存储接收到的数据，输出参数 */
    int recv_count,        /* 从源进程接收的数据项数量，输入参数 */
    MPI_Datatype recv_type, /* 接收数据的数据类型，输入参数 */
    int src_proc,          /* 源进程的rank，数据发送方，输入参数 */
    MPI_Comm comm          /* 通信域，指定散射操作的范围，输入参数 */
);

int MPI_Gather(
    void *send_buf_p,      /* 发送缓冲区指针，包含要收集的数据，输入参数 */
    int send_count,        /* 发送到接收进程的数据项数量，输入参数 */
    MPI_Datatype send_type, /* 发送数据的数据类型，输入参数 */
    void *recv_buf_p,      /* 接收缓冲区指针，存储收集的数据，输出参数 */
    int recv_count,        /* 从每个进程接收的数据项数量，输入参数 */
    MPI_Datatype recv_type, /* 接收数据的数据类型，输入参数 */
    int dest_proc,         /* 目标进程的rank，数据收集的目的地，输入参数 */
    MPI_Comm comm          /* 通信域，指定收集操作的范围，输入参数 */
);

int MPI_Allgather(
    void *send_buf_p,      /* 发送缓冲区指针，包含要收集的数据，输入参数 */
    int send_count,        /* 发送到每个进程的数据项数量，输入参数 */
    MPI_Datatype send_type, /* 发送数据的数据类型，输入参数 */
    void *recv_buf_p,      /* 接收缓冲区指针，存储所有进程收集的数据，输出参数 */
    int recv_count,        /* 从每个进程接收的数据项数量，输入参数 */
    MPI_Datatype recv_type, /* 接收数据的数据类型，输入参数 */
    MPI_Comm comm          /* 通信域，指定收集操作的范围，输入参数 */
);

Allreduce规约有树形和蝶形之分：

高级特性

创建自己的数据类型

可以用MPI做数据的聚合。创建流程：MPI_Type_create_struct → MPI_Type_commit → MPI_Type_free

计时

MPI_Wtime()可以返回自过去某个时间点以来经过的秒数。

返回两个MPI_Wtime之间的秒数。

同步

MPI_Barrier：同步所有进程计时起点。

确保在通信子中的每个进程开始调用它之前，其他调用它的进程都在这里阻塞。

MPI Safety

MPI 安全性（MPI Safety）主要关注在多线程环境中正确使用 MPI 函数，避免竞争条件和数据不一致等问题。

比如所有进程都调用Send却没人Recv就会死锁。

Pthreads编程

基本概念

POSIX线程标准，提供Unix-like系统的多线程API

操作详解

编译与运行

编译：

1

gcc -g -Wall -o pth_hello pth_hello. c -lpthread

运行：

1

./pth_hello <number of threads>

线程创建

使用pthread_create(thread_p, attr_p, start_routine, arg_p)函数可以创建线程。

1
2
3
4
5
6

int pthread_create(
    pthread_t* thread_p,       /* 输出参数：指向存储新线程ID的变量 */
    const pthread_attr_t* attr_p, /* 输入参数：线程属性（NULL表示默认属性） */
    void* (*start_routine)(void*), /* 输入参数：线程入口函数指针 */
    void* arg_p                /* 输入参数：传递给start_routine的参数（void*类型） */
);

pthread_t为线程标识符，具有系统特定的不透明属性。

线程函数原型为void* thread_function(void* args_p)，支持参数传递。

线程终止

使用pthread_join阻塞等待指定线程结束：

1

int pthread_join(pthread_t thread, void** retval);

thread：待等待的线程 ID（由 pthread_create() 返回）。

retval：输出参数，指向线程返回值的指针（类型为 void**）。若无需获取返回值，可传 NULL。

同步机制

互斥锁

POSIX线程标准包括一个特殊的类型用于互斥锁：pthread_mutex_t。

1
2
3
4

int pthread_mutex_init(
    pthread_mutex_t* mutex_p,         /* 输出参数：指向要初始化的互斥锁对象 */
    const pthread_mutexattr_t* attr_p /* 输入参数：指向互斥锁属性对象（NULL表示默认属性） */
);

核心操作：

1
2
3

pthread_mutex_lock(mutex_p)  // 获取锁
pthread_mutex_unlock(mutex_p)  // 释放锁
pthread_mutex_destroy(mutex_p)  // 销毁锁

信号量

互斥锁没法控制访问的顺序。这时候可以使用信号量。

1
2
3
4
5
6
7

#include <semaphore.h>

int sem_init(
    sem_t* semaphore_p,       /* 输出参数：指向要初始化的信号量对象 */
    int shared,               /* 输入参数：共享选项（0=线程间共享，非0=进程间共享） */
    unsigned int initial_val  /* 输入参数：信号量的初始值 */
);

核心操作：

1
2
3

int sem_destroy(sem_t* semaphore_p); /* 输入/输出：指向要销毁的信号量对象 */
int sem_post(sem_t* semaphore_p);    /* 输入/输出：指向要增加计数的信号量对象 */
int sem_wait(sem_t* semaphore_p);    /* 输入/输出：指向要等待的信号量对象 */

sem_wait可以用来阻塞当前线程，直到信号量的值大于0，解除阻塞。解除阻塞后，sem的值-1，表示公共资源被执行减少了。

sem_post用于增加信号量的值，信号量+1。当有线程阻塞在这个信号量上时，调用这个函数会使其中的一个线程不再阻塞，选择机制由线程的调度策略决定。

条件变量

条件变量用于等待某个条件成立（通常与互斥锁配合使用），而信号量、限制对资源的并发访问数量。

1
2
3
4
5
6
7
8

pthread_cond_wait(
    pthread_cond_t* cond,    /* 输入/输出：指向要等待的条件变量 */
    pthread_mutex_t* mutex   /* 输入/输出：指向与条件变量关联的互斥锁 */
);  // 原子操作：释放mutex锁，并阻塞当前线程直到条件变量cond被唤醒；唤醒后自动重新获取mutex锁

pthread_cond_signal(
    pthread_cond_t* cond     /* 输入/输出：指向要唤醒的条件变量 */
);  // 唤醒至少一个在cond条件变量上等待的线程（若有多个线程等待，唤醒策略由实现决定）

读写锁

允许多读单写模式，提升读密集型操作效率。

两种锁类型：

1
2
3
4
5
6
7

pthread_rwlock_rdlock(
    pthread_rwlock_t* rwlock  /* 输入/输出：指向读写锁对象 */
);  // 加读锁：允许多个线程同时获取读锁（共享锁），但写锁被独占时会阻塞

pthread_rwlock_wrlock(
    pthread_rwlock_t* rwlock  /* 输入/输出：指向读写锁对象 */
);  // 加写锁：写锁具有排他性，同一时间只能有一个线程获取写锁，且此时不能有其他读锁或写锁

缓存一致性

Cache coherence（缓存一致性）问题是在多处理器系统中出现的一个问题。

当多个核心同时访问同一数据时，就会产生缓存一致性问题。

缓存一致性问题的本质是由于每个核心的缓存是独立的。如果一个核心修改了某些数据，并将其写入其自己的缓存行， 那么其他核心的缓存行可能仍然包含旧的数据。

如果这些其他核心继续基于这个旧数据执行操作，就会导致错误的结果。

伪共享(False Sharing)

False-sharing是一种在多核处理器系统中出现的性能问题。它发生在多个线程访问看似不同的变量，但实际上这些变量位于同一个缓存行时。

在多核处理器中，每个核心都有自己的缓存，但是，当一个线程更新了一个变量，它可能会导致整个缓存行被刷新到主内存，这会影响到访问该缓存行中其他变量的其他线程。

线程安全

一段代码是线程安全的，如果它可以由多个线程同时执行而不引起问题。

静态变量的问题

静态变量（包括全局静态变量和局部静态变量）存储于程序的静态存储区，生命周期贯穿程序运行始终，仅在程序启动时初始化一次。

在多线程环境中，所有线程共享进程的静态存储区。若多个线程同时读写静态变量且未采取同步措施，会引发数据竞争（Data Race），导致数据污染（即数据状态不可预期）。

以下函数因内部使用静态存储类变量或共享全局状态，在多线程环境下存在数据污染风险：strtok、random、localtime。

解决方案是用可重入函数（如strtok_r），它们是线程安全的。

OpenMP

基础概念

基于编译指导指令的共享内存并行编程API，通过多线程实现并行化。

使用预处理指令#pragma扩展C语言功能，非OpenMP编译器将忽略这些指令。

假设所有线程可直接访问共享内存，视系统为多核CPU集合

操作详解

编译和运行

编译：

1

gcc -g -Wall -fopenmp -o omp_hello omp_hello. c / omp_hello 4

可以使用_OPENMP宏检测编译器是否支持OpenMP

线程管理

OpenMP通过#pragma omp parallel创建并行区域。可以加入num_threads子句指定线程数量。

1
2
3
4

#pragma omp parallel [子句列表]
{
    // 并行执行的代码
}

示例：

1
2
3
4
5
6
7
8
9
10
11

#include <stdio.h>
#include <omp.h>

int main() {
    #pragma omp parallel
    {
        printf("Hello from thread %d, nthreads %d\n", 
               omp_get_thread_num(), omp_get_num_threads());
    }
    return 0;
}

加入num_threads子句之后：

1
2
3
4
5
6
7
8
9
10
11

#include <stdio.h>
#include <omp.h>

int main() {
    #pragma omp parallel num_threads(2)
    {
        printf("Thread %d of %d\n", 
               omp_get_thread_num(), omp_get_num_threads());
    }
    return 0;
}

主线程（master）与从线程（slave）构成线程组（team）。

变量作用域

1. 共享作用域（shared）：默认作用域，所有线程可访问
2. 私有作用域（private）：线程私有存储空间

在 OpenMP 中，default 子句用于显式指定并行区域（如 #pragma omp parallel 或 #pragma omp for 等指令）内变量的默认作用域规则。它的核心目的是强制程序员明确声明并行区域中使用的外部变量的作用域，避免因隐式规则导致的意外行为（如变量共享引发的竞态条件）。

1

#pragma omp parallel [子句列表] default(none|shared)

• default(none)：禁用隐式作用域，强制显式声明所有外部变量的作用域。
• default(shared)：恢复隐式规则（外部变量默认 shared，与不使用 default 子句时行为一致）。

示例：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

#include <omp.h>
#include <stdio.h>

int x = 10; // 全局变量（外部变量）

int main() {
    int y = 20; // 局部变量（外部变量，相对于并行区域）

    #pragma omp parallel default(none) shared(x) private(y) // 必须显式声明 x 和 y 的作用域
    {
        x = omp_get_thread_num(); // x 是 shared，所有线程可见
        y = omp_get_thread_num(); // y 是 private，每个线程有独立副本
        printf("Thread %d: x = %d, y = %d\n", omp_get_thread_num(), x, y);
    }

    printf("Main: x = %d, y = %d\n", x, y); // 输出最后一个线程修改后的 x，y 为初始值 20（因主线程未进入并行区域）
    return 0;
}

reduction子句

支持+、*、&等二元操作符的规约运算。

1

reduction(<operator>: <variable list>)

示例：

1
2

# pragma omp parallel num_threads (thread_count) reduction(+:global_result)
global_result += Local_trap(double a, double b, int n);

#pragma omp parallel for

后跟一个for循环语句。OpenMp自动将迭代分配给不同线程。循环内的语句需满足迭代独立性。

Barrier

1

# pragma omp barrier

可以让各个线程在某处同步它们的进度

互斥机制

critical

# pragma omp critical保护临界区，保证单线程访问。

critical 还可以命名。借助命名，可以区分不同的临界区，从而让同一时刻只有访问相同名称临界区的线程互斥执行，访问不同名称临界区的线程则能同时执行。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32

#include <stdio.h>
#include <omp.h>

int main() {
    int a = 0, b = 0;

    #pragma omp parallel num_threads(4)
    {
        // 对a进行操作的临界区
        #pragma omp critical (update_a)
        {
            a++;
            printf("Thread %d incremented a to %d\n", omp_get_thread_num(), a);
        }

        // 对b进行操作的临界区
        #pragma omp critical (update_b)
        {
            b++;
            printf("Thread %d incremented b to %d\n", omp_get_thread_num(), b);
        }

        // 未命名的临界区
        #pragma omp critical
        {
            printf("Thread %d is in unnamed critical section\n", omp_get_thread_num());
        }
    }

    printf("Final values: a = %d, b = %d\n", a, b);
    return 0;
}

atomic

1
2

#pragma omp atomic [clause]
    expression;  // 必须是单一赋值语句

• 子句（可选）：
  • read：原子读取操作（C++11 及以上）。
  • write：原子写入操作（默认）。
  • update：原子更新操作（如自增、自减）。
  • capture：原子操作并捕获原值（如 x = atomic_fetch_add(&y, 1)）。

锁

OpenMP 提供了基础的互斥锁 API，基于 C 语言实现，包含以下函数：

调度策略

在 OpenMP 中，调度策略（Scheduling） 用于控制循环迭代如何分配给并行线程。

基本语法：

1
2
3
4

#pragma omp for schedule(类型[, 块大小])
for (int i = 0; i < N; i++) {
    // 循环体
}

• 类型：指定调度算法（如 static、dynamic、guided、runtime 等）。
• 块大小（可选）：指定每个线程一次分配的连续迭代数量，仅部分调度类型需要。

主要调度策略

static

迭代被预先、平均地分配给各线程，分配模式在编译时确定。

若未指定块大小，迭代将平均分割给线程（如线程 0 处理迭代 0~4，线程 1 处理 5~9）。

若指定块大小 chunk，迭代将按块循环分配（如线程 0 处理块 0、3、6，线程 1 处理块 1、4、7）。

1
2
3
4

#pragma omp for schedule(static, 2) // 2个迭代作为一个块分配给线程
for (int i = 0; i < 10; i++) {
   
}

dynamic

迭代在运行时动态分配，线程完成当前块后请求下一个块。

块大小 chunk 控制每次分配的迭代数量（默认 chunk=1）。

1
2
3
4

#pragma omp parallel for schedule(dynamic, 5)
for (int i = 0; i < 100; i++) {
    // 线程动态请求5个迭代的块
}

guided

分配方式：块大小随时间递减，初始块大，后续块小。

• 块大小计算公式：max(chunk_size, 剩余迭代数/线程数)。

和dynamic调度的时候一样，每个线程也执行一个块，当一个线程完成一个块时，它请求另一个块。

auto

1

schedule(auto)

系统自己选择调度策略。

runtime

1

schedule(runtime)

分配方式：调度策略在运行时由环境变量OMP_SCHEDULE决定

1

export OMP_SCHEDULE="dynamic, 4"

但其实还是在上述四种中做选择。

对比

SIMD

基本概念

3种SIMD变体：

• 向量架构：通过专用向量寄存器组实现数据并行

• 多媒体SIMD指令集扩展：面向媒体处理的数据并行

• 图形处理器单元（GPU）：大规模并行处理单元的特殊形态

向量架构

在70年代-80年代，超级计算机=向量机。Cray-1是首台向量超级计算机（1976年）。

优化技术：

• 向量链接（Chaining）：流水线级联技术（Cray-1首创），实现RAW依赖指令的零延迟衔接
• 分条带处理（Strip Mining）：通过setvl指令自动处理超长向量（循环分段执行）
• 掩码寄存器：通过谓词寄存器实现条件执行（如vpne p0,v0,f0设置非零元素掩码）

多媒体SIMD指令集拓展

许多媒体应用程序运行在比32位字大小更窄的数据类型上。

• 图形：8位颜色

• 音频样本：8-16位

通过在256位加法器内部分割进位链，处理器可以对短向量执行同时操作

• $32 \times 8$位操作数
• $16 \times 16$位操作数
• $8 \times 32$位操作数

GPU

基本概念

是异构计算模型，CPU作为主机（Host），GPU作为加速设备（Device）:

• 异构应用程序由两部分组成：

  • 主机代码

• 设备代码

• 主机代码在CPU上运行，设备代码在GPU上运行。GPU用于加速这部分数据并行的执行。被称为硬件加速器

采用SIMT（单指令多线程）编程模型，统一所有GPU并行形式为CUDA线程。

NVIDIA架构特征：

• 类似向量机：支持数据级并行、散射-聚集传输、掩码寄存器
• 关键差异：用多线程隐藏内存延迟，含大量功能单元而非深度流水线

GPU上的线程非常轻量，现代NVIDIA GPU每个多处理器可以同时支持高达1536个活动线程。

CUDA

概念

CUDA（Compute Unified Device Architecture）是NVIDIA推出的通用并行计算平台与编程模型，利用GPU并行计算引擎解决复杂计算问题。

CUDA平台组成

技术栈：

• 编程语言扩展：支持C/C++、Fortran、Python等语言的扩展语法4页5页。
• 工具链：包含编译器（nvcc）、调试器（cuda-gdb）、性能分析工具（nvprof）等8页。
• 加速库：提供数学运算库（cublas、cufft）、线性代数库（cusolver）、图像处理库（npp）等预优化模块

CUDA编程模型

概念

编程模型：

• 规定了程序组件如何共享信息和协调其活动。
• 提供了特定计算架构的逻辑视图。

• 通常，它体现在编程语言或编程环境中。

  从程序员的视角来看，可以从不同级别来观察并行计算：

• 作用域级别：如何分解数据和函数以解决问题

• 逻辑级别：如何组织你的并发线程

• 硬件级别：线程如何映射到核心可能有助于提高性能

CUDA编程模型提供：

• 通过层次结构组织GPU上线程的方法
• 通过层次结构访问GPU内存的方法

CUDA编程模型使您能够通过在C编程语言上添加一组小扩展来注释代码，从而在异构计算系统上执行应用程序。

核心架构

主机（Host）与设备（Device）

• Host：CPU及其内存（Host Memory），负责控制逻辑和轻量级计算9页20页。
• Device：GPU及其内存（Device Memory），执行高并行计算任务9页10页。
• 计算流程：Host将计算密集型任务卸载到Device执行，通过内存传输协调数据

应用程序的并行部分可以隔离成一个在设备上以许多不同线程执行的函数。这样的函数被编译成设备的指令集，得到的程序称为内核。

API层次

• Driver API：底层接口，提供对GPU硬件的精细控制，编程复杂度高。
• Runtime API：高层抽象接口，基于Driver API封装，简化开发流程。

流程

一个典型的CUDA程序结构包括5个步骤：

1. 分配GPU内存。

2. 将数据从CPU内存复制到GPU内存。

3. 调用CUDA内核执行程序特定的计算。

4. 将数据从GPU内存复制回CPU内存。

5. 销毁GPU内存。

语法

声明核函数

用__global__可以把函数标记为核函数。核函数能够被主机代码调用，并且会在 GPU 设备上执行的：

• 核函数采用特殊的调用方式：kernel_name<<<gridDim, blockDim>>>(args)。
• 所有被__global__标记的函数，其返回类型都必须是void。
• 不支持可变数量参数
• 不支持静态变量
• 不支持函数指针

一个grid由多个block组成，它代表了一次核函数调用所启动的所有线程。grid可以是一维、二维或三维的。

gridDim和blockDim都具备三维结构，可借助.x、.y、.z成员来访问。这意味着，它们能在三个维度上对线程块和线程进行组织，其定义方式如下：

1 2	>dim3 gridDim(x, y, z); // 网格维度（线程块的数量） >dim3 blockDim(x, y, z); // 块维度（线程的数量）

一个block由多个线程组成，同一block内的线程可以：

• 通过共享内存快速交换数据。
• 使用__syncthreads()实现线程同步。

与C函数调用不同，所有CUDA内核启动都是异步的。在CUDA内核被调用后，控制权会立即返回给CPU。可以调用以下函数强制主机应用程序等待所有内核执行完成。cudaDeviceSynchronize(void)

内存管理

内存类型：

• 全局内存（Global Memory）：主机与设备间主要通信媒介，高延迟，所有线程可见。
• 共享内存（Shared Memory）：块（Block）内线程共享，低延迟访问。
• 寄存器（Registers）：线程私有，速度最快。

函数：

• cudaMalloc：分配设备全局内存。
• cudaMemcpy：主机与设备间数据传输（方向包括HostToDevice、DeviceToHost等）。
• cudaFree：释放设备内存

线程层次结构

网格（Grid）与线程块（Block）

• Grid：一个内核启动生成的所有线程集合，由多个Block组成。
• Block：一组协作线程，共享数据且可同步执行（如共享内存），但不同Block间无通信。

可以用blockIdx和threadIdx获得索引

• blockIdx：Block在Grid中的索引。（三维向量，可以取其的三个坐标blockIdx.x、blockIdx.y、blockIdx.z）
• threadIdx：线程在Block中的索引。（三维）

例如：

1
2
3
4
5

__global__ void sumArraysOnGPU(float *A, float *B, float *C)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x; //blockIdx.x是线程块的索引，blockDim.x是x方向上每个块线程个数。threadIdx.x是块中线程的位置
    C[i] = A[i] + B[i];
}

函数声明符

函数类型限定符用于指定：

1. 函数是在主机上还是在设备上执行
2. 该函数是否可从主机或设备调用

分析执行时间

可以通过调用gettimeofday系统调用来获取系统的挂钟时间，从而创建一个CPU计时器。

1
2
3
4
5

double cpuSecond() {
    struct timeval tp;
    gettimeofday(&tp, NULL);
    return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6); //tv_sec是秒数，tv_usec是微妙
}

性能分析工具

nvprof是CUDA的命令行性能分析工具，它可以收集应用程序的CPU和GPU活动时间线信息，包括内核执行、内存传输以及CUDA API调用。

理论性能极限：计算设备的浮点性能（FLOPS）与内存带宽比值，判断瓶颈类型。

解释：理论性能极限

在计算密集型应用中，理解系统的理论性能极限对于优化性能至关重要。理论性能极限通常由两个关键因素决定：算术性能（Arithmetic Performance）和内存带宽（Memory Bandwidth）。我们需要确定我们的应用是受限于算术性能还是内存带宽，以便针对性地进行优化。

1. 算术性能（Arithmetic Performance）

算术性能指的是处理器每秒能够执行的浮点运算次数（FLOPS）。对于GPU等并行计算设备，算术性能可以通过以下公式计算：

$$\text{Peak FLOPS} = \text{核心时钟频率} \times \text{GPU数量} \times \text{多处理器数量} \times \text{每个多处理器的FP32核心数} \times \text{每周期操作数}$$

以Tesla K10为例：

• 核心时钟频率：745 MHz
• 每板GPU数量：2
• 每个GPU的多处理器数量：8
• 每个多处理器的FP32核心数：192
• 每周期操作数：2（假设每个核心每周期执行2次浮点运算）
  计算如下： $$\text{Peak FLOPS} = 745 \text{ MHz} \times 2 \times (8 \times 192) \times 2 = 4.58 \text{ TFLOPS}$$

  2. 内存带宽（Memory Bandwidth）

  内存带宽指的是内存系统每秒能够传输的数据量（GB/s）。对于GPU，内存带宽可以通过以下公式计算： $$\text{Peak Memory Bandwidth} = \text{GPU数量} \times \text{内存位宽} \times \text{内存时钟频率} \times \text{DDR倍数} / \text{每字节数据位数}$$

以Tesla K10为例：

• GPU数量：2
• 内存位宽：256 bit
• 内存时钟频率：2500 MHz
• DDR倍数：2（双倍数据速率）
• 每字节数据位数：8 bits/byte
  计算如下： $$\text{Peak Memory Bandwidth} = 2 \times 256 \times 2500 \times 2 / 8 = 320 \text{ GB/s}$$

3. 指令与字节的比率

为了判断应用是受限于算术性能还是内存带宽，我们需要计算指令与字节的比率：

$$\text{Ratio} = \frac{\text{Peak FLOPS}}{\text{Peak Memory Bandwidth}}$$

以Tesla K10为例：

$$\text{Ratio} = \frac{4.58 \text{ TFLOPS}}{320 \text{ GB/s}} = 13.6 \text{ instructions:1 byte}$$

4. 性能瓶颈判断

• 如果每字节访问的指令数超过13.6（即13.6 instructions:1 byte），则应用受限于算术性能。这意味着计算能力未得到充分利用，可以通过优化计算逻辑或增加计算量来提升性能。
• 如果每字节访问的指令数低于13.6，则应用受限于内存带宽。这意味着内存访问成为瓶颈，可以通过优化数据访问模式、减少内存访问次数或使用更高带宽的内存来提升性能。
• 这个比率是硬件的理论极限：在理想情况下，GPU每传输1字节数据最多能执行13.6次浮点运算。
• 如果应用的实际比率超过13.6，说明计算需求远高于数据传输能力，计算能力未被充分利用。
• 如果实际比率低于13.6，说明数据传输需求远高于计算能力，内存访问成为瓶颈。

设备管理与查询

可以使用运行时API查询GPU信息

• CUDA运行时API中提供了许多函数来帮助您管理设备。
• GPU设备的属性通过cudaDeviceProp结构体返回

比如cudaGetDeviceProperties(cudaDeviceProp* prop, int device);

还可以使用命令行工具nvidia-smi来查询 GPU 状态（如内存使用率、利用率）

多GPU管理

可以CUDA_VISIBLE_DEVICES环境变量来设置可见的GPU。

CUDA执行模型

基础概念

流式处理器

GPU架构围绕可扩展的流式多处理器（SM）阵列构建而成。 通过复制这种架构构建模块（SM），实现了GPU硬件并行化。

• GPU中的每个流式多处理器(SM)都设计用于支持数百个线程的并发执行。
• 当启动内核网格时，该内核网格的线程块会被分配到可用SM上执行。
• 线程块一旦被调度到某个SM上，其线程将仅在分配的SM上并发执行。
• 多个线程块可能被同时分配到同一个SM，并根据SM资源的可用性进行调度。

SM示意图：

一个SM可以处理多个Block，但同一时刻只能执行有限的Block。当一个Block中的所有线程执行完毕后，SM会调度其他Block继续执行。

Warp

CUDA采用单指令多线程(SIMT)架构来管理和执行线程组，每组32个线程称为Warp（线程束）。

• 一个Warp中的所有线程同时执行相同的指令
• 每个线程拥有独立的指令地址计数器和寄存器状态，并针对自身数据执行当前指令
• 每个流式多处理器(SM)会将分配给它的线程块划分为32线程的Warp，然后调度这些Warp执行。如果线程数不足32的，会填充一些不活跃的线程。
• SM通过Warp调度器动态切换活跃Warp，无上下文切换开销
• 线程块内的Warp可以按任意顺序调度
• 对于一维线程块，其唯一线程ID存储在CUDA内置变量threadIdx.x中，具有连续threadIdx.x值的线程会被分组为同一Warp
• 活跃Warp分为三类：选中（执行中）、就绪（可执行）、停滞（等待资源）。当满足以下两个条件时，一个Warp即具备执行资格：
  • 有32个CUDA核心可供执行任务
  • 当前指令的所有参数均已准备就绪

在软件层面，Block是线程组织的基本单位，而Warp是硬件执行的基本单位。一个Block中的线程会被SM分组为多个Warp，每个Warp包含32个线程。因此，Block和Warp的关系是：Block是Warp的容器，Warp是Block中线程的执行单位。

SIMT允许同一Warp内线程独立执行路径（各自指令计数器、寄存器状态）。SIMD要求向量中的所有向量元素在统一的同步组中共同执行。

资源分配

每个线程独立占用一定数量的寄存器。一个warp包含32个线程，因此单个warp总寄存器消耗量=32×每个线程寄存器数。所以SM可驻留的warp数量上限=总寄存器数/(32×每个线程寄存器数)。

分支发散（Warp Divergence）

同一Warp内线程执行不同代码路径导致串行化，性能下降。虽说SIMT允许同一Warp内线程独立执行路径，但是一个Warp中的所有线程必须在同一周期执行相同的指令。

当遇到分支时，GPU会顺序执行所有可能的分支路径，而不是并行执行。例如，如果Warp中一半线程进入if分支，另一半进入else分支，GPU会先执行if分支的指令，再执行else分支的指令。在执行某个分支时，不满足条件的线程会被“遮蔽”，即暂时不更新其状态，直到其分支被执行。

优化方法：

• 调整分支条件为Warp粒度（如(tid / warpSize) % 2 == 0）。比如如果是第0个Warp则都执行A，如果是第一个Warp则都执行B。
• 使用谓词化指令替代分支。谓词通常是一个布尔值（真/假，或1/0），用来表示某个条件是否成立。谓词化指令允许一条指令根据每个线程自身的谓词值来决定是否执行该指令，或者决定该指令对线程的影响。通过谓词，所有线程在同一周期内执行“相同的”指令（只是效果不同），避免了因分支导致的执行路径切换和延迟

延迟隐藏（Latency Hiding）

指令执行的时候可能因为访问显存等因素有延迟（指令延迟），这时候SM可以切换其他Warp执行，利用多Warp并发执行掩盖指令延迟。

为了达到延迟隐藏所需要的Warp数=指令延迟*吞吐量。也就是如果想每个周期执行6个Warp， 每个Warp延迟是5，那至少要30个Warp才能做到。

占用率

占用率 = 活跃Warp数 / SM最大Warp数。

优化技术

循环展开

可以减少分支与循环维护指令，提升指令级并行（这样线程的数目就少了，不需要频繁调度了）

动态并行

此前，所有内核均需通过主机线程调用。GPU工作负载完全受CPU控制。CUDA动态并行技术允许直接在GPU上创建和同步新的GPU内核。

父网格由主机线程配置并启动，子网格则由父网格配置并启动。在父线程中设置屏障以显式与其子网格同步。

但也有限制，比如最大嵌套深度24层

CUDA内存模型

内存模型

内存层次结构的优势

两种不同类型的局部性：

• 时间局部性：如果一个数据位置被引用过，那么它在短时间内再次被引用的可能性更高
• 空间局部性：如果某个内存位置被引用，附近的位置也很可能被引用。

层级结构

分为可编程内存和不可编程内存：

• 可编程内存：程序员可显式控制数据存储位置，包括寄存器、共享内存、局部内存、常量内存、纹理内存和全局内存。
• 不可编程内存：如L1/L2缓存，由硬件自动管理，在CPU上，内存加载和存储操作均可被缓存。在GPU上，仅内存加载操作可被缓存，内存存储操作无法被缓存

核心内存类型

寄存器：

• 最快的内存空间，线程私有，生命周期与内核相同。

本地内存：

• 内核中符合寄存器使用条件但无法适配该内核所分配寄存器空间的变量，将溢出到本地内存。
• 溢出到本地内存的数值与全局内存实际位于同一物理位置

共享内存：

• 内核中使用以下属性修饰的变量会被存储在共享内存中：__shared__
• 由于共享内存位于芯片上，其带宽远高于本地或全局内存，延迟也显著更低
• 线程块内共享，但需同步（__syncthreads()）
• 与L1缓存共享64KB片内存储，可动态配置分区比例（如48KB共享+16KB L1）

全局内存：

• 最大容量、高延迟，支持32/64/128字节粒度的事务。
• 需优化对齐（32/128字节）和合并访问以减少事务次数

只读内存：

• 常量内存（Constant Memory）：静态声明，全局可见，适合同一线程束读取相同地址（如系数）。
• 纹理内存（Texture Memory）：纹理存储器是一种通过专用只读缓存访问的全局内存类型。 针对2D空间局部性进行了优化，通过专用缓存提升性能

内存管理策略

分配与传输：

• 设备内存分配：cudaMalloc分配全局内存，cudaMemset初始化，cudaFree释放
• 主机-设备传输：
  • 使用cudaMemcpy指定方向（HostToDevice/DeviceToHost等）
  • 固定内存（Pinned Memory）：
    • 是一种在操作系统中被锁定的内存区域，不会因虚拟内存管理而被交换到磁盘上
    • 避免页锁定内存的多次拷贝，提升传输带宽
    • 分配方式：cudaMallocHost，需权衡分配成本与性能收益

• 零拷贝内存：

  • 主机与设备共享同一块物理内存，避免显式传输
  • 集成GPU（共享主存）性能更优，独立GPU仅适合小数据量

• 统一内存：

  • 单指针访问，自动迁移数据，简化编程（cudaMallocManaged）
  • 与零拷贝内存对比：零拷贝数据驻留主机，统一内存透明迁移。统一内存创建了一个托管内存池，该内存池中的每次分配均可通过相同内存地址在CPU和GPU上访问，底层系统会自动在主机与设备之间迁移统一内存中的数据。

内存访问模式

最大化利用全局内存带宽是内核性能调优的基本步骤。 CUDA执行模型的显著特征之一是以线程束(Warp)为单位发出和执行指令。 内存操作同样以线程束为单位发出。

根据线程束内内存地址的分布情况，内存访问可分为不同模式：

• 当设备内存事务的起始地址是所用缓存粒度的偶数倍时，称为对齐访问。其关注单个内存事务的起始地址是否匹配缓存粒度，影响单次访问的效率。
• 当一个 warp（32 个线程）的所有线程访问连续的内存块时，称为合并访问。其关注多个线程的访问模式是否连续，影响批量访问的吞吐量。

数据布局优化

结构体数组（AoS） vs 数组结构体（SoA）：SoA布局更利于全局内存合并访问（预对齐连续数据）。

优势有：

1. 空间局部性，当访问相同属性的数据时（如所有粒子的x坐标），数据在内存中是连续的。不需要加载多个缓存行。
2. 结构体可能需要对齐到某个字节数，导致一些字节是无效的。如果是结构体的数组就有更多无效字节。

性能优化

优化技术：

• 循环展开：增加独立内存操作以隐藏延迟。
• 执行配置优化：调整线程块/网格维度，提升并行度。

带宽与性能评估

带宽指标：

• 理论带宽：硬件极限值（如显存频率×位宽）。
• 有效带宽：实测值，通过优化访问模式逼近理论值

奇偶转置排序

参考：https://blog.csdn.net/lemon_tree12138/article/details/50605563

奇偶排序（Odd-Even Sort），也称为奇偶交换排序，是一种基于比较的排序算法。它是一种稳定的排序算法，但时间复杂度较高，通常不用于实际的大规模数据排序。奇偶排序的主要思想是通过一系列的奇偶交换步骤来将数组排序。

奇偶排序的过程如下：

1. 选取所有奇数列的元素与其右侧相邻的元素进行比较，将较小的元素排序在前面；
2. 选取所有偶数列的元素与其右侧相邻的元素进行比较，将较小的元素排序在前面；
3. 重复前面两步，直到所有序列有序为止。

相对于冒泡排序

在奇偶排序中，奇数步骤和偶数步骤的比较和交换操作是独立的。在奇数步骤中，所有奇数索引的元素与其后继元素进行比较和交换；在偶数步骤中，所有偶数索引的元素与其后继元素进行比较和交换。这些操作可以同时进行，因为它们不会相互影响。

而在冒泡排序中，每一轮排序都需要从数组的开始到结束依次进行比较和交换，每个元素的位置都可能依赖于前一个元素的位置，这使得并行化变得更加复杂。