计算CUDA数组中数字的出现次数
Counting occurrences of numbers in a CUDA array
我有一个存储在CUDA GPU上的无符号整数数组(通常是1000000
元素)。我想计算数组中每个数字出现的次数。只有几个不同的数字(关于10
),但这些数字可以从1到1000000
。关于9/10
的数字是0
,我不需要它们的计数。结果如下所示:
58458 -> 1000 occurrences
15 -> 412 occurrences
我有一个使用atomicAdd
s的实现,但它太慢了(很多线程写同一个地址)。有人知道一种快速/有效的方法吗?
您可以通过先对数字排序,然后进行键化缩减来实现直方图。
最直接的方法是先使用thrust::sort
,再使用thrust::reduce_by_key
。它通常也比基于原子的特别分组快得多。下面是一个例子
我想你可以在CUDA示例中找到帮助,特别是直方图示例。它们是GPU计算SDK的一部分。你可以在这里找到http://developer.nvidia.com/cuda-cc-sdk-code-samples#histogram。他们甚至有一个白皮书来解释算法。
我比较了在重复提问时提出的两种方法,即
- 使用
thrust::counting_iterator
和thrust::upper_bound
,遵循直方图推力示例; - 使用
thrust::unique_copy
和thrust::upper_bound
请在下面找到一个完整的工作示例。
#include <time.h> // --- time
#include <stdlib.h> // --- srand, rand
#include <iostream>
#include <thrusthost_vector.h>
#include <thrustdevice_vector.h>
#include <thrustsort.h>
#include <thrustiteratorzip_iterator.h>
#include <thrustunique.h>
#include <thrust/binary_search.h>
#include <thrustadjacent_difference.h>
#include "Utilities.cuh"
#include "TimingGPU.cuh"
//#define VERBOSE
#define NO_HISTOGRAM
/********/
/* MAIN */
/********/
int main() {
const int N = 1048576;
//const int N = 20;
//const int N = 128;
TimingGPU timerGPU;
// --- Initialize random seed
srand(time(NULL));
thrust::host_vector<int> h_code(N);
for (int k = 0; k < N; k++) {
// --- Generate random numbers between 0 and 9
h_code[k] = (rand() % 10);
}
thrust::device_vector<int> d_code(h_code);
//thrust::device_vector<unsigned int> d_counting(N);
thrust::sort(d_code.begin(), d_code.end());
h_code = d_code;
timerGPU.StartCounter();
#ifdef NO_HISTOGRAM
// --- The number of d_cumsum bins is equal to the maximum value plus one
int num_bins = d_code.back() + 1;
thrust::device_vector<int> d_code_unique(num_bins);
thrust::unique_copy(d_code.begin(), d_code.end(), d_code_unique.begin());
thrust::device_vector<int> d_counting(num_bins);
thrust::upper_bound(d_code.begin(), d_code.end(), d_code_unique.begin(), d_code_unique.end(), d_counting.begin());
#else
thrust::device_vector<int> d_cumsum;
// --- The number of d_cumsum bins is equal to the maximum value plus one
int num_bins = d_code.back() + 1;
// --- Resize d_cumsum storage
d_cumsum.resize(num_bins);
// --- Find the end of each bin of values - Cumulative d_cumsum
thrust::counting_iterator<int> search_begin(0);
thrust::upper_bound(d_code.begin(), d_code.end(), search_begin, search_begin + num_bins, d_cumsum.begin());
// --- Compute the histogram by taking differences of the cumulative d_cumsum
//thrust::device_vector<int> d_counting(num_bins);
//thrust::adjacent_difference(d_cumsum.begin(), d_cumsum.end(), d_counting.begin());
#endif
printf("Timing GPU = %fn", timerGPU.GetCounter());
#ifdef VERBOSE
thrust::host_vector<int> h_counting(d_counting);
printf("Aftern");
for (int k = 0; k < N; k++) printf("code = %in", h_code[k]);
#ifndef NO_HISTOGRAM
thrust::host_vector<int> h_cumsum(d_cumsum);
printf("nCountingn");
for (int k = 0; k < num_bins; k++) printf("element = %i; counting = %i; cumsum = %in", k, h_counting[k], h_cumsum[k]);
#else
thrust::host_vector<int> h_code_unique(d_code_unique);
printf("nCountingn");
for (int k = 0; k < N; k++) printf("element = %i; counting = %in", h_code_unique[k], h_counting[k]);
#endif
#endif
}
第一种方法已被证明是最快的。在NVIDIA GTX 960卡上,我有以下N = 1048576
数组元素的计时:
First approach: 2.35ms
First approach without thrust::adjacent_difference: 1.52
Second approach: 4.67ms
请注意,没有必要明确地计算相邻差,因为如果需要,这个操作可以在内核处理期间手动完成。
正如其他人所说,您可以使用sort & reduce_by_key
方法来计数频率。在我的情况下,我需要得到一个数组的模式(最大频率/出现),所以这是我的解决方案:
1 -首先,我们创建两个新数组,一个包含输入数据的副本,另一个包含稍后进行约简的数据(sum):
// Input: [1 3 3 3 2 2 3]
// *(Temp) dev_keys: [1 3 3 3 2 2 3]
// *(Temp) dev_ones: [1 1 1 1 1 1 1]
// Copy input data
thrust::device_vector<int> dev_keys(myptr, myptr+size);
// Fill an array with ones
thrust::fill(dev_ones.begin(), dev_ones.end(), 1);
2 -然后,我们对键进行排序,因为reduce_by_key
函数需要对数组进行排序。
// Sort keys (see below why)
thrust::sort(dev_keys.begin(), dev_keys.end());
3 -稍后,我们为(唯一的)密钥及其频率创建两个输出向量:
thrust::device_vector<int> output_keys(N);
thrust::device_vector<int> output_freqs(N);
4 -最后,我们执行键还原:
// Reduce contiguous keys: [1 3 3 3 2 2 3] => [1 3 2 1] Vs. [1 3 3 3 3 2 2] => [1 4 2]
thrust::pair<thrust::device_vector<int>::iterator, thrust::device_vector<int>::iterator> new_end;
new_end = thrust::reduce_by_key(dev_keys.begin(), dev_keys.end(), dev_ones.begin(), output_keys.begin(), output_freqs.begin());
5 -…如果我们愿意,我们可以得到最频繁的元素
// Get most frequent element
// Get index of the maximum frequency
int num_keys = new_end.first - output_keys.begin();
thrust::device_vector<int>::iterator iter = thrust::max_element(output_freqs.begin(), output_freqs.begin() + num_keys);
unsigned int index = iter - output_freqs.begin();
int most_frequent_key = output_keys[index];
int most_frequent_val = output_freqs[index]; // Frequencies
- 比较并显示使用最小值(a,b)和最大值(a、b)升序排列的4个数字
- 为什么随机数生成器不在void函数中随机化数字,而在main函数中随机化
- 检查输入是否不是整数或数字
- 编译时未启用intel oneApi CUDA支持
- 如何(从固定列表中)选择一个数字序列,该序列将与目标数字相加
- 在cuda线程之间共享大量常量数据
- 如何用数字处理log(0)
- 最高有效数字侧的第N位
- 为什么即使使用-cudart-static进行编译,库用户仍然需要链接到cuda运行时
- 如何获取一个数字的前3位
- 查找最接近的大于当前数字的数字的索引
- 找到两对数字,使它们的乘积的绝对差最小化
- 我想做一个彼此不同但重复出现的数字
- 将数字转换为字母(例如:123 转换为一二三)
- Cuda C++:设备上的Malloc类,并用来自主机的数据填充它
- C++如何计算用户输入的数字中的偶数位数
- 在添加CUDA时返回错误的数字
- 无法在 CUDA 中找到 1 到 100 个数字的简单总和
- "官方"CUDA 减少功能不能接受某些数字?
- 计算CUDA数组中数字的出现次数