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

finish hist_util.cu

Browse Source
This commit is contained in:
amdsc21 2023-03-10 05:32:38 +01:00
parent 54b076b40f
commit 911a5d8a60
5 changed files with 73 additions and 40 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

35
src/common/common.h
View File

@ -39,6 +39,41 @@
#endif // defined(__CUDACC__)
namespace dh {
#if defined(__CUDACC__)
/*
* Error handling functions
*/
#define safe_cuda(ans) ThrowOnCudaError((ans), __FILE__, __LINE__)
inline cudaError_t ThrowOnCudaError(cudaError_t code, const char *file, int line)
{
if (code != cudaSuccess) {
LOG(FATAL) << thrust::system_error(code, thrust::cuda_category(),
std::string{file} + ": " + // NOLINT
std::to_string(line)).what();
}
return code;
}
#elif defined(__HIP_PLATFORM_AMD__)
/*
* Error handling functions
*/
#define safe_cuda(ans) ThrowOnCudaError((ans), __FILE__, __LINE__)
inline hipError_t ThrowOnCudaError(hipError_t code, const char *file, int line)
{
if (code != hipSuccess) {
LOG(FATAL) << thrust::system_error(code, thrust::hip_category(),
std::string{file} + ": " + // NOLINT
std::to_string(line)).what();
}
return code;
}
#endif
} // namespace dh
namespace xgboost {
namespace common {
/*!

35
src/common/device_helpers.hip.h
View File

@ -59,41 +59,6 @@
#endif // defined(XGBOOST_USE_RMM) && XGBOOST_USE_RMM == 1
namespace dh {
#if defined(__CUDACC__)
/*
* Error handling functions
*/
#define safe_cuda(ans) ThrowOnCudaError((ans), __FILE__, __LINE__)
inline cudaError_t ThrowOnCudaError(cudaError_t code, const char *file, int line)
{
if (code != cudaSuccess) {
LOG(FATAL) << thrust::system_error(code, thrust::cuda_category(),
std::string{file} + ": " + // NOLINT
std::to_string(line)).what();
}
return code;
}
#elif defined(__HIP_PLATFORM_AMD__)
/*
* Error handling functions
*/
#define safe_cuda(ans) ThrowOnCudaError((ans), __FILE__, __LINE__)
inline hipError_t ThrowOnCudaError(hipError_t code, const char *file, int line)
{
if (code != hipSuccess) {
LOG(FATAL) << thrust::system_error(code, thrust::hip_category(),
std::string{file} + ": " + // NOLINT
std::to_string(line)).what();
}
return code;
}
#endif
} // namespace dh
namespace dh {
// FIXME(jiamingy): Remove this once we get rid of cub submodule.

29
src/common/hist_util.cu
View File

@ -19,7 +19,13 @@
#include <vector>
#include "categorical.h"
#if defined(XGBOOST_USE_CUDA)
#include "device_helpers.cuh"
#elif defined(XGBOOST_USE_HIP)
#include "device_helpers.hip.h"
#endif
#include "hist_util.cuh"
#include "hist_util.h"
#include "math.h" // NOLINT
@ -113,18 +119,35 @@ void SortByWeight(dh::device_vector<float>* weights,
dh::device_vector<Entry>* sorted_entries) {
// Sort both entries and wegihts.
dh::XGBDeviceAllocator<char> alloc;
#if defined(XGBOOST_USE_CUDA)
thrust::sort_by_key(thrust::cuda::par(alloc), sorted_entries->begin(),
sorted_entries->end(), weights->begin(),
detail::EntryCompareOp());
#elif defined(XGBOOST_USE_HIP)
thrust::sort_by_key(thrust::hip::par(alloc), sorted_entries->begin(),
sorted_entries->end(), weights->begin(),
detail::EntryCompareOp());
#endif
// Scan weights
dh::XGBCachingDeviceAllocator<char> caching;
#if defined(XGBOOST_USE_CUDA)
thrust::inclusive_scan_by_key(thrust::cuda::par(caching),
sorted_entries->begin(), sorted_entries->end(),
weights->begin(), weights->begin(),
[=] __device__(const Entry& a, const Entry& b) {
return a.index == b.index;
});
#elif defined(XGBOOST_USE_HIP)
thrust::inclusive_scan_by_key(thrust::hip::par(caching),
sorted_entries->begin(), sorted_entries->end(),
weights->begin(), weights->begin(),
[=] __device__(const Entry& a, const Entry& b) {
return a.index == b.index;
});
#endif
}
void RemoveDuplicatedCategories(
@ -192,8 +215,14 @@ void ProcessBatch(int device, MetaInfo const &info, const SparsePage &page,
sorted_entries = dh::device_vector<Entry>(host_data.begin() + begin,
host_data.begin() + end);
}
#if defined(XGBOOST_USE_CUDA)
thrust::sort(thrust::cuda::par(alloc), sorted_entries.begin(),
sorted_entries.end(), detail::EntryCompareOp());
#elif defined(XGBOOST_USE_HIP)
thrust::sort(thrust::hip::par(alloc), sorted_entries.begin(),
sorted_entries.end(), detail::EntryCompareOp());
#endif
HostDeviceVector<SketchContainer::OffsetT> cuts_ptr;
dh::caching_device_vector<size_t> column_sizes_scan;

10
src/common/hist_util.cuh
View File

@ -89,7 +89,7 @@ void GetColumnSizesScan(int device, size_t num_columns, size_t num_cuts_per_feat
cuts_ptr->DevicePointer());
thrust::exclusive_scan(thrust::hip::par(alloc), column_sizes_scan->begin(),
column_sizes_scan->end(), column_sizes_scan->begin());
#else
#elif defined(XGBOOST_USE_CUDA)
thrust::exclusive_scan(thrust::cuda::par(alloc), cut_ptr_it,
cut_ptr_it + column_sizes_scan->size(),
cuts_ptr->DevicePointer());
@ -198,7 +198,7 @@ void ProcessSlidingWindow(AdapterBatch const &batch, MetaInfo const &info,
#if defined(XGBOOST_USE_HIP)
thrust::sort(thrust::hip::par(alloc), sorted_entries.begin(),
sorted_entries.end(), detail::EntryCompareOp());
#else
#elif defined(XGBOOST_USE_CUDA)
thrust::sort(thrust::cuda::par(alloc), sorted_entries.begin(),
sorted_entries.end(), detail::EntryCompareOp());
#endif
@ -229,7 +229,7 @@ void ProcessWeightedSlidingWindow(Batch batch, MetaInfo const& info,
#if defined(XGBOOST_USE_HIP)
dh::safe_cuda(hipSetDevice(device));
#else
#elif defined(XGBOOST_USE_CUDA)
dh::safe_cuda(cudaSetDevice(device));
#endif
@ -272,7 +272,7 @@ void ProcessWeightedSlidingWindow(Batch batch, MetaInfo const& info,
batch_iter + begin,
d_temp_weights.data(), // output
is_valid);
#else
#elif defined(XGBOOST_USE_CUDA)
auto retit = thrust::copy_if(thrust::cuda::par(alloc),
weight_iter + begin, weight_iter + end,
batch_iter + begin,
@ -295,7 +295,7 @@ void ProcessWeightedSlidingWindow(Batch batch, MetaInfo const& info,
batch_iter + begin,
d_temp_weights.data(), // output
is_valid);
#else
#elif defined(XGBOOST_USE_CUDA)
auto retit = thrust::copy_if(thrust::cuda::par(alloc),
weight_iter + begin, weight_iter + end,
batch_iter + begin,

4
src/common/hist_util.hip
View File

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