カーネル内の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];
}
まず第一に、宣言された変数の多くは関数でさえ使用されていません！そして、2番目のループで「8」が何を意味するのか本当にわかりません。他のループの「4」は4x4の直交行列ブロックと関係があることを理解しています。誰かが私にここで何が起こっているのかについてより良い考えを教えてもらえますか？
とにかく、誰かが私が使用できる良いコードサンプルを持っていますか？または、CUDAカーネルでランダムなガウス数を生成する別の方法がありますか？コードサンプルは大歓迎です。
ありがとう！