Макрос CONFLICT_FREE_OFFSET, используемый в алгоритме параллельного префикса из GPU Gems 3

Question

Прежде всего, вот ссылка на алгоритм http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html

Чтобы избежать конфликтов банков, заполнение добавляется в массив разделяемой памяти каждые NUM_BANKS (т. Е. 32 для устройств вычислимости 2.х) элементы.Это делается (как на рисунке 39-5):

int ai = offset*(2*thid+1)-1
int bi = offset*(2*thid+2)-1
ai += ai/NUM_BANKS
bi += ai/NUM_BANKS
temp[bi] += temp[ai]

Я не понимаю, как ai / NUM_BANKS эквивалентно макросу:

   #define NUM_BANKS 16  
   #define LOG_NUM_BANKS 4  
   #define CONFLICT_FREE_OFFSET(n) \  
          ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))  

Не так лиравно

n >> LOG_NUM_BANKS

Любая помощь приветствуется.Спасибо

Answer 1

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

Если вы хотите выполнять сканирование простым способом, используйте thrust::inclusive_scan или thrust::exclusive_scan.

Если вы действительно хотите реализовать сканирование, обратитесь к более поздним статьям, таким как эта [1] .Или для настоящего опуса с более быстрым кодом, но для этого потребуется немного больше изучения, этот [2] .Или прочитайте учебное пособие Шона Бакстера (хотя последнее не содержит ссылок на основополагающую работу по алгоритму сканирования).

[1] Шубхабрата Сенгупта, Марк Харрис, Майкл Гарланд и ДжонД. Оуэнс.«Эффективные алгоритмы параллельного сканирования для многоядерных графических процессоров».В Jakub Kurzak, David A. Bader и Jack Dongarra, редакторы, Scientific Computing with Multicore и Accelerators, Chapman & Hall / CRC Computational Science, глава 19, страницы 413–442.Тейлор и Фрэнсис, январь 2011 г. http://www.idav.ucdavis.edu/publications/print_pub?pub_id=1041

[2] Меррилл Д. и Гримшоу А. Параллельное сканирование для потоковых архитектур.Технический отчет CS2009-14, факультет компьютерных наук, Университет Вирджинии.Декабрь 2009 г.