Производительность тяги :: Количество

Question

У меня есть следующий код в рамках реорганизации данных для последующего использования в ядре CUDA:

thrust::device_ptr<int> dev_ptr = thrust::device_pointer_cast(dev_particle_cell_indices); 
int total = 0; 
for(int i = 0; i < num_cells; i++) { 
    particle_offsets[i] = total; 
    // int num = 0; 
    int num = thrust::count(dev_ptr, dev_ptr + num_particles, i); 
    particle_counts[i] = num; 
    total += num; 
} 

Теперь, если я изложу num до 0 (раскомментировать 5-й линии, и прокомментируйте 6-е), приложение работает со скоростью более 30 кадров в секунду, что является моей целью. Однако, когда я установил num, равный вызову thrust::count, частота кадров падает примерно до 1-2 кадров в секунду. Почему это произойдет?

Мое понимание заключается в том, что тяга должна быть сборником высоко оптимизированных алгоритмов, которые используют мощь графического процессора, поэтому я удивлен, что это окажет такое влияние на производительность моей программы. Это мой первый раз, когда я использовал тягу, поэтому я могу не знать о некоторых важных деталях.

Есть ли что-то об использовании thrust::count в цикле, которое заставляет его работать так медленно? Как я могу оптимизировать свое использование?

Чтобы дать некоторые цифры, в моем текущем тесте, num_particles составляет около 2000, и num_cells составляет около 1500.

Answer 1

Производительность thrust::count просто отлично, это путь, который вы пытаетесь использовать это проблематично для производительности. Если у вас было много частиц и всего несколько ячеек, то ваша реализация с использованием thrust::count, вероятно, неплохая идея. Ваша проблема в том, что у вас 1500 ячеек. Это означает, что 1500 вызовов count и 1500 устройств для передачи памяти на хост-память каждый раз, когда вы хотите выполнить вычисление. Задержка всех запусков ядра и всех копий шины PCI-e убьет производительность, как вы уже нашли.

Лучший подход для большого количества клеток, было бы что-то вроде этого:

thrust::device_ptr<int> rawin = thrust::device_pointer_cast(dev_particle_cell_indices); 

// Sort a scratch copy of the cell indices by value 
thrust::device_vector<int> cidx(num_particles); 
thrust::copy(rawin, rawin+num_particles, cidx.begin()); 
thrust::sort(cidx.begin(), cidx.end()); 

// Use binary search to extract all the cell counts/offsets 
thrust::counting_iterator<int> cellnumber(0); 
thrust::device_vector<int> offsets(num_cells), counts(num_cells); 

// Offsets come from lower_bound of the ordered cell numbers 
thrust::lower_bound(cidx.begin(), cidx.end(), cellnumber, cellnumber+num_cells, offsets.begin()); 

// Counts come from the adjacent_difference of the upper_bound of the ordered cell numbers 
thrust::upper_bound(cidx.begin(), cidx.end(), cellnumber, cellnumber+num_cells, counts.begin()); 
thrust::adjacent_difference(counts.begin(), counts.end(), counts.begin()); 

// Copy back to the host pointer 
thrust::copy(counts.begin(), counts.end(), particle_counts); 
thrust::copy(offsets.begin(), offsets.end(), particle_offsets); 

Здесь мы сначала отсортировать локальную копию индексов клеток, а затем использовать тяговые двоичные функции поиска, чтобы выполнить то же самое как ваш код, но с гораздо меньшим количеством проходов через данные в памяти GPU и только два устройства для хранения копий памяти, чтобы вернуть все результаты на хост.

Когда я проверил вашу реализацию thrust::count с кодом, который я опубликовал выше, для нетривиального случая (10000 случайных частиц и 2000 ячеек на GeForce 320M с CUDA 4.1 на OS X), я считаю, что ваша версия занимает около 0,95 секунды для запуска, тогда как версия сортировки/поиска занимает около 0,003 секунды для запуска. Таким образом, возможно, вам потребуется несколько сотен раз быстрее, используя тягу, если вы используете более эффективную стратегию и более подходящие алгоритмы.

Answer 2

Я должен быть откровенным здесь, это будет

• критика тяги
• Хвастовство ArrayFire (из которых я являюсь главным разработчиком)

Критика тяги

Они отлично справляются с оптимизацией параллельных алгоритмов для векторных входов. Они используют параллелизм уровня данных (среди прочего) для анализа алгоритмов, которые действительно хорошо работают для big, vector Входы.Но они не могут улучшить его и пройти весь путь до совершенства true уровень параллелизма данных. то есть a большое количество мелких проблем.

Этот второй случай полезен во многих приложениях реального мира, и ArrayFire предоставляет решения в этом отношении (посмотрите gfor, параллельный цикл).

Заглушка для ArrayFire

Что должно было простой призыв к сокращению и сканирования, а не становится 4 алгоритмов (один из которых является дорогостоящим рода) и 3 копии памяти в тяге.

Вот как работает код в ArrayFire:

array cell_indices(num_particles, 1, dev_particle_cell_indices, afDevicePointer); 
array particle_counts = zeros(num_cells); 

gfor(array i, num_cells) // Parallel for loop 
     particle_counts(i) = sum(cell_indices == i); 

array particle_offsets = accum(particle_counts); // Inclusive sum 

-

Настройка

• Я использую talonmies код для ориентира против arrayfire.
• Я использую аналогичную графическую карту (gts 360m) на Linux 64 (cuda 4.1/gcc 4.7).
• Вы можете найти полный эталонный код за here.

Benchmark 1

С num_particles = 2000 и num_cells = 1500 (например, исходной задачи)

$ ./a.out 
Thrust time taken: 0.002384 
ArrayFire time taken: 0.000131 

ArrayFire является раз быстрее

Benchmark 2

С num_particles = 10000 и num_cells = 2000 (например, теста talonmies')

$ ./a.out 
Thrust time taken: 0.002920 
ArrayFire time taken: 0.000132 

ArrayFire является раз быстрее

Benchmark 3

С num_particles = 50000 и num_cells = 5000 (только больший тестовый пример)

$ ./a.out 
Thrust time taken: 0.003596 
ArrayFire time taken: 0.000157 

ArrayFire является раз быстрее

Примечания

• Thrust требует, чтобы переписать код
• Thrust обеспечивает скорость до по ~ 320X над первоначальным кодом
• ArrayFire требует мало переписывания трески е (изменение для к gfor)
• ArrayFire в 18-23 раз быстрее (эффективно ~ 7300x по сравнению с исходным кодом)
• ArrayFire масштабируется лучше (во время выполнения увеличивается на 50% за тягой, 15% для ArrayFire)

Заключение

Упорный действительно обеспечивает приличную скорость, если вы можете переписать вашу проблему. Но это не всегда возможно и нетривиально для более сложных задач. Цифры указывают на то, что существуют возможности для более высокой производительности (из-за высокой степени параллелизма данных), которые просто не достигают тяги.

ArrayFire использует параллельные ресурсы гораздо более эффективно, а время указывает, что gpu еще не насыщен.

Возможно, вы захотите написать свой собственный код cuda или использовать ArrayFire. Я просто хотел указать, что иногда использование тяги не является вариантом, потому что оно практически бесполезно при большом количестве небольших проблем.

EDIT фиксированной эталонные 1 результатов (я использую неправильные номера)