From 24ea2ef6dd90cbc74fccd6790ea8a359dc596497 Mon Sep 17 00:00:00 2001
From: Wenkai Du <wenkai.du@amd.com>
Date: Fri, 8 May 2020 15:53:51 +0000
Subject: [PATCH] Set flags when calling hipExtLaunchMultiKernelMultiDevice in
 hip-clang

---
 src/enqueue.cc | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/src/enqueue.cc b/src/enqueue.cc
index d9205ac516..3786ab2398 100644
--- a/src/enqueue.cc
+++ b/src/enqueue.cc
@@ -74,7 +74,11 @@ ncclResult_t ncclLaunchCooperativeKernelMultiDevice(hipLaunchParams *paramsList,
   if (cgMode & 0x01) {
     CUDACHECK(hipExtLaunchMultiKernelMultiDevice(paramsList, numDevices,
             // These flags are to reduce the latency of using this API
+#if __HIP__
+            hipCooperativeLaunchMultiDeviceNoPreSync|hipCooperativeLaunchMultiDeviceNoPostSync));
+#else
             0));
+#endif
     return ncclSuccess;
   }
   int savedDev;