nishio-dens's diary

Railsとかプログラミング関連の備忘録

CUDAによる並列擬似乱数生成(LCG)

線形合同法(LCG)を使った並列乱数生成のプログラムを書いてみた.ただしまだバグが残っている.出力される値になぜかマイナスが含まれてしまう.int型の値に0x7fffffffをマスクしているので符号は出力されないはずなのだが...はて

あとシードは今のところ適当.シードをどうやって与えたらいいのかは考え中.

#include <stdio.h>

#define BLOCKNUM 20
#define THREADNUM 512
//一度に生成する乱数の数
#define NUM_OF_GEN (BLOCKNUM * THREADNUM)

__global__ void lcg(unsigned int n, int seed, int* output)
{
  //anci-c rand
  unsigned int paramA = 1103515245UL;
  unsigned int paramP = 12345UL;
  unsigned int paramM = 2147483647UL;
  unsigned int paramN = n;
  unsigned int z = (seed + (threadIdx.x + blockDim.x * blockIdx.x)) % paramM;

  int i;
  for(i=0; i < paramN; i++) {
    z = (paramA * z + paramP) & 0x7fffffff; // mod 2^31
    //use z value
    output[ (blockDim.x * blockIdx.x + threadIdx.x) * i ] = (int)z;
  }

}

int main()
{
  int* buffer;
  int* deviceBuffer;
  dim3 threadNum(THREADNUM);
  dim3 blockNum(BLOCKNUM);
  int n = 500;
  int allocSize = sizeof(int) * threadNum.x * blockNum.x * n;
  
  buffer = (int*)malloc(allocSize);
  cudaMalloc(&deviceBuffer, allocSize);

  lcg<<<blockNum, threadNum>>>(n, 3, deviceBuffer);
  cudaMemcpy( buffer, deviceBuffer, allocSize, cudaMemcpyDeviceToHost);
  cudaFree(deviceBuffer);

  for(int i=0; i < threadNum.x * blockNum.x * n; i++) {
    printf("%d\n", buffer[i] );
  }

  free(buffer);

}