CUDA 直方图 reduce_by_key 失败
CUDA histogram reduce_by_key failing
我有以下 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 个元素?这是怎么回事?
如果打印出 所有 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 有 "corrupted"。
如果您解决这个问题,您将解决您描述的两个问题(第一个 bin 有 256 个元素,最后一个 bin 有 3 个。)
作为一个简单的证明,添加这行代码:
d_binResults[1023] = 255;
在您的 thrust::transform
操作后立即。结果是正确的。您如何选择更正您的 bin 计算算法取决于您。 (可能 "fixable" 通过将 1 添加到 valRange
但这可能暗示了您预期的直方图值)。
我有以下 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 个元素?这是怎么回事?
如果打印出 所有 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 有 "corrupted"。
如果您解决这个问题,您将解决您描述的两个问题(第一个 bin 有 256 个元素,最后一个 bin 有 3 个。)
作为一个简单的证明,添加这行代码:
d_binResults[1023] = 255;
在您的 thrust::transform
操作后立即。结果是正确的。您如何选择更正您的 bin 计算算法取决于您。 (可能 "fixable" 通过将 1 添加到 valRange
但这可能暗示了您预期的直方图值)。