-
Notifications
You must be signed in to change notification settings - Fork 4
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
* first code clean up * partition mtgp32 parameters into separate file * push mtgp32 parameters into separate files * Merge branch 'stashing initial edits * attempt build of threefry prng * file path fixes * clean up files * rename rng files * added kernels for threefry prng * integrated threefry prng * fixed compile * running threefry * clean threefry prng * added tester for threefry * cleaned threefry tester * complete threefry rng implementation and integration * update version
- Loading branch information
Showing
22 changed files
with
12,743 additions
and
11,648 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,87 @@ | ||
| /** | ||
| @file | ||
| Implements threefry RNG. | ||
| ******************************************************* | ||
| * Modified version of Random123 library: | ||
| * https://www.deshawresearch.com/downloads/download_random123.cgi/ | ||
| * The original copyright can be seen here: | ||
| * | ||
| * RANDOM123 LICENSE AGREEMENT | ||
| * | ||
| * Copyright 2010-2011, D. E. Shaw Research. All rights reserved. | ||
| * | ||
| * Redistribution and use in source and binary forms, with or without | ||
| * modification, are permitted provided that the following conditions | ||
| * are met: | ||
| * | ||
| * * Redistributions of source code must retain the above copyright notice, | ||
| * this list of conditions, and the following disclaimer. | ||
| * | ||
| * * Redistributions in binary form must reproduce the above copyright | ||
| * notice, this list of conditions, and the following disclaimer in the | ||
| * documentation and/or other materials provided with the distribution. | ||
| * | ||
| * Neither the name of D. E. Shaw Research nor the names of its contributors | ||
| * may be used to endorse or promote products derived from this software | ||
| * without specific prior written permission. | ||
| * | ||
| * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS | ||
| * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT | ||
| * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR | ||
| * A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT | ||
| * OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, | ||
| * SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED | ||
| * TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR | ||
| * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF | ||
| * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING | ||
| * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS | ||
| * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. | ||
| *********************************************************/ | ||
|
|
||
| /** | ||
| State of threefry RNG. We will store in global buffer as a set of uint | ||
| ** | ||
| counter: uint[4] | ||
| key: uint[4] | ||
| state: uint[4] | ||
| index: uint | ||
| typedef struct{ | ||
| uint counter[4]; | ||
| uint result[4]; | ||
| uint key[4]; | ||
| uint tracker; | ||
| } threefry_state; | ||
| **/ | ||
|
|
||
| /** | ||
| Seeds threefry RNG. | ||
| @param state Variable, that holds state of the generator to be seeded. | ||
| @param seed Value used for seeding. Should be randomly generated for each instance of generator (thread). | ||
| **/ | ||
| __kernel void | ||
| threefry_seed(__global uint __restrict *state_key, | ||
| __global uint __restrict *state_counter, | ||
| __global uint __restrict *state_result, | ||
| __global uint __restrict *state_tracker, | ||
| __global uint __restrict *seed) { | ||
| uint gid = get_global_id(0); | ||
| uint rng_count = get_global_size(0); | ||
| uint idx = gid; | ||
| uint localJ = seed[gid]; | ||
| state_key[idx] = localJ; | ||
| state_counter[idx] = 0x00000000; | ||
| state_result[idx] = 0x00000000; | ||
| state_tracker[idx] = 4; | ||
| idx += rng_count; | ||
| state_key[idx] = 0x00000000; | ||
| state_counter[idx] = 0x00000000; | ||
| state_result[idx] = 0x00000000; | ||
| idx += rng_count; | ||
| state_key[idx] = gid; | ||
| state_counter[idx] = 0x00000000; | ||
| state_result[idx] = 0x00000000; | ||
| idx += rng_count; | ||
| state_key[idx] = 0x00000000; | ||
| state_counter[idx] = 0x00000000; | ||
| state_result[idx] = 0x00000000; | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,221 @@ | ||
| /** | ||
| @file | ||
| Implements threefry RNG. | ||
| ******************************************************* | ||
| * Modified version of Random123 library: | ||
| * https://www.deshawresearch.com/downloads/download_random123.cgi/ | ||
| * The original copyright can be seen here: | ||
| * | ||
| * RANDOM123 LICENSE AGREEMENT | ||
| * | ||
| * Copyright 2010-2011, D. E. Shaw Research. All rights reserved. | ||
| * | ||
| * Redistribution and use in source and binary forms, with or without | ||
| * modification, are permitted provided that the following conditions | ||
| * are met: | ||
| * | ||
| * * Redistributions of source code must retain the above copyright notice, | ||
| * this list of conditions, and the following disclaimer. | ||
| * | ||
| * * Redistributions in binary form must reproduce the above copyright | ||
| * notice, this list of conditions, and the following disclaimer in the | ||
| * documentation and/or other materials provided with the distribution. | ||
| * | ||
| * Neither the name of D. E. Shaw Research nor the names of its contributors | ||
| * may be used to endorse or promote products derived from this software | ||
| * without specific prior written permission. | ||
| * | ||
| * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS | ||
| * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT | ||
| * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR | ||
| * A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT | ||
| * OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, | ||
| * SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED | ||
| * TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR | ||
| * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF | ||
| * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING | ||
| * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS | ||
| * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. | ||
| *********************************************************/ | ||
|
|
||
| /** | ||
| State of threefry RNG. We will store in global buffer as a set of uint | ||
| ** | ||
| typedef struct{ | ||
| uint counter[4]; | ||
| uint result[4]; | ||
| uint key[4]; | ||
| uint tracker; | ||
| } threefry_state; | ||
| **/ | ||
|
|
||
| /** | ||
| Generates a random 32-bit unsigned integer using threefry RNG. | ||
| @param state State of the RNG to use. | ||
| **/ | ||
| __kernel void | ||
| threefry_normal(__global uint __restrict *state_key, | ||
| __global uint __restrict *state_counter, | ||
| __global uint __restrict *state_result, | ||
| __global uint __restrict *state_tracker, | ||
| __global float __restrict *output, | ||
| int data_size) { | ||
| uint gid = get_global_id(0); | ||
| uint rng_count = get_global_size(0); | ||
| uint tmpIdx = gid; | ||
| threefry_state state_; | ||
| threefry_state *state = &state_; | ||
|
|
||
| // For first out of four sets... | ||
| // Read in counter | ||
| state->counter[0] = state_counter[tmpIdx]; | ||
| // Read in result | ||
| state->result[0] = state_result[tmpIdx]; | ||
| // Read in key | ||
| state->key[0] = state_key[tmpIdx]; | ||
| // Read in tracker | ||
| state->tracker = state_tracker[tmpIdx]; | ||
|
|
||
| // For second out of four sets... | ||
| tmpIdx += rng_count; | ||
| // Read in counter | ||
| state->counter[1] = state_counter[tmpIdx]; | ||
| // Read in result | ||
| state->result[1] = state_result[tmpIdx]; | ||
| // Read in key | ||
| state->key[1] = state_key[tmpIdx]; | ||
|
|
||
| // For third out of four sets... | ||
| tmpIdx += rng_count; | ||
| // Read in counter | ||
| state->counter[2] = state_counter[tmpIdx]; | ||
| // Read in result | ||
| state->result[2] = state_result[tmpIdx]; | ||
| // Read in key | ||
| state->key[2] = state_key[tmpIdx]; | ||
|
|
||
| // For last out of four sets... | ||
| tmpIdx += rng_count; | ||
| // Read in counter | ||
| state->counter[3] = state_counter[tmpIdx]; | ||
| // Read in result | ||
| state->result[3] = state_result[tmpIdx]; | ||
| // Read in key | ||
| state->key[3] = state_key[tmpIdx]; | ||
|
|
||
| for (uint outIndex = gid; outIndex < data_size / 2; outIndex += rng_count) { | ||
| uint num1[2]; | ||
| float res1[4]; | ||
| uint lidx = 0; | ||
| if (state->tracker > 3) { | ||
| threefry_round(state); | ||
| state->tracker = 1; | ||
| num1[lidx++] = state->result[0]; | ||
| } else if (state->tracker == 3) { | ||
| uint tmp = state->result[3]; | ||
| if (++state->counter[0] == 0) { | ||
| if (++state->counter[1] == 0) { | ||
| if (++state->counter[2] == 0) { | ||
| ++state->counter[3]; | ||
| } | ||
| } | ||
| } | ||
| threefry_round(state); | ||
| state->tracker = 0; | ||
| num1[lidx++] = tmp; | ||
| } else { | ||
| num1[lidx++] = state->result[state->tracker++]; | ||
| } | ||
| if (state->tracker == 3) { | ||
| uint tmp = state->result[3]; | ||
| if (++state->counter[0] == 0) { | ||
| if (++state->counter[1] == 0) { | ||
| if (++state->counter[2] == 0) { | ||
| ++state->counter[3]; | ||
| } | ||
| } | ||
| } | ||
| threefry_round(state); | ||
| state->tracker = 0; | ||
| num1[lidx] = tmp; | ||
| } else { | ||
| num1[lidx] = state->result[state->tracker++]; | ||
| } | ||
| res1[0] = uint2float(num1[0], num1[1]); | ||
| lidx = 0; | ||
| if (state->tracker == 3) { | ||
| uint tmp = state->result[3]; | ||
| if (++state->counter[0] == 0) { | ||
| if (++state->counter[1] == 0) { | ||
| if (++state->counter[2] == 0) { | ||
| ++state->counter[3]; | ||
| } | ||
| } | ||
| } | ||
| threefry_round(state); | ||
| state->tracker = 0; | ||
| num1[lidx++] = tmp; | ||
| } else { | ||
| num1[lidx++] = state->result[state->tracker++]; | ||
| } | ||
| if (state->tracker == 3) { | ||
| uint tmp = state->result[3]; | ||
| if (++state->counter[0] == 0) { | ||
| if (++state->counter[1] == 0) { | ||
| if (++state->counter[2] == 0) { | ||
| ++state->counter[3]; | ||
| } | ||
| } | ||
| } | ||
| threefry_round(state); | ||
| state->tracker = 0; | ||
| num1[lidx] = tmp; | ||
| } else { | ||
| num1[lidx] = state->result[state->tracker++]; | ||
| } | ||
| res1[1] = uint2float(num1[0], num1[1]); | ||
| res1[2] = sqrt( -2.0f * log(res1[0])) * cospi(2.0f * res1[1]); | ||
| res1[3] = sqrt( -2.0f * log(res1[0])) * sinpi(2.0f * res1[1]); | ||
| output[outIndex] = res1[2]; | ||
| output[outIndex + (data_size/2)] = res1[3]; | ||
| } | ||
|
|
||
| // For first out of four sets... | ||
| // Write out counter | ||
| tmpIdx = gid; | ||
| state_counter[tmpIdx] = state->counter[0]; | ||
| // Write out result | ||
| state_result[tmpIdx] = state->result[0]; | ||
| // Write out key | ||
| state_key[tmpIdx] = state->key[0]; | ||
| // Write out tracker | ||
| state_tracker[tmpIdx] = state->tracker; | ||
|
|
||
| // For second out of four sets... | ||
| // Write out counter | ||
| tmpIdx += rng_count; | ||
| state_counter[tmpIdx] = state->counter[1]; | ||
| // Write out result | ||
| state_result[tmpIdx] = state->result[1]; | ||
| // Write out key | ||
| state_key[tmpIdx] = state->key[1]; | ||
|
|
||
| // For third out of four sets... | ||
| // Write out counter | ||
| tmpIdx += rng_count; | ||
| state_counter[tmpIdx] = state->counter[2]; | ||
| // Write out result | ||
| state_result[tmpIdx] = state->result[2]; | ||
| // Write out key | ||
| state_key[tmpIdx] = state->key[2]; | ||
|
|
||
| // For last out of four sets... | ||
| // Write out counter | ||
| tmpIdx += rng_count; | ||
| state_counter[tmpIdx] = state->counter[3]; | ||
| // Write out result | ||
| state_result[tmpIdx] = state->result[3]; | ||
| // Write out key | ||
| state_key[tmpIdx] = state->key[3]; | ||
|
|
||
| } |
Oops, something went wrong.