_LIBCUDACXX_BEGIN_NAMESPACE_CUDA

namespace __detail_ap {

  template <typename _Property>
  __device__
  void* __associate_address_space(void* __ptr, _Property __prop) {
    if (std::is_same<_Property, access_property::shared>::value == true) {
      bool __b = __isShared(__ptr);
#if _LIBCUDACXX_DEBUG_LEVEL >= 2
      _LIBCUDACXX_DEBUG_ASSERT(__b == true);
#endif
      __builtin_assume(__b);
    } else if (std::is_same<_Property, access_property::global>::value == true ||
               std::is_same<_Property, access_property::normal>::value == true ||
               std::is_same<_Property, access_property::persisting>::value == true ||
               std::is_same<_Property, access_property::streaming>::value == true ||
               std::is_same<_Property, access_property>::value) {
      bool __b = __isGlobal(__ptr);
#if _LIBCUDACXX_DEBUG_LEVEL >= 2
      _LIBCUDACXX_DEBUG_ASSERT(__b == true);
#endif
      __builtin_assume(__b);
    }

    return __ptr;
  }

  template <typename __Prop>
  __device__
  void* __associate_descriptor(void* __ptr, __Prop __prop) {
    return __associate_descriptor(__ptr, static_cast<std::uint64_t>(access_property(__prop)));
  }

  template <>
  __device__
  void* __associate_descriptor(void* __ptr, std::uint64_t __prop) {
#if __CUDA_ARCH__ >= 800
    return __nv_associate_access_property(__ptr, __prop);
#else
    return __ptr;
#endif
  }

  template<>
  __device__
  void* __associate_descriptor(void* __ptr, access_property::shared) {
    return __ptr;
  }

  template<typename _Type, typename _Property>
  __host__ __device__
  _Type* __associate(_Type* __ptr, _Property __prop) {
#ifdef __CUDA_ARCH__
    return static_cast<_Type*>(__associate_descriptor(__associate_address_space(__ptr, __prop), __prop));
#else
    return __ptr;
#endif
  }


  template<typename _Property>
  class __annotated_ptr_base {
    using __error = typename _Property::__unknown_access_property_type;
  };

  template<>
  class __annotated_ptr_base<access_property::shared> {
    protected:
      static constexpr std::uint64_t __prop = 0;

      constexpr __annotated_ptr_base() noexcept = default;
      constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
      __device__ void* __apply_prop(void* __p) const {
        return __associate(__p, access_property::shared{});
      }
  };

  template<>
  class __annotated_ptr_base<access_property::global> {
    protected:
      static constexpr std::uint64_t __prop = __sm_80::__interleave_normal();

      constexpr __annotated_ptr_base() noexcept = default;
      constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
      __device__ void* __apply_prop(void* __p) const {
        return __associate(__p, access_property::global{});
      }
  };

  template<>
  class __annotated_ptr_base<access_property::normal> {
    protected:
      static constexpr std::uint64_t __prop = __sm_80::__interleave_normal_demote();

      constexpr __annotated_ptr_base() noexcept = default;
      constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
      __device__ void* __apply_prop(void* __p) const {
        return __associate(__p, access_property::normal{});
      }
  };

  template<>
  class __annotated_ptr_base<access_property::persisting> {
    protected:
      static constexpr std::uint64_t __prop = __sm_80::__interleave_persisting();

      constexpr __annotated_ptr_base() noexcept = default;
      constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
      __device__ void* __apply_prop(void* __p) const {
        return __associate(__p, access_property::persisting{});
      }
  };

  template<>
  class __annotated_ptr_base<access_property::streaming> {
    protected:
      static constexpr std::uint64_t __prop = __sm_80::__interleave_streaming();

      constexpr __annotated_ptr_base() noexcept = default;
      constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
      __device__ void* __apply_prop(void* __p) const {
        return __associate(__p, access_property::streaming{});
      }
  };

  template<>
  class __annotated_ptr_base<access_property> {
    protected:
      std::uint64_t __prop;

      __host__ __device__ constexpr __annotated_ptr_base() noexcept : __prop(access_property()) {}
      __host__ __device__ constexpr __annotated_ptr_base(std::uint64_t __property) noexcept : __prop(__property) {}
      constexpr __annotated_ptr_base(__annotated_ptr_base const&) = default;
      __device__ void* __apply_prop(void* __p) const {
        return __associate(__p, __prop);
      }
  };
} // namespace __detail_ap

_LIBCUDACXX_END_NAMESPACE_CUDA
