CUDA by example Chapter11 多GPU系统上的CUDA C

零拷贝主机内存

零拷贝主机内存是另一种固定内存，它也不能被CPU从物理内存交换出去或者重新定位。同时，这种内存除了可以用于主机与GPU之间的内存复制外，还可以在CUDA C核函数中直接访问这种类型的主机内存。由于这种内存不需要复制到GPU，因此也称为零拷贝内存。

通过零拷贝内存实现点积运算

1. 测试方法
我们用test性能测试的方式，编写两个函数，其中一个函数是对标准主机内存的测试，另一个函数将在GPU上执行规约运算，并使用零拷贝内存作为输入缓冲区和输出缓冲区，并比较它们的性能。
2. 分配零拷贝内存的方法
使用cudaHostAlloc()函数分配零拷贝内存，其中flags参数指定含义如下：
cudaHostAllocMapped这个标志告诉运行时将从GPU访问这块内存，即分配零拷贝内存。
cudaHostAllocWriteCombined这个标志表示，运行时应该将内存分配为"合并式写入"内存。这个标志并不会改变应用程序的功能，但却可以显著地提升GPU读取内存时的性能。然而，当CPU也要读取这块内存时，“合并式写入”会显得很低效，因此在决定是否使用这个标志之前，必须首先考虑应用程序的可能访问模式。

根据以上的说明，由于CPU不需要读取两个输入缓冲区的内容，所以它们应该使用cudaHostAllocWriteCombined标志来提升GPU读取内存时的性能；而CPU需要读取输出缓冲区的内容，因此它不应该使用cudaHostAllocWriteCombined标志，如下所示：

	//在CPU上分配内存
	cudaHostAlloc((void **)&a, size*sizeof(*a), cudaHostAllocWriteCombined|cudaHostAllocMapped);
	cudaHostAlloc((void **)&b, size*sizeof(*b), cudaHostAllocWriteCombined | cudaHostAllocMapped);
	cudaHostAlloc((void **)&partial_c, blocksPerGrid*sizeof(*partial_c), cudaHostAllocMapped);

3. 获得零拷贝内存在GPU上的有效地址
GPU的虚拟内存空间与CPU是独立的，因此在GPU上访问零拷贝内存与在CPU上访问它具有不同的地址。通过调用cudaHostAlloc()将返回这块内存在CPU上的指针，因此需要调用cudaHostGetDevicePointer()来获得这块内存在GPU上的有效指针。这些指针将被传递给核函数，并在随后由GPU对这块内存执行读取和写入操作。

	//获得主机上的零拷贝内存在GPU上的地址映射
	cudaHostGetDevicePointer(&dev_a, a, 0);
	cudaHostGetDevicePointer(&dev_b, b, 0);
	cudaHostGetDevicePointer(&dev_partial_c, partial_c, 0);

4. 需要对CPU和GPU进行同步
   因为CPU和GPU都可以访问零拷贝内存，如果在核函数中会修改零拷贝内存的内容，那么在核函数执行期间CPU不应该访问或者修改零拷贝内存的内容，因此需要对CPU和GPU进行同步。同步完成后，此时可以确信GPU已经完成了相应的工作，CPU可以安全地对零拷贝内存进行操作。

//如果核函数中会修改零拷贝内存的内容，那么在核函数的执行期间
//零拷贝内存的内容是未定义的,因此需要对CPU和GPU进行同步
cudaThreadSynchronize();

5. 将运行时置入能分配零拷贝内存的状态
必须要先将运行时置入能分配零拷贝内存的状态之后才能将主机内存映射到设备上。
通过调用cudaSetDeviceFlags()实现这个操作，并且传递标志值cudaDeviceMapHost来表示我们希望设备映射主机内存。

cudaSetDeviceFlags(cudaDeviceMapHost);

实际上，在运行完整代码时，一开始我忘记添加上述语句，发现实现的功能和性能与添加上述语句时相似的。个人猜想这可能是NVIDIA新CUDA版本做的优化。cuda by example这本书放在现在(2019年)来看确实有些古老了，但是其中的基本思想还是有借鉴之处的。
6.完整代码

// 代码11.2零拷贝主机内存
//时间：2019.08.03
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

#define imin(a,b) (a<b?a:b)


const int N = 33 * 1024 * 1024;
const int threadsPerBlock = 256;
const int blocksPerGrid = imin(32,(N+threadsPerBlock-1)/threadsPerBlock);




__global__ void dot(int size, float *a, float *b, float *partial_c)
{
	__shared__ float cache[threadsPerBlock];
	int tid = threadIdx.x + blockIdx.x*blockDim.x;
	int cacheIndex = threadIdx.x;

	float temp = 0;
	while (tid < size)
	{
		temp += a[tid] * b[tid];
		tid += gridDim.x*blockDim.x;
	}

	//设置cache中的值
	cache[cacheIndex] = temp;

	//同步这个线程块中的线程
	__syncthreads();

	//规约运算
	int i = blockDim.x / 2;
	while (i != 0)
	{
		if (cacheIndex < i)
		{
			cache[cacheIndex] += cache[cacheIndex + i];
		}
		__syncthreads();
		i /= 2;
	}

	if (cacheIndex == 0)
	{
		partial_c[blockIdx.x] = cache[0];
	}

}


float cuda_malloc_test(int size)
{
	cudaEvent_t start, stop;
	float *a, *b, c, *partial_c;
	float *dev_a, *dev_b, *dev_partial_c;
	float elapsedTime;

	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	//在CPU上分配主机内存
	a = (float *)malloc(size*sizeof(*a));
	b = (float *)malloc(size*sizeof(*b));
	partial_c = (float *)malloc(blocksPerGrid*sizeof(*partial_c));


	//在GPU上分配内存
	cudaMalloc((void **)&dev_a, size*sizeof(*dev_a));
	cudaMalloc((void **)&dev_b, size*sizeof(*dev_b));
	cudaMalloc((void **)&dev_partial_c, blocksPerGrid*sizeof(*dev_partial_c));

	//用数据填充主机内存
	for (int i = 0; i < size; i++)
	{
		a[i] = i;
		b[i] = i * 2;
	}

	//开始记录事件
	cudaEventRecord(start, 0);

	//将数组a和b复制到GPU
	cudaMemcpy(dev_a, a, size*sizeof(*a), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, size*sizeof(*b), cudaMemcpyHostToDevice);

	//执行核函数
	dot << <blocksPerGrid, threadsPerBlock >> >(size, dev_a, dev_b, dev_partial_c);

	//将GPU运算得到的结果复制到CPU上
	cudaMemcpy(partial_c, dev_partial_c, blocksPerGrid*sizeof(*dev_partial_c), cudaMemcpyDeviceToHost);

	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);

	//在CPU上把partial_c上的内容加和，结果放到c上
	c = 0;
	for (int i = 0; i < blocksPerGrid; i++)
	{
		c += partial_c[i];
	}


	//释放GPU上的内存
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_partial_c);

	//释放CPU上的内存
	free(a);
	free(b);
	free(partial_c);

	//释放事件
	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	printf("Valude calculated: %f\n", c);
	return elapsedTime;

}



float cuda_host_alloc_test(int size)
{
	cudaEvent_t start, stop;
	float *a, *b, c, *partial_c;
	float *dev_a, *dev_b, *dev_partial_c;
	float elapsedTime;

	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	//在CPU上分配内存
	cudaHostAlloc((void **)&a, size*sizeof(*a), cudaHostAllocWriteCombined|cudaHostAllocMapped);
	cudaHostAlloc((void **)&b, size*sizeof(*b), cudaHostAllocWriteCombined | cudaHostAllocMapped);
	cudaHostAlloc((void **)&partial_c, blocksPerGrid*sizeof(*partial_c), cudaHostAllocMapped);

	//用数据填充主机内存
	for (int i = 0; i < size; i++)
	{
		a[i] = i;
		b[i] = i * 2;
	}


	//获得主机上的零拷贝内存在GPU上的地址映射
	cudaHostGetDevicePointer(&dev_a, a, 0);
	cudaHostGetDevicePointer(&dev_b, b, 0);
	cudaHostGetDevicePointer(&dev_partial_c, partial_c, 0);


	//开始记录事件
	cudaEventRecord(start, 0);
	dot << <blocksPerGrid, threadsPerBlock >> >(size, dev_a, dev_b, dev_partial_c);
	cudaThreadSynchronize();//如果核函数中会修改零拷贝内存的内容，那么在核函数的执行期间，零拷贝内存的内容是未定义的
	//由于使用的是零拷贝内存，所以不需要将GPU内存拷贝到主机上


	//结束事件记录
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);

	//在CPU下对计算结果进行加和
	c = 0;
	for (int i = 0; i < blocksPerGrid; i++)
	{
		c += partial_c[i];
	}


	//释放一些内存
	cudaFreeHost(dev_a);
	cudaFreeHost(dev_b);
	cudaFreeHost(dev_partial_c);

	//释放事件
	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	printf("Value calculated:%f\n", c);
	return elapsedTime;
}


int main()
{
	cudaDeviceProp prop;
	int whichDevice;
	cudaGetDevice(&whichDevice);
	cudaGetDeviceProperties(&prop, whichDevice);
	if (prop.canMapHostMemory != 1)
	{
		printf("Device cannot map memory.\n");
		return 0;
	}

	cudaSetDeviceFlags(cudaDeviceMapHost);
	
	float elapsedTime = cuda_malloc_test(N);
	printf("Time using cudamalloc：%3.1f ms\n", elapsedTime);

	elapsedTime = cuda_host_alloc_test(N);
	printf("Time using cudaHostAlloc:%3.1f ms\n", elapsedTime);


	

	system("pause");
	return 0;

}

零拷贝内存的性能

1. 集成GPU和独立GPU的查看
使用cudaGetDeviceProperties()返回的结构来判断GPU的这个属性。该结构中有一个域integrated,如果设备是集成GPU，那么该域的值是true，否则，该域的值是false。
2. 集成GPU使用零拷贝内存带来的性能提升
集成GPU是系统芯片组中内置的图形处理器，通常与CPU共享系统内存。因此集成GPU使用零拷贝内存通常都会带来性能提升。
3. 独立GPU使用零拷贝内存带来的性能提升
当输入内存和输出内存都只能使用一次时，那么在独立GPU上使用零拷贝内存将带来性能提升。
由于GPU不会缓存零拷贝内存的内容，如果多次读取内存，性能可能会出现严重的恶化，还不如一开始就将数据复制到GPU。

使用多个GPU

1. 点积计算中所需的全部数据

struct DataStruct
{
	int deviceID;//设备标识
	int size;//缓冲区大小
	float *a;//输出缓冲区指针
	float *b;//输入缓冲区指针
	float returnValue;//点积运行结果
};

2. 获得系统中CUDA处理器的数量

	int deviceCount;
	cudaGetDeviceCount(&deviceCount);
	if (deviceCount < 2)
	{
		printf("we need at least two computr 1.0 or greater devices, but only found %d\n", deviceCount);
		system("pause");
		return 0;
	}

3. 用不同的CPU线程控制不同的GPU
每个GPU都需要由一个不同的CPU线程来控制。书中将多线程代码的大部分复杂性都移入到辅助代码文件book.h中。在精简了代码后，我们需要做的就是填充一个结构来执行计算。

	//设置一下线程执行的基本DataStruct信息
	DataStruct data[2];
	data[0].deviceID = 0;
	data[0].size = N / 2;
	data[0].a = a;
	data[0].b = b;
	
	data[1].deviceID = 1;
	data[1].size = N / 2;
	data[1].a = a + N / 2;
	data[1].b = b + N / 2;

4. 创建一个新的CPU线程调用GPU0

	CUTThread thread = start_thread( routine, &(data[0]));

函数start_thread()将创建一个新线程，这个线程将调用routine()，并将DataStruct变量作为参数传递进去。
5. 在当前CPU线程中调用GPU1

	routine(&(data[1]));

6.结束新的CPU线程

	end_thread(thread);

通过调用 end_thread(thread)， 主应用程序线程将等待另一个线程执行完成。
7.routine()函数
**其中最重要的是调用cudaSetDevice()来设定指定的GPU设备。**其他都是平凡的。

//DWORD WINAPI routine(PVOID *pvoidData)
void* routine(void *pvoidData)
{
	DataStruct *data = (DataStruct*)pvoidData;
	cudaSetDevice(data->deviceID);//设置指定的GPU设备
	int size = data->size;
	float *a, *b, c, *partial_c;
	float *dev_a, *dev_b, *dev_partial_c;

	//在CPU上分配内存
	a = data->a;
	b = data->b;
	partial_c = (float*)malloc(blocksPerGrid*sizeof(*partial_c));

	//在GPU上分配内存
	cudaMalloc((void **)&dev_a, size*sizeof(*dev_a));
	cudaMalloc((void **)&dev_b, size*sizeof(*dev_b));
	cudaMalloc((void **)&dev_partial_c, blocksPerGrid*sizeof(*dev_partial_c));

	//将数组a和b复制到GPU上
	cudaMemcpy(dev_a, a, size*sizeof(*a), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, size*sizeof(*b), cudaMemcpyHostToDevice);

	dot << <blocksPerGrid, threadsPerBlock >> >(size, dev_a, dev_b, dev_partial_c);

	//将数组partial_c从GPU复制到CPU上
	cudaMemcpy(partial_c, dev_partial_c, blocksPerGrid*sizeof(*dev_partial_c), cudaMemcpyDeviceToHost);

	//在CPU上做最后的加和处理
	c = 0;
	for (int i = 0; i < blocksPerGrid; i++)
	{
		c += partial_c[i];
	}

	//释放GPU内存
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_partial_c);

	//释放CPU内存
	free(partial_c);

	data->returnValue = c;

	return 0;

}

8. 完整代码

// 代码11.3使用多个GPU
//时间：2019.08.03
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include "book.h"

#define imin(a,b) (a<b?a:b)


const int N = 33 * 1024 * 1024;
const int threadsPerBlock = 256;
const int blocksPerGrid = imin(32, (N + threadsPerBlock - 1) / threadsPerBlock);


__global__ void dot(int size, float *a, float *b, float *partial_c)
{
	__shared__ float cache[threadsPerBlock];
	int tid = threadIdx.x + blockIdx.x*blockDim.x;
	int cacheIndex = threadIdx.x;

	float temp = 0;
	while (tid < size)
	{
		temp += a[tid] * b[tid];
		tid += gridDim.x*blockDim.x;
	}

	//设置cache中的值
	cache[cacheIndex] = temp;

	//同步这个线程块中的线程
	__syncthreads();

	//规约运算
	int i = blockDim.x / 2;
	while (i != 0)
	{
		if (cacheIndex < i)
		{
			cache[cacheIndex] += cache[cacheIndex + i];
		}
		__syncthreads();
		i /= 2;
	}

	if (cacheIndex == 0)
	{
		partial_c[blockIdx.x] = cache[0];
	}

}


struct DataStruct
{
	int deviceID;//设备标识
	int size;//缓冲区大小
	float *a;//输出缓冲区指针
	float *b;//输入缓冲区指针
	float returnValue;//点积运行结果
};

//DWORD WINAPI routine(PVOID *pvoidData)
void* routine(void *pvoidData)
{
	DataStruct *data = (DataStruct*)pvoidData;
	cudaSetDevice(data->deviceID);//设置指定的GPU设备
	int size = data->size;
	float *a, *b, c, *partial_c;
	float *dev_a, *dev_b, *dev_partial_c;

	//在CPU上分配内存
	a = data->a;
	b = data->b;
	partial_c = (float*)malloc(blocksPerGrid*sizeof(*partial_c));

	//在GPU上分配内存
	cudaMalloc((void **)&dev_a, size*sizeof(*dev_a));
	cudaMalloc((void **)&dev_b, size*sizeof(*dev_b));
	cudaMalloc((void **)&dev_partial_c, blocksPerGrid*sizeof(*dev_partial_c));

	//将数组a和b复制到GPU上
	cudaMemcpy(dev_a, a, size*sizeof(*a), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, size*sizeof(*b), cudaMemcpyHostToDevice);

	dot << <blocksPerGrid, threadsPerBlock >> >(size, dev_a, dev_b, dev_partial_c);

	//将数组partial_c从GPU复制到CPU上
	cudaMemcpy(partial_c, dev_partial_c, blocksPerGrid*sizeof(*dev_partial_c), cudaMemcpyDeviceToHost);

	//在CPU上做最后的加和处理
	c = 0;
	for (int i = 0; i < blocksPerGrid; i++)
	{
		c += partial_c[i];
	}

	//释放GPU内存
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_partial_c);

	//释放CPU内存
	free(partial_c);

	data->returnValue = c;

	return 0;

}



int main()
{
	int deviceCount;
	cudaGetDeviceCount(&deviceCount);
	if (deviceCount < 2)
	{
		printf("we need at least two computr 1.0 or greater devices, but only found %d\n", deviceCount);
		system("pause");
		return 0;
	}

	float *a = (float*)malloc(sizeof(*a)*N);
	float *b = (float*)malloc(sizeof(*b)*N);

	//用数据填充主机内存
	for (int i = 0; i < N; i++)
	{
		a[i] = i;
		b[i] = 2 * N;
	}

	//设置一下线程执行的基本DataStruct信息
	DataStruct data[2];
	data[0].deviceID = 0;
	data[0].size = N / 2;
	data[0].a = a;
	data[0].b = b;
	
	data[1].deviceID = 1;
	data[1].size = N / 2;
	data[1].a = a + N / 2;
	data[1].b = b + N / 2;

	CUTThread thread = start_thread( routine, &(data[0]));
	routine(&(data[1]));
	end_thread(thread);
	free(a);
	free(b);
	printf("Value calculated: %f\n", data[0].returnValue + data[1].returnValue);


	system("pause");
	return 0;

}