文章目录
VS2017 CUDA编程学习1:CUDA编程两变量加法运算
VS2017 CUDA编程学习2:在GPU上执行线程
VS2017 CUDA编程学习3:CUDA获取设备上属性信息
VS2017 CUDA编程学习4:CUDA并行处理初探 - 向量加法实现
VS2017 CUDA编程学习5:CUDA并行执行-线程
VS2017 CUDA编程学习6: GPU存储器架构
VS2017 CUDA编程学习7:线程同步-共享内存
VS2017 CUDA编程学习8:线程同步-原子操作
VS2017 CUDA编程学习9:常量内存
VS2017 CUDA编程学习10:纹理内存
VS2017 CUDA编程学习实例1:CUDA实现向量点乘
VS2017 CUDA编程学习11:CUDA性能测量
VS2017 CUDA编程学习12:CUDA流
VS2017 CUDA编程学习实例2:CUDA实现秩排序
前言
今天跟大家分享如何使用CUDA实现直方图统计,这个功能在图像处理中经常被用到,比如直方图归一化处理,OTSU阈值分割等。
1. C++ CUDA实现直方图统计
这里使用了共享内存,同步以及原子加法操作的概念实现直方图统计
详细代码如下所示:
#include<stdio.h>
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <device_atomic_functions.h>
#define SIZE 1000
#define NUM_BIN 256
__global__ void histogram_without_atomic(int* d_a, int* d_b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int item = d_a[tid];
if (tid < SIZE)
{
d_b[item]++;
}
}
__global__ void histogram_atomic(int* d_a, int* d_b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int item = d_a[tid];
if (tid < SIZE)
{
atomicAdd(&(d_b[item]), 1);
}
}
__global__ void histogram_shared_memory(int* d_a, int* d_b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int offset = blockDim.x * gridDim.x;
__shared__ int cache[NUM_BIN];
cache[threadIdx.x] = 0;
//等待所有线程初始化共享内存完成
__syncthreads();
while (tid < SIZE)
{
atomicAdd(&(cache[d_a[tid]]), 1);
tid += offset;
}
//等待所有线程完成直方图累加操作
__syncthreads();
//累加每个块的直方图统计结果
atomicAdd(&(d_b[threadIdx.x]), cache[threadIdx.x]);
}
int main()
{
//定义CPU变量
int h_a[SIZE];
int h_b[NUM_BIN];
for (int i = 0; i < SIZE; i++)
{
h_a[i] = i % NUM_BIN;//初始化数据,数据值范围:0~NUM_BIN-1
}
for (int i = 0; i < NUM_BIN; i++)
{
h_b[i] = 0;
}
//定义GPU变量
int* d_a, *d_b;
//分配GPU内存
cudaMalloc(&d_a, SIZE * sizeof(int));
cudaMalloc(&d_b, NUM_BIN * sizeof(int));
//拷贝数据:从CPU到GPU
cudaMemcpy(d_a, h_a, SIZE * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, NUM_BIN * sizeof(int), cudaMemcpyHostToDevice);
//调用内核函数
//histogram_without_atomic << <((SIZE + NUM_BIN - 1) / NUM_BIN), NUM_BIN >> > (d_a, d_b);
//histogram_atomic << <((SIZE + NUM_BIN - 1) / NUM_BIN), NUM_BIN >> > (d_a, d_b);
histogram_shared_memory << <((SIZE + NUM_BIN - 1) / NUM_BIN), NUM_BIN >> > (d_a, d_b);
//等待GPU所有线程执行结束
cudaDeviceSynchronize();
//拷贝数据:从GPU到CPU
cudaMemcpy(h_b, d_b, NUM_BIN * sizeof(int), cudaMemcpyDeviceToHost);
//打印结果
printf("直方图统计结果:\n");
for (int i = 0; i < NUM_BIN; i++)
{
printf("数据为 %d 的统计数量:%d\n", i, h_b[i]);
}
//释放内存
cudaFree(d_a);
cudaFree(d_b);
system("pause");
return 0;
}
3. 执行结果
因为结果较长,这里只截取部分结果图,如下所示:
总结
为了更好的理解CUDA编程的理论知识,实践是必不可少的,只有更多的实践才能更好的理解新的知识,实践出真知啊。
学习资料
《基于GPU加速的计算机视觉编程》