Генерация случайных чисел из гауссовского распределения в CUDA
Я много искал в Интернете, чтобы найти способ генерировать случайные числа на моем устройстве CUDA, в ядре. Числа должны быть из гауссовского распределения.
Лучшее, что я нашел, было от самой NVIDIA. Это алгоритм Уоллеса, который использует равномерное распределение для построения гауссовского распределения. Но примеры кода, которые они дают, не имеют объяснения, и мне действительно нужно понять, как работает алгоритм, особенно на устройстве. Например, они дают:
__device__ void generateRandomNumbers_wallace(
unsigned seed, // Initialization seed
float *chi2Corrections, // Set of correction values
float *globalPool, // Input random number pool
float *output // Output random numbers
unsigned tid=threadIdx.x;
// Load global pool into shared memory.
unsigned offset = __mul24(POOL_SIZE, blockIdx.x);
for( int i = 0; i < 4; i++ )
pool[tid+THREADS*i] = globalPool[offset+TOTAL_THREADS*i+tid];
__syncthreads();
const unsigned lcg_a=241;
const unsigned lcg_c=59;
const unsigned lcg_m=256;
const unsigned mod_mask = lcg_m-1;
seed=(seed+tid)&mod_mask ;
// Loop generating outputs repeatedly
for( int loop = 0; loop < OUTPUTS_PER_RUN; loop++ )
{
Transform();
unsigned intermediate_address;
i_a = __mul24(loop,8*TOTAL_THREADS)+8*THREADS *
blockIdx.x + threadIdx.x;
float chi2CorrAndScale=chi2Corrections[
blockIdx.x * OUTPUTS_PER_RUN + loop];
for( i = 0; i < 4; i++ )
output[i_a + i*THREADS]=chi2CorrAndScale*pool[tid+THREADS*i];
}
Во-первых, многие из объявленных переменных даже не используются в функции! И я действительно не понимаю, для чего "8" во втором цикле. Я понимаю, что "4" в других циклах как-то связано с блоком ортогональной матрицы 4x4, я прав? Кто-нибудь может дать мне лучшее представление о том, что здесь происходит?
В любом случае, есть ли у кого-нибудь хорошие примеры кода, которые я мог бы использовать? Или у кого-нибудь есть другой способ генерации случайных гауссовых чисел в ядре CUDA? Примеры кода будут высоко ценится.
Спасибо!
3 ответа
Вы можете использовать CURAND, который включен в CUDA Toolkit (версия 3.2 и выше). Это было бы намного проще!
Несколько примечаний по коду, который вы разместили:
- Генератор Уоллеса преобразует гауссову в гауссову (то есть не является равномерной гауссову)
- Код CUDA имеет две неявные переменные:
blockIdx
а такжеthreadIdx
- они определяют индекс блока и индекс потока с блоком, см. Руководство по программированию CUDA для получения дополнительной информации - В коде используется __mul24, в sm_20 и более поздних версиях это на самом деле медленнее, чем в "обычном" 32-битном умножении, поэтому я бы его избегал (даже для более старых архитектур для простоты)
Быстрое преобразование Уолша-Адамара осуществляется с помощью шаблонов сложения и вычитания. Следовательно, применима центральная предельная теорема. Массив равномерных случайных чисел, который подвергается преобразованию Уолша Адамара, будет иметь распределение Гаусса / Нормала. Есть некоторые незначительные технические детали об этом. Алгоритм не был обнаружен Уоллесом. Впервые он был опубликован в журнале Servo Magazine примерно в 1993/1994 годах мной. У меня есть код о преобразовании Уолша Адамара на www.code.google.com/p/lemontree С уважением, Шон О'Коннор