| 953 | } |
| 954 | |
| 955 | void GPUTreeLearner::ConstructHistograms(const std::vector<int8_t>& is_feature_used, bool use_subtract) { |
| 956 | std::vector<int8_t> is_sparse_feature_used(num_features_, 0); |
| 957 | std::vector<int8_t> is_dense_feature_used(num_features_, 0); |
| 958 | #pragma omp parallel for schedule(static) |
| 959 | for (int feature_index = 0; feature_index < num_features_; ++feature_index) { |
| 960 | if (!is_feature_used_[feature_index]) continue; |
| 961 | if (!is_feature_used[feature_index]) continue; |
| 962 | if (ordered_bins_[train_data_->Feature2Group(feature_index)]) { |
| 963 | is_sparse_feature_used[feature_index] = 1; |
| 964 | } else { |
| 965 | is_dense_feature_used[feature_index] = 1; |
| 966 | } |
| 967 | } |
| 968 | // construct smaller leaf |
| 969 | HistogramBinEntry* ptr_smaller_leaf_hist_data = smaller_leaf_histogram_array_[0].RawData() - 1; |
| 970 | // ConstructGPUHistogramsAsync will return true if there are availabe feature gourps dispatched to GPU |
| 971 | bool is_gpu_used = ConstructGPUHistogramsAsync(is_feature_used, |
| 972 | nullptr, smaller_leaf_splits_->num_data_in_leaf(), |
| 973 | nullptr, nullptr, |
| 974 | nullptr, nullptr); |
| 975 | // then construct sparse features on CPU |
| 976 | // We set data_indices to null to avoid rebuilding ordered gradients/hessians |
| 977 | train_data_->ConstructHistograms(is_sparse_feature_used, |
| 978 | nullptr, smaller_leaf_splits_->num_data_in_leaf(), |
| 979 | smaller_leaf_splits_->LeafIndex(), |
| 980 | &ordered_bins_, gradients_, hessians_, |
| 981 | ordered_gradients_.data(), ordered_hessians_.data(), is_constant_hessian_, |
| 982 | ptr_smaller_leaf_hist_data); |
| 983 | // wait for GPU to finish, only if GPU is actually used |
| 984 | if (is_gpu_used) { |
| 985 | if (config_->gpu_use_dp) { |
| 986 | // use double precision |
| 987 | WaitAndGetHistograms<HistogramBinEntry>(ptr_smaller_leaf_hist_data); |
| 988 | } else { |
| 989 | // use single precision |
| 990 | WaitAndGetHistograms<GPUHistogramBinEntry>(ptr_smaller_leaf_hist_data); |
| 991 | } |
| 992 | } |
| 993 | |
| 994 | // Compare GPU histogram with CPU histogram, useful for debuggin GPU code problem |
| 995 | // #define GPU_DEBUG_COMPARE |
| 996 | #ifdef GPU_DEBUG_COMPARE |
| 997 | for (int i = 0; i < num_dense_feature_groups_; ++i) { |
| 998 | if (!feature_masks_[i]) |
| 999 | continue; |
| 1000 | int dense_feature_group_index = dense_feature_group_map_[i]; |
| 1001 | size_t size = train_data_->FeatureGroupNumBin(dense_feature_group_index); |
| 1002 | HistogramBinEntry* ptr_smaller_leaf_hist_data = smaller_leaf_histogram_array_[0].RawData() - 1; |
| 1003 | HistogramBinEntry* current_histogram = ptr_smaller_leaf_hist_data + train_data_->GroupBinBoundary(dense_feature_group_index); |
| 1004 | HistogramBinEntry* gpu_histogram = new HistogramBinEntry[size]; |
| 1005 | data_size_t num_data = smaller_leaf_splits_->num_data_in_leaf(); |
| 1006 | printf("Comparing histogram for feature %d size %d, %lu bins\n", dense_feature_group_index, num_data, size); |
| 1007 | std::copy(current_histogram, current_histogram + size, gpu_histogram); |
| 1008 | std::memset(current_histogram, 0, train_data_->FeatureGroupNumBin(dense_feature_group_index) * sizeof(HistogramBinEntry)); |
| 1009 | train_data_->FeatureGroupBin(dense_feature_group_index)->ConstructHistogram( |
| 1010 | num_data != num_data_ ? smaller_leaf_splits_->data_indices() : nullptr, |
| 1011 | num_data, |
| 1012 | num_data != num_data_ ? ordered_gradients_.data() : gradients_, |
nothing calls this directly
no test coverage detected