Fluctuating (Randomized, Thermalized) LBM
- get random number from external field, add kernel to test suite @winterhalter
- physical test case to check random number correlation (auto correlation function?) @kuron
- final version should compute random numbers in the kernel itself. 'mulhi' / 'mullo' functions are required - philox RNG - write a pystencils function similar to bitoperations for these operations - on x86 there are intrinsics for mulhi/mullo, alternatively one could also use AES instead of philox since there is also a intrinsic for that - on GPU there should also be an intrinsic - C fallback as one-liner - two versions required for drawing float & double random numbers - fast versions produce 2 or more RN per step, how to handle this?
Link issues together to show that they're related. Learn more.
When these merge requests are accepted, this issue will be closed automatically.
Activity
- Maintainer
Counter-based RNG: Philox from https://doi.org/10.1145/2063384.2063405
- Download from https://www.deshawresearch.com/downloads/download_random123.cgi/Random123-1.09.tar.gz
-
std::random
interface at https://github.com/DEShawResearch/Random123-Boost - Implemented in cuRAND: https://docs.nvidia.com/cuda/curand
- Maintainer
For CUDA, see /usr/local/cuda-9.0/include/curand_philox4x32_x.h:
#define QUALIFIERS static __forceinline__ __device__ #define PHILOX_W32_0 (0x9E3779B9) #define PHILOX_W32_1 (0xBB67AE85) #define PHILOX_M4x32_0 (0xD2511F53) #define PHILOX_M4x32_1 (0xCD9E8D57) QUALIFIERS uint32 mulhilo32(uint32 a, uint32 b, uint32* hip) { #ifndef __CUDA_ARCH__ // host code uint64 product = ((uint64)a) * ((uint64)b); *hip = product >> 32; return (uint32)product; #else // device code *hip = __umulhi(a,b); return a*b; #endif } QUALIFIERS uint32* _philox4x32round(uint32* ctr, uint32* key) { uint32 hi0; uint32 hi1; uint32 lo0 = mulhilo32(PHILOX_M4x32_0, ctr[0], &hi0); uint32 lo1 = mulhilo32(PHILOX_M4x32_1, ctr[2], &hi1); ctr[0] = hi1^ctr[1]^key[0]; ctr[1] = lo1; ctr[2] = hi0^ctr[3]^key[1]; ctr[3] = lo0; return ctr; } QUALIFIERS uint32* _philox4x32bumpkey( uint32* key) { key[0] += PHILOX_W32_0; key[1] += PHILOX_W32_1; return key; } QUALIFIERS void curand_Philox4x32_10(uint32x4 ctr, uint32x2 key, uint32* out) { ctr = _philox4x32round(ctr, key); // 1 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 2 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 3 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 4 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 5 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 6 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 7 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 8 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 9 key = _philox4x32bumpkey(key); ctr = _philox4x32round(ctr, key); // 10 for(int i=0; i<4; i++) out[i] = ctr[i]; return; }
Edited by Michael Kuron - Author Reporter
Test for RNG quality https://arxiv.org/abs/0906.2425
- Michael Kuron assigned to @bauer and unassigned @winterhalter
assigned to @bauer and unassigned @winterhalter
- Maintainer
Would be nice for Walberla 4.0
- Maintainer
Random123 is fully templated, so it might be possible to just add operator overloads for the SIMD-packed datatypes to get vectorization.
- Michael Kuron mentioned in issue walberla/walberla#80 (closed)
mentioned in issue walberla/walberla#80 (closed)
- Michael Kuron changed the description
changed the description
- Developer
First WIP draft finally from my end in software/pystencils!37
- Martin Bauer marked the checklist item final version should compute random numbers in the kernel itself. 'mulhi' / 'mullo' functions are required as completed
marked the checklist item final version should compute random numbers in the kernel itself. 'mulhi' / 'mullo' functions are required as completed
- Michael Kuron mentioned in merge request !1 (closed)
mentioned in merge request !1 (closed)
- Maintainer
Missing pieces:
- seeding the RNG (make the second key settable from the constructor)
- correct passthrough of AST objects everywhere. The assignment handling works with the changes in the test_martin branch, but CSE does not.
- Michael Kuron mentioned in merge request walberla/walberla!220 (merged)
mentioned in merge request walberla/walberla!220 (merged)
- Martin Bauer marked the checklist item get random number from external field, add kernel to test suite @winterhalter as completed
marked the checklist item get random number from external field, add kernel to test suite @winterhalter as completed
- Martin Bauer marked the checklist item get random number from external field, add kernel to test suite @winterhalter as incomplete
marked the checklist item get random number from external field, add kernel to test suite @winterhalter as incomplete
- Martin Bauer marked the checklist item physical test case to check random number correlation (auto correlation function?) @kuron as completed
marked the checklist item physical test case to check random number correlation (auto correlation function?) @kuron as completed
- Martin Bauer marked the checklist item physical test case to check random number correlation (auto correlation function?) @kuron as incomplete
marked the checklist item physical test case to check random number correlation (auto correlation function?) @kuron as incomplete
- Felix Winterhalter mentioned in merge request !2 (closed)
mentioned in merge request !2 (closed)
- Michael Kuron mentioned in merge request lbmpy_walberla!1 (merged)
mentioned in merge request lbmpy_walberla!1 (merged)
- Michael Kuron mentioned in merge request !3 (merged)
mentioned in merge request !3 (merged)
- Martin Bauer closed via merge request !3 (merged)
closed via merge request !3 (merged)
- Martin Bauer mentioned in commit 1c4dfcf8
mentioned in commit 1c4dfcf8