用写一段cuda执行的直方图计算代码

以下是一个简单的 CUDA 直方图计算代码示例：

__global__ void histogram_kernel(int* data, int data_size, int* histogram, int histogram_size) {
  int tid = blockIdx.x * blockDim.x + threadIdx.x;
  if (tid < data_size) {
    atomicAdd(&histogram[data[tid]], 1);
  }
}

void histogram(int* data, int data_size, int* histogram, int histogram_size) {
  int* d_data;
  int* d_histogram;

  cudaMalloc(&d_data, sizeof(int) * data_size);
  cudaMemcpy(d_data, data, sizeof(int) * data_size, cudaMemcpyHostToDevice);

  cudaMalloc(&d_histogram, sizeof(int) * histogram_size);
  cudaMemset(d_histogram, 0, sizeof(int) * histogram_size);

  int block_size = 256;
  int grid_size = (data_size + block_size - 1) / block_size;

  histogram_kernel<<<grid_size, block_size>>>(d_data, data_size, d_histogram, histogram_size);

  cudaMemcpy(histogram, d_histogram, sizeof(int) * histogram_size, cudaMemcpyDeviceToHost);

  cudaFree(d_data);
  cudaFree(d_histogram);
}

这个代码包含两个函数：`histogram` 和 `histogram_kernel`。`histogram_kernel` 是在 GPU 上执行的实际计算。`histogram` 函数分配和释放 GPU 内存，并将数据从主机内存复制到 GPU 内存，然后调用 `histogram_kernel` 函数进行计算，并将结果从 GPU 内存复制回主机内存。

`histogram_kernel` 函数接受四个参数：`data` 指向数据数组的指针，`data_size` 是数据数组的大小，`histogram` 指向直方图数组的指针，`histogram_size` 是直方图数组的大小。该函数计算每个数据元素出现的次数，并使用原子操作将其增加到直方图数组的正确位置。

`histogram` 函数分配和释放两个 GPU 内存数组，并将数据从主机内存复制到 GPU 内存。然后计算需要的线程块和线程数量，调用 `histogram_kernel` 函数执行计算。最后，将计算结果从 GPU 内存复制回主机内存。