Создание случайных чисел из гауссовского распределения в CUDA

Question

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

The best thing I found был из самой 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? Образцы кода будут высоко оценены.

Спасибо!

Answer 1

Вы можете использовать CURAND, который входит в комплект CUDA Toolkit (версия 3.2 и более поздняя версия). Это было бы намного проще!

Несколько замечаний по коду вы публикуемую:

• Генератор Wallace преобразует Gaussian к гауссовой (т.е. не Uniform к гауссовой)
• CUDA код имеет две неявные переменные: blockIdx и threadIdx - Они определяют индекс блока и индекс потока с блоком, см. CUDA Programming Guide для получения дополнительной информации
• Код использует __mul24, на sm_20, а позже это медленнее, чем «обычное» 32-битное умножение, поэтому я бы избегал его (даже на старых архитектурах для простота)

Answer 2

Box-Muller method также хорош.

Answer 3

Преобразование Fast Walsh Hadamard осуществляется путем сложения и вычитания. Следовательно, применяется центральная предельная теорема. Массив однородных случайных чисел, который претерпевает преобразование Уолша Адамара, будет иметь гауссовское/нормальное распределение. Есть некоторые незначительные технические подробности об этом. Уоллес алгоритм не был обнаружен. Он был впервые опубликован в журнале Servo около 1993/1994 годов. У меня есть код о Уолша Адамара на www.code.google.com/p/lemontree С уважением, Шон О'Коннор