Наскоро тествах алгоритъма за намаляване с помощта на CUDA (този, който можете да намерите например на http://www.cuvilib.com/Reduction.pdf, страница 16). Но в края на това се натъкнах на проблеми, когато не използвах атомарност. Така че основно правя сбора на всеки блок и го съхранявам в споделен масив. След това го връщам в глобалния масив x (tdx е threadIndex.x, а i е глобален индекс).
if(i==0){
*sum = 0.; // Initialize to 0
}
__syncthreads();
if (tdx == 0){
x[blockIdx.x] = s_x[tdx]; //get the shared sums in global memory
}
__syncthreads();
След това искам да сумирам първите x елемента (колкото имам блокове). Когато правя с атомарност, работи добре (същия резултат като процесора), но когато използвам коментирания ред по-долу, не работи и често дава "nan":
if(i == 0){
for(int k = 0; k < gridDim.x; k++){
atomicAdd(sum, x[k]); //Works good
//sum[0] += x[k]; //or *sum += x[k]; //Does not work, often results in nan
}
}
Сега всъщност използвам atomicadd директно, за да сумирам споделените суми, но бих искал да разбера защо това не работи. Атомното добавяне е пълна глупост, когато ограничавате операцията до една нишка. И простата сума трябва да работи добре!
__syncthreads()синхронизира нишки само в един и същ блок, а не в различни блокове. Мисля, че неправилният резултат се дължи на проблем със синхронизацията. ЧрезatomicAddвие налагате синхронизацията между различните блокове, които ви липсват от__syncthreads(). - person Vitality schedule 24.07.2013x[k]са резултатите от изчисленията от различни блокове:x[0]е резултатът от блок0,x[1]е резултатът от блок1и т.н. Подозирам, че нишката0може да започне да ги добавя, преди някои блокове наистина да са приключили своите изчисления. Опитайте следното. Поставете втория кодов фрагмент в друго ядро, така че да се наложи синхронизация, и след това опитайте дали редътsum[0] += x[k];работи. - person Vitality schedule 24.07.2013