From 9b2f41330ac1a38ac2869d815cea1cdb538f0aee Mon Sep 17 00:00:00 2001 From: amduser Date: Fri, 24 Nov 2023 08:17:49 +0000 Subject: [PATCH] Fix comments --- include/mscclpp/fifo_device.hpp | 2 +- include/mscclpp/packet_device.hpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/include/mscclpp/fifo_device.hpp b/include/mscclpp/fifo_device.hpp index 9ef6fab9c..5806ca510 100644 --- 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 index 8fdcad91e..91246067e 100644 --- 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(®); - // 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(®);