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
",*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
",sum,time_used);
sum = 0;
for(int i = 0; i < DATA_SIZE; i++)
{
sum += data[i] * data[i];
}
printf("sum (CPU): %d
", 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
",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
", 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
",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
",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
",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
",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
",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
",final_sum);
system("pause");
}声明：该文观点仅代表作者本人，牛骨文系教育信息发布平台，牛骨文仅提供信息存储空间服务。
- 上一篇： php字符串函数相关试题场景
- 下一篇： CUDA常用函数介绍
