Merge pull request #274 from gonzalobg/bugfix/annotated_ptr_odr

ODR bugfixes for annotated_ptr and access_property
NVIDIA · May 10, 2022 · 0edbb68 · 0edbb68
2 parents 017cb64 + f2f08bf
commit 0edbb68
Show file tree

Hide file tree

Showing 3 changed files with 11 additions and 8 deletions.
diff --git a/include/cuda/annotated_ptr b/include/cuda/annotated_ptr
@@ -170,6 +170,7 @@ void apply_access_property(const volatile void* __ptr, const _Shape __shape, acc
 #endif
 }
 
+inline
 __host__ __device__
 void discard_memory(volatile void* __ptr, std::size_t __nbytes) noexcept {
 #if __CUDA_ARCH__ >= 800

diff --git a/include/cuda/std/detail/__access_property b/include/cuda/std/detail/__access_property
@@ -155,6 +155,7 @@ namespace __detail_ap {
           std::uint64_t(__ap_reserved3)           << 63;
       }
 
+      inline
       __host__ __device__
       std::uint64_t __get_descriptor_non_cexpr() const noexcept { return *reinterpret_cast<const std::uint64_t*>(this); }
 
@@ -256,6 +257,7 @@ namespace __detail_ap {
           std::uint64_t(__ap_reserved2)           << 63;
       }
 
+      inline
       __host__ __device__
       std::uint64_t __get_descriptor_non_cexpr() const noexcept { return *reinterpret_cast<const std::uint64_t*>(this); }
 

diff --git a/include/cuda/std/detail/__annotated_ptr b/include/cuda/std/detail/__annotated_ptr
@@ -86,7 +86,7 @@ namespace __detail_ap {
   }
 
   template <>
-  __device__
+  inline __device__
   void* __associate_descriptor(void* __ptr, std::uint64_t __prop) {
 #if __CUDA_ARCH__ >= 800
     return __nv_associate_access_property(__ptr, __prop);
@@ -96,7 +96,7 @@ namespace __detail_ap {
   }
 
   template<>
-  __device__
+  inline __device__
   void* __associate_descriptor(void* __ptr, access_property::shared) {
     return __ptr;
   }
@@ -125,7 +125,7 @@ namespace __detail_ap {
       constexpr __annotated_ptr_base() noexcept = default;
       constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
       _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __annotated_ptr_base& operator=(const __annotated_ptr_base&) = default;
-      __device__ void* __apply_prop(void* __p) const {
+      inline __device__ void* __apply_prop(void* __p) const {
         return __associate(__p, access_property::shared{});
       }
   };
@@ -138,7 +138,7 @@ namespace __detail_ap {
       constexpr __annotated_ptr_base() noexcept = default;
       constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
       _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __annotated_ptr_base& operator=(const __annotated_ptr_base&) = default;
-      __device__ void* __apply_prop(void* __p) const {
+      inline __device__ void* __apply_prop(void* __p) const {
         return __associate(__p, access_property::global{});
       }
   };
@@ -151,7 +151,7 @@ namespace __detail_ap {
       constexpr __annotated_ptr_base() noexcept = default;
       constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
       _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __annotated_ptr_base& operator=(const __annotated_ptr_base&) = default;
-      __device__ void* __apply_prop(void* __p) const {
+      inline __device__ void* __apply_prop(void* __p) const {
         return __associate(__p, access_property::normal{});
       }
   };
@@ -164,7 +164,7 @@ namespace __detail_ap {
       constexpr __annotated_ptr_base() noexcept = default;
       constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
       _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __annotated_ptr_base& operator=(const __annotated_ptr_base&) = default;
-      __device__ void* __apply_prop(void* __p) const {
+      inline __device__ void* __apply_prop(void* __p) const {
         return __associate(__p, access_property::persisting{});
       }
   };
@@ -177,7 +177,7 @@ namespace __detail_ap {
       constexpr __annotated_ptr_base() noexcept = default;
       constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
       _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __annotated_ptr_base& operator=(const __annotated_ptr_base&) = default;
-      __device__ void* __apply_prop(void* __p) const {
+      inline __device__ void* __apply_prop(void* __p) const {
         return __associate(__p, access_property::streaming{});
       }
   };
@@ -191,7 +191,7 @@ namespace __detail_ap {
       __host__ __device__ constexpr __annotated_ptr_base(std::uint64_t __property) noexcept : __prop(__property) {}
       constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
       _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __annotated_ptr_base& operator=(const __annotated_ptr_base&) = default;
-      __device__ void* __apply_prop(void* __p) const {
+      inline __device__ void* __apply_prop(void* __p) const {
         return __associate(__p, __prop);
       }
   };