CUDA-CODE5-并行编程(2)

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

/*
线程块Block的数量限制是：65535
我的机器上maxThreadPerBlock = 1024.
也就是一个Block中最多1024个线程
介于前一篇中，是对长度很小的向量进行计算，所以直接在一个block中开几个并行线程即可
（前两篇分别是多Blcok单Thread--MBST和单Block多Thread--SBMT）
因此，如果我们要对一个长度超过1024的向量进行运算，该如何做呢？
在这种情况下，要将线程和block结合起来才行,本篇将会使用MBMT--多Block和多线程的方法
修改两个地方：核函数中的索引计算方法和核函数的调用方法
*/

/*
注意，计算核函数运行了几次要用原子操作
*/

/*
此例子是线程少，数据多的情况，具有典型性！！！
*/

#define N (33 * 1024)



__global__ void  add(int *a , int *b , int* c ,int *n)
{
	int tid = threadIdx.x + blockIdx.x * blockDim.x;//获得线程索引
	//			printf("tid = %d\n",tid);
	atomicAdd(n,1);	//计算函数运行了几次	 
	while( tid < N)	//必须检查是否在正确的偏移范围类内
	{
			c[ tid ] = a[ tid ] + b[ tid ];
			tid += blockDim.x * gridDim.x; // 这一句不是很懂
	}
}

int main()
{
	int a[N] , b[N] , c[N];
	int *dev_a, *dev_b , *dev_c;
	int *n = 0;
	int count = 2;	//存放最后的计算此数结果
	//在GPU上分配内存
	cudaMalloc( (void**)&dev_a , N * sizeof(int ) );
	cudaMalloc( (void**)&dev_b , N * sizeof(int ) );
	cudaMalloc( (void**)&dev_c , N * sizeof(int ) );
	cudaMalloc( (void**)&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, n);//那么将会有128 * 128 个核函数运行
	//计算完毕，将数组C从GPU复制到CPU
	cudaMemcpy( c , dev_c , N * sizeof(int) , cudaMemcpyDeviceToHost);
	//计算完毕，将n从GPU复制到CPU
	cudaMemcpy( &count , n , sizeof(int) , cudaMemcpyDeviceToHost);
	//验证GPU确实完成了我们要求的工作
	bool success = true;
	//显示结果
	for(int  i = 0; i < N ; i++)
	{
		if( ( a[i] + b[i]) != c[i] )
		{
			printf("error\n");
			success = false;
		}
		//printf(" %d + %d = %d \n", a[i] , b[i] , c[i]);
	}
	if(success)
		printf("we did it ! \n");
	//最后记得释放在GPU上分配的内存
		printf("n = %d", count);

	cudaFree( dev_a );
	cudaFree( dev_b );
	cudaFree( dev_c );
	cudaFree( n );

	getchar();
	return 0;
}
/*
	使用线程与使用线程块相比存在哪些优势呢？几乎没有任何优势，
	但是，线程块中的并行线程能够完成线程块无法完成的工作。
*/