Генерация случайных чисел из гауссовского распределения в 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 С уважением, Шон О'Коннор

Метод Бокса-Мюллера тоже хорош.

Другие вопросы по тегам