CUDA 内核，用于查找一维数组中大于特定阈值的最小和最大索引值

Question

我正在尝试编写一个 CUDA 内核，用于查找一维数组中大于特定阈值的最小和最大索引值 下面是 CPU 中执行相同操作的伪代码

int min_index = 0, max_index = length;
    for(int i = 0; i < length; i++)
    {
        if(arr[i] >= threshold)
        {
            min_index = i;
            break;
        }
    }

    for(int i = length - 1 ; i >= min_index; i--)
    {
        if(arr[i] >= threshold)
        {
            max_index = i;
            break;
        }
    }

做到这一点的一种方法是 GPU 使用原子操作，但我不喜欢使用它们。我正在考虑使用归约，但无法弄清楚我需要进行归约的逻辑。对于最小减少，可以遵循如下逻辑

    for(int i = 8; i >= 1; i /= 2)
    {
        if(threadIdx.x < i)
            sh_mem[threadIdx.x] = min(sh_mem[threadIdx.x], sh_mem[threadIdx.x + i]);
        __syncthreads();
    }

但是我可以用什么逻辑来处理我的情况，我需要将共享内存中的值与阈值进行比较并计算满足此条件的共享内存中的最小索引？

非常感谢有关如何解决上述问题的任何想法

Answer 1

以下是如何在 Thrust 中使用

transform_reduce

过滤掉变换中高于阈值的值的示例。缩减部分是对过滤索引的基本

minmax

缩减。虽然使用问题下面的评论中讨论的自定义内核可能可以获得更好的性能，但这应该作为一个很好的参考解决方案（即，这不一定是最终/接受的答案）。

#include <iostream>

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/transform_reduce.h>
#include <thrust/zip_function.h>

int main() {
    thrust::host_vector<float> h_vals{0.0f, 100.0f, 0.0f, 100.0f, 42.0f, 0.0, 0.0};
    thrust::device_vector<float> vals(h_vals);
    float threshold = 42.0f;

    const auto enumerate_iter = thrust::make_zip_iterator(
        thrust::make_tuple(
            thrust::make_counting_iterator(0),
            vals.cbegin()));
    const auto filter_minmax_input = thrust::make_zip_function(
            [threshold] __host__ __device__ (int idx, float val) {
                if (val >= threshold) {
                    return thrust::make_tuple(idx, idx);
                }
                else {
                    return thrust::make_tuple(INT_MAX, -1);
                }
            });
    const auto minmax_f = 
        [] __host__ __device__ (thrust::tuple<int, int> left,
                                thrust::tuple<int, int> right) {
            const int left_min_idx = thrust::get<0>(left);
            const int left_max_idx = thrust::get<1>(left);
            const int right_min_idx = thrust::get<0>(right);
            const int right_max_idx = thrust::get<1>(right);
            return thrust::make_tuple(
                thrust::min(left_min_idx, right_min_idx),
                thrust::max(left_max_idx, right_max_idx));
        };
    const auto res = thrust::transform_reduce(
        enumerate_iter, enumerate_iter + vals.size(),
        filter_minmax_input,
        thrust::make_tuple(INT_MAX, -1),
        minmax_f);
    
    std::cout << thrust::get<0>(res) << ' ' << thrust::get<1>(res) << '\n';
    return 0;
}