From 654bcfa6ba403cc0c40f31cf2224b3f2fa524569 Mon Sep 17 00:00:00 2001
From: Binyang Li <binyli@microsoft.com>
Date: Fri, 8 May 2026 03:54:32 +0000
Subject: [PATCH] update

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

diff --git a/include/mscclpp/switch_channel_device.hpp b/include/mscclpp/switch_channel_device.hpp
index 7b749f7a..841b7f32 100644
--- a/include/mscclpp/switch_channel_device.hpp
+++ b/include/mscclpp/switch_channel_device.hpp
@@ -84,7 +84,9 @@ struct SwitchChannelDeviceHandle {
           : "=r"(val.words[0]), "=r"(val.words[1]), "=r"(val.words[2]), "=r"(val.words[3])
           : "l"(ptr)
           : "memory");
-    } else if constexpr (std::is_same_v<VectorType, f8_e4m3x4>) {
+    }
+#if (defined(__CUDA_ARCH_SPECIFIC__) || defined(__CUDA_ARCH_FAMILY_SPECIFIC__)) && (__CUDA_ARCH__ >= 1000)
+    else if constexpr (std::is_same_v<VectorType, f8_e4m3x4>) {
       if constexpr (std::is_same_v<AccumT, __half>) {
         asm("multimem.ld_reduce.relaxed.sys.global.add.acc::f16.e4m3x4 %0, [%1];"
             : "=r"(val.words[0])
@@ -150,7 +152,9 @@ struct SwitchChannelDeviceHandle {
             : "l"(ptr)
             : "memory");
       }
-    } else {
+    }
+#endif
+    else {
       static_assert(dependentFalse<VectorType>, "Not supported type");
     }
     return val;