4 #ifndef XGBOOST_COMMON_TRANSFORM_H_ 5 #define XGBOOST_COMMON_TRANSFORM_H_ 8 #include <dmlc/common.h> 13 #include <type_traits> 20 #if defined (__CUDACC__) 21 #include "device_helpers.cuh" 22 #endif // defined (__CUDACC__) 31 #if defined(__CUDACC__) 32 template <
typename Functor,
typename... SpanType>
33 __global__
void LaunchCUDAKernel(Functor _func,
Range _range,
35 for (
auto i : dh::GridStrideRange(*_range.
begin(), *_range.
end())) {
39 #endif // defined(__CUDACC__) 57 template <
bool CompiledWithCuda = WITH_CUDA()>
60 template <
typename Functor>
63 Evaluator(Functor func,
Range range,
int device,
bool shard) :
64 func_(func), range_{std::move(range)},
74 template <
typename... HDV>
75 void Eval(HDV... vectors)
const {
76 bool on_device = device_ >= 0;
79 LaunchCUDA(func_, vectors...);
81 LaunchCPU(func_, vectors...);
103 template <
typename T>
109 template <
typename T>
113 template <
typename Head,
typename... Rest>
114 void UnpackShard(
int device,
118 UnpackShard(device, _vectors...);
121 #if defined(__CUDACC__) 122 template <typename std::enable_if<CompiledWithCuda>::type* =
nullptr,
124 void LaunchCUDA(Functor _func, HDV*... _vectors)
const {
126 UnpackShard(device_, _vectors...);
128 size_t range_size = *range_.end() - *range_.begin();
133 size_t shard_size = range_size;
135 dh::safe_cuda(cudaSetDevice(device_));
137 static_cast<int>(
DivRoundUp(*(range_.end()), kBlockThreads));
141 detail::LaunchCUDAKernel<<<kGrids, kBlockThreads>>>(
142 _func, shard_range, UnpackHDVOnDevice(_vectors)...);
146 template <typename std::enable_if<!CompiledWithCuda>::type* =
nullptr,
148 void LaunchCUDA(Functor _func, HDV*... _vectors)
const {
149 LOG(FATAL) <<
"Not part of device code. WITH_CUDA: " <<
WITH_CUDA();
151 #endif // defined(__CUDACC__) 153 template <
typename... HDV>
154 void LaunchCPU(Functor func, HDV*... vectors)
const {
156 dmlc::OMPException omp_exc;
157 #pragma omp parallel for schedule(static) 158 for (
omp_ulong idx = 0; idx < end; ++idx) {
159 omp_exc.Run(func, idx, UnpackHDV(vectors)...);
188 template <
typename Functor>
189 static Evaluator<Functor>
Init(Functor func,
Range const range,
191 bool const shard =
true) {
192 return Evaluator<Functor> {func, std::move(range), device, shard};
199 #endif // XGBOOST_COMMON_TRANSFORM_H_ XGBOOST_DEVICE Iterator begin() const
Definition: common.h:121
std::size_t index_type
Definition: span.h:394
Definition: host_device_vector.h:85
The input data structure of xgboost.
T * HostPointer()
Definition: host_device_vector.h:107
dmlc::omp_ulong omp_ulong
define unsigned long for openmp loop
Definition: base.h:244
int64_t DifferenceType
Definition: common.h:85
A device-and-host vector abstraction layer.
XGBOOST_DEVICE Iterator end() const
Definition: common.h:122
common::Span< const T > ConstDeviceSpan() const
const T * ConstHostPointer() const
Definition: host_device_vector.h:108
span class implementation, based on ISO++20 span<T>. The interface should be the same.
Definition: span.h:115
common::Span< T > DeviceSpan()
XGBOOST_DEVICE T1 DivRoundUp(const T1 a, const T2 b)
Definition: common.h:76
void SetDevice(int device) const
#define WITH_CUDA()
Definition: common.h:27
constexpr size_t kBlockThreads
Definition: transform.h:27
namespace of xgboost
Definition: base.h:102