diff --git a/CMakeLists.txt b/CMakeLists.txt
index 6e9e993823..ea901485c1 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -836,6 +836,7 @@ endif()
 target_compile_options(rccl PRIVATE -Werror=uninitialized)
 target_compile_options(rccl PRIVATE -Werror=sometimes-uninitialized)
 target_compile_options(rccl PRIVATE -Wall)
+target_compile_options(rccl PRIVATE -Werror=deprecated-copy-with-user-provided-copy)
 target_compile_options(rccl PRIVATE -Wno-format-nonliteral)
 target_compile_options(rccl PRIVATE -fgpu-rdc)
 
diff --git a/src/device/op128.h b/src/device/op128.h
index 04e410cf7d..647360ac5f 100644
--- a/src/device/op128.h
+++ b/src/device/op128.h
@@ -125,7 +125,11 @@ union alignas(16) BytePack<16> {
   uint64_t u64[2];
   ulong2 ul2, native;
 #if !defined(USE_INDIRECT_FUNCTION_CALL) || defined(__gfx942__) || defined(__gfx950__)
-  inline __device__ BytePack<16>& operator=(BytePack<16> other) {
+  inline __device__ BytePack<16>() = default;
+  inline __device__ BytePack<16>(const BytePack<16>& other) {
+    *this = other;
+  }
+  inline __device__ BytePack<16>& operator=(const BytePack<16>& other) {
     u64[0] = other.u64[0];
     u64[1] = other.u64[1];
     return *this;
diff --git a/src/include/rccl_float8.h b/src/include/rccl_float8.h
index 5db894394a..91e08f96d6 100644
--- a/src/include/rccl_float8.h
+++ b/src/include/rccl_float8.h
@@ -393,6 +393,8 @@ struct rccl_float8
     // default constructor
     HIP_HOST_DEVICE rccl_float8() = default;
 
+    constexpr inline HIP_HOST_DEVICE rccl_float8(const rccl_float8& a) : data(a.data) {}
+
 #if defined(__gfx942__) || defined(__gfx950__)
     // device specific optimized F8 down-conversion code
 
@@ -541,7 +543,7 @@ struct rccl_float8
     }
 
     // assignment overloading only from the same F8 types
-    inline __host__ __device__ rccl_float8& operator=(const rccl_float8& a)
+    inline HIP_HOST_DEVICE rccl_float8& operator=(const rccl_float8& a)
     {
         data = a.data;
         return *this;
@@ -560,6 +562,8 @@ struct rccl_bfloat8
     // default constructor
     HIP_HOST_DEVICE rccl_bfloat8() = default;
 
+    constexpr inline HIP_HOST_DEVICE rccl_bfloat8(const rccl_bfloat8& a) : data(a.data) {}
+
 #if defined(__gfx942__) || defined(__gfx950__)
     // device specific optimized F8 down-conversion code
 
@@ -708,7 +712,7 @@ struct rccl_bfloat8
     }
 
     // assignment overloading only from the same F8 types
-    inline __host__ __device__ rccl_bfloat8& operator=(const rccl_bfloat8& a)
+    inline HIP_HOST_DEVICE rccl_bfloat8& operator=(const rccl_bfloat8& a)
     {
         data = a.data;
         return *this;
@@ -733,11 +737,11 @@ namespace std
     {
         return rccl_bfloat8(cosf(float(a)));
     }
-    __device__ __host__ constexpr rccl_float8 real(const rccl_float8& a)
+    HIP_HOST_DEVICE constexpr rccl_float8 real(const rccl_float8& a)
     {
         return a;
     }
-    __device__ __host__ constexpr rccl_bfloat8 real(const rccl_bfloat8& a)
+    HIP_HOST_DEVICE constexpr rccl_bfloat8 real(const rccl_bfloat8& a)
     {
         return a;
     }