Как генерировать случайные перестановки с CUDA
Какие параллельные алгоритмы я мог бы использовать для генерации случайных перестановок из заданного набора?
Особенно полезны предложения или ссылки на документы, подходящие для CUDA.
Последовательной версией этого будет перетасовка Фишера-Йейтса.
Пример:
Пусть S = {1, 2, …, 7} будет набором исходных индексов.
Цель состоит в том, чтобы генерировать n случайных перестановок параллельно.
Каждая из n перестановок содержит каждый из исходных индексов ровно один раз,
например {7, 6, …, 1}.
Решение
Перестановка Фишера-Йейтса может быть распараллелена. Например, для 4 одновременных рабочих требуется всего 3 итерации, чтобы перемешать вектор из 8 элементов. На первой итерации они меняются местами 0<-> 1, 2<-> 3, 4<-> 5, 6<-> 7; на второй итерации 0<-> 2, 1<-> 3, 4<-> 5, 6<-> 7; и на последней итерации 0<-> 4, 1<-> 5, 2<-> 6, 3<-> 7.
Это может быть легко реализовано как CUDA __device__ код (вдохновленный стандартом минимальное / максимальное снижение):
const int id = threadIdx.x;
__shared__ int perm_shared[2 * BLOCK_SIZE];
perm_shared[2 * id] = 2 * id;
perm_shared[2 * id + 1] = 2 * id + 1;
__syncthreads();
unsigned int shift = 1;
unsigned int pos = id * 2;
while(shift <= BLOCK_SIZE)
{
if (curand(&curand_state) & 1) swap(perm_shared, pos, pos + shift);
shift = shift << 1;
pos = (pos & ~shift) | ((pos & shift) >> 1);
__syncthreads();
}
Здесь код инициализации curand опущен, а метод swap(int *p, int i, int j) обменивается ценностями p[i] а также p[j],
Обратите внимание, что приведенный выше код имеет следующие предположения:
- Длина перестановки равна 2 * BLOCK_SIZE, где BLOCK_SIZE — степень 2.
- 2 * BLOCK_SIZE целые числа вписываются в
__shared__память устройства CUDA - BLOCK_SIZE — допустимый размер блока CUDA (обычно между 32 и 512)
Чтобы сгенерировать более одной перестановки, я бы предложил использовать разные блоки CUDA. Если цель состоит в том, чтобы сделать перестановку из 7 элементов (как это было упомянуто в первоначальном вопросе), то я считаю, что это будет быстрее сделать это в одном потоке.
Другие решения
Если длина s = s_L, очень грубый способ сделать это можно реализовать в толчке:
Сначала создайте вектор val длиной s_L x n, который повторяется s n раз.
Создать вектор val_keys, связать n уникальных ключей, повторенных s_L раз с каждым элементом val, например,
val = {1,2,...,7,1,2,...,7,....,1,2,...7}
val_keys = {0,0,0,0,0,0,0,1,1,1,1,1,1,2,2,2,...., n,n,n}
Теперь самое интересное. создать вектор длиной s_L x n равномерно распределенных случайных величин
U = {0.24, 0.1, .... , 0.83}
тогда вы можете сделать итератор zip над val, val_keys и отсортировать их по U:
http://codeyarns.com/2011/04/04/thrust-zip_iterator/
оба val, val_keys будут повсюду, поэтому вы должны снова собрать их вместе с помощью thrust :: stable_sort_by_key (), чтобы убедиться, что если val [i] и val [j] оба принадлежат key [k] и val [i] предшествует val [j] после случайной сортировки, тогда в окончательной версии val [i] все еще должно предшествовать val [j]. Если все идет по плану, val_keys должен выглядеть так же, как и раньше, но val должен отражать тасование.
надеюсь, эта ссылка поможет вам http://www.codeproject.com/Articles/380399/Permutations-with-CUDA-and-OpenCL
Для больших наборов использование примитива сортировки на векторе случайных ключей может быть достаточно эффективным для ваших нужд. Сначала настройте несколько векторов:
const int N = 65535;
thrust:device_vector<uint16_t> d_cards(N);
thrust:device_vector<uint16_t> d_keys(N);
thrust::sequence(d_cards.begin(), d_cards.end());
Затем каждый раз, когда вы хотите перетасовать карты d_cards, вызывайте пару:
thrust::tabulate(d_keys.begin(), d_keys.end(), PRNFunc(rand()*rand());
thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_cards.begin());
// d_cards now freshly shuffled
Случайные ключи генерируются из функтора, который использует начальное число (вычисляется в коде узла и копируется в ядро во время запуска) и номер ключа (который передается в виде таблицы во время создания потока):
struct PRNFunc
{
uint32_t seed;
PRNFunc(uint32_t s) { seed = s; }
__device__ __host__ uint32_t operator()(uint32_t kn) const
{
thrust::minstd_rand randEng(seed);
randEng.discard(kn);
return randEnd();
}
};
Я обнаружил, что производительность могла бы быть улучшена (вероятно, на 30%), если бы я мог выяснить, как кэшировать выделения, которые внутренне выполняет thrust :: sort_by_key.
Любые исправления или предложения приветствуются.