Resolve compiler warnings when building with CUDA by Sichao25 · Pull Request #251 · SCOREC/pcms

Sichao25 · 2025-12-19T16:38:04Z

No description provided.

jacobmerson · 2025-12-19T17:08:38Z

test/test_field_communication.cpp

-  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) {


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

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.

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

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?

I see. I'll update accordingly.

jacobmerson · 2025-12-19T17:09:20Z

test/test_field_communication.cpp

-  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) {


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

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.

Similar warnings about calling host functions from device functions.

jacobmerson · 2025-12-19T17:13:35Z

src/pcms/adapter/omega_h/omega_h_field2.cpp

@@ -34,7 +34,7 @@ class MeshFieldBackendImpl : public MeshFieldBackend
    Kokkos::deep_copy(data_d, Kokkos::View<const Real*, HostMemorySpace>(
                                data.data_handle(), data.size()));
    Omega_h::parallel_for(


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

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

jacobmerson · 2025-12-19T17:14:36Z

src/pcms/adapter/omega_h/omega_h_field2.cpp

    // 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);


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

jacobmerson · 2025-12-21T21:36:47Z

test/test_field_copy.cpp

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


use std::iota

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.

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

jacobmerson · 2025-12-21T21:38:21Z

test/test_field_copy.cpp

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


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

Similar warnings about calling host functions from device functions.

jacobmerson · 2025-12-21T21:39:32Z

test/test_interpolation_on_ltx_mesh.cpp


  // ------------------ Verification ------------------ //
-  double tol = 10.0 / 100.0; // 10 percent tolerance
+  // double tol = 10.0 / 100.0; // 10 percent tolerance


If this is unused, just delete it.

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)); } */

Sichao25 added 4 commits December 9, 2025 17:14

fix compiler warnings

fb48823

Merge branch 'develop' into yus/debug_gpu_warning

e34a560

Merge branch 'develop' into yus/debug_gpu_warning

c8f91cc

comment unused code

f8a58b6

Sichao25 marked this pull request as ready for review December 19, 2025 16:58

jacobmerson reviewed Dec 21, 2025

View reviewed changes

remove KOKKOS_LAMBDA

477c325

Sichao25 had a problem deploying to asics December 24, 2025 02:04 — with GitHub Actions Failure

Trigger workflow

95cd841

Sichao25 had a problem deploying to asics January 30, 2026 17:28 — with GitHub Actions Failure

Sichao25 had a problem deploying to asics February 2, 2026 17:22 — with GitHub Actions Error

Sichao25 temporarily deployed to asics February 3, 2026 17:14 — with GitHub Actions Inactive

Conversation

Sichao25 commented Dec 19, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants