Skip to content

chore(rocrand): removal of host device attributes #8806

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
umfranzw merged 2 commits into from
Jul 18, 2026
Merged

chore(rocrand): removal of host device attributes #8806

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
397ce13
chore(rocRAND): removal of host device attributes from a deduction guide
borysborys May 12, 2026
fd37157
Formatting changes for CI checks
umfranzw Jul 17, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions projects/rocrand/library/src/rng/threefry.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -164,8 +164,8 @@ struct generate_threefry
size_t index = thread_id;
while(index < vec_n)
{
const auto v = engine.next_leap(stride);
cpp_utils::vec_wrapper vs(v);
const auto v = engine.next_leap(stride);
cpp_utils::vec_wrapper<std::remove_cv_t<decltype(v)>> vs(v);
for(unsigned int s = 0; s < output_per_thread; s++)
{
for(unsigned int i = 0; i < input_width; i++)
Expand Down
27 changes: 18 additions & 9 deletions projects/rocrand/library/src/rng/utils/cpp_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2023-2026 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
Expand Down Expand Up @@ -107,7 +107,8 @@ constexpr auto numeric_combinations(const std::array<T, Ns>... inputs)
}

/// \brief Calculates greatest common divisor.
__host__ __device__ constexpr unsigned int gcd(const unsigned int a, const unsigned int b)
__host__ __device__
constexpr unsigned int gcd(const unsigned int a, const unsigned int b)
{
if(a == 0)
return b;
Expand All @@ -117,7 +118,8 @@ __host__ __device__ constexpr unsigned int gcd(const unsigned int a, const unsig
}

/// \brief Calculates least common multiple
__host__ __device__ constexpr unsigned int lcm(const unsigned int a, const unsigned int b)
__host__ __device__
constexpr unsigned int lcm(const unsigned int a, const unsigned int b)
{
if(a == 0 || b == 0)
return 0;
Expand Down Expand Up @@ -192,10 +194,14 @@ struct vec_wrapper
{
static_assert(is_vector_type_v<V>, "vec_wrapper can only be used with vector types");

__host__ __device__ explicit vec_wrapper(V vec) : m_vec(vec) {}
__host__ __device__
explicit vec_wrapper(V vec)
: m_vec(vec)
{}

template<class U = V, std::enable_if_t<vector_size_v<U> == 2, int> = 0>
__host__ __device__ auto& operator[](int idx)
__host__ __device__
auto& operator[](int idx)
{
switch(idx)
{
Expand All @@ -206,7 +212,8 @@ struct vec_wrapper
}

template<class U = V, std::enable_if_t<vector_size_v<U> == 2, int> = 0>
__host__ __device__ const auto& operator[](int idx) const
__host__ __device__
const auto& operator[](int idx) const
{
switch(idx)
{
Expand All @@ -217,7 +224,8 @@ struct vec_wrapper
}

template<class U = V, std::enable_if_t<vector_size_v<U> == 4, int> = 0>
__host__ __device__ auto& operator[](int idx)
__host__ __device__
auto& operator[](int idx)
{
switch(idx)
{
Expand All @@ -230,7 +238,8 @@ struct vec_wrapper
}

template<class U = V, std::enable_if_t<vector_size_v<U> == 4, int> = 0>
__host__ __device__ const auto& operator[](int idx) const
__host__ __device__
const auto& operator[](int idx) const
{
switch(idx)
{
Expand All @@ -246,7 +255,7 @@ struct vec_wrapper
};

template<class V>
__host__ __device__ vec_wrapper(V) -> vec_wrapper<V>;
vec_wrapper(V) -> vec_wrapper<V>;

/// \brief Returns the maximum of its arguments.
/// \note This function must be the choice in `__host__ __device__` and preferably on
Expand Down
Loading