#include <iostream>
#include "cuda_runtime.h"
#include "time.h"
#include "device_launch_parameters.h"
#include "sm_20_atomic_functions.h"
#include "device_functions.h"
using namespace std;
#define num (1000*10000)//256 * 1024*1000
// 核函数
// 注意,为了方便验证GPU的统计结果,这里采用了"逆直方图",
// 即每发现一个数字,就从CPU的统计结果中减1
__global__ void hist(unsigned char* inputdata, int* outputhist, long size)
{
// 开辟共享内存,否则在全局内存采用原子操作会非常慢(因为冲突太多)
__shared__ int temp[256];
temp[threadIdx.x] = 0;
__syncthreads();
// 计算线程索引及线程偏移量
int ids = blockIdx.x * blockDim.x + threadIdx.x;
int offset = blockDim.x * gridDim.x;
while (ids < size)
{
//采用原子操作对一个block中的数据进行直方图统计
atomicAdd(&temp[inputdata[ids]], 1);
ids += offset;
}
// 等待统计完成,减去统计结果
__syncthreads();
atomicAdd(&outputhist[threadIdx.x], temp[threadIdx.x]);
}
v
cuda矩阵加法
最新推荐文章于 2023-08-01 22:00:00 发布