Skip to content

frontier: need host/device fn, use kokkos array#297

Merged
cwsmith merged 2 commits intodevelopfrom
cws/bboxUseKokkosArray
Mar 16, 2026
Merged

frontier: need host/device fn, use kokkos array#297
cwsmith merged 2 commits intodevelopfrom
cws/bboxUseKokkosArray

Conversation

@cwsmith
Copy link
Contributor

@cwsmith cwsmith commented Mar 16, 2026

This change fixes a compile error (pasted below) on Frontier using hipcc (/opt/rocm-6.2.4/lib/llvm/bin/clang++) regarding a device function calling the host only simplex_bbox(...) function. To that end, the use of std::array was replaced with Kokkos::Array in the functions that use the AABBox struct.

All tests are passing in a local build (with petsc disabled).

~/pcms/src/pcms/point_search.cpp:241:21: error: reference to __host__ function 'simplex_bbox<3U>' in __host__ __device__ function                                                                                                                                                                                                                             
  241 |   return intersects(simplex_bbox<dim>(coords), bbox);                                                                                                                                  
      |                     ^                                                                                                                                                                  
~/pcms/src/pcms/point_search.cpp:340:11: note: called by 'operator()'                                                                                                                                                                                                                                                                                         
  340 |       if (simplex_intersects_bbox<3>(vertex_coords, grid_cell_bbox)) {                                                                                                                                                                                                                                                                                                                
      |           ^                                                                                                                                                                                                                                                                                                                                                                           
~/pcmsDev/builds/AMD_GFX90A/kokkos/install/include/Kokkos_Crs.hpp:316:19: note: called by 'operator()'                                                                                                                                                                                                                                                     
  316 |     m_counts(i) = m_functor(i, nullptr);                                                                                                                                                                                                                                                                                                                                              
      |                   ^                                                                                                                                                                    
~/pcmsDev/builds/AMD_GFX90A/kokkos/install/include/HIP/Kokkos_HIP_ParallelFor_Range.hpp:50:5: note: called by 'exec_range<Kokkos::CountAndFillBase<Kokkos::Crs<int, Kokkos::HIP, void, int>, pcms::detail::GridTriIntersectionFunctor3D>::Count>'                                                                                                          
   50 |     m_functor(TagType(), i);                                                                                                                                                           
      |     ^                                                                                                                                                                                                                                                                                                                                                                                 
~/pcmsDev/builds/AMD_GFX90A/kokkos/install/include/HIP/Kokkos_HIP_ParallelFor_Range.hpp:70:22: note: called by 'operator()'                                                                                                                                                                                                                                
   70 |       this->template exec_range<WorkTag>(iwork);                                                                                                                                                                                                                                                                                                                                      
      |                      ^                                                                                                                                                                                                                                                                                                                                                                
~/pcmsDev/builds/AMD_GFX90A/kokkos/install/include/HIP/Kokkos_HIP_KernelLaunch.hpp:76:3: note: called by 'hip_parallel_launch_local_memory<Kokkos::Impl::ParallelFor<Kokkos::CountAndFill<Kokkos::Crs<int, Kokkos::HIP, void, int>, pcms::detail::GridTriIntersectionFunctor3D>, Kokkos::RangePolicy<int, Kokkos::HIP, Kokkos::CountAndFillBase<Kokkos::Crs
<int, Kokkos::HIP, void, int>, pcms::detail::GridTriIntersectionFunctor3D>::Count>>>'                                                                                                          
   76 |   driver();                                                                                                                                                                            
      |   ^                                                                                                                                                                                    
~/pcms/src/pcms/point_search.cpp:47:13: note: 'simplex_bbox<3U>' declared here                                                                                 
   47 | AABBox<dim> simplex_bbox(const Omega_h::Matrix<dim, dim + 1>& coords)                                                                                                                  
      |             ^                                                                                                                                                                                                                                                                                                                                                                         
1 error generated when compiling for gfx90a.                                                                

@jacobmerson
Copy link
Collaborator

Thanks.

this is using clang-format from llvm 16 instead of 18 the CI suggested
@cwsmith cwsmith merged commit a3ccc35 into develop Mar 16, 2026
6 checks passed
@cwsmith cwsmith deleted the cws/bboxUseKokkosArray branch March 16, 2026 23:47
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