| 950 | |
| 951 | template<typename T> |
| 952 | __global__ void normalThreefry(T *out, uint hi, uint lo, uint hic, uint loc, |
| 953 | uint elementsPerBlock, uint elements) { |
| 954 | uint index = blockIdx.x * elementsPerBlock + threadIdx.x; |
| 955 | uint key[2] = {lo, hi}; |
| 956 | uint ctr[2] = {loc, hic}; |
| 957 | ctr[0] += index; |
| 958 | ctr[1] += (ctr[0] < loc); |
| 959 | uint o[4]; |
| 960 | |
| 961 | threefry(key, ctr, o); |
| 962 | uint step = elementsPerBlock / 2; |
| 963 | ctr[0] += step; |
| 964 | ctr[1] += (ctr[0] < step); |
| 965 | threefry(key, ctr, o + 2); |
| 966 | |
| 967 | if (blockIdx.x != (gridDim.x - 1)) { |
| 968 | boxMullerWriteOut128Bytes(out, index, o[0], o[1], o[2], o[3]); |
| 969 | } else { |
| 970 | partialBoxMullerWriteOut128Bytes(out, index, o[0], o[1], o[2], o[3], |
| 971 | elements); |
| 972 | } |
| 973 | } |
| 974 | |
| 975 | template<typename T> |
| 976 | __global__ void normalMersenne(T *const out, uint *const gState, |
nothing calls this directly
no test coverage detected