Files
derek.holloway 7502bf64a4
CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
Flatten subprojects into a single base monorepo
2026-09-02 16:34:47 -07:00

147 lines
4.1 KiB
C++

// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#ifndef GUARD_RANDOM_GEN_
#define GUARD_RANDOM_GEN_
#include <miopen/env.hpp>
#include <miopen/kernel_tuning_mode.hpp>
#include <cassert>
#include <iostream>
#include <random>
MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_DRIVER_PRNG_SEED, 12345678)
namespace env = miopen::env;
namespace prng {
namespace details {
using glibc_gen = std::linear_congruential_engine<std::uint32_t, 1103515245, 12345, 2147483648>;
inline std::random_device::result_type get_default_seed()
{
static std::random_device::result_type seed{[] {
auto external_seed = env::value(MIOPEN_DEBUG_DRIVER_PRNG_SEED);
auto seed_ = external_seed == 0
? std::random_device{}()
: static_cast<std::random_device::result_type>(external_seed);
// Check if JSON mode is enabled in MIOPEN_PERFORMANCE_LOGS
if(miopen::IsPerformanceLoggingEnabled())
{
std::cout << "{\"prng_seed\":" << seed_ << "}" << std::endl;
}
else
{
std::cout << "PRNG seed: " << seed_ << "\n";
}
return seed_;
}()};
return seed;
}
inline glibc_gen& get_prng()
{
static thread_local glibc_gen gen{get_default_seed()};
return gen;
}
template <class, class = void>
struct has_digits : std::false_type
{
};
template <class T>
struct has_digits<T, std::void_t<decltype(std::numeric_limits<T>::digits)>> : std::true_type
{
};
} // namespace details
inline void reset_seed(std::random_device::result_type seed = 0)
{
details::get_prng().seed(seed + details::get_default_seed());
}
// similar to std::generate_canonical, but simpler and faster
template <typename T>
inline T gen_canonical()
{
if constexpr(std::is_floating_point_v<T>) // native fp
{
static constexpr T range =
static_cast<T>(1) /
static_cast<T>(details::glibc_gen::max() - details::glibc_gen::min() + 1);
return range * static_cast<T>(details::get_prng()() - details::glibc_gen::min());
}
else if constexpr(std::is_integral_v<T>)
{
auto val = details::get_prng()();
return static_cast<T>(((val >> 4) + (val >> 16)) & 0x1);
}
else
{
return static_cast<T>(gen_canonical<float>());
}
}
template <typename T>
inline T gen_0_to_B(T B)
{
if constexpr(std::is_floating_point_v<T>) // native fp
{
return gen_canonical<T>() * B;
}
else if constexpr(std::is_integral_v<T>)
{
// can only generate 27bit range, so it may not be suitable
// for huge 64 bit ranges, but we do not expect such ranges
return static_cast<T>((details::get_prng()() >> 4) % B);
}
else // half/bfloat/etc
{
return static_cast<T>(gen_0_to_B(static_cast<float>(B)));
}
}
template <typename T>
inline T gen_A_to_B(T A, T B)
{
assert(B > A);
return gen_0_to_B(B - A) + A;
}
template <typename T>
inline T gen_off_range(T offset, T range)
{
static_assert(std::is_integral_v<T>);
return prng::gen_0_to_B(range) + offset;
}
template <typename T, bool Signed = false>
inline T gen_subnorm()
{
T denorm_val = static_cast<T>(0);
if constexpr(!std::is_integral_v<T> && !std::is_same_v<T, double> &&
std::is_trivially_copyable<T>::value && details::has_digits<T>::value)
{
using BitType = std::conditional_t<sizeof(T) == 1,
uint8_t,
std::conditional_t<sizeof(T) == 2, uint16_t, uint32_t>>;
static_assert(sizeof(T) == sizeof(BitType));
// -1 because ::digits counts the first implicit digit
static constexpr auto mantissa_bits = std::numeric_limits<T>::digits - 1;
BitType denorm_bits = static_cast<BitType>(gen_0_to_B(1 << mantissa_bits));
denorm_bits |= Signed ? (gen_canonical<BitType>() << (sizeof(T) * 8 - 1)) : 0;
// the proper way to do a type punning
std::memcpy(&denorm_val, &denorm_bits, sizeof(T));
}
return denorm_val;
}
} // namespace prng
#endif // GUARD_RANDOM_GEN_