CUDA统计直方图

#include <iostream>
#include <cuda.h>
#include <curand_kernel.h>

const int N = 1 << 16; // 数据点数量
const int bins = 256;  // 直方图的箱数

__global__ void histogramKernel(unsigned char *data, int *hist, int n) {
    // 每个线程计算一个数据点的直方图
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        atomicAdd(&hist[data[idx]], 1); //假设data[idx] = 5，hist[5]代表第6个数数量加1，
    }
}

void generateRandomData(unsigned char *data, int n) {
    for (int i = 0; i < n; i++) {
        data[i] = rand() % bins; // 生成随机数
    }
}

int main() {
    unsigned char *h_data = new unsigned char[N];
    int *h_hist = new int[bins]();
    unsigned char *d_data;
    int *d_hist;

    // 生成随机数据
    generateRandomData(h_data, N);

    // 分配设备内存
    cudaMalloc(&d_data, N * sizeof(unsigned char));
    cudaMalloc(&d_hist, bins * sizeof(int));
    cudaMemset(d_hist, 0, bins * sizeof(int));

    // 将数据从主机复制到设备
    cudaMemcpy(d_data, h_data, N * sizeof(unsigned char), cudaMemcpyHostToDevice);

    // 启动CUDA内核
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    histogramKernel<<<numBlocks, blockSize>>>(d_data, d_hist, N);

    // 将结果从设备复制到主机
    cudaMemcpy(h_hist, d_hist, bins * sizeof(int), cudaMemcpyDeviceToHost);

    // 打印直方图结果
    for (int i = 0; i < bins; i++) {
        std::cout << "Bin " << i << ": " << h_hist[i] << std::endl;
    }

    // 清理内存
    delete[] h_data;
    delete[] h_hist;
    cudaFree(d_data);
    cudaFree(d_hist);

    return 0;
}