如何并行化包含Thrust变换的for循环？CUDA技术求助

我需要并行化一个包含多个Thrust变换的for循环，目标是计算列向量r的每个元素r2[i]（i从0到N），所有元素可独立计算。原始串行循环代码如下：

for(int i=0; i < N; i++) {
   thrust::device_vector<float> P(N, 0.0);  
   thrust::device_vector<int> corr_col_indices_d(col_indices.begin() + row_begin[i], col_indices.begin() + row_begin[i+1]); // 列索引
   thrust::device_vector<float> corr_values_d(values_d.begin() + row_begin[i], values_d.begin() + row_begin[i+1]); // 列对应的值

    // P[j] = corr_values_d[k] 如果j在corr_col_indices_d中，否则为0（j在索引中时k自增）
    thrust::scatter(corr_values_d.begin(), corr_values_d.end(), corr_col_indices_d.begin(), P.begin());
    r2[i] = thrust::inner_product(P.begin(), P.end(), r1.begin(), 0.0f);
}

我尝试了三种并行实现方法，但均无法编译：

1) 整合到含Lambda的transform中
尝试将所有操作放到一个包含循环变量i的transform里，用Lambda表达式实现：

auto counting_iter = thrust::make_counting_iterator(0);
thrust::transform(counting_iter, counting_iter + N, r2.begin(), [&](int i) {
    thrust::device_vector<float> P(N, 0.0);  
    thrust::device_vector<int> corr_col_indices_d(col_indices.begin() + row_begin[i], col_indices.begin() + row_begin[i+1]); 
    thrust::device_vector<float> corr_values_d(values_d.begin() + row_begin[i], values_d.begin() + row_begin[i+1]); 
    thrust::scatter(corr_values_d.begin(), corr_values_d.end(), corr_col_indices_d.begin(), P.begin());
    thrust::transform(P.begin(), P.end(), r1.begin(), P.begin(), thrust::multiplies<float>());
    return thrust::reduce(P.begin(), P.end());
});

此方案无法运行，推测是该变换方式不被支持或存在语法错误。

2) 接收原始指针的Functor
创建接收device_vector原始指针的Functor：

struct loop {
    // 构造函数接收向量参数
    __host__ __device__
    loop(int *t_row_begin, int *t_col_indices, float*t_values, float *r1): 
        t_row_begin_(t_row_begin), t_col_indices_(t_col_indices), t_values_(t_values), r1_(r1) {}

    // 存储向量的成员变量
    int *t_row_begin_;
    int *t_col_indices_;
    float *t_values_;
    float *r1_;

    __host__ __device__
    float operator()(int i) const {
        thrust::device_vector<float> P(N, 0.0);  
        thrust::device_vector<int> corr_col_indices_d(t_col_indices_ + t_row_begin_[i], t_col_indices_ + t_row_begin_[i + 1]); // 列索引
        thrust::device_vector<float> corr_values_d(t_values_ + t_row_begin_[i], t_values_ + t_row_begin_[i+1]); // 列对应的值

        thrust::scatter(corr_values_d.begin(), corr_values_d.end(), corr_col_indices_d.begin(), P.begin());
        return thrust::inner_product(P.begin(), P.end(), r1.begin(), 0.0f);
    }
};

调用代码：

loop lp(thrust::raw_pointer_cast(row_begin_d.data()), 
            thrust::raw_pointer_cast(col_indices_d.data()), 
            thrust::raw_pointer_cast(values_d.data()), 
            thrust::raw_pointer_cast(r1.data()));
auto iter = thrust::make_counting_iterator(0);
// 用transform执行循环的每个迭代操作
thrust::transform(iter, iter + N, r2.begin(), lp);

该方案仍无法编译。

3) 向operator传递参数的Functor
尝试通过operator参数而非构造函数传入数据：

struct loop {
    __host__ __device__
    float operator()(int i, thrust::device_vector<int>& col_indices, thrust::device_vector<float>& values_d, thrust::device_vector<int>& row_begin, thrust::device_vector<float>& r1) const {
        thrust::device_vector<float> P(N, 0.0);  
        thrust::device_vector<int> corr_col_indices_d(col_indices.begin() + row_begin[i], col_indices.begin() + row_begin[i+1]); // 列索引
        thrust::device_vector<float> corr_values_d(values_d.begin() + row_begin[i], values_d.begin() + row_begin[i+1]); // 列对应的值
        thrust::scatter(corr_values_d.begin(), corr_values_d.end(), corr_col_indices_d.begin(), P.begin());
        return thrust::inner_product(P.begin(), P.end(), r1.begin(), 0.0f);
    }
};

调用代码：

auto iter = thrust::make_counting_iterator(0);
thrust::transform(iter, iter + N, r2.begin(), 
   thrust::make_transform_iterator(iter, loop()), 
   thrust::make_zip_iterator(thrust::make_tuple(col_indices, values_d, row_begin, r1)));

此方案同样无法编译，复杂的错误信息未提供有效调试线索。

环境信息 #

• CUDA版本：11.2
• Thrust版本：1.10.0

补充说明 #

上述向量对应CSR矩阵的组成部分：

vector<int> row_begin;
vector<float> values;
vector<int> col_indices;

更新 #

• 已按照@paleonix的建议将transform和reduce融合为inner_product。

内容的提问来源于stack exchange，提问作者Muhteva