numba cuda random device function

xorshift.py

MAX32 = uint32(0xffffffff)

@cuda.jit("(uint64[::1], uint64)", device=True)
def cuda_xorshift(states, id):
    x = states[id]
    x ^= x >> 12
    x ^= x << 25
    x ^= x >> 27
    states[id] = x
    return uint64(x) * uint64(2685821657736338717)


@cuda.jit("float32(uint64[::1], uint64)", device=True)
def cuda_xorshift_float(states, id):
    return float32(float32(MAX32 & cuda_xorshift(states, id)) / float32(MAX32))

Could you provide a very brief example of how to use this generator? It's not clear what the input variables states and id should be...

The states is an array of random number states.
If you assume one state per thread, the id will be the global thread id.

from numba import cuda
import numpy as np

@cuda.jit
def example_usage(states, out):
    tid = cuda.grid(1)
    out[tid] = cuda_xorshift(states, tid)


states = np.arange(4 * 5, dtype=np.uint64) + 1  # seed states with anything but 0
out = np.zeros_like(states)

example_usage[4, 5](states, out)

print(states)
# output:
# [ 33554433  67108866 100663299 134217733 167772164 201326599 234881030
#   268435466 301989899 335544328 369098761 402653199 436207630 469762061
#   503316492 536870932 570425365 603979798 637534231 671088657]
print(out)
# output:
# [ 5180492295206395165 10360984590412790330 15541476885619185495
#  4961046764852367761  4769895744586085492 15322031355265158091
#15130880334998875822 9922093529704735522 15102585824911130687
#9539791489172170984 14720283784378566149 14883140294557103283
#14691989274290821014 14500838254024538745 14309687233758256476
#  1397442985699919428  6577935280906314593 11758427576112709758
#16938919871319104923  3318660562371129069]
'''