diff --git a/build_linux64/libamsculib3.linux64.a b/build_linux64/libamsculib3.linux64.a index 9087993..17b0407 100644 Binary files a/build_linux64/libamsculib3.linux64.a and b/build_linux64/libamsculib3.linux64.a differ diff --git a/build_linux64/objstore/amscu_comp128.o b/build_linux64/objstore/amscu_comp128.o index bf17260..d7d4407 100644 Binary files a/build_linux64/objstore/amscu_comp128.o and b/build_linux64/objstore/amscu_comp128.o differ diff --git a/build_linux64/objstore/amscu_comp64.o b/build_linux64/objstore/amscu_comp64.o index 1459322..4e0ef13 100644 Binary files a/build_linux64/objstore/amscu_comp64.o and b/build_linux64/objstore/amscu_comp64.o differ diff --git a/build_linux64/objstore/amscu_cudafunctions.o b/build_linux64/objstore/amscu_cudafunctions.o index 6f6e6af..f996fbb 100644 Binary files a/build_linux64/objstore/amscu_cudafunctions.o and b/build_linux64/objstore/amscu_cudafunctions.o differ diff --git a/build_linux64/objstore/amscu_random.o b/build_linux64/objstore/amscu_random.o deleted file mode 100644 index 975a066..0000000 Binary files a/build_linux64/objstore/amscu_random.o and /dev/null differ diff --git a/build_linux64/objstore/amscuarray.o b/build_linux64/objstore/amscuarray.o index 5a1469c..c0177d3 100644 Binary files a/build_linux64/objstore/amscuarray.o and b/build_linux64/objstore/amscuarray.o differ diff --git a/build_linux64/objstore/amscuarray_dops.o b/build_linux64/objstore/amscuarray_dops.o index 84dd042..55fc2b2 100644 Binary files a/build_linux64/objstore/amscuarray_dops.o and b/build_linux64/objstore/amscuarray_dops.o differ diff --git a/build_linux64/objstore/amscufhash.o b/build_linux64/objstore/amscufhash.o new file mode 100644 index 0000000..1d4827a Binary files /dev/null and b/build_linux64/objstore/amscufhash.o differ diff --git a/build_linux64/objstore/amscugeom.o b/build_linux64/objstore/amscugeom.o index 910a439..69e1bc2 100644 Binary files a/build_linux64/objstore/amscugeom.o and b/build_linux64/objstore/amscugeom.o differ diff --git a/build_linux64/objstore/amsculib3.o b/build_linux64/objstore/amsculib3.o index 4118168..adc331e 100644 Binary files a/build_linux64/objstore/amsculib3.o and b/build_linux64/objstore/amsculib3.o differ diff --git a/build_linux64/objstore/amscumath.o b/build_linux64/objstore/amscumath.o index cd9e9ea..96cb9dd 100644 Binary files a/build_linux64/objstore/amscumath.o and b/build_linux64/objstore/amscumath.o differ diff --git a/build_linux64/objstore/amscupcg.o b/build_linux64/objstore/amscupcg.o index 1b39e49..1bfa65a 100644 Binary files a/build_linux64/objstore/amscupcg.o and b/build_linux64/objstore/amscupcg.o differ diff --git a/build_linux64/objstore/amscurandlcg.o b/build_linux64/objstore/amscurandlcg.o new file mode 100644 index 0000000..db1dbaf Binary files /dev/null and b/build_linux64/objstore/amscurandlcg.o differ diff --git a/build_linux64/objstore/amscurandom1.o b/build_linux64/objstore/amscurandom1.o index 3da73e8..6f57ad2 100644 Binary files a/build_linux64/objstore/amscurandom1.o and b/build_linux64/objstore/amscurandom1.o differ diff --git a/build_linux64/objstore/amscurandom1_dbuff.o b/build_linux64/objstore/amscurandom1_dbuff.o index 1e4d71c..5291cba 100644 Binary files a/build_linux64/objstore/amscurandom1_dbuff.o and b/build_linux64/objstore/amscurandom1_dbuff.o differ diff --git a/build_linux64/objstore/amscurandom1_hbuff.o b/build_linux64/objstore/amscurandom1_hbuff.o index aea0e8a..4f03534 100644 Binary files a/build_linux64/objstore/amscurandom1_hbuff.o and b/build_linux64/objstore/amscurandom1_hbuff.o differ diff --git a/build_linux64/objstore/amscurandom_tests1.o b/build_linux64/objstore/amscurandom_tests1.o new file mode 100644 index 0000000..1c8bb86 Binary files /dev/null and b/build_linux64/objstore/amscurandom_tests1.o differ diff --git a/build_linux64/objstore/amscurarray.o b/build_linux64/objstore/amscurarray.o index 2f2abd8..6df5b4d 100644 Binary files a/build_linux64/objstore/amscurarray.o and b/build_linux64/objstore/amscurarray.o differ diff --git a/build_linux64/objstore/amscusplitmix.o b/build_linux64/objstore/amscusplitmix.o index ed3492e..9df9f59 100644 Binary files a/build_linux64/objstore/amscusplitmix.o and b/build_linux64/objstore/amscusplitmix.o differ diff --git a/build_linux64/objstore/amsxoroshiro.o b/build_linux64/objstore/amsxoroshiro.o index a9e51d3..2c7743c 100644 Binary files a/build_linux64/objstore/amsxoroshiro.o and b/build_linux64/objstore/amsxoroshiro.o differ diff --git a/build_linux64/objstore/cuvec2.o b/build_linux64/objstore/cuvec2.o index 33d3deb..ed52394 100644 Binary files a/build_linux64/objstore/cuvec2.o and b/build_linux64/objstore/cuvec2.o differ diff --git a/build_linux64/objstore/cuvec2f.o b/build_linux64/objstore/cuvec2f.o index 1b01a7a..22ae19e 100644 Binary files a/build_linux64/objstore/cuvec2f.o and b/build_linux64/objstore/cuvec2f.o differ diff --git a/build_linux64/objstore/cuvec2i.o b/build_linux64/objstore/cuvec2i.o index ce98493..59929a8 100644 Binary files a/build_linux64/objstore/cuvec2i.o and b/build_linux64/objstore/cuvec2i.o differ diff --git a/build_linux64/objstore/cuvec3.o b/build_linux64/objstore/cuvec3.o index 4e37970..d323ec4 100644 Binary files a/build_linux64/objstore/cuvec3.o and b/build_linux64/objstore/cuvec3.o differ diff --git a/build_linux64/objstore/cuvec3f.o b/build_linux64/objstore/cuvec3f.o index bb676bd..cd4bef3 100644 Binary files a/build_linux64/objstore/cuvec3f.o and b/build_linux64/objstore/cuvec3f.o differ diff --git a/build_linux64/objstore/cuvec3i.o b/build_linux64/objstore/cuvec3i.o index 670ee6c..d337f37 100644 Binary files a/build_linux64/objstore/cuvec3i.o and b/build_linux64/objstore/cuvec3i.o differ diff --git a/build_linux64/objstore/cuvec4.o b/build_linux64/objstore/cuvec4.o index 860f897..16abc9b 100644 Binary files a/build_linux64/objstore/cuvec4.o and b/build_linux64/objstore/cuvec4.o differ diff --git a/build_linux64/objstore/cuvec4f.o b/build_linux64/objstore/cuvec4f.o index a5ac0b9..b810182 100644 Binary files a/build_linux64/objstore/cuvec4f.o and b/build_linux64/objstore/cuvec4f.o differ diff --git a/build_linux64/objstore/cuvec4i.o b/build_linux64/objstore/cuvec4i.o index 7ea34eb..d78340a 100644 Binary files a/build_linux64/objstore/cuvec4i.o and b/build_linux64/objstore/cuvec4i.o differ diff --git a/build_linux64/test b/build_linux64/test index d132678..ff34099 100644 Binary files a/build_linux64/test and b/build_linux64/test differ diff --git a/include/amsculib3/amscu_random.hpp b/include/amsculib3/amscu_random.hpp deleted file mode 100644 index e974e47..0000000 --- a/include/amsculib3/amscu_random.hpp +++ /dev/null @@ -1,55 +0,0 @@ -#ifndef __AMSCU_RANDOM_HPP__ -#define __AMSCU_RANDOM_HPP__ - -namespace amscuda -{ - -// Random Number Gerneators - - -// faster floating point hash function used in fractal generators -__device__ __host__ float fhash1d_su(float x); - -__device__ __host__ float fhash3d_su(float x, float y, float z); - -__device__ __host__ float fhash4d_su(float x, float y, float z, float w); - - -////////////////////////////////////////////////// -// Deterministic Pseudorandom int32_t Generator // -////////////////////////////////////////////////// - -//Next seed in simple 32 bit integer deterministic psuedo-rand generator -__host__ __device__ void dpr32_nextseed(int32_t *rseed_inout); - -//Simple 32 bit integer deterministic pseudo-random generator -// *not* for cryptography -// Frequency of generated floats should be uniform [0,1) -__host__ __device__ float dpr32_randf(int32_t *rseed_inout); - -//box muller standard normal pseudorandom variable -__host__ __device__ float dpr32_randnf(int32_t *rseed_inout); - -////////////////////////////////////////////////// -// Deterministic Pseudorandom int64_t Generator // -////////////////////////////////////////////////// - -//operates without side-effects on explicit seed for threaded use -//deterministic pseudorandom number generator - takes seed and returns next seed -__host__ __device__ void dpr64_nextseed(int64_t *seedinout); - -//deterministic pseudorandom number generator - takes seed and returns next seed -//returns uniformly distributed double -__host__ __device__ double dpr64_randd(int64_t *seedinout); - -__host__ __device__ float dpr64_randf(int64_t *seedinout); - - -void test_dprg64(); -void test_dprg32(); - - -}; //end namespace amscuda - -#endif - diff --git a/include/amsculib3/amscuarray_dops.hpp b/include/amsculib3/amscuarray_dops.hpp index eff18a6..be5af92 100644 --- a/include/amsculib3/amscuarray_dops.hpp +++ b/include/amsculib3/amscuarray_dops.hpp @@ -39,8 +39,6 @@ namespace amscuda //random device buffer functions void dbuff_rand_dpr32(float *devbuffer, int N, int32_t *rseedinout, int nblocks, int nthreads); // void dbuff_rand_dpr32n(float *devbuffer, int N, int32_t *rseedinout, int nblocks, int nthreads); // - - void dbuff_rand_dpr64(float *devbuffer, int N, int64_t *rseedinout, int nblocks, int nthreads); // //Elementwise device-buffer vector binary operation diff --git a/include/amsculib3/amsculib3.hpp b/include/amsculib3/amsculib3.hpp index a29f9eb..3d4d8c3 100644 --- a/include/amsculib3/amsculib3.hpp +++ b/include/amsculib3/amsculib3.hpp @@ -50,7 +50,6 @@ namespace amscuda #include #include -#include #include diff --git a/include/amsculib3/random/amscufhash.cuh b/include/amsculib3/random/amscufhash.cuh index 0f652aa..d771ce3 100644 --- a/include/amsculib3/random/amscufhash.cuh +++ b/include/amsculib3/random/amscufhash.cuh @@ -11,7 +11,12 @@ namespace fhash { //Floating point hash functions +// legacy fast floating point hash function used in fractal generators +__device__ __host__ float fhash1d_su(float x); +__device__ __host__ float fhash3d_su(float x, float y, float z); + +__device__ __host__ float fhash4d_su(float x, float y, float z, float w); }; diff --git a/include/amsculib3/random/amsculcg.cuh b/include/amsculib3/random/amsculcg.cuh index 9877415..14229d9 100644 --- a/include/amsculib3/random/amsculcg.cuh +++ b/include/amsculib3/random/amsculcg.cuh @@ -11,7 +11,38 @@ namespace lcg { //Legacy linear congruential generators +////////////////////////////////////////////////// +// Deterministic Pseudorandom int32_t Generator // +////////////////////////////////////////////////// +//Next seed in simple 32 bit integer deterministic psuedo-rand generator +__host__ __device__ void dpr32_nextseed(int32_t *rseed_inout); + +//Simple 32 bit integer deterministic pseudo-random generator +// *not* for cryptography +// Frequency of generated floats should be uniform [0,1) +__host__ __device__ float dpr32_randf(int32_t *rseed_inout); + +//box muller standard normal pseudorandom variable +__host__ __device__ float dpr32_randnf(int32_t *rseed_inout); + +////////////////////////////////////////////////// +// Deterministic Pseudorandom int64_t Generator // +////////////////////////////////////////////////// + +//operates without side-effects on explicit seed for threaded use +//deterministic pseudorandom number generator - takes seed and returns next seed +__host__ __device__ void dpr64_nextseed(int64_t *seedinout); + +//deterministic pseudorandom number generator - takes seed and returns next seed +//returns uniformly distributed double +__host__ __device__ double dpr64_randd(int64_t *seedinout); + +__host__ __device__ float dpr64_randf(int64_t *seedinout); + + +void test_dprg64(); +void test_dprg32(); }; }; diff --git a/include/amsculib3/random/amscurandom.cuh b/include/amsculib3/random/amscurandom.cuh index 157552e..68f52af 100644 --- a/include/amsculib3/random/amscurandom.cuh +++ b/include/amsculib3/random/amscurandom.cuh @@ -48,6 +48,10 @@ __host__ int dbuff_randn(double *hbuffer, int64_t size, randstate_t *state = NUL __host__ int dbuff_randint(int *hbuffer, int64_t size, int low, int high, randstate_t *state = NULL); +//Tests +__host__ void amscurand_tests1(); //test basic random functions + + }; }; diff --git a/src/amsculib3/amscuarray_dops.cu b/src/amsculib3/amscuarray_dops.cu index 024ff7d..cd12dc2 100644 --- a/src/amsculib3/amscuarray_dops.cu +++ b/src/amsculib3/amscuarray_dops.cu @@ -26,7 +26,7 @@ __global__ void dbuff_rand_dpr32_kf(float *devbuffer, int N, int32_t *seeds) lseed = seeds[I0]; for(I=I0;I + +namespace amscuda +{ +namespace random +{ +namespace fhash +{ + +__device__ __host__ float fhash1d_su(float x) +{ + float ret; + ret = x*(x>0.0f) + -x*(x<0.0f); //sign without conditionals? + ret = fmodf(ret,10000.0f); //restrain domain + ret = fmodf(ret*(ret+3678.453f)+7890.453f,10000.0f); + ret = fmodf(ret*(ret+8927.2134f),10000.0f); + ret = fmodf(ret*(ret+3656.234f),10000.0f); + //ret = fmodf(ret*(ret+892.2134f),1000.0f); + //ret = fmodf(ret,1000.0f); + ret = ret/10000.0f; + return ret; +} + + +__device__ __host__ float fhash3d_su(float x, float y=0.0f, float z=0.0f) +{ + float ret = 0.0f; + + ret = fhash1d_su(z); + ret = fhash1d_su(1000.0f*ret*ret + y); + ret = fhash1d_su(1000.0f*ret*ret + x); + + return ret; +} + +__device__ __host__ float fhash4d_su(float x, float y=0.0f, float z=0.0f, float w=0.0f) +{ + float ret = 0.0f; + + ret = fhash1d_su(w); + ret = fhash1d_su(1000.0f*ret*ret + z); + ret = fhash1d_su(1000.0f*ret*ret + y); + ret = fhash1d_su(1000.0f*ret*ret + x); + + return ret; +} + + +}; +}; +}; //end namespaces \ No newline at end of file diff --git a/src/amsculib3/amscu_random.cu b/src/amsculib3/random/amscurandlcg.cu similarity index 81% rename from src/amsculib3/amscu_random.cu rename to src/amsculib3/random/amscurandlcg.cu index 8fabb1f..0bb9ac8 100644 --- a/src/amsculib3/amscu_random.cu +++ b/src/amsculib3/random/amscurandlcg.cu @@ -2,44 +2,10 @@ namespace amscuda { - -__device__ __host__ float fhash1d_su(float x) +namespace random { - float ret; - ret = x*(x>0.0f) + -x*(x<0.0f); //sign without conditionals? - ret = fmodf(ret,10000.0f); //restrain domain - ret = fmodf(ret*(ret+3678.453f)+7890.453f,10000.0f); - ret = fmodf(ret*(ret+8927.2134f),10000.0f); - ret = fmodf(ret*(ret+3656.234f),10000.0f); - //ret = fmodf(ret*(ret+892.2134f),1000.0f); - //ret = fmodf(ret,1000.0f); - ret = ret/10000.0f; - return ret; -} - - -__device__ __host__ float fhash3d_su(float x, float y=0.0f, float z=0.0f) +namespace lcg { - float ret = 0.0f; - - ret = fhash1d_su(z); - ret = fhash1d_su(1000.0f*ret*ret + y); - ret = fhash1d_su(1000.0f*ret*ret + x); - - return ret; -} - -__device__ __host__ float fhash4d_su(float x, float y=0.0f, float z=0.0f, float w=0.0f) -{ - float ret = 0.0f; - - ret = fhash1d_su(w); - ret = fhash1d_su(1000.0f*ret*ret + z); - ret = fhash1d_su(1000.0f*ret*ret + y); - ret = fhash1d_su(1000.0f*ret*ret + x); - - return ret; -} ////////////////////////////////////////////////// // Deterministic Pseudorandom int32_t Generator // @@ -218,5 +184,6 @@ void test_dprg32() // return; } - -}; //namespace amscuda \ No newline at end of file +}; +}; +}; //end namespaces \ No newline at end of file diff --git a/src/amsculib3/random/amscurandom_tests1.cu b/src/amsculib3/random/amscurandom_tests1.cu new file mode 100644 index 0000000..2db04e2 --- /dev/null +++ b/src/amsculib3/random/amscurandom_tests1.cu @@ -0,0 +1,21 @@ +#include + +namespace amscuda +{ +namespace random +{ + + __host__ void amscurand_tests1() + { + using namespace random; + printf("Random number generator basic function test:\n"); + + + rand_seed(0); + + + return; + } + +}; //end namespaces +}; \ No newline at end of file diff --git a/src/main.cu b/src/main.cu index 948347a..1c6254e 100644 --- a/src/main.cu +++ b/src/main.cu @@ -23,5 +23,7 @@ int main(int argc, char* argv[]) //test_amscurarray1(); + random::amscurand_tests1(); + return 0; } \ No newline at end of file