Optimisations for gpu_hist.

[RAMitchell]

• Use streams to overlap operations.

• Reduce redundant calls to cudaSetDevice().

• ColumnSampler now uses HostDeviceVector to prevent repeatedly copying feature vectors to the device.

[RAMitchell]

Running tests/benchmark/benchmark_tree.py 5 times on each branch and averaging the result. Old: 13.476s, New: 11.54s, 14.4% improvement

This was on my Windows development machine with a single 1080ti.

[trivialfis]

@RAMitchell Could you elaborate on why moving cudaSetDevice from device shard methods into ExecuteIndexShards?

And I recently did some small testings for CUDA APIs to see if it's possible to add some checks for issues like #4245, and other things like whether we are using the right device etc... cudaSetDevice is really cheap to call, I guess it's just a simple global integer variable indicating device id, redundant call sometimes makes debugging easier...

[trivialfis]

One more, please don't merge until the clang-tidy PR is done.

[RAMitchell]

It seemed a little more consistent to do it this way, each device shard performs work for strictly one GPU so if we set the device once before each shard performs work it should be safe.

I am seeing the cuda API calls come up on the profiler taking some nontrivial amount of time, I'm not actually sure if they take this long.

[hcho3]

@RAMitchell It looks like the multi-GPU test is failing due to memory error: https://xgboost-ci.net/blue/organizations/jenkins/xgboost/detail/PR-4248/5/pipeline/51#step-88-log-1317

[2019-03-12T22:18:57.198Z] tests/python-gpu/test_gpu_updaters.py::TestGPU::test_gpu_hist_mgpu Training on dataset: Boston
[2019-03-12T22:18:57.198Z] Using parameters: {'n_gpus': -1, 'eval_metric': 'rmse', 'max_bin': 2, 'objective': 'reg:linear', 'grow_policy': 'lossguide', 'tree_method': 'gpu_hist', 'max_depth': 2, 'max_leaves': 255}
[2019-03-12T22:18:57.582Z] terminate called after throwing an instance of 'thrust::system::system_error'
[2019-03-12T22:18:57.582Z]   what():  device free failed: an illegal memory access was encountered
[2019-03-12T22:18:58.403Z] tests/ci_build/test_mgpu.sh: line 7: 35100 Aborted                 (core dumped) pytest -v -s --fulltrace -m "(not slow) and mgpu" tests/python-gpu
[2019-03-12T22:18:59.338Z] Terminated

[trivialfis]

I am seeing the cuda API calls come up on the profiler taking some nontrivial amount of time

That's weird

[RAMitchell]

This PR is ready to be merged, just having some difficulty with R test failures on Travis.