c++ - CUDA 直方图 reduce_by_key 失败

Question

我有以下 CUDA Thrust 代码，它使用 reduce_by_key 将值 [0, 1024) 的直方图绘制到 256 个桶中。我希望每个桶的计数 = 4，但我看到桶 0 有 256 个，桶 255 有 3 个，其余桶有 4 个。

#include <stdio.h>
#include <stdlib.h>

#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#include <thrust/device_vector.h>
#include <thrust/extrema.h>
#include <thrust/pair.h>

#define SIZE 1024

struct binFunc {
    const float minVal;
    const float valRange;
    const int numBins;
    binFunc(float _minVal, float _valRange, int _numBins) :
        minVal(_minVal), valRange(_valRange), numBins(_numBins) {}

    __host__ __device__
    int operator()(float v) const {
        int b = int((v - minVal) / valRange * float(numBins));
        return b;
    }
};

int main() {
    thrust::device_vector<float> d_vec(SIZE);
    for (int i = 0; i < SIZE; ++i)
        d_vec[i] = float(i);

    thrust::device_vector<float>::iterator min;
    thrust::device_vector<float>::iterator max;
    thrust::pair<thrust::device_vector<float>::iterator,
            thrust::device_vector<float>::iterator> minmax =
            thrust::minmax_element(d_vec.begin(), d_vec.end());
    min = minmax.first;
    max = minmax.second;
    float minVal = *min;
    float maxVal = *max;

    std::cout << "The minimum value is " << minVal
            << " and the maximum value is " << maxVal << "." << std::endl;

    float valRange = maxVal - minVal;

    std::cout << "The range is " << valRange << "." << std::endl;

    int numBins = 256;

    thrust::device_vector<int> d_binResults(SIZE);
    thrust::transform(d_vec.begin(), d_vec.end(), d_binResults.begin(),
            binFunc(minVal, valRange, numBins));

    thrust::device_vector<int>::iterator d_binResults_iter =
            d_binResults.begin();
    for (int i = 0; i < 10; ++i) {
        int b = *d_binResults_iter;
        printf("d_binResults[%d]=%d\n", i, b);
        d_binResults_iter++;
    }

    std::cout << "The numBins is " << numBins << "." << std::endl;

    thrust::device_vector<int> d_binsKeys(numBins);
    thrust::device_vector<int> d_binsValues(numBins);

    thrust::pair<thrust::device_vector<int>::iterator,
            thrust::device_vector<int>::iterator> keys_and_values =
            thrust::reduce_by_key(d_binResults.begin(), d_binResults.end(),
                    thrust::constant_iterator<int>(1), d_binsKeys.begin(),
                    d_binsValues.begin());

    thrust::device_vector<int>::iterator d_binsKeys_begin_iter =
            d_binsKeys.begin();
    thrust::device_vector<int>::iterator d_binsValues_begin_iter =
            d_binsValues.begin();
    for (int i = 0; i < numBins; ++i) {
        int key = *d_binsKeys_begin_iter;
        int val = *d_binsValues_begin_iter;
        printf("d_binsValues[%d]=(%d,%d)\n", i, key, val);
        d_binsKeys_begin_iter++;
        d_binsValues_begin_iter++;
    }
    return 0;
}

输出的显着部分是:

d_binsValues[0]=(0,256)
d_binsValues[1]=(1,4)
d_binsValues[2]=(2,4)
...
d_binsValues[254]=(254,4)
d_binsValues[255]=(255,3)

那么，0 号桶有 256 个元素，255 号桶有 3 个元素？这是怎么回事？

Answer 1

如果打印出 所有 d_binResults[] 值而不是前 10 个值，您会发现最后一个元素 (d_binResults[1023]) 的值为 256!但那是一个无效的 bin 索引。对于 numBins = 256，有效索引为 0..255。

它是由于你的仿函数中的计算算法而发生的:

    int b = int((v - minVal) / valRange * float(numBins));

插入最后一个元素的相关值，我们有:

(1023 - 0)/1023*256 = 256

但是 256 是无效的 bin 索引。事实证明，这会破坏 reduce_by_key 操作，导致最后一个 bin 有 3 个元素，第一个 bin 被“损坏”。

如果您解决这个问题，您将解决您描述的两个问题(第一个 bin 有 256 个元素，最后一个 bin 有 3 个。)

作为一个简单的证明，添加这行代码:

d_binResults[1023] = 255;

紧接在您的thrust::transform 操作之后。结果是正确的。您如何选择更正您的 bin 计算算法取决于您。 (可能通过将 1 添加到 valRange 来“修复”，但这可能暗示了您预期的直方图值)。