我有几个块，每个块在整数数组的单独部分上执行。例如：从 array[0] 到 array[9] 的块 1 和从 array[10] 到 array[20] 的块 2。
我可以获得每个块的数组最大值索引的最佳方法是什么？
示例块一 a[0] 到 a[10] 具有以下值：5 10 2 3 4 34 56 3 9 10
所以 56 是索引 6 处的最大值。
我无法使用共享内存，因为数组的大小可能非常大。因此它不适合。有没有库可以让我做到这么快？
我知道约简算法，但我认为我的情况不同，因为我想获取最大元素的索引。
【问题讨论】：
CUDA setup。
标签： cuda
如果我完全理解您想要的是：获取其中最大值的数组 A 的索引。
如果是这样，那么我建议您使用推力库：
你会怎么做：
#include <thrust/device_vector.h>
#include <thrust/tuple.h>
#include <thrust/reduce.h>
#include <thrust/fill.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>
#include <cstdlib>
#include <time.h>
using namespace thrust;
// return the biggest of two tuples
template <class T>
struct bigger_tuple {
__device__ __host__
tuple<T,int> operator()(const tuple<T,int> &a, const tuple<T,int> &b)
{
if (a > b) return a;
else return b;
}
};
template <class T>
int max_index(device_vector<T>& vec) {
// create implicit index sequence [0, 1, 2, ... )
counting_iterator<int> begin(0); counting_iterator<int> end(vec.size());
tuple<T,int> init(vec[0],0);
tuple<T,int> smallest;
smallest = reduce(make_zip_iterator(make_tuple(vec.begin(), begin)), make_zip_iterator(make_tuple(vec.end(), end)),
init, bigger_tuple<T>());
return get<1>(smallest);
}
int main(){
thrust::host_vector<int> h_vec(1024);
thrust::sequence(h_vec.begin(), h_vec.end()); // values = indices
// transfer data to the device
thrust::device_vector<int> d_vec = h_vec;
int index = max_index(d_vec);
std::cout << "Max index is:" << index <<std::endl;
std::cout << "Value is: " << h_vec[index] <<std::endl;
return 0;
}
【讨论】：
这不会使原始海报受益，但对于那些来到此页面寻找答案的人，我会建议使用已经具有功能的推力推力::max_element 正是这样做的 - 返回最大的索引元素。还提供了 min_element 和 minmax_element 函数。有关详细信息，请参阅推力文档here。
【讨论】：
除了使用 Thrust 的建议外，您还可以使用 CUBLAS cublasIsamax 函数。
【讨论】：
与共享内存相比，数组的大小几乎无关紧要，因为每个块中的线程数是限制因素，而不是数组的大小。一种解决方案是让每个线程块在与线程块大小相同的数组大小上工作。也就是说，如果您有 512 个线程，那么块 n 将查看数组 [n] 到数组 [n + 511]。每个块都会进行归约以找到数组该部分中的最高成员。然后将每个部分的最大值带回主机并进行简单的线性搜索以定位整个数组中的最大值。 GPU 每次减少都会将线性搜索减少 512 倍。根据数组的大小，您可能希望在将数据带回之前进行更多减少。 （如果您的数组大小为 3*512^10，您可能希望对 gpu 进行 10 次缩减，并让主机搜索剩余的 3 个数据点。）
【讨论】：
在进行最大值加索引缩减时要注意的一点是，如果数组中有多个相同值的最大元素，即在您的示例中，如果有 2 个或多个等于 56 的值，则返回的索引不会是唯一的，并且可能在每次运行代码时都不同，因为 GPU 上线程排序的时间不是确定性的。
要解决此类问题，您可以使用唯一的排序索引，例如 threadid + threadsperblock * blockid，或者元素索引位置（如果这是唯一的）。然后最大测试是沿着这些路线：
if(a>max_so_far || a==max_so_far && order_a>order_max_so_far)
{
max_so_far = a;
index_max_so_far = index_a;
order_max_so_far = order_a;
}
（索引和顺序可以是同一个变量，具体取决于应用程序。）
【讨论】：