самостоятельные поиски на GPU - как синхронизировать его финиш? - PullRequest
Новая
       92

самостоятельные поиски на GPU - как синхронизировать его финиш?

1 голос
/ 06 января 2012

Предположим, у меня есть некоторый алгоритм generateRandomNumbersAndTestThem (), который возвращает true с вероятностью p и false с вероятностью 1-p. Обычно p очень мало, например р = 0,000001.

Я пытаюсь построить программу в JOCL, которая оценивает p следующим образом: generateRandomNumbersAndTestThem () выполняется параллельно на всех доступных ядрах шейдеров (предпочтительно из нескольких графических процессоров), пока не будет найдено как минимум 100 истин. Тогда оценка для p равна 100 / n, где n - это общее число раз, которое было выполнено generateRandomNumbersAndTestThem ().

Для p = 0,0000001 это означает примерно 10 ^ 9 независимых попыток, что должно сделать очевидным, почему я пытаюсь сделать это на графических процессорах. Но я немного изо всех сил пытаюсь правильно выполнить условие остановки. Моя идея состояла в том, чтобы иметь что-то вроде ядра: 

__kernel void sampleKernel(all_the_input, __global unsigned long *totAttempts) {
    int gid = get_global_id(0);
    //here code that localizes all_the_input for faster access
    while (lessThan100truesFound) {
        totAttempts[gid]++;
        if (generateRandomNumbersAndTestThem()) 
            reportTrue();
    }
}

Как мне реализовать это без серьезной потери производительности, учитывая, что

  • срабатывание «если» будет очень редким событием, и поэтому не является проблемой, если все потоки должны ждать, пока выполняется reportTrue ()
  • lessThan100truesFound должен быть изменен только один раз (от true до false), когда reportTrue () вызывается в сотый раз (поэтому я даже не знаю, является ли логическое значение правильным)
  • план состоит в том, чтобы купить для этого новое оборудование для графических процессоров, так что вы можете использовать новейший графический процессор, например, несколько ATI Radeon HD7970. Но было бы неплохо, если бы я смог протестировать его на моем текущем HD5450.

Я предполагаю, что что-то можно сделать аналогично модификатору "synchronized" в Java, но я не могу найти точный способ сделать это. Каков «правильный» способ сделать это, то есть любой способ, который работает без существенной потери производительности?

1 Ответ

1 голос
/ 06 января 2012

Я бы предложил не использовать глобальный флаг для остановки ядра, а запустить ядро ​​для выполнения определенного количества попыток, проверить на хосте, накопили ли вы достаточно «успехов», и при необходимости повторить.Использование цикла неопределенной длины в ядре является плохим, поскольку драйвер GPU может быть уничтожен сторожевым таймером.Кроме того, проверка некоторой глобальной переменной на каждой итерации, безусловно, приведет к снижению производительности ядра.

Таким образом, reportTrue может быть реализовано как atomic_inc для некоторого счетчика, находящегося в глобальной памяти. 

__kernel void sampleKernel(all_the_input, __global unsigned long *successes) {
    int gid = get_global_id(0);
    //here code that localizes all_the_input for faster access
    for (int i = 0; i < ATT_PER_THREAD; ++i) {
        if (generateRandomNumbersAndTestThem()) 
            atomic_inc(successes);
    }
}

ATT_PER_THREAD настраивается в зависимости от того, сколько времени требуется для выполнения generateRandomNumbersAndTestThem().Затраты на запуск ядра довольно малы, поэтому обычно нет необходимости заставлять ваше ядро ​​работать дольше 0,1–1 секунды

...