Skip to content
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

Add HIP logic, move to consistent naming convention #53

Merged
merged 2 commits into from
Feb 13, 2025
Merged

Add HIP logic, move to consistent naming convention #53

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
0b0388e
Add HIP logic, move to same naming convention as other traccc project…
Feb 11, 2025
1e8d0ed
Formatting/linting
Feb 11, 2025
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
9 changes: 2 additions & 7 deletions benchmarks/common/covfie/benchmark/lorentz.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,15 +10,10 @@
#include <covfie/benchmark/types.hpp>
#include <covfie/core/definitions.hpp>
#include <covfie/core/field_view.hpp>

#ifdef __CUDACC__
#define HOST_DEVICE __host__ __device__
#else
#define HOST_DEVICE
#endif
#include <covfie/core/qualifiers.hpp>

template <typename propagator_t, typename backend_t>
HOST_DEVICE inline __attribute__((always_inline)) void lorentz_step(
COVFIE_HOST_DEVICE inline __attribute__((always_inline)) void lorentz_step(
const covfie::field_view<backend_t> & f,
covfie::benchmark::lorentz_agent<3> & o,
float s
Expand Down
8 changes: 1 addition & 7 deletions benchmarks/common/covfie/benchmark/propagate.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,12 +10,6 @@
#include <covfie/core/definitions.hpp>
#include <covfie/core/field_view.hpp>

#ifdef __CUDACC__
#define HOST_DEVICE __host__ __device__
#else
#define HOST_DEVICE
#endif

class Wide
{
};
Expand All @@ -33,7 +27,7 @@ class RungeKutta4
};

template <typename propagator_t, typename backend_t>
HOST_DEVICE inline __attribute__((always_inline)) void propagation_step(
COVFIE_HOST_DEVICE inline __attribute__((always_inline)) void propagation_step(
const covfie::field_view<backend_t> & f,
covfie::benchmark::propagation_agent<3> & o,
float s
Expand Down
12 changes: 7 additions & 5 deletions lib/core/covfie/core/algebra/affine.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,12 +22,13 @@ struct affine : public matrix<N, N + 1, T, I> {
affine & operator=(const affine &) = default;
affine & operator=(affine &&) = default;

COVFIE_DEVICE affine(const matrix<N, N + 1, T, I> & o)
COVFIE_HOST_DEVICE affine(const matrix<N, N + 1, T, I> & o)
: matrix<N, N + 1, T, I>(o)
{
}

COVFIE_DEVICE vector<N, T, I> operator*(const vector<N, T, I> & v) const
COVFIE_HOST_DEVICE vector<N, T, I> operator*(const vector<N, T, I> & v
) const
{
vector<N + 1, T, I> r;

Expand All @@ -40,7 +41,8 @@ struct affine : public matrix<N, N + 1, T, I> {
return matrix<N, N + 1, T, I>::operator*(r);
}

COVFIE_DEVICE affine<N, T, I> operator*(const affine<N, T, I> & m) const
COVFIE_HOST_DEVICE affine<N, T, I> operator*(const affine<N, T, I> & m
) const
{
matrix<N + 1, N + 1, T, I> m1, m2;

Expand Down Expand Up @@ -74,7 +76,7 @@ struct affine : public matrix<N, N + 1, T, I> {
}

template <typename... Args>
COVFIE_DEVICE static affine<N, T, I> translation(const Args &... args)
COVFIE_HOST_DEVICE static affine<N, T, I> translation(const Args &... args)
{
static_assert(
(std::is_convertible_v<Args, T> && ...),
Expand All @@ -99,7 +101,7 @@ struct affine : public matrix<N, N + 1, T, I> {
}

template <typename... Args>
COVFIE_DEVICE static affine<N, T, I> scaling(const Args &... args)
COVFIE_HOST_DEVICE static affine<N, T, I> scaling(const Args &... args)
{
static_assert(
(std::is_convertible_v<Args, T> && ...),
Expand Down
12 changes: 6 additions & 6 deletions lib/core/covfie/core/algebra/matrix.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,11 +18,11 @@ template <
typename T = float,
typename I = std::size_t>
struct matrix {
COVFIE_DEVICE matrix()
COVFIE_HOST_DEVICE matrix()
{
}

COVFIE_DEVICE matrix(array::array<array::array<T, M>, N> l)
COVFIE_HOST_DEVICE matrix(array::array<array::array<T, M>, N> l)
{
for (I i = 0; i < l.size(); ++i) {
for (I j = 0; j < l[i].size(); ++j) {
Expand All @@ -33,18 +33,18 @@ struct matrix {

matrix(const matrix<N, M, T, I> &) = default;

COVFIE_DEVICE T operator()(const I i, const I j) const
COVFIE_HOST_DEVICE T operator()(const I i, const I j) const
{
return m_elems[i][j];
}

COVFIE_DEVICE T & operator()(const I i, const I j)
COVFIE_HOST_DEVICE T & operator()(const I i, const I j)
{
return m_elems[i][j];
}

template <std::size_t P>
COVFIE_DEVICE matrix<N, P, T, I> operator*(const matrix<M, P, T, I> & o
COVFIE_HOST_DEVICE matrix<N, P, T, I> operator*(const matrix<M, P, T, I> & o
) const
{
matrix<N, P, T, I> r;
Expand All @@ -64,7 +64,7 @@ struct matrix {
return r;
}

COVFIE_DEVICE static matrix<N, M, T, I> identity()
COVFIE_HOST_DEVICE static matrix<N, M, T, I> identity()
{
matrix<N, M, T, I> result;

Expand Down
12 changes: 6 additions & 6 deletions lib/core/covfie/core/algebra/vector.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,20 +18,20 @@
namespace covfie::algebra {
template <std::size_t N, typename T = float, typename I = std::size_t>
struct vector : public matrix<N, 1, T, I> {
COVFIE_DEVICE vector()
COVFIE_HOST_DEVICE vector()
: matrix<N, 1, T, I>()
{
}

COVFIE_DEVICE vector(array::array<T, N> l)
COVFIE_HOST_DEVICE vector(array::array<T, N> l)
: matrix<N, 1, T, I>()
{
for (I i = 0; i < N; ++i) {
matrix<N, 1, T, I>::operator()(i, 0) = l[i];
}
}

COVFIE_DEVICE vector(const matrix<N, 1, T, I> & o)
COVFIE_HOST_DEVICE vector(const matrix<N, 1, T, I> & o)
: matrix<N, 1, T, I>(o)
{
}
Expand All @@ -43,17 +43,17 @@ struct vector : public matrix<N, 1, T, I> {
std::conjunction_v<std::is_convertible<Args, T>...> &&
sizeof...(Args) == N,
bool> = true>
COVFIE_DEVICE vector(Args... args)
COVFIE_HOST_DEVICE vector(Args... args)
: vector(array::array<T, N>{std::forward<Args>(args)...})
{
}

COVFIE_DEVICE T operator()(const I & i) const
COVFIE_HOST_DEVICE T operator()(const I & i) const
{
return matrix<N, 1, T, I>::operator()(i, 0);
}

COVFIE_DEVICE T & operator()(const I & i)
COVFIE_HOST_DEVICE T & operator()(const I & i)
{
return matrix<N, 1, T, I>::operator()(i, 0);
}
Expand Down
37 changes: 19 additions & 18 deletions lib/core/covfie/core/array.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,99 +19,100 @@ requires(_size > 0) struct array {
using value_type = _scalar_t;
static constexpr std::size_t dimensions = _size;

COVFIE_DEVICE array() = default;
COVFIE_HOST_DEVICE array() = default;

COVFIE_DEVICE array(const scalar_t (&arr)[dimensions])
COVFIE_HOST_DEVICE array(const scalar_t (&arr)[dimensions])
requires(dimensions > 1)
: array(arr, std::make_index_sequence<dimensions>())
{
}

COVFIE_DEVICE array(const scalar_t & val)
COVFIE_HOST_DEVICE array(const scalar_t & val)
: array(val, std::make_index_sequence<dimensions>())
{
}

template <typename... Ts>
requires(sizeof...(Ts) == dimensions) COVFIE_DEVICE array(Ts... args)
requires(sizeof...(Ts) == dimensions) COVFIE_HOST_DEVICE array(Ts... args)
: m_data{std::forward<Ts>(args)...}
{
}

COVFIE_DEVICE constexpr scalar_t & at(const std::size_t & n)
COVFIE_HOST_DEVICE constexpr scalar_t & at(const std::size_t & n)
{
assert(n < dimensions);

return m_data[n];
}

COVFIE_DEVICE constexpr const scalar_t & at(const std::size_t & n) const
COVFIE_HOST_DEVICE constexpr const scalar_t & at(const std::size_t & n
) const
{
assert(n < dimensions);

return m_data[n];
}

COVFIE_DEVICE constexpr scalar_t & operator[](const std::size_t & n)
COVFIE_HOST_DEVICE constexpr scalar_t & operator[](const std::size_t & n)
{
assert(n < dimensions);

return m_data[n];
}

COVFIE_DEVICE constexpr const scalar_t & operator[](const std::size_t & n
) const
COVFIE_HOST_DEVICE constexpr const scalar_t &
operator[](const std::size_t & n) const
{
assert(n < dimensions);

return m_data[n];
}

COVFIE_DEVICE constexpr std::size_t size() const
COVFIE_HOST_DEVICE constexpr std::size_t size() const
{
return dimensions;
}

COVFIE_DEVICE constexpr scalar_t * begin()
COVFIE_HOST_DEVICE constexpr scalar_t * begin()
{
return m_data + 0;
}

COVFIE_DEVICE constexpr const scalar_t * begin() const
COVFIE_HOST_DEVICE constexpr const scalar_t * begin() const
{
return m_data + 0;
}

COVFIE_DEVICE constexpr const scalar_t * cbegin() const
COVFIE_HOST_DEVICE constexpr const scalar_t * cbegin() const
{
return m_data + 0;
}

COVFIE_DEVICE constexpr scalar_t * end()
COVFIE_HOST_DEVICE constexpr scalar_t * end()
{
return m_data + dimensions;
}

COVFIE_DEVICE constexpr const scalar_t * end() const
COVFIE_HOST_DEVICE constexpr const scalar_t * end() const
{
return m_data + dimensions;
}

COVFIE_DEVICE constexpr const scalar_t * cend() const
COVFIE_HOST_DEVICE constexpr const scalar_t * cend() const
{
return m_data + dimensions;
}

private:
template <std::size_t... Is>
COVFIE_DEVICE
COVFIE_HOST_DEVICE
array(const scalar_t (&arr)[dimensions], std::index_sequence<Is...>)
: m_data{arr[Is]...}
{
}

template <std::size_t... Is>
COVFIE_DEVICE array(const scalar_t & val, std::index_sequence<Is...>)
COVFIE_HOST_DEVICE array(const scalar_t & val, std::index_sequence<Is...>)
: m_data{((void)Is, val)...}
{
}
Expand Down
2 changes: 1 addition & 1 deletion lib/core/covfie/core/backend/primitive/array.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -207,7 +207,7 @@ struct array {
{
}

COVFIE_DEVICE typename covariant_output_t::vector_t
COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
at(typename contravariant_input_t::vector_t i) const
{
assert(i < m_size);
Expand Down
2 changes: 1 addition & 1 deletion lib/core/covfie/core/backend/primitive/constant.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -96,7 +96,7 @@ struct constant {
{
}

COVFIE_DEVICE typename covariant_output_t::vector_t
COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
at(typename contravariant_input_t::vector_t) const
{
return m_value;
Expand Down
2 changes: 1 addition & 1 deletion lib/core/covfie/core/backend/transformer/affine.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -136,7 +136,7 @@ struct affine {
{
}

COVFIE_DEVICE typename covariant_output_t::vector_t
COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
at(typename contravariant_input_t::vector_t c) const
{
covfie::algebra::vector<
Expand Down
2 changes: 1 addition & 1 deletion lib/core/covfie/core/backend/transformer/backup.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -171,7 +171,7 @@ struct backup {
{
}

COVFIE_DEVICE typename covariant_output_t::vector_t
COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
at(typename contravariant_input_t::vector_t coord) const
{
for (std::size_t i = 0; i < contravariant_input_t::dimensions; ++i)
Expand Down
4 changes: 2 additions & 2 deletions lib/core/covfie/core/backend/transformer/clamp.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -144,14 +144,14 @@ struct clamp {
}

template <std::size_t... Is>
COVFIE_DEVICE typename contravariant_input_t::vector_t
COVFIE_HOST_DEVICE typename contravariant_input_t::vector_t
adjust(typename contravariant_input_t::vector_t coord, std::index_sequence<Is...>)
const
{
return {std::clamp(coord[Is], m_min[Is], m_max[Is])...};
}

COVFIE_DEVICE typename covariant_output_t::vector_t
COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
at(typename contravariant_input_t::vector_t coord) const
{
return m_backend.at(adjust(
Expand Down
4 changes: 2 additions & 2 deletions lib/core/covfie/core/backend/transformer/covariant_cast.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -101,14 +101,14 @@ struct covariant_cast {
}

template <std::size_t... Is>
COVFIE_DEVICE typename covariant_output_t::vector_t
COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
at_helper(typename contravariant_input_t::vector_t c, std::index_sequence<Is...>)
const
{
return {static_cast<target_type>(m_backend.at(c)[Is])...};
}

COVFIE_DEVICE typename covariant_output_t::vector_t
COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
at(typename contravariant_input_t::vector_t c) const
{
return at_helper(
Expand Down
2 changes: 1 addition & 1 deletion lib/core/covfie/core/backend/transformer/dereference.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -99,7 +99,7 @@ struct dereference {
{
}

COVFIE_DEVICE typename covariant_output_t::vector_t
COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
at(typename contravariant_input_t::vector_t c) const
{
return m_backend.at(c);
Expand Down
Loading
Loading