Skip to content

Commit

Permalink
Merge pull request #123 from CKegel/const_compatibility
Browse files Browse the repository at this point in the history
Use `constexpr` in place of `__constant__` when using the kokkos backend.

Variables with the __constant__ qualifier cannot be read in Kokkos host functions, causing the compiler to throw a warning. This branch replaces the __constant__ qualifier with constexpr on non-CUDA platforms to resolve this warning. The changes were tested on Perlmutter and Frontier using the Delta Wing test, showing a less than 1% increase in completion time.
  • Loading branch information
cwsmith authored Jan 7, 2025
2 parents 7e79631 + 3eac0f8 commit 5f30588
Showing 1 changed file with 6 additions and 2 deletions.
8 changes: 6 additions & 2 deletions src/Omega_h_macros.h
Original file line number Diff line number Diff line change
Expand Up @@ -47,19 +47,23 @@
#define OMEGA_H_INLINE_BIG OMEGA_H_INLINE
#define OMEGA_H_DEVICE __host__ __device__ inline
#define OMEGA_H_LAMBDA [=] __host__ __device__
#if defined(__CUDA_ARCH__)
#define OMEGA_H_CONSTANT_DATA __constant__
#else
#define OMEGA_H_CONSTANT_DATA
#endif
#elif defined(OMEGA_H_USE_KOKKOS)
#define OMEGA_H_INLINE KOKKOS_INLINE_FUNCTION
#define OMEGA_H_INLINE_BIG OMEGA_H_INLINE
#define OMEGA_H_DEVICE KOKKOS_INLINE_FUNCTION
#define OMEGA_H_LAMBDA KOKKOS_LAMBDA
#define OMEGA_H_CONSTANT_DATA
#define OMEGA_H_CONSTANT_DATA constexpr
#elif defined(OMEGA_H_USE_CUDA)
#define OMEGA_H_INLINE __host__ __device__ inline
#define OMEGA_H_INLINE_BIG OMEGA_H_INLINE
#define OMEGA_H_DEVICE __host__ __device__ inline
#define OMEGA_H_LAMBDA [=] __host__ __device__
#define OMEGA_H_CONSTANT_DATA __constant__
#define OMEGA_H_CONSTANT_DATA constexpr
#elif defined(_MSC_VER)
#define OMEGA_H_INLINE __forceinline
#define OMEGA_H_INLINE_BIG inline
Expand Down

0 comments on commit 5f30588

Please sign in to comment.