CUDA向量运算

向量相加

为了理解GPU上的向量运算，我们将首先在CPU上编写一个向量加法程序，然后利用GPU的并行结构对其进行修改。我们取两个含有一些数值的数组，按元素计算它们的加法，然后将结果保存在第三个数组。
cpuAdd函数应该很好理解。但其中的tid变量使用可能是理解的难点。我们在这里包含它，是用来在CPU上模仿GPU的写法，而在GPU中,tid代表特定的某个线程的ID。这里请注意，如果你的CPU是双核的，你可以在每个核心上运行一个线程，分别将tid初始化为0和1，然后每次循环的时候+2。这样一个CPU核心将计算偶数元素的和，另外一个CPU核心将计算奇数元素的和。
程序中有两个函数: main和 cpuAdd。在main 函数中，我们首先定义了两个输入数组，并将其中的元素初始化为一些随意的值。然后，我们将这两个数组作为输人传递给cpuAdd函数。cpuAdd 函数将计算结果存储在第三个数组中。

#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()
{
	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;
}

在cpuAdd函数中，对tid用法的解释,会给你一些在多核的GPU 上，如何编写同样的函数的提示。如果我们用GPU上的每个核心的ID，作为这里的tid，那么我们就可以并行的对所有的元素进行相加。
在gpuAdd内核函数中，每个块中的这个线程，用当前块的ID来初始化tid变量。然后根据tid变量，每个线程将一对元素进行相加。这样如果块的总数等于每个数组中元素的总数，那么所有的加法操作将并行完成。

#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()
{
	//Defining host arrays
	int h_a[N], h_b[N], h_c[N];
	//Defining device pointers
	int *d_a, *d_b, *d_c;
	// allocate the memory
	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);
	//Calling kernels with N blocks and one thread per block, passing device pointers as parameters
	gpuAdd << <N, 1 >> >(d_a, d_b, d_c);
	//Copy result back to host memory from device memory
	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;
}

向量各元素求平方

你可能会问:既然我们能并行启动N个块，每个块只有1个线程。那么我们能否只启动1个块，然后里面有N个线程?答案是肯定的，我们可以在1个块里并行启动N个线程。来做第二个例子:将一个向量数组里面的元素进行平方。我们以一个数值数组作为输入，然后输出一个每个元素都平方后的数组。
gpuSquare内核函数具有两个指针参数。第一个指针 d_in指向存储输入数据的显存。第二个指针d_out则指向输出结果的显存。在这程序里，我们并行启动了多个线程，而不是启动多个块，所以每个线程特有的tid被初始化为threadldx.x。

#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() 
{
	//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 cpu
	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 卡，每个块最大1 024个线程)。因此，N的值应该小于或等于这个值。