Fix comments

microsoft · Nov 24, 2023 · 9b2f413 · 9b2f413
1 parent 2022df4
commit 9b2f413
Show file tree

Hide file tree

Showing 2 changed files with 3 additions and 3 deletions.
diff --git a/include/mscclpp/fifo_device.hpp b/include/mscclpp/fifo_device.hpp
@@ -68,7 +68,7 @@ struct FifoDeviceHandle {
 #if defined(MSCCLPP_DEVICE_CUDA)
     asm volatile("st.global.release.sys.v2.u64 [%0], {%1,%2};" ::"l"(triggerPtr), "l"(trigger.fst), "l"(trigger.snd));
 #else   // !defined(MSCCLPP_DEVICE_CUDA)
-    // TODO: revisit after changing clang++ to hipcc
+    // TODO: both atomic and clang built-ins are buggy here
     triggerPtr->fst = trigger.fst;
     triggerPtr->snd = trigger.snd;
 #endif  // !defined(MSCCLPP_DEVICE_CUDA)

diff --git a/include/mscclpp/packet_device.hpp b/include/mscclpp/packet_device.hpp
@@ -39,7 +39,7 @@ union alignas(16) LLPacket {
 #else  // !defined(MSCCLPP_DEVICE_CUDA)
     uint4 reg = make_uint4(val1, flag, val2, flag);
     ulonglong2* p = reinterpret_cast<ulonglong2*>(&reg);
-    // TODO: revisit after changing clang++ to hipcc
+    // TODO: clang built-ins are buggy here
     atomicStore(&(raw_.x), p->x, memoryOrderRelaxed);
     atomicStore(&(raw_.y), p->y, memoryOrderRelaxed);
 #endif
@@ -63,7 +63,7 @@ union alignas(16) LLPacket {
     return (flag1 != flag) || (flag2 != flag);
 #else  // !defined(MSCCLPP_DEVICE_CUDA)
     ulonglong2 reg;
-    // TODO: revisit after changing clang++ to hipcc
+    // TODO: clang built-ins are buggy here
     reg.x = atomicLoad(&(raw_.x), memoryOrderRelaxed);
     reg.y = atomicLoad(&(raw_.y), memoryOrderRelaxed);
     uint4* ptr = reinterpret_cast<uint4*>(&reg);