WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/curand/curandTools.h

   1 #ifndef CURAND_TOOLS_H_
   2 #define CURAND_TOOLS_H_
   3
   4 #include <curand_kernel.h>
   5 #include "Indice1D.h"
   6
   7 /*----------------------------------------------------------------------*\
   8  |*                     Declaration                                     *|
   9  \*---------------------------------------------------------------------*/
  10
  11 /*--------------------------------------*\
  12  |*             Public                  *|
  13  \*-------------------------------------*/
  14
  15 /**
  16  *  Utilisation:
  17  *      (1) Fabriquer sur le device les generateurs parallels de nomber aleatoire : setup_kernel_rand
  18  *      (2) Passer au kernel métier ces générateurs : ptrDevTabGeneratorThread
  19  *      (3) Dans le kernel métier
  20  *               curandState generatorThread = ptrDevTabGeneratorThread[tid]; //Optimisation
  21  *               xAlea = curand_uniform(&generatorThread);
  22  * Hyp
  23  *      (H1) Grid et block monodimensionnelle
  24  */
  25 __global__ void setup_kernel_rand(curandState* ptrDevTabGeneratorThread, int deviceId);
  26
  27 /*--------------------------------------*\
  28  |*             Private                 *|
  29  \*-------------------------------------*/
  30
  31 /*----------------------------------------------------------------------*\
  32  |*                     Implementation                                  *|
  33  \*---------------------------------------------------------------------*/
  34
  35 /*--------------------------------------*\
  36  |*             Public                  *|
  37  \*-------------------------------------*/
  38
  39 __global__ void setup_kernel_rand(curandState* ptrDevTabGeneratorThread, int deviceId)
  40     {
  41     int tid = Indice1D::tid();
  42
  43     // Customisation du generator: Proposition (au lecteur de faire mieux)
  44     // Contrainte : Doit etre différent d'un GPU à l'autre
  45     int deltaSeed = deviceId * INT_MAX;
  46     int deltaSequence = deviceId * 100;
  47     int deltaOffset = deviceId * 100;
  48
  49     int seed = 1234 + deltaSeed;    // deviceId+tid;
  50     int sequenceNumber = tid + deltaSequence;    // + tid;
  51     int offset = deltaOffset;
  52
  53     //Each thread gets same seed , a different sequence number , no offset
  54     curand_init(seed, sequenceNumber, offset, &ptrDevTabGeneratorThread[tid]);
  55     }
  56
  57 /*--------------------------------------*\
  58  |*             Private                 *|
  59  \*-------------------------------------*/
  60
  61 #endif
  62
  63 /*----------------------------------------------------------------------*\
  64  |*                     End                                             *|
  65  \*---------------------------------------------------------------------*/