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

finished updater_gpu_coordinate.cu

Browse Source
This commit is contained in:
amdsc21 2023-03-09 20:53:54 +01:00
parent c875f0425f
commit 5044713388
2 changed files with 40 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

36
src/linear/updater_gpu_coordinate.cu
View File

@ -11,7 +11,13 @@
#include "coordinate_common.h"
#include "../common/common.h"
#if defined(XGBOOST_USE_CUDA)
#include "../common/device_helpers.cuh"
#elif defined(XGBOOST_USE_HIP)
#include "../common/device_helpers.hip.h"
#endif
#include "../common/timer.h"
#include "./param.h"
@ -60,7 +66,12 @@ class GPUCoordinateUpdater : public LinearUpdater { // NOLINT
return;
}
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaSetDevice(ctx_->gpu_id));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipSetDevice(ctx_->gpu_id));
#endif
// The begin and end indices for the section of each column associated with
// this device
std::vector<std::pair<bst_uint, bst_uint>> column_segments;
@ -86,10 +97,18 @@ class GPUCoordinateUpdater : public LinearUpdater { // NOLINT
for (size_t fidx = 0; fidx < batch.Size(); fidx++) {
auto col = page[fidx];
auto seg = column_segments[fidx];
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaMemcpy(
data_.data().get() + row_ptr_[fidx],
col.data() + seg.first,
sizeof(Entry) * (seg.second - seg.first), cudaMemcpyHostToDevice));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipMemcpy(
data_.data().get() + row_ptr_[fidx],
col.data() + seg.first,
sizeof(Entry) * (seg.second - seg.first), hipMemcpyHostToDevice));
#endif
}
}
@ -170,7 +189,12 @@ class GPUCoordinateUpdater : public LinearUpdater { // NOLINT
// This needs to be public because of the __device__ lambda.
GradientPair GetBiasGradient(int group_idx, int num_group) {
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaSetDevice(ctx_->gpu_id));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipSetDevice(ctx_->gpu_id));
#endif
auto counting = thrust::make_counting_iterator(0ull);
auto f = [=] __device__(size_t idx) {
return idx * num_group + group_idx;
@ -194,7 +218,12 @@ class GPUCoordinateUpdater : public LinearUpdater { // NOLINT
// This needs to be public because of the __device__ lambda.
GradientPair GetGradient(int group_idx, int num_group, int fidx) {
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaSetDevice(ctx_->gpu_id));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipSetDevice(ctx_->gpu_id));
#endif
common::Span<xgboost::Entry> d_col = dh::ToSpan(data_).subspan(row_ptr_[fidx]);
size_t col_size = row_ptr_[fidx + 1] - row_ptr_[fidx];
common::Span<GradientPair> d_gpair = dh::ToSpan(gpair_);
@ -227,10 +256,17 @@ class GPUCoordinateUpdater : public LinearUpdater { // NOLINT
}
void UpdateGpair(const std::vector<GradientPair> &host_gpair) {
#if defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaMemcpyAsync(
gpair_.data().get(),
host_gpair.data(),
gpair_.size() * sizeof(GradientPair), cudaMemcpyHostToDevice));
#elif defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipMemcpyAsync(
gpair_.data().get(),
host_gpair.data(),
gpair_.size() * sizeof(GradientPair), hipMemcpyHostToDevice));
#endif
}
// training parameter

4
src/linear/updater_gpu_coordinate.hip
View File

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