minor fix

microsoft · Oct 20, 2023 · f8899ef · f8899ef
1 parent 6f43282
commit f8899ef
Showing 1 changed file with 2 additions and 2 deletions.
diff --git a/include/mscclpp/sm_channel_device.hpp b/include/mscclpp/sm_channel_device.hpp
@@ -97,7 +97,7 @@ __forceinline__ __device__ void store<longlong2>(longlong2* p, const longlong2&
 template <>
 __forceinline__ __device__ void load<int4>(int4& v, const int4* p) {
   asm volatile("ld.volatile.global.v4.u32 {%0,%1,%2,%3}, [%4];"
-               : "=r"(v.w), "=r"(v.x), "=r"(v.y), "=r"(v.z)
+               : "=r"(v.x), "=r"(v.y), "=r"(v.z), "=r"(v.w)
                : "l"(p)
                : "memory");
 }
@@ -106,7 +106,7 @@ template <>
 __forceinline__ __device__ void store<int4>(int4* p, const int4& v) {
   asm volatile("st.volatile.global.v4.u32 [%0], {%1,%2,%3,%4};"
                :
-               : "l"(p), "r"(v.w), "r"(v.x), "r"(v.y), "r"(v.z)
+               : "l"(p), "r"(v.x), "r"(v.y), "r"(v.z), "r"(v.w)
                : "memory");
 }