X-Git-Url: http://git.euphorik.ch/index.cgi?a=blobdiff_plain;f=WCudaMSE%2FBilatTools_Cuda%2Fsrc%2Fcore%2Fcudatools%2Fheader%2Fdevice%2Fcurand%2FcurandTools.h;fp=WCudaMSE%2FBilatTools_Cuda%2Fsrc%2Fcore%2Fcudatools%2Fheader%2Fdevice%2Fcurand%2FcurandTools.h;h=8ac7bd4173def42e07cad6cfb529a3b97cc7a639;hb=8d08c12b29c2a14684f35c023ee39e694bb80d25;hp=0000000000000000000000000000000000000000;hpb=226de81f7e1f1fbf4ac79d0d089e8a05ec7159a0;p=GPU.git

diff --git a/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/curand/curandTools.h b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/curand/curandTools.h
new file mode 100755
index 0000000..8ac7bd4
--- /dev/null
+++ b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/curand/curandTools.h
@@ -0,0 +1,65 @@
+#ifndef CURAND_TOOLS_H_
+#define CURAND_TOOLS_H_
+
+#include <curand_kernel.h>
+#include "Indice1D.h"
+
+/*----------------------------------------------------------------------*\
+ |*			Declaration 					*|
+ \*---------------------------------------------------------------------*/
+
+/*--------------------------------------*\
+ |*		Public			*|
+ \*-------------------------------------*/
+
+/**
+ *  Utilisation:
+ *  	(1) Fabriquer sur le device les generateurs parallels de nomber aleatoire : setup_kernel_rand
+ *  	(2) Passer au kernel métier ces générateurs : ptrDevTabGeneratorThread
+ *  	(3) Dans le kernel métier
+ *  		 curandState generatorThread = ptrDevTabGeneratorThread[tid]; //Optimisation
+ *  		 xAlea = curand_uniform(&generatorThread);
+ * Hyp
+ * 	(H1) Grid et block monodimensionnelle
+ */
+__global__ void setup_kernel_rand(curandState* ptrDevTabGeneratorThread, int deviceId);
+
+/*--------------------------------------*\
+ |*		Private			*|
+ \*-------------------------------------*/
+
+/*----------------------------------------------------------------------*\
+ |*			Implementation 					*|
+ \*---------------------------------------------------------------------*/
+
+/*--------------------------------------*\
+ |*		Public			*|
+ \*-------------------------------------*/
+
+__global__ void setup_kernel_rand(curandState* ptrDevTabGeneratorThread, int deviceId)
+    {
+    int tid = Indice1D::tid();
+
+    // Customisation du generator: Proposition (au lecteur de faire mieux)
+    // Contrainte : Doit etre différent d'un GPU à l'autre
+    int deltaSeed = deviceId * INT_MAX;
+    int deltaSequence = deviceId * 100;
+    int deltaOffset = deviceId * 100;
+
+    int seed = 1234 + deltaSeed;    // deviceId+tid;
+    int sequenceNumber = tid + deltaSequence;    // + tid;
+    int offset = deltaOffset;
+
+    //Each thread gets same seed , a different sequence number , no offset
+    curand_init(seed, sequenceNumber, offset, &ptrDevTabGeneratorThread[tid]);
+    }
+
+/*--------------------------------------*\
+ |*		Private			*|
+ \*-------------------------------------*/
+
+#endif 
+
+/*----------------------------------------------------------------------*\
+ |*			End	 					*|
+ \*---------------------------------------------------------------------*/