Лучший способ поддерживать состояние RNG на нескольких устройствах в openCL

Question

Поэтому я пытаюсь использовать эту пользовательскую библиотеку RNG для openCL: http://cas.ee.ic.ac.uk/people/dt10/research/rngs-gpu-mwc64x.html

Библиотека определяет структуру состояния:

//! Represents the state of a particular generator
typedef struct{ uint x; uint c; } mwc64x_state_t;

И чтобы сгенерировать случайную uint, вы переходите в состояние в следующую функцию:

uint MWC64X_NextUint(mwc64x_state_t *s)

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

Для проекта, который я создаю, мне нужно иметь возможность генерировать случайные числа не только в разных рабочих группах / элементах, но и на нескольких устройствах одновременно, и у меня возникают проблемы с поиском наилучшего способа разработки этого. Как я должен создать 1 объект mwc64x_state_t на устройство / очередь команд и передать это состояние в качестве глобальной переменной? Или возможно создать 1 объект состояния для всех устройств одновременно? Или я даже не передаю ее как глобальную переменную и не объявляю новое состояние локально в каждой функции ядра?

Библиотека также поставляется с этой функцией:

void MWC64X_SeedStreams(mwc64x_state_t *s, ulong baseOffset, ulong perStreamOffset)

Который предположительно должен разделить ГСЧ на несколько "потоков", но включение этого в мое ядро ​​делает его невероятно медленным. Например, если я сделаю что-то очень простое, например, следующее:

__kernel void myKernel()
{
    mwc64x_state_t rng;
    MWC64X_SeedStreams(&rng, 0, 10000);
}

Тогда вызов ядра становится примерно в 40 раз медленнее.

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

Так что, если кто-то знаком с RNG в openCL или вы уже использовали эту конкретную библиотеку, я буду очень признателен за ваш совет.

Answer 1

Функция MWC64X_SeedStreams действительно относительно медленная, по крайней мере по сравнению с вызовом MWC64X_NextUint, но это верно для большинства параллельных групп RNG, которые пытаются разделить большой глобальный поток на множество подпотоков, которые могут использоваться параллельно. Предполагается, что вы будете вызывать NextUint много раз в ядре (например, сто или больше), но SeedStreams находится только наверху.

Это аннотированная версия примера EstimatePi, поставляемого с библиотекой (mwc64x/test/ valu_pi.cpp и mwc64x/test/test_mwc64x.cl).

__kernel void EstimatePi(ulong n, ulong baseOffset, __global ulong *acc)
{
    // One RNG state per work-item
    mwc64x_state_t rng;

    // This calculates the number of samples that each work-item uses
    ulong samplesPerStream=n/get_global_size(0);

    // And then skip each work-item to their part of the stream, which
    // will from stream offset:
    //   baseOffset+2*samplesPerStream*get_global_id(0)
    // up to (but not including):
    //   baseOffset+2*samplesPerStream*(get_global_id(0)+1)
    //
    MWC64X_SeedStreams(&rng, baseOffset, 2*samplesPerStream);


    // Now use the numbers
    uint count=0;
    for(uint i=0;i<samplesPerStream;i++){
        ulong x=MWC64X_NextUint(&rng);
        ulong y=MWC64X_NextUint(&rng);
        ulong x2=x*x;
        ulong y2=y*y;
        if(x2+y2 >= x2)
            count++;
    }
    acc[get_global_id(0)] = count;
}

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

Если вам нужно несколько ядер на нескольких устройствах, то вам нужно добавить еще один уровень иерархии для разделения потока, например, если у вас есть:

• K: количество устройств (возможно, на параллельных машинах)
• W: количество рабочих элементов на устройство
• C: количество обращений к NextUint на один рабочий элемент

В итоге вы получите N = KW C общих вызовов к NextUint по всем рабочим элементам. Если ваши устройства определены как k=0..(K-1), то в каждом ядре вы должны сделать:

MWC64X_SeedStreams(&rng, W*C*k, C);

Тогда индексы в потоке будут:

[0             .. N ) : Parts of stream used across all devices
[k*(W*C)       .. (k+1)*(W*C) )    : Used within device k
[k*(W*C)+(i*C) .. (k*W*C)+(i+1)*C ) : Used by work-item i in device k.

Хорошо, если каждое ядро ​​использует меньше образцов C, вы можете переоценить при необходимости.

(Я автор библиотеки).