【CUDA】直方图 Histogram

最新推荐文章于 2025-05-29 09:43:06 发布

原创

最新推荐文章于 2025-05-29 09:43:06 发布 · 1.3k 阅读

23 ·

CC 4.0 BY-SA版权

文章标签：

#算法 #c语言 #c++ #开发语言 #深度学习

直方图

直方图计算是计算在给定范围内数据出现次数的过程。数据可以是一组数字或一组字符。范围也可以是一组数字或一组字符。直方图计算是图像处理和计算机视觉中非常常见的操作。

例如，下面是句子“programming massively parallel processors”的直方图。

范围的线程分配与原子操作的使用

为了并行计算直方图，我们需要将给定范围内数据分成块，并将每个块分配给一个线程。每个线程将使用原子操作来更新直方图。

使用这种方法，我们可以并行地计算一组数据的直方图。然而，这种方法有一个主要缺点。全局内存上的原子操作具有较高的延迟。这意味着线程将等待原子操作完成，然后才能继续。这将导致GPU的利用率降低。

一个解决方案是在共享内存上使用原子操作。这将减少原子操作的延迟。每个块将在共享内存中初始化和更新自己的直方图。然后线程将使用全局内存上的原子操作来更新最终直方图。这种技术称为私有化（privatization）。

Code

Host代码使用随机值初始化输入数据，并调用kernel执行直方图计算。

#include <iostream>
#include <cstdio>
#include <ctime>
#include <cmath>

#include <cuda_runtime.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include "histogram.cuh"
#include "helper_cuda.h"
#include "error.cuh"

const int FORTIME = 1;

int main(void) {

    int n, n_bins, t_x;

    n = 16384;
    n_bins = 4096;
    t_x = 1024;

    thrust::host_vector<uint32_t> h_array(n);
    thrust::host_vector<uint32_t> h_bins(n_bins);

    srand(time(NULL));
    for (int i = 0; i < n; ++i)
        h_array[i] = rand() % n_bins;

    for (int i = 0; i < n; ++i)
        ++h_bins[h_array[i]];

    thrust::device_vector<uint32_t> array = h_array;
    thrust::device_vector<uint32_t>