drand48 on OpenCL Kernel

Question

Does anyone have an implementation of drand48() or an equivalent that can work in an OpenCL kernel?

I have been sending random numbers generated on the host through a buffer but I need random numbers generated on the device if there is any way to do this.

What will you use the random numbers for? If those numbers are intended to further information security in any way (e.g., they serve as random encryption keys, nonces, or passwords), then drand48 is not an appropriate choice of RNG. Also, have you confirmed that random number generation is the performance bottleneck in your application? — Peter O.
@PeterO. They are being used in a ray tracing application to generate realistic output from a material simulation algorithm, therefore the more random the number the more realistic the output is. In order to prevent random numbers from repeating though, I have to send a very large buffer, while I'm not focused on the performance of this, I thought there was likely a better way to do it natively? — Carter Roeser

faken faken · Accepted Answer · 2019-08-02T10:54:04

Here's an OpenCL device function which you can call from an OpenCL kernel:

uint rng_next(__global ulong *states, uint index) {

    /* Assume 32 bits */
    uint bits = 32;

    /* Get current state */
    ulong state = states[index];

    /* Update state */
    state = (state * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);

    /* Keep new state */
    states[index] = state;

    /* Return value */
    return (uint) (state >> (48 - bits));
}

The states array contains the state of the PRNG for each work-item and the index is basically - but not necessarily - the work-item ID (which you can get with get_global_id()).

The states array can be generated in the host (using another PRNG) and copied to the device, or it can be initialized in the device using some kind of hash function applied to the work-item global IDs. If you use the work-item global IDs as initial seeds, the random streams for each work-item will be very low quality (due to high correlation between them). Here's a kernel to apply a hash function to decorrelate the initial seeds (note you need a main initial seed, passed by the host):

__kernel void rng_init(
        const ulong main_seed,
        __global clo_statetype *seeds) {

    /* Get initial seed for this workitem. */
    ulong seed = get_global_id(0) + main_seed;

    /* Apply basic xor-shift hash, better ones probably exist. */
    seed = ((seed >> 16) ^ seed) * 0x45d9f3b;
    seed = ((seed >> 16) ^ seed) * 0x45d9f3b;
    seed = ((seed >> 16) ^ seed);

    /* Update seeds array. */
    seeds[get_global_id(0)] = seed;
}

Note that, as pointed out in the comments, the drand48 is of very low quality, and if you use a lot of work-items you will see artifacts in your rendering. This post explains this in more detail.

This code is taken from the cl_ops library, which I'm the author of.

drand48 on OpenCL Kernel

1 Answers