random: improve gegl-random implementation

- introduce a GeglRandom structure instead of accessing the LUT each time
- make a larger cycle for the seed
- avoid segfault when a negative seed is given
- use g(u)int64 instead of long to avoid plaform-dependant behavior
- make opencl and operations follow that api change
- build the GeglRandom structure in the gegl-chant machinery when using
  gegl_chant_seed
- make sure the pointer gegl_random_data is 32bits aligned when used with
  CL_MEM_USE_HOST_PTR

4 files changed, 146 insertions, 142 deletions

diff --git a/opencl/noise-hurl.cl b/opencl/noise-hurl.cl
index 6cc04122..ccb44be2 100644
--- a/opencl/noise-hurl.cl
+++ b/opencl/noise-hurl.cl
@@ -16,16 +16,15 @@
  * Copyright 2013 Carlos Zubieta <czubieta.dev@gmail.com>
  */
 
-__kernel void cl_noise_hurl(__global       float4 *src,
-                            __global const int    *random_data,
-                            __global const long   *random_primes,
-                                           int     x_offset,
-                                           int     y_offset,
-                                           int     roi_width,
-                                           int     whole_region_width,
-                                           int     seed,
-                                           float   pct_random,
-                                           int     offset)
+__kernel void cl_noise_hurl(__global       float4    *src,
+                            __global const int       *random_data,
+                                           int        x_offset,
+                                           int        y_offset,
+                                           int        roi_width,
+                                           int        whole_region_width,
+                                           GeglRandom rand,
+                                           float      pct_random,
+                                           int        offset)
 {
   int gid  = get_global_id(0);
   int gidy = gid / roi_width;
@@ -37,14 +36,14 @@ __kernel void cl_noise_hurl(__global       float4 *src,
 
   float4 src_v = src[gid];
 
-  float pc          = gegl_cl_random_float_range (random_data, random_primes,
-                                                  seed, x, y, 0, n, 0, 100);
-  float red_noise   = gegl_cl_random_float (random_data, random_primes,
-                                            seed, x, y, 0, n+1);
-  float green_noise = gegl_cl_random_float (random_data, random_primes,
-                                            seed, x, y, 0, n+2);
-  float blue_noise  = gegl_cl_random_float (random_data, random_primes,
-                                            seed, x, y, 0, n+3);
+  float pc          = gegl_cl_random_float_range (random_data,
+                                                  rand, x, y, 0, n, 0, 100);
+  float red_noise   = gegl_cl_random_float (random_data,
+                                            rand, x, y, 0, n+1);
+  float green_noise = gegl_cl_random_float (random_data,
+                                            rand, x, y, 0, n+2);
+  float blue_noise  = gegl_cl_random_float (random_data,
+                                            rand, x, y, 0, n+3);
 
   if(pc <= pct_random)
     {
diff --git a/opencl/noise-hurl.cl.h b/opencl/noise-hurl.cl.h
index 9cc9f78a..b57b32bc 100644
--- a/opencl/noise-hurl.cl.h
+++ b/opencl/noise-hurl.cl.h
@@ -17,16 +17,15 @@ static const char* noise_hurl_cl_source =
 " * Copyright 2013 Carlos Zubieta <czubieta.dev@gmail.com>                     \n"
 " */                                                                           \n"
 "                                                                              \n"
-"__kernel void cl_noise_hurl(__global       float4 *src,                       \n"
-"                            __global const int    *random_data,               \n"
-"                            __global const long   *random_primes,             \n"
-"                                           int     x_offset,                  \n"
-"                                           int     y_offset,                  \n"
-"                                           int     roi_width,                 \n"
-"                                           int     whole_region_width,        \n"
-"                                           int     seed,                      \n"
-"                                           float   pct_random,                \n"
-"                                           int     offset)                    \n"
+"__kernel void cl_noise_hurl(__global       float4    *src,                    \n"
+"                            __global const int       *random_data,            \n"
+"                                           int        x_offset,               \n"
+"                                           int        y_offset,               \n"
+"                                           int        roi_width,              \n"
+"                                           int        whole_region_width,     \n"
+"                                           GeglRandom rand,                   \n"
+"                                           float      pct_random,             \n"
+"                                           int        offset)                 \n"
 "{                                                                             \n"
 "  int gid  = get_global_id(0);                                                \n"
 "  int gidy = gid / roi_width;                                                 \n"
@@ -38,14 +37,14 @@ static const char* noise_hurl_cl_source =
 "                                                                              \n"
 "  float4 src_v = src[gid];                                                    \n"
 "                                                                              \n"
-"  float pc          = gegl_cl_random_float_range (random_data, random_primes, \n"
-"                                                  seed, x, y, 0, n, 0, 100);  \n"
-"  float red_noise   = gegl_cl_random_float (random_data, random_primes,       \n"
-"                                            seed, x, y, 0, n+1);              \n"
-"  float green_noise = gegl_cl_random_float (random_data, random_primes,       \n"
-"                                            seed, x, y, 0, n+2);              \n"
-"  float blue_noise  = gegl_cl_random_float (random_data, random_primes,       \n"
-"                                            seed, x, y, 0, n+3);              \n"
+"  float pc          = gegl_cl_random_float_range (random_data,                \n"
+"                                                  rand, x, y, 0, n, 0, 100);  \n"
+"  float red_noise   = gegl_cl_random_float (random_data,                      \n"
+"                                            rand, x, y, 0, n+1);              \n"
+"  float green_noise = gegl_cl_random_float (random_data,                      \n"
+"                                            rand, x, y, 0, n+2);              \n"
+"  float blue_noise  = gegl_cl_random_float (random_data,                      \n"
+"                                            rand, x, y, 0, n+3);              \n"
 "                                                                              \n"
 "  if(pc <= pct_random)                                                        \n"
 "    {                                                                         \n"
diff --git a/opencl/random.cl b/opencl/random.cl
index e6e33405..b900b379 100644
--- a/opencl/random.cl
+++ b/opencl/random.cl
@@ -17,77 +17,85 @@
  */
 
 /* XXX: this file should be kept in sync with gegl-random. */
-#define XPRIME     103423
-#define YPRIME     101359
-#define NPRIME     101111
+__constant const long XPRIME = 103423;
+__constant const long YPRIME = 101359;
+__constant const long NPRIME = 101111;
 
-#define RANDOM_DATA_SIZE (15083+15091+15101)
-#define PRIME_SIZE 533u
+#define RANDOM_DATA_SIZE (15101 * 3)
 
-inline uint _gegl_cl_random_int  (__global const int  *cl_random_data,
-                                  __global const long *cl_random_primes,
-                                  int seed, int x, int y, int z, int n);
+typedef ushort4 GeglRandom;
 
-uint gegl_cl_random_int          (__global const int  *cl_random_data,
-                                  __global const long *cl_random_primes,
-                                  int seed, int x, int y, int z, int n);
+unsigned int _gegl_cl_random_int (__global const int *cl_random_data,
+                                  const GeglRandom  rand,
+                                  int x,
+                                  int y,
+                                  int z,
+                                  int n);
+
+unsigned int gegl_cl_random_int  (__global const int *cl_random_data,
+                                  const GeglRandom  rand,
+                                  int x,
+                                  int y,
+                                  int z,
+                                  int n);
 
 int gegl_cl_random_int_range     (__global const int  *cl_random_data,
-                                  __global const long *cl_random_primes,
-                                  int seed, int x, int y, int z, int n,
-                                  int min, int max);
+                                  const GeglRandom  rand,
+                                  int x,
+                                  int y,
+                                  int z,
+                                  int n,
+                                  int min,
+                                  int max);
 
 float gegl_cl_random_float       (__global const int  *cl_random_data,
-                                  __global const long *cl_random_primes,
-                                  int seed, int x, int y, int z, int n);
+                                  const GeglRandom  rand,
+                                  int x,
+                                  int y,
+                                  int z,
+                                  int n);
 
 float gegl_cl_random_float_range (__global const int  *cl_random_data,
-                                  __global const long *cl_random_primes,
-                                  int seed, int x, int y, int z, int n,
-                                  float min, float max);
+                                  const GeglRandom  rand,
+                                  int x,
+                                  int y,
+                                  int z,
+                                  int n,
+                                  float min,
+                                  float max);
 
-inline uint
+unsigned int
 _gegl_cl_random_int (__global const int  *cl_random_data,
-                     __global const long *cl_random_primes,
-                     int                 seed,
+                     const GeglRandom    rand,
                      int                 x,
                      int                 y,
                      int                 z,
                      int                 n)
 {
-  unsigned long idx = x * XPRIME + 
-                      y * YPRIME * XPRIME + 
+  unsigned long idx = x * XPRIME +
+                      y * YPRIME * XPRIME +
                       n * NPRIME * YPRIME * XPRIME;
-#define ROUNDS 3
-    /* 3 rounds gives a reasonably high cycle for */
-    /*   our synthesized larger random set. */
-  unsigned long seed_idx = seed % (PRIME_SIZE - 1 - ROUNDS);
-  int prime0 = cl_random_primes[seed_idx],
-      prime1 = cl_random_primes[seed_idx+1],
-      prime2 = cl_random_primes[seed_idx+2];
-  int r0 = cl_random_data[idx % prime0],
-      r1 = cl_random_data[prime0 + (idx % (prime1))],
-      r2 = cl_random_data[prime0 + prime1 + (idx % (prime2))];
+
+  int r0 = cl_random_data[idx % rand.x],
+      r1 = cl_random_data[rand.x + (idx % rand.y)],
+      r2 = cl_random_data[rand.x + rand.y + (idx % rand.z)];
   return r0 ^ r1 ^ r2;
 }
 
-uint
-gegl_cl_random_int (__global const int  *cl_random_data,
-                    __global const long *cl_random_primes,
-                    int                 seed,
+unsigned int
+gegl_cl_random_int (__global const int *cl_random_data,
+                    const GeglRandom    rand,
                     int                 x,
                     int                 y,
                     int                 z,
                     int                 n)
 {
-  return _gegl_cl_random_int (cl_random_data, cl_random_primes,
-                              seed, x, y, z, n);
+  return _gegl_cl_random_int (cl_random_data, rand, x, y, z, n);
 }
 
 int
 gegl_cl_random_int_range (__global const int  *cl_random_data,
-                          __global const long *cl_random_primes,
-                          int                 seed,
+                          const GeglRandom    rand,
                           int                 x,
                           int                 y,
                           int                 z,
@@ -95,8 +103,7 @@ gegl_cl_random_int_range (__global const int  *cl_random_data,
                           int                 min,
                           int                 max)
 {
-  uint ret = _gegl_cl_random_int (cl_random_data, cl_random_primes,
-                                  seed, x, y, z, n);
+  int ret = _gegl_cl_random_int (cl_random_data, rand, x, y, z, n);
   return (ret % (max-min)) + min;
 }
 
@@ -104,23 +111,20 @@ gegl_cl_random_int_range (__global const int  *cl_random_data,
 #define G_RAND_FLOAT_TRANSFORM  0.00001525902189669642175f
 
 float
-gegl_cl_random_float (__global const int  *cl_random_data,
-                      __global const long *cl_random_primes,
-                      int                 seed,
+gegl_cl_random_float (__global const int *cl_random_data,
+                      const GeglRandom    rand,
                       int                 x,
                       int                 y,
                       int                 z,
                       int                 n)
 {
-  uint u = _gegl_cl_random_int (cl_random_data, cl_random_primes,
-                                seed, x, y, z, n);
+  int u = _gegl_cl_random_int (cl_random_data, rand, x, y, z, n);
   return (u & 0xffff) * G_RAND_FLOAT_TRANSFORM;
 }
 
 float
-gegl_cl_random_float_range (__global const int  *cl_random_data,
-                            __global const long *cl_random_primes,
-                            int                 seed,
+gegl_cl_random_float_range (__global const int *cl_random_data,
+                            const GeglRandom    rand,
                             int                 x,
                             int                 y,
                             int                 z,
@@ -128,7 +132,6 @@ gegl_cl_random_float_range (__global const int  *cl_random_data,
                             float               min,
                             float               max)
 {
-  float f = gegl_cl_random_float (cl_random_data, cl_random_primes,
-                                  seed, x, y, z, n);
+  float f = gegl_cl_random_float (cl_random_data, rand, x, y, z, n);
   return f * (max - min) + min;
 }
diff --git a/opencl/random.cl.h b/opencl/random.cl.h
index 6b789514..d48e087e 100644
--- a/opencl/random.cl.h
+++ b/opencl/random.cl.h
@@ -18,39 +18,56 @@ static const char* random_cl_source =
 " */                                                                           \n"
 "                                                                              \n"
 "/* XXX: this file should be kept in sync with gegl-random. */                 \n"
-"#define XPRIME     103423                                                     \n"
-"#define YPRIME     101359                                                     \n"
-"#define NPRIME     101111                                                     \n"
+"__constant const long XPRIME = 103423;                                        \n"
+"__constant const long YPRIME = 101359;                                        \n"
+"__constant const long NPRIME = 101111;                                        \n"
 "                                                                              \n"
-"#define RANDOM_DATA_SIZE (15083+15091+15101)                                  \n"
-"#define PRIME_SIZE 533u                                                       \n"
+"#define RANDOM_DATA_SIZE (15101 * 3)                                          \n"
 "                                                                              \n"
-"inline uint _gegl_cl_random_int  (__global const int  *cl_random_data,        \n"
-"                                  __global const long *cl_random_primes,      \n"
-"                                  int seed, int x, int y, int z, int n);      \n"
+"typedef ushort4 GeglRandom;                                                   \n"
 "                                                                              \n"
-"uint gegl_cl_random_int          (__global const int  *cl_random_data,        \n"
-"                                  __global const long *cl_random_primes,      \n"
-"                                  int seed, int x, int y, int z, int n);      \n"
+"unsigned int _gegl_cl_random_int (__global const int *cl_random_data,         \n"
+"                                  const GeglRandom  rand,                     \n"
+"                                  int x,                                      \n"
+"                                  int y,                                      \n"
+"                                  int z,                                      \n"
+"                                  int n);                                     \n"
+"                                                                              \n"
+"unsigned int gegl_cl_random_int  (__global const int *cl_random_data,         \n"
+"                                  const GeglRandom  rand,                     \n"
+"                                  int x,                                      \n"
+"                                  int y,                                      \n"
+"                                  int z,                                      \n"
+"                                  int n);                                     \n"
 "                                                                              \n"
 "int gegl_cl_random_int_range     (__global const int  *cl_random_data,        \n"
-"                                  __global const long *cl_random_primes,      \n"
-"                                  int seed, int x, int y, int z, int n,       \n"
-"                                  int min, int max);                          \n"
+"                                  const GeglRandom  rand,                     \n"
+"                                  int x,                                      \n"
+"                                  int y,                                      \n"
+"                                  int z,                                      \n"
+"                                  int n,                                      \n"
+"                                  int min,                                    \n"
+"                                  int max);                                   \n"
 "                                                                              \n"
 "float gegl_cl_random_float       (__global const int  *cl_random_data,        \n"
-"                                  __global const long *cl_random_primes,      \n"
-"                                  int seed, int x, int y, int z, int n);      \n"
+"                                  const GeglRandom  rand,                     \n"
+"                                  int x,                                      \n"
+"                                  int y,                                      \n"
+"                                  int z,                                      \n"
+"                                  int n);                                     \n"
 "                                                                              \n"
 "float gegl_cl_random_float_range (__global const int  *cl_random_data,        \n"
-"                                  __global const long *cl_random_primes,      \n"
-"                                  int seed, int x, int y, int z, int n,       \n"
-"                                  float min, float max);                      \n"
+"                                  const GeglRandom  rand,                     \n"
+"                                  int x,                                      \n"
+"                                  int y,                                      \n"
+"                                  int z,                                      \n"
+"                                  int n,                                      \n"
+"                                  float min,                                  \n"
+"                                  float max);                                 \n"
 "                                                                              \n"
-"inline uint                                                                   \n"
+"unsigned int                                                                  \n"
 "_gegl_cl_random_int (__global const int  *cl_random_data,                     \n"
-"                     __global const long *cl_random_primes,                   \n"
-"                     int                 seed,                                \n"
+"                     const GeglRandom    rand,                                \n"
 "                     int                 x,                                   \n"
 "                     int                 y,                                   \n"
 "                     int                 z,                                   \n"
@@ -59,36 +76,27 @@ static const char* random_cl_source =
 "  unsigned long idx = x * XPRIME +                                            \n"
 "                      y * YPRIME * XPRIME +                                   \n"
 "                      n * NPRIME * YPRIME * XPRIME;                           \n"
-"#define ROUNDS 3                                                              \n"
-"    /* 3 rounds gives a reasonably high cycle for */                          \n"
-"    /*   our synthesized larger random set. */                                \n"
-"  unsigned long seed_idx = seed % (PRIME_SIZE - 1 - ROUNDS);                  \n"
-"  int prime0 = cl_random_primes[seed_idx],                                    \n"
-"      prime1 = cl_random_primes[seed_idx+1],                                  \n"
-"      prime2 = cl_random_primes[seed_idx+2];                                  \n"
-"  int r0 = cl_random_data[idx % prime0],                                      \n"
-"      r1 = cl_random_data[prime0 + (idx % (prime1))],                         \n"
-"      r2 = cl_random_data[prime0 + prime1 + (idx % (prime2))];                \n"
+"                                                                              \n"
+"  int r0 = cl_random_data[idx % rand.x],                                      \n"
+"      r1 = cl_random_data[rand.x + (idx % rand.y)],                           \n"
+"      r2 = cl_random_data[rand.x + rand.y + (idx % rand.z)];                  \n"
 "  return r0 ^ r1 ^ r2;                                                        \n"
 "}                                                                             \n"
 "                                                                              \n"
-"uint                                                                          \n"
-"gegl_cl_random_int (__global const int  *cl_random_data,                      \n"
-"                    __global const long *cl_random_primes,                    \n"
-"                    int                 seed,                                 \n"
+"unsigned int                                                                  \n"
+"gegl_cl_random_int (__global const int *cl_random_data,                       \n"
+"                    const GeglRandom    rand,                                 \n"
 "                    int                 x,                                    \n"
 "                    int                 y,                                    \n"
 "                    int                 z,                                    \n"
 "                    int                 n)                                    \n"
 "{                                                                             \n"
-"  return _gegl_cl_random_int (cl_random_data, cl_random_primes,               \n"
-"                              seed, x, y, z, n);                              \n"
+"  return _gegl_cl_random_int (cl_random_data, rand, x, y, z, n);              \n"
 "}                                                                             \n"
 "                                                                              \n"
 "int                                                                           \n"
 "gegl_cl_random_int_range (__global const int  *cl_random_data,                \n"
-"                          __global const long *cl_random_primes,              \n"
-"                          int                 seed,                           \n"
+"                          const GeglRandom    rand,                           \n"
 "                          int                 x,                              \n"
 "                          int                 y,                              \n"
 "                          int                 z,                              \n"
@@ -96,8 +104,7 @@ static const char* random_cl_source =
 "                          int                 min,                            \n"
 "                          int                 max)                            \n"
 "{                                                                             \n"
-"  uint ret = _gegl_cl_random_int (cl_random_data, cl_random_primes,           \n"
-"                                  seed, x, y, z, n);                          \n"
+"  int ret = _gegl_cl_random_int (cl_random_data, rand, x, y, z, n);           \n"
 "  return (ret % (max-min)) + min;                                             \n"
 "}                                                                             \n"
 "                                                                              \n"
@@ -105,23 +112,20 @@ static const char* random_cl_source =
 "#define G_RAND_FLOAT_TRANSFORM  0.00001525902189669642175f                    \n"
 "                                                                              \n"
 "float                                                                         \n"
-"gegl_cl_random_float (__global const int  *cl_random_data,                    \n"
-"                      __global const long *cl_random_primes,                  \n"
-"                      int                 seed,                               \n"
+"gegl_cl_random_float (__global const int *cl_random_data,                     \n"
+"                      const GeglRandom    rand,                               \n"
 "                      int                 x,                                  \n"
 "                      int                 y,                                  \n"
 "                      int                 z,                                  \n"
 "                      int                 n)                                  \n"
 "{                                                                             \n"
-"  uint u = _gegl_cl_random_int (cl_random_data, cl_random_primes,             \n"
-"                                seed, x, y, z, n);                            \n"
+"  int u = _gegl_cl_random_int (cl_random_data, rand, x, y, z, n);             \n"
 "  return (u & 0xffff) * G_RAND_FLOAT_TRANSFORM;                               \n"
 "}                                                                             \n"
 "                                                                              \n"
 "float                                                                         \n"
-"gegl_cl_random_float_range (__global const int  *cl_random_data,              \n"
-"                            __global const long *cl_random_primes,            \n"
-"                            int                 seed,                         \n"
+"gegl_cl_random_float_range (__global const int *cl_random_data,               \n"
+"                            const GeglRandom    rand,                         \n"
 "                            int                 x,                            \n"
 "                            int                 y,                            \n"
 "                            int                 z,                            \n"
@@ -129,8 +133,7 @@ static const char* random_cl_source =
 "                            float               min,                          \n"
 "                            float               max)                          \n"
 "{                                                                             \n"
-"  float f = gegl_cl_random_float (cl_random_data, cl_random_primes,           \n"
-"                                  seed, x, y, z, n);                          \n"
+"  float f = gegl_cl_random_float (cl_random_data, rand, x, y, z, n);          \n"
 "  return f * (max - min) + min;                                               \n"
 "}                                                                             \n"
 ;