#include "device_functions.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "stdlib.h"
#include <iostream>
using namespace std;
#define LENG 117//数组长度
//返回thread和block
int getThreadNum()
{
cudaDeviceProp prop;//cudaDeviceProp的一个对象
int count = 0;//GPU的个数
cudaGetDeviceCount(&count);
std::cout << "gpu 的个数:" << count << '\n';
cudaGetDeviceProperties(&prop, 0);//第二参数为那个gpu
cout << "最大线程数:" << prop.maxThreadsPerBlock << endl;
cout << "最大网格类型:" << prop.maxGridSize[0] << '\t' << prop.maxGridSize[1] << '\t' << prop.maxGridSize[2] << endl;
return prop.maxThreadsPerBlock;
}
//形参:枚举类型
//判断是否成功分配内存
void GetCudaCalError(cudaError err)
{
if (err != cudaSuccess)
{
cout << "分配内存失败!程序结束!";
}
return;
}
//这个函数有局限性:不能放到多个block上运行
__global__ void sumReduced(float*aGpu, float*sumGpu, int countNum)
{
const int id = threadIdx.x;
//定义一个共享内存
__shared__ float sData[LENG];
//为其赋值
sData[id] = aGpu[id];
//等待每个线程赋值完成
__syncthreads();
//实现归约求和
/*
1、每个数和当前索引加上总数一半的值相加,如果超出最大索引数就加0
2、等到所有线程计算完毕
3、线程数减半,直到减到1
*/
int leng = LENG;
for (int i = countNum / 2.0 + 0.5; i > 1; i = i / 2.0 + 0.5)
{
if (id < i)
{
if (id + i < leng)
{
sData[id] += sData[id + i];
}
//sData[id] += sData[id + i];
}
__syncthreads();
leng = leng / 2.0 + 0.5;
}
if (id == 0)
{
sumGpu[0] = sData[0]+ sData[1];
}
}
int main()
{
float a[LENG];
float asum = 0;
for (int i = 0; i < LENG; ++i)
{
a[i] = i * (i + 1);
//cout << a[i] << '\t';
}
//定义Device上的内存
float *aGpu = 0;
float *sumGpu = 0;
//为其开辟内存
GetCudaCalError(cudaMalloc(&aGpu, LENG * sizeof(float)));
GetCudaCalError(cudaMalloc(&sumGpu, 1 * sizeof(float)));
//给aGpu 赋值
cudaMemcpy(aGpu, a, LENG * sizeof(float), cudaMemcpyHostToDevice);
//开一个block,每个block里面有16个thread
sumReduced << <1, LENG >> > (aGpu, sumGpu, LENG);
//将结果传回host
cudaMemcpy(&asum, sumGpu, 1 * sizeof(float), cudaMemcpyDeviceToHost);
cout << "cuda数组和为:" << asum << endl;
cudaFree(aGpu);
cudaFree(sumGpu);
float testSum = 0;
for (int i = 0; i < LENG; ++i)
{
testSum += a[i];
}
cout << "for数组和为:" << testSum << endl;
}
cuda归约求和最终版
最新推荐文章于 2024-06-03 10:18:50 发布