support cuda 10.1 (#4223)

* support cuda 10.1

* add cuda 10.1 to jenkins build matrix

src/common/host_device_vector.cu

@@ -53,7 +53,8 @@ struct HostDeviceVectorImpl {
       if (vec_ == nullptr) { vec_ = vec; }
       CHECK_EQ(vec, vec_);
       device_ = device;
-      LazyResize(vec_->Size());
+      // TODO(rongou): remove pointer dereference once CUDA 10.1 is fixed.
+      LazyResize((*vec_).Size());
       perm_d_ = vec_->perm_h_.Complementary();
     }
 

src/tree/updater_gpu.cu

@@ -140,7 +140,11 @@ __global__ void CubScanByKeyL1(
   // in order to pass on the partial scan values.
   // this statement MUST appear before the checks below!
   // else, the result of this shuffle operation will be undefined
+#if (__CUDACC_VER_MAJOR__ >= 9)
+  int previousKey = __shfl_up_sync(0xFFFFFFFF, myKey, 1);
+#else
   int previousKey = __shfl_up(myKey, 1);
+#endif
   // Collectively compute the block-wide exclusive prefix sum
   BlockScan(temp_storage)
       .ExclusiveScan(threadData, threadData, rootPair, AddByKey());