"sims/mem/basicmem/rules.mk" did not exist on "32a74b8eb0183e6911e87292795f8a8968d214af"
random_gen.hpp 2.02 KB
Newer Older
1
2
3
4
5
6
7
8
9
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.

#pragma once

namespace ck {

// Pseudo random number generator
// version for fp32
10
template <typename T, uint32_t seed_t, ck::enable_if_t<ck::is_same<float, T>{}, bool> = false>
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
__host__ __device__ uint32_t prand_generator(index_t id, T val, uint32_t seed = seed_t)
{
    uint32_t x         = *(reinterpret_cast<uint32_t*>(&val));
    uint32_t drop_bits = uint32_t(x) & 0xFFFFu;
    drop_bits ^= x >> 16;
    drop_bits = ((drop_bits & 31) << 11) | (drop_bits >> 5);
    drop_bits *= 0x7000149;
    // NOTE: If id is in 64 bit, we are only using lower 32 bit.
    //       So, it can have an effect of using same id for multiple elements when the id is very
    //       large!
    uint32_t rng = (drop_bits ^ 0x13371337 ^ (id * 229791) ^ seed);
    return rng;
}

// version for fp16
26
template <typename T, uint32_t seed_t, ck::enable_if_t<ck::is_same<half_t, T>{}, bool> = false>
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
__host__ __device__ uint32_t prand_generator(index_t id, T val, uint32_t seed = seed_t)
{
    uint16_t x         = *(reinterpret_cast<uint16_t*>(&val));
    uint32_t drop_bits = uint32_t(x) & 0xFFFFu;
    drop_bits          = ((drop_bits & 31) << 11) | (drop_bits >> 5);
    drop_bits *= 0x7000149;
    // NOTE: If id is in 64 bit, we are only using lower 32 bit.
    //       So, it can have an effect of using same id for multiple elements when the id is very
    //       large!
    uint32_t rng = (drop_bits ^ 0x13371337 ^ (id * 229791) ^ seed);
    return rng;
}

// return 0 if data is not fp16 or fp32
template <typename T,
          uint32_t seed_t,
43
          ck::enable_if_t<!(ck::is_same<float, T>{} || ck::is_same<half_t, T>{}), bool> = false>
44
45
__host__ __device__ uint32_t prand_generator(int id, T val, uint32_t seed = seed_t)
{
46
47
48
49
50
    #ifdef __HIPCC_RTC__
    static_cast<void>(id);
    static_cast<void>(val);
    static_cast<void>(seed);
    #else
51
52
53
    std::ignore = id;
    std::ignore = val;
    std::ignore = seed;
54
    #endif
55
56
57
58
59

    return 0;
}

} // namespace ck