Skip to content

Conversation

@Sichao25
Copy link
Contributor

No description provided.

@Sichao25 Sichao25 marked this pull request as ready for review December 19, 2025 16:58
Kokkos::parallel_for(
"id 0", Kokkos::RangePolicy<pcms::HostMemorySpace::execution_space>(0, n),
KOKKOS_LAMBDA(int i) { ids[i] = 0; });
for (size_t i = 0; i < static_cast<size_t>(n); ++i) {
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

since we know this is on the host use std::fill

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This would use OpenMP threads for host parallel if available. What is the specific warning that is being resolved with this? Basicly I want to decide if it makes sense to swich this to a host for loop or stay with kokkos host parallel_for.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The warning is:

warning #20011-D: calling a __host__ function("Omega_h::HostWrite<double> ::operator [](int) const") from a __host__ __device__ function("server(int,  ::Omega_h::Mesh &,    ::std::__cxx11::basic_string<char,     ::std::char_traits<char> , ::std::allocator<char> > , int, const     ::std::map<   ::std::__cxx11::basic_string<char,     ::std::char_traits<char> , ::std::allocator<char> > ,    ::std::__cxx11::basic_string<char,     ::std::char_traits<char> , ::std::allocator<char> > ,     ::std::less<   ::std::__cxx11::basic_string<char,     ::std::char_traits<char> , ::std::allocator<char> > > , ::std::allocator<    ::std::pair<const    ::std::__cxx11::basic_string<char,     ::std::char_traits<char> , ::std::allocator<char> > ,    ::std::__cxx11::basic_string<char,     ::std::char_traits<char> , ::std::allocator<char> > > > >  &, const    ::std::__cxx11::basic_string<char,     ::std::char_traits<char> , ::std::allocator<char> >  &)::[lambda(int) (instance 1)]::operator () const") is not allowed

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see...since it's in the HostSpace using KOKKOS_LAMBDA is marking that as __host__ __device__. Probably just using [=](int i){} for the lambda will resolve this. Alternatively there may be a KOKKOS_HOST_LAMBDA?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see. I'll update accordingly.

Kokkos::parallel_for(
"id gid", Kokkos::RangePolicy<pcms::HostMemorySpace::execution_space>(0, n),
KOKKOS_LAMBDA(int i) { ids[i] = gids[i]; });
for (size_t i = 0; i < static_cast<size_t>(n); ++i) {
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

since we know this is on the host, use std::copy

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This would use OpenMP threads for host parallel if available. What is the specific warning that is being resolved with this? Basicly I want to decide if it makes sense to swich this to a host for loop or stay with kokkos host parallel_for.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Similar warnings about calling host functions from device functions.

"data_d", data.size());
Kokkos::deep_copy(data_d, Kokkos::View<const Real*, HostMemorySpace>(
data.data_handle(), data.size()));
Omega_h::parallel_for(
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

would it make sense to swap out the Omega_h::parallel_for with Kokkos::parallel_for if we are using the Kokkos macros?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That makes sense. I think we should be consistent and use only one of them if possible.

// element should be inside the domain (positive)
PCMS_ALWAYS_ASSERT(elem_idx >= 0 && elem_idx < mesh_.nelems());
// PCMS_ALWAYS_ASSERT(elem_idx >= 0 && elem_idx < mesh_.nelems());
LO count = Kokkos::atomic_sub_fetch(&elem_counts_(elem_idx), 1);
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should we define a Host macro that only evaluates on host?

KOKKOS_LAMBDA(int i) { ids[i] = i; });
for (size_t i = 0; i < static_cast<size_t>(ndata); ++i) {
ids[i] = i;
}
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

use std::iota

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This would use OpenMP threads for host parallel if available. What is the specific warning that is being resolved with this? Basicly I want to decide if it makes sense to swich this to a host for loop or stay with kokkos host parallel_for.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The warning is:

warning #20011-D: calling a __host__ function("Omega_h::HostWrite<double> ::operator [](int) const") from a __host__ __device__ function("test_copy(    ::std::shared_ptr< ::Omega_h::Comm> , int, int, int)::[lambda(int) (instance 1)]::operator () const") is not allowed

if (std::abs(ids[i] - copied_array[i]) < 1e-12) {
sum++;
}
}
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This would use OpenMP threads for host parallel if available. What is the specific warning that is being resolved with this?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Similar warnings about calling host functions from device functions.


// ------------------ Verification ------------------ //
double tol = 10.0 / 100.0; // 10 percent tolerance
// double tol = 10.0 / 100.0; // 10 percent tolerance
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If this is unused, just delete it.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think this is intended to be used with the commented-out verification code. Would you like me to remove that entire block?

 // ------------------ Verification ------------------ //
  double tol = 10.0 / 100.0; // 10 percent tolerance
  // XGC Originated Values
  /*
  for (int i = 0; i < xgc_num_nodes; ++i) {
    CHECK_THAT(interpolated_back_density_at_xgc_nodes[i],
               Catch::Matchers::WithinRel(density_at_xgc_nodes[i], tol) ||
                 Catch::Matchers::WithinAbs(density_at_xgc_nodes[i], tol));
    CHECK_THAT(interpolated_back_temp_at_xgc_nodes[i],
               Catch::Matchers::WithinRel(temp_at_xgc_nodes[i], tol) ||
                 Catch::Matchers::WithinAbs(temp_at_xgc_nodes[i], tol));
  }
  */

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants