From 346fc49514741fe97a2707a7daffcf4e1785e0bc Mon Sep 17 00:00:00 2001
From: Christian Sigg <csigg@google.com>
Date: Thu, 13 Dec 2018 16:09:12 +0100
Subject: [PATCH] Two temporary workarounds for cuda-clang issues.

---
 src/collectives/device/functions.cu | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/src/collectives/device/functions.cu b/src/collectives/device/functions.cu
index 16f1865d8e..1fb8108166 100644
--- a/src/collectives/device/functions.cu
+++ b/src/collectives/device/functions.cu
@@ -56,9 +56,17 @@
 
 // Must be consistent with the ncclFuncSet enum
 __device__ ncclKern_t ncclFuncs[ncclCollCount*ncclNumOps*ncclNumTypes*2] = {
+// Don't try to initialize the host shadow copy of this device-side global
+// variable. There is no host pointer to a device-side function, which
+// confuses clang. This will be fixed in the next clang release.
+#if __CUDA_ARCH__
   NCCL_FUNCS2B(ncclBroadcast),
   NCCL_FUNCS2A(ncclReduce),
   NCCL_FUNCS2B(ncclAllGather),
   NCCL_FUNCS2A(ncclReduceScatter),
   NCCL_FUNCS2A(ncclAllReduce)
+#endif
 };
+
+// Workaround for https://reviews.llvm.org/D55580
+__device__ void ncclWorkaroundClangD55580() {}