diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst
index 481ed39230813..8f473c21e1918 100644
--- a/clang/docs/HIPSupport.rst
+++ b/clang/docs/HIPSupport.rst
@@ -286,6 +286,26 @@ Example Usage
       basePtr->virtualFunction(); // Allowed since obj is constructed in device code
    }
 
+Host and Device Attributes of Default Destructors
+===================================================
+
+If a default destructor does not have explicit host or device attributes,
+clang infers these attributes based on the destructors of its data members
+and base classes. If any conflicts are detected among these destructors,
+clang diagnoses the issue. Otherwise, clang adds an implicit host or device
+attribute according to whether the data members's and base classes's
+destructors can execute on the host or device side.
+
+For explicit template classes with virtual destructors, which must be emitted,
+the inference adopts a conservative approach. In this case, implicit host or
+device attributes from member and base class destructors are ignored. This
+precaution is necessary because, although a constexpr destructor carries
+implicit host or device attributes, a constexpr function may call a
+non-constexpr function, which is by default a host function.
+
+Users can override the inferred host and device attributes of default
+destructors by adding explicit host and device attributes to them.
+
 C++ Standard Parallelism Offload Support: Compiler And Runtime
 ==============================================================
 
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 473956c37bb51..0e5fc5e1a40b4 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -437,7 +437,9 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
     if (!SMOR.getMethod())
       continue;
 
-    CUDAFunctionTarget BaseMethodTarget = IdentifyTarget(SMOR.getMethod());
+    CUDAFunctionTarget BaseMethodTarget =
+        IdentifyTarget(SMOR.getMethod(), IsExpVDtor);
+
     if (!InferredTarget) {
       InferredTarget = BaseMethodTarget;
     } else {
@@ -481,7 +483,9 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
     if (!SMOR.getMethod())
       continue;
 
-    CUDAFunctionTarget FieldMethodTarget = IdentifyTarget(SMOR.getMethod());
+    CUDAFunctionTarget FieldMethodTarget =
+        IdentifyTarget(SMOR.getMethod(), IsExpVDtor);
+
     if (!InferredTarget) {
       InferredTarget = FieldMethodTarget;
     } else {
diff --git a/clang/test/SemaCUDA/dtor.cu b/clang/test/SemaCUDA/dtor.cu
index c266e51f5c29e..cc37837e70791 100644
--- a/clang/test/SemaCUDA/dtor.cu
+++ b/clang/test/SemaCUDA/dtor.cu
@@ -32,22 +32,24 @@ public:
 template class B<float>;
 }
 
-// The implicit host/device attrs of virtual dtor B<float>::~B() is inferred to
-// have implicit device attr since dtors of its members and parent classes can
-// be executed on device. This causes a diagnostic since B<float>::~B() must
-// be emitted, and it eventually causes host_fun() called on device side.
+// The implicit host/device attrs of virtual dtor ~B() should be
+// conservatively inferred, where constexpr member dtor's should
+// not be considered device since they may call host functions.
+// Therefore B<float>::~B() should not have implicit device attr.
+// However C<float>::~C() should have implicit device attr since
+// it is trivial.
 namespace ExplicitInstantiationDtorNoAttr {
-void host_fun() // dev-note {{'host_fun' declared here}}
+void host_fun()
 {}
 
 template <unsigned>
 constexpr void hd_fun() {
-  host_fun(); // dev-error{{reference to __host__ function 'host_fun' in __host__ __device__ function}}
+  host_fun();
 }
 
 struct A {
-  constexpr ~A() { // dev-note {{called by '~B'}}
-     hd_fun<8>(); // dev-note {{called by '~A'}}
+  constexpr ~A() {
+     hd_fun<8>();
   }
 };