Thrust: суммирование элементов массива, проиндексированных другим массивом [Синтаксическая сумма Matlab (x (индексы))]

Question

Я пытаюсь суммировать элементы массива, индексированные другим массивом, используя библиотеку Thrust, но я не мог " t найти пример. Другими словами, я хочу реализовать синтаксис Matlab в

sum(x(indices)) 

Вот руководство код пытается указать на то, что мне нравится, чтобы достичь:

#define N 65536 

// device array copied using cudaMemcpyToSymbol 
__device__ int global_array[N]; 

// function to implement with thrust 
__device__ int support(unsigned short* _memory, unsigned short* _memShort) 
{ 
    int support = 0; 

    for(int i=0; i < _memSizeShort; i++) 
     support += global_array[_memory[i]]; 

    return support;  
} 

Кроме того, от хоста кода, можно использовать global_array [N] без копирования с cudaMemcpyFromSymbol?

Каждый комментарий/ответ ценится :)

Благодарности

Answer 1

Это очень поздно ответ, представленный здесь, чтобы удалить этот вопрос из списка без ответа. Я уверен, что OP уже нашел решение (с мая 2012 года :-)), но я считаю, что следующее может быть полезным для других пользователей.

Как указано в @talonmies, проблема может быть решена с помощью плавного сбора-сокращения. Решение действительно является применением Thurst's permutation_iterator и reduce. permutation_iterator позволяет (неявно) изменять порядок целевого массива x в соответствии с индексами в массиве indices. reduce выполняет сумму (неявно) переупорядоченного массива.

Это приложение является частью Thrust's documentation ниже сообщил для удобства

#include <thrust/iterator/permutation_iterator.h> 
#include <thrust/reduce.h> 
#include <thrust/device_vector.h> 

// this example fuses a gather operation with a reduction for 
// greater efficiency than separate gather() and reduce() calls 

int main(void) 
{ 
    // gather locations 
    thrust::device_vector<int> map(4); 
    map[0] = 3; 
    map[1] = 1; 
    map[2] = 0; 
    map[3] = 5; 

    // array to gather from 
    thrust::device_vector<int> source(6); 
    source[0] = 10; 
    source[1] = 20; 
    source[2] = 30; 
    source[3] = 40; 
    source[4] = 50; 
    source[5] = 60; 

    // fuse gather with reduction: 
    // sum = source[map[0]] + source[map[1]] + ... 
    int sum = thrust::reduce(thrust::make_permutation_iterator(source.begin(), map.begin()), 
          thrust::make_permutation_iterator(source.begin(), map.end())); 

    // print sum 
    std::cout << "sum is " << sum << std::endl; 

    return 0; 
} 

В приведенном выше примере, map играет роль indices, в то время как source играет роль x.

Что касается дополнительного вопроса в своем комментарии (итерацию уменьшенного числа членов), будет достаточно, чтобы изменить следующую строку

int sum = thrust::reduce(thrust::make_permutation_iterator(source.begin(), map.begin()), 
         thrust::make_permutation_iterator(source.begin(), map.end())); 

в

int sum = thrust::reduce(thrust::make_permutation_iterator(source.begin(), map.begin()), 
         thrust::make_permutation_iterator(source.begin(), map.begin()+N)); 

если вы хотите перебрать только за первые N терминов индексационного массива map.

Наконец, что касается возможности использования global_array с хоста, вы должны заметить, что это вектор, который находится на устройстве, поэтому вам необходимо установить cudaMemcpyFromSymbol, чтобы перенести его на хост.