CUDA系列二：向量的相加

      本文主要介绍下cuda的向量加法的实现，该代码只是为了熟悉概念和初步上手cuda，过程也比较简单。cuda c上手其实还是挺简单的，但是要精通还是需要通过大量实践才能达到的，有兴趣的同学需要多加练习。本文的向量相加，由简逐渐深入。

初步实现只使用了一个block：

#include <iostream>


#define N 10

static void HandleError( cudaError_t err,
                         const char *file,
                         int line ) {
    if (err != cudaSuccess) {
        printf( "%s in %s at line %d\n", cudaGetErrorString( err ),
                file, line );
        exit( EXIT_FAILURE );
    }
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

//核函数具体实现
__global__ void add(int *a,int *b,int *c)
{
	int tid = threadIdx.x; //获取数据索引位置，每个线程对应一个位置
	if(tid < N)
		c[tid] = a[tid] + b[tid];
}

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

    //设备端内存分配
	HANDLE_ERROR(cudaMalloc((void**)&dev_a,N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b,N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c,N * sizeof(int)));

    //生成实验数据
	for(int i = 0;i < N;i++)
	{
		a[i] = i;
		b[i] = i * i;
	}

    //拷贝数据，从主机到设备
	HANDLE_ERROR(cudaMemcpy(dev_a,
				a,
				N * sizeof(int),
				cudaMemcpyHostToDevice));

	HANDLE_ERROR(cudaMemcpy(dev_b,
				b,
				N * sizeof(int),
				cudaMemcpyHostToDevice));

    //调用核函数，这里仅仅采用了一个线程块，对于数据很大的时候会出现问题
	add<<<1,N>>>(dev_a,dev_b,dev_c);

    //拷贝数据，从设备端到主机端
	HANDLE_ERROR(cudaMemcpy(c,
				dev_c,
				N * sizeof(int),
				cudaMemcpyDeviceToHost));

	//show result
	for(int i = 0;i < N;i++)
	{
		printf("%d + %d = %d\n",a[i],b[i],c[i]);
	}

	
	return 0;

}

运行结果：

 

相比上面的实现，下面这种方式实现更加高效，对资源使用更加充分

#include <iostream>


#define N (33 * 1024)


static void HandleError( cudaError_t err,
                         const char *file,
                         int line ) {
    if (err != cudaSuccess) {
        printf( "%s in %s at line %d\n", cudaGetErrorString( err ),
                file, line );
        exit( EXIT_FAILURE );
    }
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))


//核函数的具体实现
__global__ void add(int *a,int *b,int *c)
{
	int tid = threadIdx.x + blockIdx.x * blockDim.x;//具体位置的索引
	
	while(tid < N)
	{
		c[tid] = a[tid] + b[tid];
		tid += blockDim.x * gridDim.x;
	}
}

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

	//设备端内存分配
	HANDLE_ERROR(cudaMalloc((void**)&dev_a,N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b,N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c,N * sizeof(int)));

	//生成数据
	for(int i = 0;i < N;i++)
	{
		a[i] = i;
		b[i] = i * i;
	}

	//拷贝数据，从主机到设备
	HANDLE_ERROR(cudaMemcpy(dev_a,
				a,
				N * sizeof(int),
				cudaMemcpyHostToDevice));
	
	HANDLE_ERROR(cudaMemcpy(dev_b,
				b,
				N * sizeof(int),
				cudaMemcpyHostToDevice));
	
	add<<<128,128>>>(dev_a,dev_b,dev_c); //采用128个block，每个block中128个线程

	//拷贝数据，从设备到主机
	HANDLE_ERROR(cudaMemcpy(c,
				dev_c,
				N * sizeof(int),
				cudaMemcpyDeviceToHost));
	
	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");
	
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);


	return 0;
}

运行结果：

水平有限，如有错去，请指教，谢谢！