CUDA:求平方和

CUDA:求平方和

1.求平方和(1)

#include<stdio.h>
#include<stdlib.h>
#include<cuda_runtime.h>
#define DATA_SIZE 1048576
int data[DATA_SIZE];

void GenerateNumbers(int *number,int size)
{
	//这个函式会产生一大堆 0 ~ 9 之间的随机数
	for(int i=0;i<size;i++)
	{
		number[i]=rand()%10;
	}

}

//在 CUDA 中,在函式前面加上__global__表示这个函式是要在显示芯片上执行的。因此,加入以下的函式:
__global__ static void sumOfSquares(int *num,int *result,clock_t *time)
{
	int sum=0;
	int i;
	clock_t start=clock();
	//CUDA 提供了一个clock 函式,可以取得目前的 timestamp
	for(i=0;i<DATA_SIZE;i++)
	{
		sum+=num[i]*num[i];
	}
	*result=sum;
	printf("%d\n",*result);
	*time=clock()-start;
}
int main()
{
	//要利用 CUDA 进行计算之前,要先把数据复制到显卡内存中,才能让显示芯片使用。
	GenerateNumbers(data, DATA_SIZE);
	int* gpudata, *result;
	clock_t *time;
	cudaMalloc((void **)&gpudata, sizeof(int)*DATA_SIZE);
	cudaMalloc((void **)&result, sizeof(int));
	cudaMalloc((void**) &time, sizeof(clock_t));
	cudaMemcpy(gpudata,data,sizeof(int)*DATA_SIZE,cudaMemcpyHostToDevice);
	//上面这段程序会先呼叫GenerateNumbers 产生随机数,并呼叫cudaMalloc 取得一块显卡内存
	//(result 则是用来存取计算结果,在稍后会用到),并透过 cudaMemcpy将产生的随机数复制到显卡内存中

	//函式名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
	//因为这个程序只使用一个 thread,所以 block 数目、thread 数目都是 1。我们也没有使用到任何 shared memory,所以设为 0
	sumOfSquares<<<1,1,0>>>(gpudata,result,time);
	int sum;
	clock_t time_used;
	cudaMemcpy(&sum,result,sizeof(int),cudaMemcpyDeviceToHost);
	cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost);
	cudaFree(gpudata);
	cudaFree(result);
	printf("sum:%d,time:%d\n",sum,time_used);
	sum = 0; 
	for(int i = 0; i < DATA_SIZE; i++) 
	{
		sum += data[i] * data[i]; 
	} 
	printf("sum (CPU): %d\n", sum);
	system("pause");
}

2.求平方和(2)——使用多线程

//要怎么把计算平方和的程序并行化呢?最简单的方法,似乎就是把数字分成若干组,
//把各组数字分别计算平方和后,最后再把每组的和加总起来就可以了。
//一开始,我们可以把最后加总的动作,由 CPU 来进行。
#include<stdio.h>
#include<stdlib.h>
#include<cuda_runtime.h>
#include<thread>
#define data_size 1048576
#define thread_num 256 //设定thread数目
int data[data_size];

//产生数组元素值
void generatenumbers(int *number,int size)
{
	for (int i=0;i<size;i++)
	{
		number[i]=rand()%10; //产生0~9的随机数
	}
}
__global__ static void sumofsquares(int *num,int *result,clock_t* time)
{

	const int tid = threadIdx.x; //取得线程号
	printf("%d\n",tid);
	//程序里的threadidx 是 cuda 的一个内建的变量,
	//表示目前的 thread 是第几个 thread(由 0 开始计算)
	const int size = data_size / thread_num;//将数据分成多少组
	int sum=0;
	int i;
	clock_t start;
	if(tid==0)
		start = clock(); //开始时间
	//for(i==0;i<DATA_SIZE;i+=thread_num)
	//我们应该要让 thread 0 读取第一个数字,thread 1 读取第二个数字…依此类推。
	for(i = tid * size; i < (tid + 1) * size; i++)
	{ 
		sum += num[i] * num[i]; 
	}
	result[tid] = sum; 
	if(tid == 0) 
		*time = clock() - start;
}

int main()
{
	//...cuda初始化
	generatenumbers(data,data_size); //产生随机数

	int *gpudata,*result; //gpu设备上的数
	clock_t* time; //计算时间(以gpu时钟为基准)

	cudaMalloc((void**) &gpudata, sizeof(int) * data_size); 
	cudaMalloc((void**) &result, sizeof(int) * thread_num);
	cudaMalloc((void**) &time, sizeof(clock_t)); 
	cudaMemcpy(gpudata, data, sizeof(int) * data_size, cudaMemcpyHostToDevice);

	// 函式名称<<<block 数目,thread 数目,shared memory 大小>>>(参数...)
	sumofsquares<<<1, thread_num, 0>>>(gpudata, result, time); 
	int sum[thread_num]; 
	clock_t time_used;
	cudaMemcpy(&sum, result, sizeof(int) * thread_num, cudaMemcpyDeviceToHost); 
	cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost);
	cudaFree(gpudata); 
	cudaFree(result); 
	cudaFree(time);

	//计算每个线程计算出来和的总和
	int final_sum = 0; 
	for(int i = 0; i < thread_num; i++) 
	{ final_sum += sum[i]; }
	printf("sum: %d time: %d\n", final_sum, time_used); 
	
}

3.求平方和(3)

//在 CUDA 中,thread 是可以分组的,也就是 block。
//一个 block 中的 thread,具有一个共享的 shared memory,也可以进行同步工作。
//不同 block 之间的 thread 则不行。在我们的程序中,其实不太需要进行 thread 的同步动作,
//因此我们可以使用多个 block 来进一步增加 thread 的数目。


#include<stdio.h>
#include<stdlib.h>
#include<cuda_runtime.h>
#include<thread>
#define DATA_SIZE 1048576
#define BLOCK_NUM 32 //块数量
#define THREAD_NUM 256 //每个块中线程数
int data[DATA_SIZE];
//这表示我们会建立 32 个 blocks,每个 blocks 有 256 个 threads,总共有 32*256 = 8192 个 threads。

//产生数组元素值
void GenerateNumbers(int *number,int size)
{
	for (int i=0;i<size;i++)
	{
		number[i]=rand()%10; //产生0~9的随机数
	}
}
__global__ static void sumOfSquares(int *num,int *result,clock_t* time)
{

	const int tid = threadIdx.x; //取得线程号
	const int bid = blockIdx.x; //获得块号
	int sum=0;
	int i;
	if(tid==0)
		time[bid]=clock(); //开始时间
	for(i=bid*THREAD_NUM+tid;i<DATA_SIZE;i+=BLOCK_NUM*THREAD_NUM) //注意循环内变化
	{
		sum += num[i]*num[i];
	}
	result[bid*THREAD_NUM+tid] = sum; //第bid个块内第tid个线程计算的结果
	if(tid==0)
		time[bid+BLOCK_NUM] = clock(); //运行时间
	//把计算时间的方式改成每个 block 都会记录开始时间及结束时间。
}

int main()
{
	//...CUDA初始化
	GenerateNumbers(data,DATA_SIZE); //产生随机数

	int *gpudata,*result; //gpu设备上的数
	clock_t* time; //计算时间(以GPU时钟为基准)

	cudaMalloc((void**)&gpudata,sizeof(int)*DATA_SIZE); //分配显存,Allocate memory on the device
	cudaMalloc((void**)&result,sizeof(int)*BLOCK_NUM*THREAD_NUM);
	cudaMalloc((void**)&time,sizeof(clock_t)*BLOCK_NUM*2);
	// Copies data between host and device
	cudaMemcpy(gpudata,data,sizeof(int)*DATA_SIZE,cudaMemcpyHostToDevice); //Host->Device

	// 函式名称<<<block 数目,thread 数目,shared memory 大小>>>(参数...)
	sumOfSquares<<<BLOCK_NUM,THREAD_NUM,0>>>(gpudata,result,time);

	int sum[THREAD_NUM*BLOCK_NUM];
	clock_t time_used[BLOCK_NUM*2]; //运行时间
	cudaMemcpy(&sum,result,sizeof(int)*THREAD_NUM*BLOCK_NUM,cudaMemcpyDeviceToHost);
	cudaMemcpy(&time_used,time,sizeof(clock_t)*BLOCK_NUM*2,cudaMemcpyDeviceToHost);
	cudaFree(gpudata); //释放显存
	cudaFree(result);
	cudaFree(time);

	//计算每个线程计算出来和的总和
	int final_sum=0;
	for(int i=0;i<THREAD_NUM*BLOCK_NUM;i++)
	{
		final_sum += sum[i];
	}
	clock_t min_start,max_end;
	min_start=time_used[0];
	max_end=time_used[BLOCK_NUM];
	//计算GPU上总运行时间
	//找到计算时间最长的block,
	//把每个 block 最早的开始时间,和最晚的结束时间相减,取得总运行时间。
	for (int i=0;i<BLOCK_NUM;i++)
	{
		if(min_start>time_used[i])
			min_start=time_used[i];
		if(max_end<time_used[i+BLOCK_NUM])
			max_end=time_used[i+BLOCK_NUM];
	}
	printf("sum:%d  time:%d\n",final_sum,max_end-min_start);

	//在上面的代码后
	//在cpu上计算验证
	final_sum=0;
	for(int i=0;i<DATA_SIZE;i++)
	{
		final_sum += data[i]*data[i];
	}
	printf("(CPU) sum:%d\n",final_sum);
	system("pause");
}




4.求平方和(4)——Thread的同步

//一个block内的thread可以有共享的内存,也可以进行同步。
//我们可以利用这一点,让每个block内的所有thread把自己计算的结果加起来。

#include<stdio.h>
#include<stdlib.h>
#include<cuda_runtime.h>
#include<thread>
#define DATA_SIZE 1048576
#define BLOCK_NUM 32 //块数量
#define THREAD_NUM 256 //每个块中线程数
int data[DATA_SIZE];
//这表示我们会建立 32 个 blocks,每个 blocks 有 256 个 threads,总共有 32*256 = 8192 个 threads。

//产生数组元素值
void GenerateNumbers(int *number,int size)
{
	for (int i=0;i<size;i++)
	{
		number[i]=rand()%10; //产生0~9的随机数
	}
}
__global__ static void sumOfSquares(int *num,int *result,clock_t* time){
	extern __shared__ int shared[];
	//利用__shared__ 声明的变量表示这是 shared memory,是一个 block 中每个 thread 都共享的内存。
	//它会使用在 GPU 上的内存,所以存取的速度相当快,不需要担心 latency 的问题。
	const int tid = threadIdx.x; //取得线程号
	const int bid = blockIdx.x; //获得块号
	int i;
	if(tid==0)time[bid]=clock();
	shared[tid]=0;
	for(i=bid*THREAD_NUM+tid;i<DATA_SIZE;i+=BLOCK_NUM*THREAD_NUM)
	{
		shared[tid]+=num[i]*num[i];
	}

	__syncthreads();
	//__syncthreads() 是一个 CUDA 的内部函数,
	//表示 block 中所有的 thread 都要同步到这个点,才能继续执行。
	//在我们的例子中,由于之后要把所有 thread 计算的结果进行加总,
	//所以我们需要确定每个 thread 都已经把结果写到 shared[tid] 里面了。
	if(tid==0){
		for(i=1;i<THREAD_NUM;i++){
			shared[0]+=shared[i];
		}
		result[bid]=shared[0];
	}
	if(tid == 0) time[bid + BLOCK_NUM] = clock();
}

int main()
{
	//...CUDA初始化
	GenerateNumbers(data,DATA_SIZE); //产生随机数

	int *gpudata,*result; //gpu设备上的数
	clock_t* time; //计算时间(以GPU时钟为基准)

	cudaMalloc((void**)&gpudata,sizeof(int)*DATA_SIZE); //分配显存,Allocate memory on the device
	cudaMalloc((void**)&result,sizeof(int)*BLOCK_NUM);
	cudaMalloc((void**)&time,sizeof(clock_t)*BLOCK_NUM*2);
	
	cudaMemcpy(gpudata,data,sizeof(int)*DATA_SIZE,cudaMemcpyHostToDevice); //Host->Device

	// 函式名称<<<block 数目,thread 数目,shared memory 大小>>>(参数...)
	sumOfSquares<<<BLOCK_NUM,THREAD_NUM,THREAD_NUM*sizeof(int)>>>(gpudata,result,time);

	int sum[BLOCK_NUM];
	clock_t time_used[BLOCK_NUM*2]; //运行时间
	cudaMemcpy(&sum,result,sizeof(int)*BLOCK_NUM,cudaMemcpyDeviceToHost);
	cudaMemcpy(&time_used,time,sizeof(clock_t)*BLOCK_NUM*2,cudaMemcpyDeviceToHost);
	cudaFree(gpudata); //释放显存
	cudaFree(result);
	cudaFree(time);

	//计算每个线程计算出来和的总和
	int final_sum=0;
	for(int i=0;i<BLOCK_NUM;i++)
	{
		final_sum += sum[i];
	}
	clock_t min_start,max_end;
	min_start=time_used[0];
	max_end=time_used[BLOCK_NUM];
	//计算GPU上总运行时间
	//找到计算时间最长的block,
	//把每个 block 最早的开始时间,和最晚的结束时间相减,取得总运行时间。
	for (int i=0;i<BLOCK_NUM;i++)
	{
		if(min_start>time_used[i])
			min_start=time_used[i];
		if(max_end<time_used[i+BLOCK_NUM])
			max_end=time_used[i+BLOCK_NUM];
	}
	printf("sum:%d  time:%d\n",final_sum,max_end-min_start);

	//在上面的代码后
	//在cpu上计算验证
	final_sum=0;
	for(int i=0;i<DATA_SIZE;i++)
	{
		final_sum += data[i]*data[i];
	}
	printf("(CPU) sum:%d\n",final_sum);
	system("pause");
}



5.求平方和(5)

//求平方和(4)中最后加总的工作,只由每个 block 的 thread 0 来进行,但这并不是最有效率的方法。
//理论上,把 256 个数字加总的动作,是可以并行化的。最常见的方法,是透过树状的加法:
#include<stdio.h>
#include<stdlib.h>
#include<cuda_runtime.h>
#include<thread>
#define DATA_SIZE 1048576
#define BLOCK_NUM 32 //块数量
#define THREAD_NUM 256 //每个块中线程数
int data[DATA_SIZE];
//这表示我们会建立 32 个 blocks,每个 blocks 有 256 个 threads,总共有 32*256 = 8192 个 threads。

//产生数组元素值
void GenerateNumbers(int *number,int size)
{
	for (int i=0;i<size;i++)
	{
		number[i]=rand()%10; //产生0~9的随机数
	}
}
__global__ static void sumOfSquares(int *num,int *result,clock_t* time){
	extern __shared__ int shared[];
	//利用__shared__ 声明的变量表示这是 shared memory,是一个 block 中每个 thread 都共享的内存。
	//它会使用在 GPU 上的内存,所以存取的速度相当快,不需要担心 latency 的问题。
	const int tid = threadIdx.x; //取得线程号
	const int bid = blockIdx.x; //获得块号
	int i;
	int offset=1,mask=1;
	if(tid==0)time[bid]=clock();
	shared[tid]=0;
	for(i=bid*THREAD_NUM+tid;i<DATA_SIZE;i+=BLOCK_NUM*THREAD_NUM)
	{
		shared[tid]+=num[i]*num[i];
	}

	__syncthreads();
	//__syncthreads() 是一个 CUDA 的内部函数,
	//表示 block 中所有的 thread 都要同步到这个点,才能继续执行。
	//在我们的例子中,由于之后要把所有 thread 计算的结果进行加总,
	//所以我们需要确定每个 thread 都已经把结果写到 shared[tid] 里面了。

	while(offset<THREAD_NUM)
	{// while 循环就是进行树状加法。
		if((tid & mask)==0){
			shared[tid]+=shared[tid+offset];
		}
		offset+=offset;
		mask=offset+mask;
		__syncthreads();
	}
	//树状加法是一般的写法,但是它在 GPU 上执行的时候,会有 share memory 的 bank conflict 的问题
	//采用下面的方法,可以避免这个问题:
    //offset = THREAD_NUM / 2;
	//while(offset > 0) { if(tid < offset) { shared[tid] += shared[tid + offset]; } offset >>= 1; __syncthreads(); }

	//这样同时也省去了mask 变数。因此,这个版本的执行的效率就可以再提高一些
	//如果还要再提高效率,可以把树状加法整个展开:
	//if(tid < 128) { shared[tid] += shared[tid + 128]; } __syncthreads(); 
	//if(tid < 64) { shared[tid] += shared[tid + 64]; } __syncthreads(); 
	//if(tid < 32) { shared[tid] += shared[tid + 32]; } __syncthreads(); 
	//if(tid < 16) { shared[tid] += shared[tid + 16]; } __syncthreads(); 
	//if(tid < 8) { shared[tid] += shared[tid + 8]; } __syncthreads();
	//if(tid < 4) { shared[tid] += shared[tid + 4]; } __syncthreads(); 
	//if(tid < 2) { shared[tid] += shared[tid + 2]; } __syncthreads(); 
	//if(tid < 1) { shared[tid] += shared[tid + 1]; }__syncthreads();
	if(tid == 0) 
	{ 
		result[bid] = shared[0];
		time[bid + BLOCK_NUM] = clock();
	}
}

int main()
{
	//...CUDA初始化
	GenerateNumbers(data,DATA_SIZE); //产生随机数

	int *gpudata,*result; //gpu设备上的数
	clock_t* time; //计算时间(以GPU时钟为基准)

	cudaMalloc((void**)&gpudata,sizeof(int)*DATA_SIZE); //分配显存,Allocate memory on the device
	cudaMalloc((void**)&result,sizeof(int)*BLOCK_NUM);
	cudaMalloc((void**)&time,sizeof(clock_t)*BLOCK_NUM*2);
	
	cudaMemcpy(gpudata,data,sizeof(int)*DATA_SIZE,cudaMemcpyHostToDevice); //Host->Device

	// 函式名称<<<block 数目,thread 数目,shared memory 大小>>>(参数...)
	sumOfSquares<<<BLOCK_NUM,THREAD_NUM,THREAD_NUM*sizeof(int)>>>(gpudata,result,time);

	int sum[BLOCK_NUM];
	clock_t time_used[BLOCK_NUM*2]; //运行时间
	cudaMemcpy(&sum,result,sizeof(int)*BLOCK_NUM,cudaMemcpyDeviceToHost);
	cudaMemcpy(&time_used,time,sizeof(clock_t)*BLOCK_NUM*2,cudaMemcpyDeviceToHost);
	cudaFree(gpudata); //释放显存
	cudaFree(result);
	cudaFree(time);

	//计算每个线程计算出来和的总和
	int final_sum=0;
	for(int i=0;i<BLOCK_NUM;i++)
	{
		final_sum += sum[i];
	}
	clock_t min_start,max_end;
	min_start=time_used[0];
	max_end=time_used[BLOCK_NUM];
	//计算GPU上总运行时间
	//找到计算时间最长的block,
	//把每个 block 最早的开始时间,和最晚的结束时间相减,取得总运行时间。
	for (int i=0;i<BLOCK_NUM;i++)
	{
		if(min_start>time_used[i])
			min_start=time_used[i];
		if(max_end<time_used[i+BLOCK_NUM])
			max_end=time_used[i+BLOCK_NUM];
	}
	printf("sum:%d  time:%d\n",final_sum,max_end-min_start);

	//在上面的代码后
	//在cpu上计算验证
	final_sum=0;
	for(int i=0;i<DATA_SIZE;i++)
	{
		final_sum += data[i]*data[i];
	}
	printf("(CPU) sum:%d\n",final_sum);
	system("pause");
}



评论 2
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值