Enable -Wdeprecated-copy-with-user-provided-copy (#1643)

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)
 

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;

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;
     }