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

SumReduction logging

Browse Source
This commit is contained in:
Hendrik Groove 2024-10-20 23:27:50 +02:00
parent bf2ef6c586
commit db66fad9e9
1 changed files with 21 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -665,21 +665,31 @@ struct PinnedMemory {
*/
template <typename T>
typename std::iterator_traits<T>::value_type SumReduction(T in, int nVals) {
std::cerr << "Entering SumReduction, nVals: " << nVals << std::endl;
using ValueT = typename std::iterator_traits<T>::value_type;
size_t tmpSize {0};
ValueT *dummy_out = nullptr;
dh::safe_cuda(hipcub::DeviceReduce::Sum(nullptr, tmpSize, in, dummy_out, nVals));
try {
dh::safe_cuda(hipcub::DeviceReduce::Sum(nullptr, tmpSize, in, dummy_out, nVals));
std::cerr << "Temporary storage size: " << tmpSize << std::endl;
TemporaryArray<char> temp(tmpSize + sizeof(ValueT));
auto ptr = reinterpret_cast<ValueT *>(temp.data().get()) + 1;
dh::safe_cuda(hipcub::DeviceReduce::Sum(
reinterpret_cast<void *>(ptr), tmpSize, in,
reinterpret_cast<ValueT *>(temp.data().get()),
nVals));
ValueT sum;
dh::safe_cuda(hipMemcpy(&sum, temp.data().get(), sizeof(ValueT),
hipMemcpyDeviceToHost));
return sum;
TemporaryArray<char> temp(tmpSize + sizeof(ValueT));
auto ptr = reinterpret_cast<ValueT *>(temp.data().get()) + 1;
dh::safe_cuda(hipcub::DeviceReduce::Sum(
reinterpret_cast<void *>(ptr), tmpSize, in, reinterpret_cast<ValueT *>(temp.data().get()), nVals));
ValueT sum;
dh::safe_cuda(hipMemcpy(&sum, temp.data().get(), sizeof(ValueT), hipMemcpyDeviceToHost));
std::cerr << "SumReduction completed successfully" << std::endl;
return sum;
} catch (const std::exception& e) {
std::cerr << "Exception in SumReduction: " << e.what() << std::endl;
throw;
}
}
constexpr std::pair<int, int> CUDAVersion() {