Более быстрый способ вычислить массив в Thrust-Cuda

Question

У меня есть четыре массива:

thrust::device_vector<int> vertices;
thrust::device_vector<int> adjacency:
thrust::device_vector<int> degree_summation;
thrust::device_vector<int> degree;

Массив индексов вершин представляет вершины графа, а его содержимое указывает на первого соседа каждой вершины в массиве смежности. Степень массива хранит степень каждой вершины в вершинах.

Я хочу, чтобы массив deg_summation содержал степень каждой вершины плюс суммирование степени каждой окрестности вершины.

Ex: Given the graph:
 0
/ \
1--2

vertices = {0,2,4}
adjacency = {1,2,0,2,0,1}
degree = {2,2,2}

Что я хочу:

degree_summation = {6,6,6}

В настоящее время я использую цикл for для вычисления этого значения, но я думаю, что, возможно, есть более быстрый способ вычисления этих значений с использованием примитивов, заданных thrust. Ядро, которое я использую для вычисления массива степени:

__global__ void computeDegreeSummation(int* vertices,int* adjacency,unsigned int* degree_summation, unsigned int * degree,unsigned int num_vertices, unsigned int num_edges){
unsigned long int tid = (gridDim.y*blockIdx.y+blockIdx.x)*blockDim.x+threadIdx.x;
if(tid < num_vertices){
    int pos_first_neighbor = vertices[tid];
    int pos_last_neighbor;
    if(tid != num_vertices - 1){
        int it = tid;
        pos_last_neighbor = vertices[it+1];
        while(pos_last_neighbor == 0){
            it++;
            pos_last_neighbor = vertices[it+1];
        }
        pos_last_neighbor--;
    }//if
    else{
        pos_last_neighbor = num_edges - 1;
    }//else
    for(int nid = pos_first_neighbor; nid <= pos_last_neighbor; nid++){
        if(adjacency[nid]!=tid){
            degree_summation[tid]+=degrees[adjacency[nid]];
        }//if
    }//for
}//if
}//kernel

Внутри for есть if, так как я инициализировал степень_суммации степенью каждой вершины до вычисления ядра.

Answer 1

Я не знаю, будет ли это быстрее - поскольку я не знаю, как быстро работает ваша эталонная реализация. Но это один из возможных подходов:

1. создать вектор, который определяет окрестности в adjacency для каждой вершины.

   vertices =     {0,  2,  4}
   adjacency =    {1,2,0,2,0,1}
   neighborhood = {1,0,1,0,1,0}  (after scatter)
   neighborhood = {1,1,2,2,3,3}  (after inclusive scan)
   

   Для этого мы будем толкать:: рассеивать набор 1 ( thrust:: constant_iterator) в neighborhood массив в соответствии с индексами, предоставленными vertices, а затем сделать толчок:: inclusive_scan на neighborhood,

2. Затем мы можем использовать thrust:: redu_by_key, используя neighborhood массив в качестве ключей, и с помощью thrust::permutation_iterator для выбора значения из degree соответствующие каждой вершине, перечисленной в каждом neighborhood из adjacency,

3. Тогда нам просто нужно добавить degree до промежуточного degree_summation результат, полученный на шаге 2. Для этого мы можем использовать thrust:: transform.

Вот полностью проработанный пример:

$ cat t612.cu
#include <thrust/device_vector.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/scatter.h>
#include <thrust/scan.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/reduce.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/copy.h>
#include <iostream>


int main(){

  // input data setup

  int h_vertices[] = {0,2,4};
  int h_adjacency[] = {1,2,0,2,0,1};
  int h_degree[] = {2,2,2};
  int vertices_size = sizeof(h_vertices)/sizeof(int);
  int adjacency_size = sizeof(h_adjacency)/sizeof(int);
  int degree_size = sizeof(h_degree)/sizeof(int);
  thrust::device_vector<int> vertices(h_vertices, h_vertices+vertices_size);
  thrust::device_vector<int> adjacency(h_adjacency, h_adjacency+adjacency_size);
  thrust::device_vector<int> degree(h_degree, h_degree+degree_size);
  thrust::device_vector<int> degree_summation(degree_size);

  // create neighborhood array

  thrust::device_vector<int> neighborhood(adjacency_size);
  thrust::constant_iterator<int> first(1);
  thrust::scatter(first, first+adjacency_size, vertices.begin(), neighborhood.begin());
  thrust::inclusive_scan(neighborhood.begin(), neighborhood.end(), neighborhood.begin());

  // sum degree over the neighborhoods

  thrust::reduce_by_key(neighborhood.begin(), neighborhood.end(), thrust::make_permutation_iterator(degree.begin(), adjacency.begin()), thrust::make_discard_iterator(), degree_summation.begin());

  // add the vertex degrees

  thrust::transform(degree.begin(), degree.end(), degree_summation.begin(), degree_summation.begin(), thrust::plus<int>());

  // display results

  thrust::copy(degree_summation.begin(), degree_summation.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}
$ nvcc -arch=sm_35 -o t612 t612.cu
$ ./t612
6,6,6,
$

(удаляя оставшуюся часть моей предыдущей ошибки. Альтернативный подход, который я предложил, был недействительным.)