[4]CUDA中的向量计算与并行通信模式

CUDA中的向量计算与并行通信模式

• 本节开始，我们将利用GPU的并行能力，对其执行向量和数组操作
• 讨论每个通信模式，将帮助你识别通信模式相关的应用程序，以及如何编写代码

1.两个向量加法程序

• 先写一个通过cpu实现向量加法的程序
• 如下所示，向量相加实际上是模仿GPU的写法，在GPU中，tid 代表特定的某个线程的ID。
• 如果你的cpu是双核的，可以在每个核心上运行一个线程，分别将tid初始化为0和1，然后每次循环的时候+2，这样的话可以实现一个核激素那偶数元素的和，一个核计算基数元素的和，通过两个线程的实现并行计算

#include "stdio.h"
#include<iostream>
//Defining Number of elements in Array
#define N	5
//Defining vector addition function for CPU
void cpuAdd(int *h_a, int *h_b, int *h_c) {
	int tid = 0;	
	while (tid < N)
	{
		h_c[tid] = h_a[tid] + h_b[tid];
		tid += 1;
	}
}

int main(void) {
	int h_a[N], h_b[N], h_c[N];
		//Initializing two arrays for addition
	for (int i = 0; i < N; i++) {
		h_a[i] = 2 * i*i;
		h_b[i] = i;
	}
	//Calling CPU function for vector addition
	cpuAdd (h_a, h_b, h_c);
	//Printing Answer
	printf("Vector addition on CPU\n");
	for (int i = 0; i < N; i++) {
		printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]);
	}
	return 0;
}

• 而总所周知，NVIDIA GPU包含多个块，每个块又包含多个线程，因此可以通过GPU实现更多线程并行计算向量的和，最大程度提高速度
• 可以将代码修改为核函数如下：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>

//Defining number of elements in Array
#define N	5

//Defining Kernel function for vector addition
__global__ void gpuAdd(int* d_a, int* d_b, int* d_c) {
	//Getting block index of current kernel
	int tid = blockIdx.x;	// handle the data at this index
	if (tid < N)
		d_c[tid] = d_a[tid] + d_b[tid];
}

int main(void) {

	//定义主机数组变量
	int h_a[N], h_b[N], h_c[N];

	//定义设备指针变量
	int* d_a, * d_b, * d_c;
	//分配显卡内存
	cudaMalloc((void**)&d_a, N * sizeof(int));
	cudaMalloc((void**)&d_b, N * sizeof(int));
	cudaMalloc((void**)&d_c, N * sizeof(int));
	//Initializing Arrays
	for (int i = 0; i < N; i++) {
		h_a[i] = 2 * i * i;
		h_b[i] = i;
	}
	// Copy input arrays from host to device memory
	cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);

	//设置内核参数为5个块，每个块一个线程 ，并像核函数传递参数
	gpuAdd << <N, 1 >> > (d_a, d_b, d_c);

	//将计算结果从显卡拷贝到主机
	cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);

	printf("Vector addition on GPU \n");
	//Printing result on console
	for (int i = 0; i < N; i++) {
		printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]);
	}
	//Free up memory
	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);
	return 0;
}

• 以上可以发现，通过GPU并行运算，或者多个线程的并行计算，明显的减少了数组的处理时间。比起CPU上的串行计算，提高了吞吐率
• 在此说一下吞吐量的含义：只对网络、设备端口、虚电路或者其他设施，单位时间内成功的传送数据的数量（以比特、字节、分贝等测量）

2. 对比CPU代码和GPU代码的延迟

• CPU的加法程序和GPU的加法程序都是以一个模块化的方式来编写的
• N的值较小时，看不出cpu与GPU的差异，但是当N值很大时，会发现两者计算效率的显著差异
• 下边将展示如何为并行计算计时并对两者时间进行比较

clock_t start_d = clock();
	printf("Doing GPU Vector add\n");
	gpuAdd << <N, 1 >> > (d_a, d_b, d_c);

	cudaThreadSynchronize();
	clock_t end_d = clock();
	double time_d = double(end_d - start_d) / CLOCKS_PER_SEC;
	printf("No of elements in Array: %d \n Device time %f second \n Host time %f second \n ",
		N, time_d, time_h);

3. 对向量的每个元素进行平方

• 前边调用内核函数时启用了N个块，每个块一个线程执行计算；另一种也可以只启动1个块，块里边有N个线程，淡然也可以启用N个块，每个块M个线程
• 下边通过启用一个块中的N个线程来执行向量每个元素的平方运算

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining number of elements in Array
#define N	5
//Kernel function for squaring number
__global__ void gpuSquare(float *d_in, float *d_out) {
	//Getting thread index for current kernel
	int tid = threadIdx.x;	// handle the data at this index
	float temp = d_in[tid];
	d_out[tid] = temp*temp;
}

int main(void) {
	//Defining Arrays for host
	float h_in[N], h_out[N];
	//Defining Pointers for device
	float *d_in, *d_out;

	// allocate the memory on the gpu
	cudaMalloc((void**)&d_in, N * sizeof(float));
	cudaMalloc((void**)&d_out, N * sizeof(float));
	//Initializing Array
	for (int i = 0; i < N; i++) {
		h_in[i] = i;
	}
	//Copy Array from host to device
	cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
	//Calling square kernel with one block and N threads per block
	gpuSquare << <1, N >> >(d_in, d_out);
	//Coping result back to host from device memory
	cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
	//Printing result on console
	printf("Square of Number on GPU \n");
	for (int i = 0; i < N; i++) {
		printf("The square of %f is %f\n", h_in[i], h_out[i]);
	}
	//Free up memory
	cudaFree(d_in);
	cudaFree(d_out);
	return 0;
}

• 需注意

  • 每当使用这种方式启动N个线程并行的时候，需要注意每个块的最大线程不超过 512 或 1024
  • 现在所有计算能力/显卡算力在 3.0 - 7.5 的GPU卡，每个块最大1024个线程
  • 如果N是2000，而你的GPU卡线程的最大数量是512，那么不能写成 << <12 000 > >>，而应该使用<< <4,500 > >>，应该理性的选择合适数量的块和每个块具有的线程数量

4. 并行通信模式

• 当多个线程并行执行时，它们遵循一定的通信模式，知道它们在显存里哪里输入，哪里输出

4.1 映射

• 一对一操作，每个线程或任务读取单一输入，产生单一输出，就是Map模式

d_out[i] = d_in[i] * 2

4.2 收集

• 此模式下，每个线程或者任务，具有多个输入，并产生单个输出，保存到存储器的单一位置，即Gather模式：

out[i] = (in[i-1] + in[i] + in[i+1]) / 3

4.3 分散式

• Scatter 模式，线程或者任务读取单一输入，单项存储器产生多个输出，比如数组排序：

out[i-1] += 2 * in[i] and out[i+1] += 3 * in[i]

4.4 蒙版

• 当线程或者任务要从数组中读取固定形状的相邻元素时，这叫stencil模式，在图像处理中非常有用。比如想用一个3X3或者5X5的窗口进行滑动滤波
• 代码类似Gather

4.5 转置

• 当想要输入矩阵行主序，输出矩阵想要列主序，或者有一个结构数组（SoA)，想转换成一个数组结构(AoS)，它是特别有用的。Transpose模式如下：

out[i+j*128] = in[j + i*128]