Ajout de l'ensemble du workspace.

[GPU.git] / WCudaMSE / BilatTools_Cuda / src / core / cudatools / header / device / curand / curandTools.h

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 (executable)
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_\r
+#define CURAND_TOOLS_H_\r
+\r
+#include <curand_kernel.h>\r
+#include "Indice1D.h"\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    Declaration                                     *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
+/*--------------------------------------*\\r
+ |*            Public                  *|\r
+ \*-------------------------------------*/\r
+\r
+/**\r
+ *  Utilisation:\r
+ *     (1) Fabriquer sur le device les generateurs parallels de nomber aleatoire : setup_kernel_rand\r
+ *     (2) Passer au kernel métier ces générateurs : ptrDevTabGeneratorThread\r
+ *     (3) Dans le kernel métier\r
+ *              curandState generatorThread = ptrDevTabGeneratorThread[tid]; //Optimisation\r
+ *              xAlea = curand_uniform(&generatorThread);\r
+ * Hyp\r
+ *     (H1) Grid et block monodimensionnelle\r
+ */\r
+__global__ void setup_kernel_rand(curandState* ptrDevTabGeneratorThread, int deviceId);\r
+\r
+/*--------------------------------------*\\r
+ |*            Private                 *|\r
+ \*-------------------------------------*/\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    Implementation                                  *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
+/*--------------------------------------*\\r
+ |*            Public                  *|\r
+ \*-------------------------------------*/\r
+\r
+__global__ void setup_kernel_rand(curandState* ptrDevTabGeneratorThread, int deviceId)\r
+    {\r
+    int tid = Indice1D::tid();\r
+\r
+    // Customisation du generator: Proposition (au lecteur de faire mieux)\r
+    // Contrainte : Doit etre différent d'un GPU à l'autre\r
+    int deltaSeed = deviceId * INT_MAX;\r
+    int deltaSequence = deviceId * 100;\r
+    int deltaOffset = deviceId * 100;\r
+\r
+    int seed = 1234 + deltaSeed;    // deviceId+tid;\r
+    int sequenceNumber = tid + deltaSequence;    // + tid;\r
+    int offset = deltaOffset;\r
+\r
+    //Each thread gets same seed , a different sequence number , no offset\r
+    curand_init(seed, sequenceNumber, offset, &ptrDevTabGeneratorThread[tid]);\r
+    }\r
+\r
+/*--------------------------------------*\\r
+ |*            Private                 *|\r
+ \*-------------------------------------*/\r
+\r
+#endif \r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    End                                             *|\r
+ \*---------------------------------------------------------------------*/\r