NCCL fixes (#592)

* Fix an allgather bug
* Fix the `nccl.h` installation path (now goes under `mscclpp/`)

Author: chhwang

apps/nccl/CMakeLists.txt

@@ -2,7 +2,7 @@
 # Licensed under the MIT license.
 
 file(GLOB_RECURSE SOURCES CONFIGURE_DEPENDS src/*)
-file(GLOB_RECURSE HEADERS CONFIGURE_DEPENDS include/nccl.h)
+file(GLOB_RECURSE HEADERS CONFIGURE_DEPENDS include/*.h)
 
 if(MSCCLPP_USE_ROCM)
     set_source_files_properties(${SOURCES} PROPERTIES LANGUAGE CXX)
@@ -30,12 +30,8 @@ target_link_libraries(mscclpp_nccl_static PUBLIC mscclpp_obj mscclpp_nccl_obj)
 set_target_properties(mscclpp_nccl_static PROPERTIES VERSION ${MSCCLPP_VERSION} SOVERSION ${MSCCLPP_SOVERSION})
 
 install(TARGETS mscclpp_nccl_obj
-    FILE_SET HEADERS DESTINATION ${INSTALL_PREFIX}/include)
+    FILE_SET HEADERS DESTINATION ${INSTALL_PREFIX})
 install(TARGETS mscclpp_nccl
     LIBRARY DESTINATION ${INSTALL_PREFIX}/lib)
 install(TARGETS mscclpp_nccl_static
     ARCHIVE DESTINATION ${INSTALL_PREFIX}/lib)
-
-if(MSCCLPP_BUILD_TESTS)
-    add_subdirectory(test)
-endif()

apps/nccl/include/nccl.h renamed to apps/nccl/include/mscclpp/nccl.h

@@ -104,7 +104,7 @@ __global__ void __launch_bounds__(1024, 1)
   deviceSyncer.sync(gridDim.x);
 
   if (threadIdx.x < nPeer) {
-    memChans[threadIdx.x].relaxedSignal();
+    memChans[threadIdx.x].signal();
     memChans[threadIdx.x].wait();
   }
 }

apps/nccl/src/allgather.hpp

@@ -20,12 +20,12 @@
 #include <mscclpp/npkit/npkit.hpp>
 #endif
 #include <dlfcn.h>
+#include <mscclpp/nccl.h>
 
 #include "allgather.hpp"
 #include "allreduce.hpp"
 #include "broadcast.hpp"
 #include "debug.h"
-#include "nccl.h"
 
 #define NCCL_API extern "C" __attribute__((visibility("default")))