Class initMersenneKernel

src/backend/oneapi/kernel/random_engine_mersenne.hpp:105–144

Initialization

// Initialization
class initMersenneKernel {
   public:
    initMersenneKernel(write_accessor<uint> state, read_accessor<uint> tbl,
                       sycl::local_accessor<uint, 1> lstate, uintl seed)
        : state_(state), tbl_(tbl), lstate_(lstate), seed_(seed) {}

    void operator()(sycl::nd_item<1> it) const {
        sycl::group g = it.get_group();

        const uint *ltbl =
            tbl_.get_pointer() + (TABLE_SIZE * g.get_group_id(0));
        uint hidden_seed = ltbl[4] ^ (ltbl[8] << 16);
        uint tmp         = hidden_seed;
        tmp += tmp >> 16;
        tmp += tmp >> 8;
        tmp &= 0xff;
        tmp |= tmp << 8;
        tmp |= tmp << 16;
        lstate_[it.get_local_id(0)] = tmp;
        it.barrier();
        if (it.get_local_id(0) == 0) {
            lstate_[0] = seed_;
            lstate_[1] = hidden_seed;
            for (int i = 1; i < N; ++i) {
                lstate_[i] ^= ((uint)(1812433253) *
                                   (lstate_[i - 1] ^ (lstate_[i - 1] >> 30)) +
                               i);
            }
        }
        it.barrier();
        state_[N * g.get_group_id(0) + it.get_local_id(0)] =
            lstate_[it.get_local_id(0)];
    }

   protected:
    write_accessor<uint> state_;
    read_accessor<uint> tbl_;
    sycl::local_accessor<uint, 1> lstate_;
    uintl seed_;
};

void initMersenneState(Param<uint> state, const Param<uint> tbl, uintl seed) {
    sycl::nd_range<1> ndrange({BLOCKS * N}, {N});