From 5516bdbb6be2b307754053898db6b940c83cb011 Mon Sep 17 00:00:00 2001
From: Binyang Li <binyli@microsoft.com>
Date: Fri, 8 May 2026 04:22:50 +0000
Subject: [PATCH] fix

---
 include/mscclpp/switch_channel_device.hpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/include/mscclpp/switch_channel_device.hpp b/include/mscclpp/switch_channel_device.hpp
index 841b7f32..4e0396dd 100644
--- a/include/mscclpp/switch_channel_device.hpp
+++ b/include/mscclpp/switch_channel_device.hpp
@@ -155,7 +155,7 @@ struct SwitchChannelDeviceHandle {
     }
 #endif
     else {
-      static_assert(dependentFalse<VectorType>, "Not supported type");
+      assert(false && "Unsupported vector type for multimemLoadReduce");
     }
     return val;
   };
@@ -219,7 +219,7 @@ struct SwitchChannelDeviceHandle {
                    "r"(val.words[1]), "r"(val.words[2]), "r"(val.words[3])
                    : "memory");
     } else {
-      static_assert(dependentFalse<VectorType>, "Not supported type");
+      assert(false && "Unsupported vector type for multimemStore");
     }
   };
 
@@ -244,7 +244,7 @@ struct SwitchChannelDeviceHandle {
     } else if constexpr (std::is_same_v<TValue, uint1> && std::is_same_v<T, __half2>) {
       asm volatile("multimem.red.relaxed.sys.global.add.f16x2 [%0], {%1};" ::"l"(ptr), "r"(val.x) : "memory");
     } else {
-      static_assert(dependentFalse<T>, "Not supported type");
+      assert(false && "Unsupported vector type for multimemStoreReduce");
     }
   };
 #endif  // defined(MSCCLPP_DEVICE_CUDA)