From f5f800c80d7c6387b2b33ee039f3cb859c6ec280 Mon Sep 17 00:00:00 2001 From: amdsc21 <96135754+amdsc21@users.noreply.github.com> Date: Wed, 8 Mar 2023 01:33:38 +0100 Subject: [PATCH] add HIP flags --- src/data/array_interface.h | 18 +++++++++--------- src/data/ellpack_page_source.h | 4 ++-- src/data/iterative_dmatrix.h | 4 ++-- src/data/proxy_dmatrix.h | 8 ++++---- src/data/sparse_page_source.h | 2 +- src/data/validation.h | 2 +- src/tree/split_evaluator.h | 2 +- 7 files changed, 20 insertions(+), 20 deletions(-) diff --git a/src/data/array_interface.h b/src/data/array_interface.h index e9045899b..997bc4788 100644 --- a/src/data/array_interface.h +++ b/src/data/array_interface.h @@ -302,12 +302,12 @@ class ArrayInterfaceHandler { template struct ToDType; // float -#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 +#if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(__HIP_PLATFORM_AMD__) template <> struct ToDType<__half> { 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 <> struct ToDType { static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF4; @@ -356,10 +356,10 @@ struct ToDType { 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 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 @@ -458,11 +458,11 @@ class ArrayInterface { CHECK(sizeof(long double) == 16) << "128-bit floating point is not supported on current platform."; } 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; #else 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') { type = T::kF4; } else if (typestr[1] == 'f' && typestr[2] == '8') { @@ -497,12 +497,12 @@ class ArrayInterface { using T = ArrayInterfaceHandler::Type; switch (type) { 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)); #else SPAN_CHECK(false); return func(reinterpret_cast(data)); -#endif // defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600 +#endif // (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600) || defined(__HIP_PLATFORM_AMD__) } case T::kF4: return func(reinterpret_cast(data)); @@ -555,7 +555,7 @@ class ArrayInterface { static_assert(sizeof...(index) <= D, "Invalid index."); return this->DispatchCall([=](auto const *p_values) -> T { 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 using Type = std::conditional_t< std::is_same<__half, diff --git a/src/data/ellpack_page_source.h b/src/data/ellpack_page_source.h index dc0802472..9ac513ec3 100644 --- a/src/data/ellpack_page_source.h +++ b/src/data/ellpack_page_source.h @@ -43,14 +43,14 @@ class EllpackPageSource : public PageSourceIncMixIn { void Fetch() final; }; -#if !defined(XGBOOST_USE_CUDA) +#if !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP) inline void EllpackPageSource::Fetch() { // silent the warning about unused variables. (void)(row_stride_); (void)(is_dense_); common::AssertGPUSupport(); } -#endif // !defined(XGBOOST_USE_CUDA) +#endif // !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP) } // namespace data } // namespace xgboost diff --git a/src/data/iterative_dmatrix.h b/src/data/iterative_dmatrix.h index 28c4087c4..d3ee62696 100644 --- a/src/data/iterative_dmatrix.h +++ b/src/data/iterative_dmatrix.h @@ -121,7 +121,7 @@ void GetCutsFromRef(std::shared_ptr ref_, bst_feature_t n_features, Bat */ 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) { // silent the warning about unused variables. (void)(proxy_); @@ -138,7 +138,7 @@ inline BatchSet IterativeDMatrix::GetEllpackBatches(const BatchPara inline void GetCutsFromEllpack(EllpackPage const &, common::HistogramCuts *) { common::AssertGPUSupport(); } -#endif // !defined(XGBOOST_USE_CUDA) +#endif // !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_HIP) } // namespace data } // namespace xgboost diff --git a/src/data/proxy_dmatrix.h b/src/data/proxy_dmatrix.h index fa55a481f..fa2901c47 100644 --- a/src/data/proxy_dmatrix.h +++ b/src/data/proxy_dmatrix.h @@ -47,10 +47,10 @@ class DMatrixProxy : public DMatrix { dmlc::any batch_; Context ctx_; -#if defined(XGBOOST_USE_CUDA) +#if defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP) void FromCudaColumnar(StringView interface_str); void FromCudaArray(StringView interface_str); -#endif // defined(XGBOOST_USE_CUDA) +#endif // defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP) public: int DeviceIdx() const { return ctx_.gpu_id; } @@ -58,7 +58,7 @@ class DMatrixProxy : public DMatrix { void SetCUDAArray(char const* c_interface) { common::AssertGPUSupport(); CHECK(c_interface); -#if defined(XGBOOST_USE_CUDA) +#if defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP) StringView interface_str{c_interface}; Json json_array_interface = Json::Load(interface_str); if (IsA(json_array_interface)) { @@ -66,7 +66,7 @@ class DMatrixProxy : public DMatrix { } else { this->FromCudaArray(interface_str); } -#endif // defined(XGBOOST_USE_CUDA) +#endif // defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP) } void SetArrayData(char const* c_interface); diff --git a/src/data/sparse_page_source.h b/src/data/sparse_page_source.h index 088f1e98c..f35ccd07c 100644 --- a/src/data/sparse_page_source.h +++ b/src/data/sparse_page_source.h @@ -206,7 +206,7 @@ class SparsePageSourceImpl : public BatchIteratorImpl { } }; -#if defined(XGBOOST_USE_CUDA) +#if defined(XGBOOST_USE_CUDA) || defined(XGBOOST_USE_HIP) void DevicePush(DMatrixProxy* proxy, float missing, SparsePage* page); #else inline void DevicePush(DMatrixProxy*, float, SparsePage*) { common::AssertGPUSupport(); } diff --git a/src/data/validation.h b/src/data/validation.h index 6d3701114..914a2d740 100644 --- a/src/data/validation.h +++ b/src/data/validation.h @@ -13,7 +13,7 @@ namespace xgboost { namespace data { struct LabelsCheck { XGBOOST_DEVICE bool operator()(float y) { -#if defined(__CUDA_ARCH__) +#if defined(__CUDA_ARCH__) || defined(__HIP_PLATFORM_AMD__) return ::isnan(y) || ::isinf(y); #else return std::isnan(y) || std::isinf(y); diff --git a/src/tree/split_evaluator.h b/src/tree/split_evaluator.h index c036cc3ed..b6625339d 100644 --- a/src/tree/split_evaluator.h +++ b/src/tree/split_evaluator.h @@ -121,7 +121,7 @@ class TreeEvaluator { // Fast floating point division instruction on device XGBOOST_DEVICE float Divide(float a, float b) const { -#ifdef __CUDA_ARCH__ +#if defined(__CUDA_ARCH__) || defined(__HIP_PLATFORM_AMD__) return __fdividef(a, b); #else return a / b;