《GPU高性能编程CUDA实战》 《学习笔记二》

在GPU上对任意长度的矢量求和

//**************在GPU上对任意长度的矢量求和**********//
/* 				时间：2017-8-9     页数：46
				参考书：《GPU高性能编程CUDA实战》	*/
#include <iostream>
#include <stdio.h>

#define N 33*1024

using namespace std;

__global__ void add(int *a,int *b,int *c);

int main()
{
	int a[N],b[N],c[N];
	int *dev_a,*dev_b,*dev_c;

	//在GPU上分配内存
	cudaMalloc((void**)&dev_a,N*sizeof(int));
	cudaMalloc((void**)&dev_b,N*sizeof(int));
	cudaMalloc((void**)&dev_c,N*sizeof(int));

	//在CPU上为数组"a""b"赋值
	for(int i=0;i<N;i++)
	{
		a[i]=i;
		b[i]=i*i;
	}

	//将数组“a”"b"复制到GPU
	cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice);

	add<<<128,128>>>(dev_a,dev_b,dev_c);

	//将数组“c”从GPU赋值到CPU
	cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);

	//验证GPU确实完成了我们要求的工作
	bool success = true;
	for(int i=0;i<N;i++)
	{
		if((a[i]+b[i])!=c[i])
		{
			printf("Error:%d + %d != %d\n",a[i],b[i],c[i]);
			success = false;
		}
	}
	if(success)
		printf("We did it!\n");

	//释放在GPU上分配的内存
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);

	return 0;
}

__global__ void add(int *a,int *b,int *c)
{
	int tid=threadIdx.x+blockIdx.x*blockDim.x;//线程的起始索引,block是一维的
	while(tid<N)//这里要用while，表示循环迭代，不能用if，不然运行完这一次，就直接跳出循环
	{
		c[tid]=a[tid]+b[tid];
		tid += blockDim.x*gridDim.x;
		//索引的递增步长为线程格中正在运行的线程数量，加上这一句是因为：总的thread数量为128*128=16384,而N=33*1024=33792,如果每个线程计算一个数，那么我们的线程数量不够，所以要加上递增步长。
	}
}