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

add HIP flags

Browse Source
This commit is contained in:
amdsc21 2023-03-08 01:33:38 +01:00
parent 6b7be96373
commit f5f800c80d
7 changed files with 20 additions and 20 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

18
src/data/array_interface.h
View File

@ -302,12 +302,12 @@ class ArrayInterfaceHandler {
template <typename T, typename E = void> template <typename T, typename E = void>
struct ToDType; struct ToDType;
// float // float
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 #if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(__HIP_PLATFORM_AMD__)
template <> template <>
struct ToDType<__half> { struct ToDType<__half> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF2; static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF2;
}; };
#endif // defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 #endif // (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(__HIP_PLATFORM_AMD__)
template <> template <>
struct ToDType<float> { struct ToDType<float> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF4; static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF4;
@ -356,10 +356,10 @@ struct ToDType<int64_t> {
static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI8; static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI8;
}; };
#if !defined(XGBOOST_USE_CUDA) #if !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP)
inline void ArrayInterfaceHandler::SyncCudaStream(int64_t) { common::AssertGPUSupport(); } inline void ArrayInterfaceHandler::SyncCudaStream(int64_t) { common::AssertGPUSupport(); }
inline bool ArrayInterfaceHandler::IsCudaPtr(void const *) { return false; } inline bool ArrayInterfaceHandler::IsCudaPtr(void const *) { return false; }
#endif // !defined(XGBOOST_USE_CUDA) #endif // !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP)
/** /**
* \brief A type erased view over __array_interface__ protocol defined by numpy * \brief A type erased view over __array_interface__ protocol defined by numpy
@ -458,11 +458,11 @@ class ArrayInterface {
CHECK(sizeof(long double) == 16) CHECK(sizeof(long double) == 16)
<< "128-bit floating point is not supported on current platform."; << "128-bit floating point is not supported on current platform.";
} else if (typestr[1] == 'f' && typestr[2] == '2') { } else if (typestr[1] == 'f' && typestr[2] == '2') {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 #if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(XGBOOST_USE_HIP)
type = T::kF2; type = T::kF2;
#else #else
LOG(FATAL) << "Half type is not supported."; LOG(FATAL) << "Half type is not supported.";
#endif // defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 #endif // (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(XGBOOST_USE_HIP)
} else if (typestr[1] == 'f' && typestr[2] == '4') { } else if (typestr[1] == 'f' && typestr[2] == '4') {
type = T::kF4; type = T::kF4;
} else if (typestr[1] == 'f' && typestr[2] == '8') { } else if (typestr[1] == 'f' && typestr[2] == '8') {
@ -497,12 +497,12 @@ class ArrayInterface {
using T = ArrayInterfaceHandler::Type; using T = ArrayInterfaceHandler::Type;
switch (type) { switch (type) {
case T::kF2: { case T::kF2: {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 #if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(__HIP_PLATFORM_AMD__)
return func(reinterpret_cast<__half const *>(data)); return func(reinterpret_cast<__half const *>(data));
#else #else
SPAN_CHECK(false); SPAN_CHECK(false);
return func(reinterpret_cast<float const *>(data)); return func(reinterpret_cast<float const *>(data));
#endif // defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 #endif // (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(__HIP_PLATFORM_AMD__)
} }
case T::kF4: case T::kF4:
return func(reinterpret_cast<float const *>(data)); return func(reinterpret_cast<float const *>(data));
@ -555,7 +555,7 @@ class ArrayInterface {
static_assert(sizeof...(index) <= D, "Invalid index."); static_assert(sizeof...(index) <= D, "Invalid index.");
return this->DispatchCall([=](auto const *p_values) -> T { return this->DispatchCall([=](auto const *p_values) -> T {
std::size_t offset = linalg::detail::Offset<0ul>(strides, 0ul, index...); std::size_t offset = linalg::detail::Offset<0ul>(strides, 0ul, index...);
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 #if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(__HIP_PLATFORM_AMD__)
// No operator defined for half -> size_t // No operator defined for half -> size_t
using Type = std::conditional_t< using Type = std::conditional_t<
std::is_same<__half, std::is_same<__half,

4
src/data/ellpack_page_source.h
View File

@ -43,14 +43,14 @@ class EllpackPageSource : public PageSourceIncMixIn<EllpackPage> {
void Fetch() final; void Fetch() final;
}; };
#if !defined(XGBOOST_USE_CUDA) #if !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP)
inline void EllpackPageSource::Fetch() { inline void EllpackPageSource::Fetch() {
// silent the warning about unused variables. // silent the warning about unused variables.
(void)(row_stride_); (void)(row_stride_);
(void)(is_dense_); (void)(is_dense_);
common::AssertGPUSupport(); common::AssertGPUSupport();
} }
#endif // !defined(XGBOOST_USE_CUDA) #endif // !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP)
} // namespace data } // namespace data
} // namespace xgboost } // namespace xgboost

4
src/data/iterative_dmatrix.h
View File

@ -121,7 +121,7 @@ void GetCutsFromRef(std::shared_ptr<DMatrix> ref_, bst_feature_t n_features, Bat
*/ */
void GetCutsFromEllpack(EllpackPage const &page, common::HistogramCuts *cuts); void GetCutsFromEllpack(EllpackPage const &page, common::HistogramCuts *cuts);
#if !defined(XGBOOST_USE_CUDA) #if !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP)
inline void IterativeDMatrix::InitFromCUDA(DataIterHandle, float, std::shared_ptr<DMatrix>) { inline void IterativeDMatrix::InitFromCUDA(DataIterHandle, float, std::shared_ptr<DMatrix>) {
// silent the warning about unused variables. // silent the warning about unused variables.
(void)(proxy_); (void)(proxy_);
@ -138,7 +138,7 @@ inline BatchSet<EllpackPage> IterativeDMatrix::GetEllpackBatches(const BatchPara
inline void GetCutsFromEllpack(EllpackPage const &, common::HistogramCuts *) { inline void GetCutsFromEllpack(EllpackPage const &, common::HistogramCuts *) {
common::AssertGPUSupport(); common::AssertGPUSupport();
} }
#endif // !defined(XGBOOST_USE_CUDA) #endif // !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP)
} // namespace data } // namespace data
} // namespace xgboost } // namespace xgboost

8
src/data/proxy_dmatrix.h
View File

@ -47,10 +47,10 @@ class DMatrixProxy : public DMatrix {
dmlc::any batch_; dmlc::any batch_;
Context ctx_; Context ctx_;
#if defined(XGBOOST_USE_CUDA) #if defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP)
void FromCudaColumnar(StringView interface_str); void FromCudaColumnar(StringView interface_str);
void FromCudaArray(StringView interface_str); void FromCudaArray(StringView interface_str);
#endif // defined(XGBOOST_USE_CUDA) #endif // defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP)
public: public:
int DeviceIdx() const { return ctx_.gpu_id; } int DeviceIdx() const { return ctx_.gpu_id; }
@ -58,7 +58,7 @@ class DMatrixProxy : public DMatrix {
void SetCUDAArray(char const* c_interface) { void SetCUDAArray(char const* c_interface) {
common::AssertGPUSupport(); common::AssertGPUSupport();
CHECK(c_interface); CHECK(c_interface);
#if defined(XGBOOST_USE_CUDA) #if defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP)
StringView interface_str{c_interface}; StringView interface_str{c_interface};
Json json_array_interface = Json::Load(interface_str); Json json_array_interface = Json::Load(interface_str);
if (IsA<Array>(json_array_interface)) { if (IsA<Array>(json_array_interface)) {
@ -66,7 +66,7 @@ class DMatrixProxy : public DMatrix {
} else { } else {
this->FromCudaArray(interface_str); this->FromCudaArray(interface_str);
} }
#endif // defined(XGBOOST_USE_CUDA) #endif // defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP)
} }
void SetArrayData(char const* c_interface); void SetArrayData(char const* c_interface);

2
src/data/sparse_page_source.h
View File

@ -206,7 +206,7 @@ class SparsePageSourceImpl : public BatchIteratorImpl<S> {
} }
}; };
#if defined(XGBOOST_USE_CUDA) #if defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP)
void DevicePush(DMatrixProxy* proxy, float missing, SparsePage* page); void DevicePush(DMatrixProxy* proxy, float missing, SparsePage* page);
#else #else
inline void DevicePush(DMatrixProxy*, float, SparsePage*) { common::AssertGPUSupport(); } inline void DevicePush(DMatrixProxy*, float, SparsePage*) { common::AssertGPUSupport(); }

2
src/data/validation.h
View File

@ -13,7 +13,7 @@ namespace xgboost {
namespace data { namespace data {
struct LabelsCheck { struct LabelsCheck {
XGBOOST_DEVICE bool operator()(float y) { XGBOOST_DEVICE bool operator()(float y) {
#if defined(__CUDA_ARCH__) #if defined(__CUDA_ARCH__) || defined(__HIP_PLATFORM_AMD__)
return ::isnan(y) || ::isinf(y); return ::isnan(y) || ::isinf(y);
#else #else
return std::isnan(y) || std::isinf(y); return std::isnan(y) || std::isinf(y);

2
src/tree/split_evaluator.h
View File

@ -121,7 +121,7 @@ class TreeEvaluator {
// Fast floating point division instruction on device // Fast floating point division instruction on device
XGBOOST_DEVICE float Divide(float a, float b) const { XGBOOST_DEVICE float Divide(float a, float b) const {
#ifdef __CUDA_ARCH__ #if defined(__CUDA_ARCH__) || defined(__HIP_PLATFORM_AMD__)
return __fdividef(a, b); return __fdividef(a, b);
#else #else
return a / b; return a / b;