Hendrik/xgboost
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

finish iterative_dmatrix.cu

Browse Source
This commit is contained in:
amdsc21 2023-03-10 03:47:00 +01:00
parent ec9f500a49
commit 49732359ef
4 changed files with 45 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
src/common/hist_util.cuh
View File

@ -12,7 +12,13 @@
#include <cstddef> // for size_t #include <cstddef> // for size_t
#include "../data/device_adapter.cuh" #include "../data/device_adapter.cuh"
#if defined(XGBOOST_USE_CUDA)
#include "device_helpers.cuh" #include "device_helpers.cuh"
#elif defined(XGBOOST_USE_HIP)
#include "device_helpers.hip.h"
#endif
#include "hist_util.h" #include "hist_util.h"
#include "quantile.cuh" #include "quantile.cuh"
#include "timer.h" #include "timer.h"

6
src/common/quantile.cuh
View File

@ -5,7 +5,13 @@
#include "xgboost/span.h" #include "xgboost/span.h"
#include "xgboost/data.h" #include "xgboost/data.h"
#if defined(XGBOOST_USE_CUDA)
#include "device_helpers.cuh" #include "device_helpers.cuh"
#elif defined(XGBOOST_USE_HIP)
#include "device_helpers.hip.h"
#endif
#include "quantile.h" #include "quantile.h"
#include "timer.h" #include "timer.h"
#include "categorical.h" #include "categorical.h"

29
src/data/iterative_dmatrix.cu
View File

@ -44,7 +44,13 @@ void IterativeDMatrix::InitFromCUDA(DataIterHandle iter_handle, float missing,
bst_feature_t cols = 0; bst_feature_t cols = 0;
int32_t current_device; int32_t current_device;
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaGetDevice(&current_device)); dh::safe_cuda(cudaGetDevice(&current_device));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipGetDevice(&current_device));
#endif
auto get_device = [&]() -> int32_t { auto get_device = [&]() -> int32_t {
int32_t d = (ctx_.gpu_id == Context::kCpuId) ? current_device : ctx_.gpu_id; int32_t d = (ctx_.gpu_id == Context::kCpuId) ? current_device : ctx_.gpu_id;
CHECK_NE(d, Context::kCpuId); CHECK_NE(d, Context::kCpuId);
@ -59,7 +65,13 @@ void IterativeDMatrix::InitFromCUDA(DataIterHandle iter_handle, float missing,
// We use do while here as the first batch is fetched in ctor // We use do while here as the first batch is fetched in ctor
ctx_.gpu_id = proxy->DeviceIdx(); ctx_.gpu_id = proxy->DeviceIdx();
CHECK_LT(ctx_.gpu_id, common::AllVisibleGPUs()); CHECK_LT(ctx_.gpu_id, common::AllVisibleGPUs());
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaSetDevice(get_device())); dh::safe_cuda(cudaSetDevice(get_device()));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipSetDevice(get_device()));
#endif
if (cols == 0) { if (cols == 0) {
cols = num_cols(); cols = num_cols();
collective::Allreduce<collective::Operation::kMax>(&cols, 1); collective::Allreduce<collective::Operation::kMax>(&cols, 1);
@ -83,7 +95,13 @@ void IterativeDMatrix::InitFromCUDA(DataIterHandle iter_handle, float missing,
row_stride = std::max(row_stride, Dispatch(proxy, [=](auto const& value) { row_stride = std::max(row_stride, Dispatch(proxy, [=](auto const& value) {
return GetRowCounts(value, row_counts_span, get_device(), missing); return GetRowCounts(value, row_counts_span, get_device(), missing);
})); }));
#if defined(XGBOOST_USE_CUDA)
nnz += thrust::reduce(thrust::cuda::par(alloc), row_counts.begin(), row_counts.end()); nnz += thrust::reduce(thrust::cuda::par(alloc), row_counts.begin(), row_counts.end());
#elif defined(XGBOOST_USE_HIP)
nnz += thrust::reduce(thrust::hip::par(alloc), row_counts.begin(), row_counts.end());
#endif
batches++; batches++;
} while (iter.Next()); } while (iter.Next());
iter.Reset(); iter.Reset();
@ -91,7 +109,12 @@ void IterativeDMatrix::InitFromCUDA(DataIterHandle iter_handle, float missing,
auto n_features = cols; auto n_features = cols;
CHECK_GE(n_features, 1) << "Data must has at least 1 column."; CHECK_GE(n_features, 1) << "Data must has at least 1 column.";
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaSetDevice(get_device())); dh::safe_cuda(cudaSetDevice(get_device()));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipSetDevice(get_device()));
#endif
if (!ref) { if (!ref) {
HostDeviceVector<FeatureType> ft; HostDeviceVector<FeatureType> ft;
common::SketchContainer final_sketch( common::SketchContainer final_sketch(
@ -130,7 +153,13 @@ void IterativeDMatrix::InitFromCUDA(DataIterHandle iter_handle, float missing,
size_t n_batches_for_verification = 0; size_t n_batches_for_verification = 0;
while (iter.Next()) { while (iter.Next()) {
init_page(); init_page();
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaSetDevice(get_device())); dh::safe_cuda(cudaSetDevice(get_device()));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipSetDevice(get_device()));
#endif
auto rows = num_rows(); auto rows = num_rows();
dh::caching_device_vector<size_t> row_counts(rows + 1, 0); dh::caching_device_vector<size_t> row_counts(rows + 1, 0);
common::Span<size_t> row_counts_span(row_counts.data().get(), row_counts.size()); common::Span<size_t> row_counts_span(row_counts.data().get(), row_counts.size());

4
src/data/iterative_dmatrix.hip
View File

@ -0,0 +1,4 @@
#if defined(XGBOOST_USE_HIP)
#include "iterative_dmatrix.cu"
#endif