Address NVCC warning #20012-D (#528)

Author: chhwang

include/mscclpp/concurrency_device.hpp

@@ -15,10 +15,10 @@ namespace mscclpp {
 struct DeviceSyncer {
  public:
   /// Construct a new DeviceSyncer object.
-  DeviceSyncer() = default;
+  MSCCLPP_INLINE DeviceSyncer() = default;
 
   /// Destroy the DeviceSyncer object.
-  ~DeviceSyncer() = default;
+  MSCCLPP_INLINE ~DeviceSyncer() = default;
 
 #if defined(MSCCLPP_DEVICE_COMPILE)
   /// Synchronize all threads inside a kernel. Guarantee that all previous work of all threads in cooperating blocks is

include/mscclpp/device.hpp

@@ -11,6 +11,7 @@
 #if (defined(__NVCC__) || defined(__HIP_PLATFORM_AMD__))
 
 #define MSCCLPP_DEVICE_COMPILE
+#define MSCCLPP_INLINE __forceinline__
 #define MSCCLPP_DEVICE_INLINE __forceinline__ __device__
 #define MSCCLPP_HOST_DEVICE_INLINE __forceinline__ __host__ __device__
 #if defined(__HIP_PLATFORM_AMD__)
@@ -22,6 +23,7 @@
 #else  // !(defined(__NVCC__) || defined(__HIP_PLATFORM_AMD__))
 
 #define MSCCLPP_HOST_COMPILE
+#define MSCCLPP_INLINE inline
 #define MSCCLPP_HOST_DEVICE_INLINE inline
 
 #endif  // !(defined(__NVCC__) || defined(__HIP_PLATFORM_AMD__))

include/mscclpp/memory_channel_device.hpp

@@ -15,7 +15,7 @@ namespace mscclpp {
 struct BaseMemoryChannelDeviceHandle {
   MemoryDevice2DeviceSemaphoreDeviceHandle semaphore_;
 
-  MSCCLPP_HOST_DEVICE_INLINE BaseMemoryChannelDeviceHandle() = default;
+  MSCCLPP_INLINE BaseMemoryChannelDeviceHandle() = default;
 
   MSCCLPP_HOST_DEVICE_INLINE BaseMemoryChannelDeviceHandle(MemoryDevice2DeviceSemaphoreDeviceHandle semaphore)
       : semaphore_(semaphore) {}
@@ -65,7 +65,7 @@ struct MemoryChannelDeviceHandle : public BaseMemoryChannelDeviceHandle {
   void* src_;
   void* packetBuffer_;
 
-  MSCCLPP_HOST_DEVICE_INLINE MemoryChannelDeviceHandle() = default;
+  MSCCLPP_INLINE MemoryChannelDeviceHandle() = default;
 
   MSCCLPP_HOST_DEVICE_INLINE MemoryChannelDeviceHandle(MemoryDevice2DeviceSemaphoreDeviceHandle semaphore, void* dst,
                                                        void* src, void* packetBuffer)

include/mscclpp/packet_device.hpp

@@ -29,7 +29,7 @@ union alignas(16) LL16Packet {
 #if defined(MSCCLPP_DEVICE_COMPILE)
   ulonglong2 raw_;
 
-  MSCCLPP_DEVICE_INLINE LL16Packet() = default;
+  MSCCLPP_INLINE LL16Packet() = default;
 
   MSCCLPP_DEVICE_INLINE LL16Packet(uint2 val, uint32_t flag) : data1(val.x), flag1(flag), data2(val.y), flag2(flag) {}
 
@@ -107,7 +107,7 @@ union alignas(8) LL8Packet {
   using Payload = uint32_t;
 #if defined(MSCCLPP_DEVICE_COMPILE)
 
-  MSCCLPP_DEVICE_INLINE LL8Packet() = default;
+  MSCCLPP_INLINE LL8Packet() = default;
 
   MSCCLPP_DEVICE_INLINE LL8Packet(uint32_t val, uint32_t flag) : data(val), flag(flag) {}
 

include/mscclpp/port_channel_device.hpp

@@ -49,7 +49,7 @@ union ChannelTrigger {
 
 #if defined(MSCCLPP_DEVICE_COMPILE)
   /// Default constructor.
-  MSCCLPP_DEVICE_INLINE ChannelTrigger() = default;
+  MSCCLPP_INLINE ChannelTrigger() = default;
 
   /// Copy constructor.
   MSCCLPP_DEVICE_INLINE ChannelTrigger(ProxyTrigger value) : value(value) {}
@@ -92,7 +92,7 @@ struct BasePortChannelDeviceHandle {
   // can produce for and the sole proxy thread consumes it.
   FifoDeviceHandle fifo_;
 
-  MSCCLPP_HOST_DEVICE_INLINE BasePortChannelDeviceHandle() = default;
+  MSCCLPP_INLINE BasePortChannelDeviceHandle() = default;
 
   MSCCLPP_HOST_DEVICE_INLINE BasePortChannelDeviceHandle(SemaphoreId semaphoreId,
                                                          Host2DeviceSemaphoreDeviceHandle semaphore,
@@ -190,7 +190,7 @@ struct PortChannelDeviceHandle : public BasePortChannelDeviceHandle {
   MemoryId dst_;
   MemoryId src_;
 
-  MSCCLPP_HOST_DEVICE_INLINE PortChannelDeviceHandle() = default;
+  MSCCLPP_INLINE PortChannelDeviceHandle() = default;
 
   MSCCLPP_HOST_DEVICE_INLINE PortChannelDeviceHandle(SemaphoreId semaphoreId,
                                                      Host2DeviceSemaphoreDeviceHandle semaphore, FifoDeviceHandle fifo,