第8章 多CPU与GPU解决方案 摘录

8.1  简介 

在现代计算机系统中通常由多个设备：CPU和GPU。在CPU方面将讨论卡槽和核。卡槽是指主板用来放置CPU的物理插槽。一个CPU可能包含一个或多个核。每个核实际上是一个独立的实体。许多CPU和GPU卡槽都位于单个节点或计算机系统。

了解核、卡槽已经节点的物理布局，会使得我们更有效地调度和分配任务。

8.2 局部性

局部性原理在GPU和CPU上是相当不错的。接近设备的内存(GPU上共享内存或CPU上的高速缓存)会更快地访问。一个卡槽内(比如核之间)的通信远远快于和不同卡槽上的另一核的通信。访问另一个节点核的通信方式至少比该节点内的访问慢一个数量级。

这种卡槽感知软件可以按照硬件布局分隔数据，确保一个核在一致的数据集上工作，并且需要多个核合作。

8.3 多GPU系统

任何多处理器系统的主要问题之一是内存的一致性。CPU和GPU将内存分配在单个设备。对于GPU，这是每个GPU卡上的全局内存。对于CPU，这是主板上的系统内存。

当需要两个核相互合作的时候，就会遇见问题。为了加快内存的访问速度，CPU大量使用高速缓存。一个参数的值被更新时(++x)，x真的被写入内存吗？假设两个核都需要更新x，则两个核写入存储参数x的内存位置内容未必是一致的。

这就是缓存一致性问题，限制了可以在当个节点上实际合作的核的最大数量。实际上当核1写入x时，它会通知所有其他的核x的值已经改变，然后缓慢地写回主存中，而不是快速写到高速缓存。

在一个简单的一致性模型中，其他核标记条目表明x在它们的缓存中无效。下一次访问x时，会从缓慢的主存中读取x值。随后的核写入x时，重复该过程，下一个核访问参数x必须再次从主存中读取并且重新写回。实际上，参数x没有被缓存，这对于CPU意味着巨大的性能损失。

在一个更复杂的一致性模型中，用更新请求替换无效请求，而不是替换无效的参数x。因此，每一个写操作被分配到N个高速缓存中。随着N的增长，同步这些缓存所用的时间变得不切实际。这往往限制了可以放入对称多处理器(SMP)系统的实际数目。

详情参考写直达与写回

现在请记住，高速缓存都应该高速运行。在一个单独的卡槽，这是不难的。然而，只要你用到多个卡槽，高时钟速率难以维持，因此一切都开始放缓。用的卡槽越多，保持一切同步就变得更加困难。

下一个重要的问题是内存访问时间。为了使在这样的机器中编程更容易，往往内存在逻辑上被安排成一个巨大的线性地址空间。然而，只要卡槽1的核试图访问卡槽2的一个内存地址，就必须要卡槽2(作为唯一可以实际访问内存地址的卡槽)提供服务。这就是所谓的非一致内存访问(Nonuniform Memory Access, NUMA)。尽管在概念上，它使程序员工作更轻松，但是在实践中，你还是需要考虑内存地址的问题，否则你写的程序执行速度会非常缓慢。

8.4 多GPU系统

添加一个额外的GPU卡，你通常会看到性能水平增加了一倍，当前执行时间减半。你很少能这么容易获得这样的加速。

8.5 多GPU算法

多GPU会带来一个明显问题----系统中的数据移动。

在这种情况下，我们重组算法使它可以被分解成独立的数据块。然而这并不总是简易的，许多类型的问题至少会从其他GPU获得少量数据。当你需要另一个GPU数据，你必须显式地共享这些数据，并明确GPU之间的数据访问顺序。CUDA4,0之后，提供了GPU对等通信模型，或者也可以在CPU层使用CPU级原语以进行合作。前者不适合所有的操作系统。

8.6 按需选用GPU

CUDA运行时会自动选择基于最高级别计算设备的二进制代码执行。

某些功能，比如原子操作，寄存器数目，每个线程块的最大值在不同架构可能会有一些区别。

8.7 单节点系统

在CUDA4.0之前，单节点系统是唯一支持多GPU模型的系统。如上图所示，一个基于CPU的任务将与单GPU上下文相关联。在该上下文中的是一个进程或线程。后台的CUDA运行时将CPU进程/线程ID绑定到GPU上下文。因此，其随后所有CUDA调用将在绑定到该上下文的设备中分配内存。

上述方法有许多缺点，但也有优点。从编程的角度看，在主机端的进程/线程模型被操作系统切分了。进程作为一个独立CPU调度单位运行，并且有其自己的数据空间的程序。为了节省内存，通常在同一进程的多个实例共享代码空间，并且操作系统内每个进程保留了一组寄存器(或者上下文)。

相比而言，线程是一种更轻便的CPU调度元素。多个所属线程共享父进程的代码及数据空间。然而，与进程不同的是，每个线程需要操作系统保持一个状态(指令指针，堆栈指针和寄存器等)。

线程可以与同一进程内其他线程沟通和合作。进程可以通过进程间通信与其他进程沟通和合作。这样进程之间的通信可以是在一个CPU核内，一个CPU插槽内、一个CPU节点内、一个机架内、一个计算机系统内甚至是在不同计算机系统之间。

执行内核时，CPU线程与GPU线程是相似的，唯一的不同是它不像GPU一样划分为组或线程束执行。GPU线程通过共享内存和显式同步以确保每一个线程对该内存进行了读/写。共享内存对于SM来说是本地的，这意味着线程只能(理论上)与同一个SM上的其他线程进行通信。因为线程块是SM的调度单元，线程间通信实际上是限制在一个线程块内的。

从线程调度的角度而言，SM的行为很像CPU核。

请注意，在一个主机节点内通信时，它们都利用了共享内存传输。然而，支持线程的模型(pthread，ZeroMQ)执行基于线程间的通信比那些基于进程(MPI)更快。

8.8 流

流是GPU上虚拟工作队列。它们用于异步操作，也就是你希望GPU独立于CPU进行操作。一些操作会隐式地引发一个同步点，例如主机与设备间默认的内存传输。

通过创建一个流，你可以将任务和事件压入流，然后按照他们会被压入流的顺序执行它们。流和事件与它们被创建时所在的GPU上下文关联。

一个多流测试案例：

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdio.h>
#include <cmath>

typedef unsigned long int u32;

// 简单使用0~num_element对数组进行填充。
void fill_array(u32* data, const u32 num_elements)
{
	for (u32 i = 0; i < num_elements; ++i)
		data[i] = i;
}

// 单纯检查GPU的结果是否是我们所期待的
void check_array(char* device_prefix, u32* data, const u32 num_elements)
{
	bool error_found = false;

	for (u32 i = 0; i < num_elements; ++i)
	{
		if (data[i] != (i * 2))
		{
			printf("%s Error: %u %u", device_prefix, i, data[i]);
		}

		error_found = true;
	}

	if (error_found == false)
		printf("%s Array check passed", device_prefix);
}

// 核函数功能为让每一个数据元素乘以2
__global__ void gpu_test_kernel(u32* data)
{
	const int tid = blockIdx.x * blockDim.x + threadIdx.x;
	data[tid] *= 2;
	printf("OK!");
}

// 定义最大利用的设备数，已经数组元素个数
const int Max_Num_Devices = 2;
const u32 Num_Elem = 1024 * 1024 * 1024;

__host__ void gpu_kernel(void)
{
	// 定义最大支持的GPU数
	int num_devices = 0;

	// 创建流数组
	cudaStream_t streams[Max_Num_Devices];

	// 为每个GPU定义一个字符串输出
	char device_prefix[Max_Num_Devices][300];

	// 我们需要添加一些一些计时代码，以查看每个内核实际花费时间。
	// 我们需要将事件添加到工作队列中。现在这些事件是很特殊的。
	// 因为我们可以查询事件而无需考虑当前选择的GPU。
	// 要做到这一点，我们需要声明一个启动和停止事件。
	cudaEvent_t kernel_start_event[Max_Num_Devices];
	cudaEvent_t memcpy_to_start_event[Max_Num_Devices];
	cudaEvent_t memcpy_from_start_event[Max_Num_Devices];
	cudaEvent_t memcpy_from_stop_event[Max_Num_Devices];

	// 创建事件

	// 无需动态收集共享内存
	const int shared_memoery_usage = 0;

	// 每个GPU上的字节数大小
	const size_t single_gpu_chunk_size = Num_Elem * sizeof(u32);

	// 定义执行环境
	const int num_threads = 256;

	const int num_blocks = ceil(Num_Elem / num_threads);

	// 定义每个GPU上的数列
	u32* gpu_data[Max_Num_Devices];

	// 定义每个CPU的源数列GPU传输的目的段数列
	u32* cpu_src_data[Max_Num_Devices];
	u32* cpu_dest_data[Max_Num_Devices];

	cudaGetDeviceCount(&num_devices);
	printf("是否成功找到的设备数目： %s\n", cudaGetErrorString(cudaGetLastError()));
	
	if (num_devices > Max_Num_Devices)
		num_devices = Max_Num_Devices;

	// 分配GPU
	for (int device_idx = 0; device_idx < num_devices; ++device_idx)
	{
		cudaSetDevice(device_idx);
		cudaDeviceProp device_prop;
		cudaGetDeviceProperties(&device_prop, device_idx);
		sprintf(&device_prefix[device_idx][0], "\n ID: %d %s: ", device_idx, device_prop.name);

		// 在相应的GPU上创建流
		cudaStreamCreate(&streams[device_idx]);
		printf("是否成功创建流： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 请注意主机上的锁业内存使用，使用cudaMallocHost分配内存，而不要使用常规的C语言的malloc函数。
		// 因为锁业内存是不能被交换到磁盘的。
		// 由于内存复制操作通过在PEI-E总线上的直接内存访问DMA实现，因此CPU端的内存一定总是存在与物理内存。
		// 用malloc分配的内存可以交换到磁盘上，如果DMA试图对其进行操作，那么就会导致失败。
		// 另外，cudaMallocHost分配，则cudaFreeHost释放。
		// 分配GPU侧内存
		cudaMalloc((void**)&gpu_data[device_idx], single_gpu_chunk_size);
		printf("是否成功创建gpu_data： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 分配CPU侧内存
		cudaMallocHost((void**)&cpu_src_data[device_idx], single_gpu_chunk_size);
		printf("是否成功创建锁业内存cpu_src_data： %s\n", cudaGetErrorString(cudaGetLastError()));

		cudaMallocHost((void**)&cpu_dest_data[device_idx], single_gpu_chunk_size);
		printf("是否成功创建锁业内存cpu_dest_data： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 初始化CPU侧数据
		fill_array(cpu_src_data[device_idx], Num_Elem);

		// 我们将主机CPU的内存复制操作到GPU全局内存的操作和紧接着内核调用
		// 还有数据传回CPU的操作放入到流中。它们将按照数据执行，
		// 这样内核只有在前面的内存复制操作完成才会执行。
		
		// 创建传输数据传输H2D事件
		cudaEventCreate(&memcpy_to_start_event[device_idx]);
		cudaEventCreate(&kernel_start_event[device_idx]);

		// 开始计时传输数据传输H2D事件开始
		cudaEventRecord(memcpy_to_start_event[device_idx], streams[device_idx]);

		// 初始化GPU侧数据--异步Host2Device
		cudaMemcpyAsync(gpu_data[device_idx], 
						cpu_src_data[device_idx], 
						single_gpu_chunk_size, 
						cudaMemcpyHostToDevice, 
						streams[device_idx]);
		printf("是否成功传输cpu_src_data： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 开始计时内核函数gpu_test_kernel开始
		cudaEventCreate(&kernel_start_event[device_idx]);
		cudaEventRecord(kernel_start_event[device_idx], streams[device_idx]);

		// 开启内核
		gpu_test_kernel << <num_blocks, num_threads, shared_memoery_usage, streams[device_idx] >> > (gpu_data[device_idx]);
		printf("是否成功开启内核gpu_test_kernel： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 开始计时传输数据传输D2H事件开始
		cudaEventCreate(&memcpy_from_start_event[device_idx]);
		cudaEventRecord(memcpy_from_start_event[device_idx], streams[device_idx]);

		// 将GPU侧数据写回CPU--异步Device2Host
		cudaMemcpyAsync(cpu_dest_data[device_idx],
						gpu_data[device_idx],
						single_gpu_chunk_size,
						cudaMemcpyHostToDevice,
						streams[device_idx]);
		printf("是否成功异步传输D2H： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 计时传输数据传输D2H事件结束
		cudaEventCreate(&memcpy_from_stop_event[device_idx]);
		cudaEventRecord(memcpy_from_stop_event[device_idx], streams[device_idx]);

		// 统计数据传输事件传输H2D事件花费时间
		float time_copy_to_ms = 0;
		cudaEventElapsedTime(&time_copy_to_ms, memcpy_to_start_event[device_idx], kernel_start_event[device_idx]);

		float time_kernel_ms = 0;
		cudaEventElapsedTime(&time_copy_to_ms, kernel_start_event[device_idx], memcpy_from_start_event[device_idx]);

		float time_from_ms = 0;
		cudaEventElapsedTime(&time_copy_to_ms, memcpy_from_start_event[device_idx], memcpy_from_stop_event[device_idx]);

		// cudaEventSynchronize(); 此处是否需要同步事件？
		
		// 删除创建的计时Event
		cudaEventDestroy(memcpy_from_stop_event[device_idx]);
		cudaEventDestroy(memcpy_from_start_event[device_idx]);
		cudaEventDestroy(kernel_start_event[device_idx]);
		cudaEventDestroy(memcpy_to_start_event[device_idx]);

		// 输出统计的三个事件花费的时间
		printf("H2D: %.6f\n", time_copy_to_ms);
		printf("Kernel: %.6f\n", time_kernel_ms);
		printf("D2H: %.6f\n", time_from_ms);

	}

	for (int device_idx = 0; device_idx < num_devices; ++device_idx)
	{
		// 选择相应的设备
		cudaSetDevice(device_idx);
		printf("是否成功选择设备： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 等待每个流内的操作都完成
		cudaStreamSynchronize(streams[device_idx]);
		printf("是否成功全部完成相应流中所有操作： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 目前相应编号的GPU data和流都已经使用完，可以删除了。
		cudaStreamDestroy(streams[device_idx]);
		printf("是否成功删除相应流： %s\n", cudaGetErrorString(cudaGetLastError()));

		cudaFree(gpu_data[device_idx]);
		printf("是否成功释放GPU侧数据： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 数据已经到达CPU侧，检测数据
		check_array(device_prefix[device_idx], cpu_dest_data[device_idx], Num_Elem);

		// 检测完后，删除CPU侧数据
		cudaFreeHost(cpu_src_data[device_idx]);
		cudaFreeHost(cpu_dest_data[device_idx]);
		printf("是否成功释放CPU侧数据： %s\n", cudaGetErrorString(cudaGetLastError()));

		// 重设设备
		cudaDeviceReset();
	}
}


int main(void)
{
	gpu_kernel();
	return 0;
}

目前PCI-E链路的速度甚至比最慢的设备内存还要慢得多，所以我们的传输速度实际上是被PCI-E总线的速度限制了。

有些设备处理能力是不一样的。比如rtx 3090与rtx 2080，在3090完成时候，就可以去处理其他的任务。

我们可以通过查询结束事件来解决该问题，而不是简单地等待事件结束。也就是说，我们可以检查内核是否完成，如果没有就转移下一个设备，然后在回到较慢设备。这可以通过cudaEventQuery函数来实现。

该函数接受并判断某一特定的事件。如果该事件已经发生就返回cudaSuccess，如果事件尚未发生，则返回cudaErrorNotReady。由于cudaErrorNotReady的状态并不是一个真正的错误状态，它只是状态信息。我们还需要通过下面的调用说明CUDA如何跟踪处于等待的GPU任务。

cudaSetDeviceFlags(cudaDeviceScheduleYield);

该调用是在所有其他CUDA调用之前完成，并且告诉驱动程序，在任何情况下等待操作的时候，该CPU线程应该给其他CPU让步。这意味着额外的延迟，因为驱动程序要求CPU线程在CPU工作队列中等待它的调度机会，但允许其他CPU任务继续进行。另一种方法是，驱动程序只围绕该CPU线程进行检测(轮训设备cudaDeviceScheduleSpin， polling)，当有设备就绪时，这肯定不是我们所期待的。

为了避免对事件队列进行轮询从而导致程序在涉及其他CPU任务时表现不佳，程序需要把自己休眠，过一段时间唤醒并且再次检查事件队列。

我们将重新排列数据的处理顺序，以移除cudaStreamSynchronize调用，并且把此代码放置在一个函数中。把代码放在主循环之外。这种特定操作是很重要的，因为在循环内这样做，会导致串行的驱动程序调用。（文中此处代码晦涩，暂时还未理解）

重新排布数据的处理顺序是为了让处理能力的强GPU在处理好当前任务后，立马去处理其他任务。不必等待其他处理能力的不强的GPU。

在大多数情况下，在最快与最慢的两个GPU都执行完任务的期间，CPU线程都处于闲置状态，它本来可以做一些有用的事情，比如当GPU完成任务时向它们发布更多的任务。

8.9 多节点系统

单台计算机形成网络上的一个节点，在连接许多机器后，就得到了一个机器集群。通常，这种集群将由一套机架式节点组成，这个机架本身可以与一个或多个机架相互连接。

目前大多数研究和商业组织通常购买16~48端口的千兆交换机上节点上。

如果CPU负载很大，那么它很有可能限制吞吐量。为了克服这个问题，我们需要为每个CPU核分配更少的GPU，变成应用程序需要的1:2或者1:1的比例，最简单、最可扩展的方法是为节点上的每一组CPU/GPU分配一个进程。一旦我们使用了这个模型，它将允许更大的扩展。

如果问题被分解为翻倍个节点上的块，那么其性能理论是节点的翻倍。事实上，由于通信开销的存在，以及随着节点数的增加，网络通信对问题的影响也会增加。因此，你通常会发现每个节点具有更多数量GPU的网络，其性能会超过同样数量的GPU但是分布在更多节点上的网络。本地节点上的资源(磁盘、内存、CPU)会对给定问题的最佳拓扑产生重大影响。

可使用的通信机制有ZeroMQ，MPI，套接字等等。

我们从调度和工作负载角度而言，一个主要的问题在于网络加载和通信开销。我们在网络上发送数据的数量会对系统的性能产生巨大的影响。收到任意入站数据，在CPU上改变它并在CPU上再次将它发送出去的总体时间必须比GPU内核执行所用的时间要少。否则，应用程序可能是CPU密集的或网络密集的。

在服务器发送客户端一系列的数据而言。你需要尽可能的做法是简单地将数据本身发送给客户端。要利用节点上的本地资源，该资源可能是CPU、本机内存、本地存储空间或任何可能的资源。另外，输出的数据被全部运送回服务器。这个问题可能是：输出数据不是一个巨大的数据块，而仅仅是单个值，比如来自规约操作。通常情况下输入空间是很大的。然而，如果输入空间可以划分为N个本地磁盘。这时网络流量是非常小的，通过使用多个GPU节点你可以看到很大的扩展性。

8.10 本章小结

单节点内多GPU、多节点多GPU案例。多节点的ZeroMQ使用。

使用流来实现双缓冲系统，这意味着CPU在准备下一个数据块和处理前一个数据块时，GPU总是忙碌的。我们将流的使用从双流延伸到多流，以允许我们对同一个节点中不同速度的GPU设备的工作进行平衡。

为了进一步提升GPU的吞吐量，你需要使用多个节点并估计将要在网络上的通信的数据量。