Повышение эффективности компактного/рассеяния в CUDA

Question

Резюме:

Любые идеи о том, как еще больше улучшить базовую работу рассеяния в CUDA? Особенно, если известно, что он будет использоваться только для компактного массива большего размера в меньший размер? или почему приведенные ниже методы векторизации операций памяти и разделяемой памяти не сработали? Я чувствую, что может быть что-то фундаментальное, чего я не вижу, и любая помощь будет оценена по достоинству.

EDIT 03/09/15: Таким образом, я нашел это Parallel For All Blog post «Оптимизированная фильтрация с разбитой арифметикой атома». Я предположил, что атомистика будет по своей сути медленнее для этой цели, однако я ошибался - тем более, что я не думаю, что мне нужно поддерживать порядок элементов в массиве во время моей симуляции. Мне нужно подумать об этом еще и затем реализовать его, чтобы посмотреть, что произойдет!

EDIT 01/04/16: Я понял, что никогда не писал о моих результатах. К сожалению, в этой статье Parallel for All Blog они сравнили глобальный атомный метод для компактного метода с префикс-суммой Thrust, который на самом деле довольно медленный. Устройства CUB :: IF намного быстрее, чем Thrust - как и префиксная версия, которую я написал с использованием кода CUB для Device :: Scan +. Глобальный атомный метод warp-aggregrate все еще быстрее примерно на 5-10%, но нигде почти в 3-4 раза быстрее, чем я надеялся, основываясь на результатах в блоге. Я все еще использую метод префикс-sum, так как при сохранении порядка элементов не требуется, я предпочитаю согласованность результатов префикс-суммы, и преимущество от атомистики не очень велико. Я по-прежнему пытаюсь использовать различные методы для улучшения компактности, но до сих пор только незначительные улучшения (2%) в лучшем случае значительно увеличили сложность кода.

---

Деталь:

Я пишу моделирование в CUDA, где я COMPACT из элементов я больше не заинтересован в моделировании каждые 40-60 шагов по времени. Из профилирования кажется, что разброс op занимает наибольшее количество времени при уплотнении - более того, чем ядро ​​фильтра или сумма префикса. Сейчас я использую довольно основную функцию рассеяния:

__global__ void scatter_arrays(float * new_freq, const float * const freq, const int * const flag, const int * const scan_Index, const int freq_Index){ 
      int myID = blockIdx.x*blockDim.x + threadIdx.x; 
      for(int id = myID; id < freq_Index; id+= blockDim.x*gridDim.x){ 
       if(flag[id]){ 
        new_freq[scan_Index[id]] = freq[id]; 
       } 
      } 
    } 

freq_Index является количество элементов в старом массиве. Массив флага является результатом фильтра. Scan_ID - результат суммы префикса в массиве флагов.

Попытки, которые я попытался улучшить, - сначала считывать помеченные частоты в разделяемую память, а затем записывать из общей памяти в глобальную память - идея состоит в том, что записи в глобальную память будут более объединены между искажениями (например, вместо потока 0, записывающего в позицию 0 и поток 128, записывающий в позицию 1, поток 0 будет записывать в 0, а поток 1 будет записывать в 1). Я также пробовал векторизовать чтения и записи - вместо чтения и записи float/ints я читал/писал float4/int4 из глобальных массивов, когда это было возможно, поэтому по четыре числа за раз. Я думал, что это ускорит разброс за счет меньшего количества операций с памятью, передающих большие объемы памяти. «Кухонная раковина» код как с Векторизованной нагрузкой памяти/хранят и разделяемая память ниже:

const int compact_threads = 256; 
    __global__ void scatter_arrays2(float * new_freq, const float * const freq, const int * const flag, const int * const scan_Index, const int freq_Index){ 
     int gID = blockIdx.x*blockDim.x + threadIdx.x; //global ID 
     int tID = threadIdx.x; //thread ID within block 
     __shared__ float row[4*compact_threads]; 
     __shared__ int start_index[1]; 
     __shared__ int end_index[1]; 
     float4 myResult; 
     int st_index; 
     int4 myFlag; 
     int4 index; 
     for(int id = gID; id < freq_Index/4; id+= blockDim.x*gridDim.x){ 
      if(tID == 0){ 
       index = reinterpret_cast<const int4*>(scan_Index)[id]; 
       myFlag = reinterpret_cast<const int4*>(flag)[id]; 
       start_index[0] = index.x; 
       st_index = index.x; 
       myResult = reinterpret_cast<const float4*>(freq)[id]; 
       if(myFlag.x){ row[0] = myResult.x; } 
       if(myFlag.y){ row[index.y-st_index] = myResult.y; } 
       if(myFlag.z){ row[index.z-st_index] = myResult.z; } 
       if(myFlag.w){ row[index.w-st_index] = myResult.w; } 
      } 
      __syncthreads(); 
      if(tID > 0){ 
       myFlag = reinterpret_cast<const int4*>(flag)[id]; 
       st_index = start_index[0]; 
       index = reinterpret_cast<const int4*>(scan_Index)[id]; 
       myResult = reinterpret_cast<const float4*>(freq)[id]; 
       if(myFlag.x){ row[index.x-st_index] = myResult.x; } 
       if(myFlag.y){ row[index.y-st_index] = myResult.y; } 
       if(myFlag.z){ row[index.z-st_index] = myResult.z; } 
       if(myFlag.w){ row[index.w-st_index] = myResult.w; } 
       if(tID == blockDim.x -1 || gID == mutations_Index/4 - 1){ end_index[0] = index.w + myFlag.w; } 
      } 
      __syncthreads(); 
      int count = end_index[0] - st_index; 

      int rem = st_index & 0x3; //equivalent to modulo 4 
      int offset = 0; 
      if(rem){ offset = 4 - rem; } 

      if(tID < offset && tID < count){ 
       new_mutations_freq[population*new_array_Length+st_index+tID] = row[tID]; 
      } 

      int tempID = 4*tID+offset; 
      if((tempID+3) < count){ 
       reinterpret_cast<float4*>(new_freq)[tID] = make_float4(row[tempID],row[tempID+1],row[tempID+2],row[tempID+3]); 
      } 

      tempID = tID + offset + (count-offset)/4*4; 
      if(tempID < count){ new_freq[st_index+tempID] = row[tempID]; } 
     } 
     int id = gID + freq_Index/4 * 4; 
     if(id < freq_Index){ 
      if(flag[id]){ 
       new_freq[scan_Index[id]] = freq[id]; 
      } 
     } 
    } 

Очевидно, это становится немного сложнее. :) Хотя указанное ядро ​​кажется стабильным, когда в массиве есть сотни тысяч элементов, я заметил состояние гонки, когда числа массивов в десятках миллионов. Я все еще пытаюсь отследить ошибку.

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

Answer 1

Я нашел алгоритм, упомянутый в этом poster (аналогичный алгоритм, также обсуждаемый в этом paper) работает очень хорошо, особенно для уплотнения больших массивов. Он использует меньше памяти для этого и немного быстрее, чем мой предыдущий метод (5-10%). Я приложил несколько настроек к алгоритму плаката: 1) устранение окончательного сокращения тасования в фазе 1, может просто суммировать элементы по мере их вычисления, 2) дать функции возможность работать больше, чем просто массивы, размер которых равен кратное 1024 + добавление сетчатых петель и 3) позволяет каждому потоку загружать свои регистры одновременно на фазе 3 вместо одного за раз. Я также использую CUB вместо Thrust для Inclusive sum для более быстрого сканирования. Могу быть больше хитростей, которые я могу сделать, но пока это хорошо.

//kernel phase 1 
int myID = blockIdx.x*blockDim.x + threadIdx.x; 
//padded_length is nearest multiple of 1024 > true_length 
for(int id = myID; id < (padded_length >> 5); id+= blockDim.x*gridDim.x){ 
    int lnID = threadIdx.x % warp_size; 
    int warpID = id >> 5; 

    unsigned int mask; 
    unsigned int cnt=0;//;// 

    for(int j = 0; j < 32; j++){ 
     int index = (warpID<<10)+(j<<5)+lnID; 

     bool pred; 
     if(index > true_length) pred = false; 
     else pred = predicate(input[index]); 
     mask = __ballot(pred); 

     if(lnID == 0) { 
      flag[(warpID<<5)+j] = mask; 
      cnt += __popc(mask); 
     } 
    } 

    if(lnID == 0) counter[warpID] = cnt; //store sum 
} 

//kernel phase 2 -> CUB Inclusive sum transforms counter array to scan_Index array 

//kernel phase 3 
int myID = blockIdx.x*blockDim.x + threadIdx.x; 

for(int id = myID; id < (padded_length >> 5); id+= blockDim.x*gridDim.x){ 
    int lnID = threadIdx.x % warp_size; 
    int warpID = id >> 5; 

    unsigned int predmask; 
    unsigned int cnt; 

    predmask = flag[(warpID<<5)+lnID]; 
    cnt = __popc(predmask); 

    //parallel prefix sum 
#pragma unroll 
    for(int offset = 1; offset < 32; offset<<=1){ 
     unsigned int n = __shfl_up(cnt, offset); 
     if(lnID >= offset) cnt += n; 
    } 

    unsigned int global_index = 0; 
    if(warpID > 0) global_index = scan_Index[warpID - 1]; 

    for(int i = 0; i < 32; i++){ 
     unsigned int mask = __shfl(predmask, i); //broadcast from thread i 
     unsigned int sub_group_index = 0; 
     if(i > 0) sub_group_index = __shfl(cnt, i-1); 
     if(mask & (1 << lnID)){ 
      compacted_array[global_index + sub_group_index + __popc(mask & ((1 << lnID) - 1))] = input[(warpID<<10)+(i<<5)+lnID]; 
     } 
    } 
} 

}

EDIT: Существует новый article подмножества авторов плаката, где они EXAMINE быстрее вариации компактна, чем то, что написано выше. Однако их новая версия не является сохранением заказа, поэтому она не полезна для меня, и я ее не реализовал, чтобы проверить ее. Тем не менее, если ваш проект не полагается на порядок объектов, их новая компактная версия, вероятно, может ускорить ваш алгоритм.