并行编程实战——CUDA编程的统一内存的预取

一、内存优化与数据预取

有过编程开发经验的大多都明白，内存优化中，数据预取的重要性。当然，数据预取也是有一定规则的，否则就可能导致预取的数据不合适而重新准备数据，反而会使操作变得更复杂。
在CUDA中，也有这种数据预取的优化方式。目的同样，也是为了将需要的统一内存中的数据提前在CPU和GPU中进行加载。CUDA中提供了API接口：

cudaError_t cudaMemPrefetchAsync(const void* devPtr, size_t count, int dstDevice, cudaStream_t stream = 0);
devPtr: 需要预取的数据指针
count: 需要预取数据的字节数
dstDevice: 预取的目标设备（CPU或GPU）
- cudaCpuDeviceId: CPU设备ID
- GPU设备ID (可使用cudaGetDevice(&deviceId)得到)
stream: CUDA流（可选）

数据预取的主要优势在于可以提高数据访问的速度（减少内存延迟），提高性能，减少因数据加载问题导致的内存页面问题并提供稳定的的性能表现。当然，事物都是有两面性的，过度预取数据以及不当时机的预取数据，都可能导致性能的下降；同时，大量的预取也可能导致内存的压力。

二、CUDA的应用

在CUDA的编程中，数据预取在科学计算等数据流协作要求紧密的情况下，非常重要。它的应用场景主要包括：

1. 在数据计算和数据传递前进行
   比如在GPU启动计算任务前将数据从统一内存中预取；或者在计算任务完成后，将计算数据预取到CPU中准备使用
2. 批次预取
   在针对一些大数据计算任务时，可以根据实际情况，分批预取，形成类似流水线的机制
3. 多流预取
   这有些类似于多线程中的不同线程分块加载不同的内存中的数据。而在CUDA中则是不同的流加载不同的数据块

在CUDA编程中，和预取操作相关一些API还有：

1. cudaMemAdvise函数
   用于对内存页面设置访问建议，从而优化GPU和主机内存之间的数据迁移策略。比如开发者可以通过其设定内存数据保存的位置（GPU或主机内存），从而减少内存数据的迁移开销
2. cudaMemRangeGetAttribute函数
   查询指定内存范围的属性，用于获取内存使用模式、访问频率等信息。从而可以优化内存和获取相应的内存调度相关信息

当然其实还有一些可配合使用的相关API，这里就不再一一分析说明了。有兴趣可以查阅CUDA官网的文档说明。

三、例程

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <algorithm>

#define N (1 << 20)  // 1M
#define BLOCK_NUM 256
#define ITLIMIT 100

__global__ void vecProcess(float* data, float factor, int n) {
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	if (i < n) {
		data[i] = data[i] * factor + sinf(data[i]) * cosf(data[i]);
	}
}

void initMemData(float* data, int n) {
	for (int i = 0; i < n; i++) {
		data[i] = static_cast<float>(rand()) / RAND_MAX * 100.0f;
	}
}


int main() {
	printf("-------------- CUDA Unified Memory prefetch test -----------\n");

	const size_t size = N * sizeof(float);
	int bkPerGrid = (N + BLOCK_NUM - 1) / BLOCK_NUM;
	int devID = 0;

	// get device info
	cudaGetDevice(&devID);
	cudaDeviceProp prop;
	cudaGetDeviceProperties(&prop, devID);
	printf("used device: %s\n", prop.name);
	printf("mem data size: %zu MB\n", size / (1024 * 1024));

	// malloc unified mem
	float *data= NULL;
	cudaMallocManaged(&data, size);

	// prefetch GPU
	printf("\n---------------GPU prefetch--------------\n");
	initMemData(data, N);


	// prefetch GPU
	cudaMemPrefetchAsync(data, size, devID);
	cudaDeviceSynchronize();

	for (int it = 0; it < ITLIMIT; it++) {
		vecProcess << <bkPerGrid, BLOCK_NUM >> >(data, 1.1f, N);
		cudaDeviceSynchronize();
	}

	// batch prefetch
	printf("\n------------------batch prefetch---------------------\n");

	const int chunkSize = N / 4;  
	const size_t chunkCount = chunkSize * sizeof(float);

	initMemData(data, N);

	for (int it = 0; it < ITLIMIT; it++) {
		for (int chunk = 0; chunk < 4; chunk++) {
			float* pChunk = data + chunk * chunkSize;

			cudaMemPrefetchAsync(pChunk, chunkCount, devID);

			vecProcess << <bkPerGrid / 4, BLOCK_NUM >> >(
				pChunk, 1.1f, chunkSize);

			if (chunk < 3) {
				float* nextpChunk = data + (chunk + 1) * chunkSize;
				cudaMemPrefetchAsync(nextpChunk, chunkCount, devID);
			}
		}
		cudaDeviceSynchronize();
	}

	// stream prefetch
	printf("\n----------stream prefetch-----------------------------\n");

	cudaStream_t stream1, stream2;
	cudaStreamCreate(&stream1);
	cudaStreamCreate(&stream2);

	initMemData(data, N);

	for (int it = 0; it < ITLIMIT; it++) {
		for (int chunk = 0; chunk < 4; chunk += 2) {
			float* c1 = data + chunk * chunkSize;
			cudaMemPrefetchAsync(c1, chunkCount, devID, stream1);
			vecProcess << <bkPerGrid / 4, BLOCK_NUM, 0, stream1 >> >(
				c1, 1.1f, chunkSize);

			if (chunk + 1 < 4) {
				float* c2 = data + (chunk + 1) * chunkSize;
				cudaMemPrefetchAsync(c2, chunkCount, devID, stream2);
				vecProcess << <bkPerGrid / 4, BLOCK_NUM, 0, stream2 >> >(
					c2, 1.1f, chunkSize);
			}
		}
	}

	cudaStreamSynchronize(stream1);
	cudaStreamSynchronize(stream2);

	// pretch CPU
	printf("\-------------------cpu prefetch----------------------------\n");

	cudaMemPrefetchAsync(data, size, devID);
	cudaDeviceSynchronize();

	cudaMemPrefetchAsync(data, size, cudaCpuDeviceId);
	cudaDeviceSynchronize();

	for (int i = 0; i < N; i++) {
		data[i] = data[i] * 1.5f;
	}


	// clean
	cudaFree(data);
	cudaStreamDestroy(stream1);
	cudaStreamDestroy(stream2);

	printf("\n---------------end---------------------------\n");

	return 0;
}

大家可以根据代码对三种预取方式进行分析，其实通过相关时间查看，在本机的测试上三者的整体效率没有较大差别。原因可能在于本机测试的整体场景不是太匹配。

四、总结

通过学习CUDA中的技术，大家是不是发现虽然GPU与CPU的应用任务领域不同，但它们是分工协作解决实际的问题的关系而不是某一方替代另外一方的情况。虽然从抽象问题和解决问题的具体形式不同，甚至可以说区别非常大，但仍然保持在计算机开发的思想范围之内。