CUDA Atomic операция над массивом в глобальной памяти

Question

У меня есть программа CUDA, ядро ​​которой в основном делает следующее.

• Я предоставляю список из n точек в декартовых координатах, например (x_i,y_i) в плоскости измерения dim_x * dim_y. Я призываю ядро ​​соответственно.
• Для каждой точки на этой плоскости (x_p,y_p) я вычисляю по формуле время, которое потребуется для достижения каждой из этих n точек; учитывая, что эти n точек движутся с определенной скоростью.
• Я упорядочиваю эти времена в возрастающем порядке t_0,t_1,...t_n, где точность t_i установлена ​​в 1. т.е. если t'_i=2.3453, то я бы использовал только t_i=2.3.
• Предполагая, что времена получены из нормального распределения, я моделирую 3 самых быстрых времени, чтобы найти процент времени, когда эти 3 точки достигли самого раннего. Следовательно, предположим, что в результате случайного эксперимента prob_0 = 0,76,prob_1=0,20 и prob_2=0,04. Поскольку t_0 достигает первого наибольшего значения среди трех, я также возвращаю исходный индекс (до сортировки раз) точки. Скажите idx_0 = 5 (целое число).
• Следовательно, для каждой точки на этой плоскости я получаю пару (prob,idx).

Предположим, что n/2 из этих точек относятся к одному виду, а остальные - к другому. Образец сгенерированного изображения выглядит следующим образом.

Особенно, когда точность времени была установлена ​​на 1, я заметил, что количество уникальных трех наборов времени (t_0, t_1, t_2) составило всего лишь 2,5% от общего количества точек данных, то есть количества точек на плоскости. Это означало, что в большинстве случаев ядро ​​бесполезно имитировало, когда оно могло просто использовать значения из предыдущих симуляций. Следовательно, я мог бы использовать словарь, имеющий ключ как 3-х кратный раз и значение как index и prob. Поскольку, насколько я знаю и тестировал, доступ к STL внутри ядра невозможен, я создал массив чисел с плавающей точкой размером 201000000. Этот выбор был сделан экспериментально, поскольку ни один из первых 3 раз не превышал 20 секунд. Следовательно, t_0 может принимать любое значение из {0.0,0.1,0.2,...,20.0}, таким образом, имея 201 выбор. Я мог бы построить ключ для такого словаря, как следующий

• Ключ = t_o * 10^6 + t_1 * 10^3 + t_2

Что касается значения, я мог бы сделать это как (prob+idx). Поскольку idx является целым числом и 0,0<=prob<=1,0, я мог бы получить оба эти значения позже

• проб = ДИКТ [ключ]-Пол (ДИКТ [ключ])
• idx = floor (dict [ключ])

Так что теперь мое ядро ​​выглядит следующим образом

__global__ my_kernel(float* points,float* dict,float *p,float *i,size_t w,...){
  unsigned int col = blockIdx.y*blockDim.y + threadIdx.y;
  unsigned int row = blockIdx.x*blockDim.x + threadIdx.x;
  //Calculate time taken for each of the points to reach a particular point on the plane
  //Order the times in increasing order t_0,t_1,...,t_n
  //Calculate Key = t_o * 10^6 + t_1 * 10^3 + t_2
  if(dict[key]>0.0){
    prob=dict[key]-floor(dict[key])
    idx = floor(dict[key])
  }
  else{
    //Simulate and find prob and idx

    dict[key]=(prob+idx)
  }
  p[row*width+col]=prob;
  i[row*width+col]=idx;
} 

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

Я совершенно уверен, что это связано с состоянием гонки. Обратите внимание, что dict был инициализирован со всеми нулями. Основная идея заключается в том, чтобы структура данных "читала много раз за раз" в определенном месте.

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

Answer 1

Если только ваша симуляция в else ветвление очень длинное (~100 с операций с плавающей запятой), таблица поиска в глобальной памяти, вероятно, будет медленнее, чем выполнение вычислений. Глобальный доступ к памяти очень дорогой!

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

редактировать:

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

Возможно, я что-то неправильно понял, но если вы хотите рассчитать вероятность события x, предполагая нормальное распределение и учитывая среднее mu и стандартное отклонение sigma, нет необходимости генерировать нагрузку случайных чисел и аппроксимировать гауссову кривую. Вы можете напрямую рассчитать вероятность:

p = exp(-((x - mu) * (x - mu) / (2.0f * sigma * sigma))) / 
    (sigma * sqrt(2.0f * M_PI));