| 819 | } |
| 820 | |
| 821 | bool GPUTreeLearner::BeforeFindBestSplit(const Tree* tree, int left_leaf, int right_leaf) { |
| 822 | int smaller_leaf; |
| 823 | data_size_t num_data_in_left_child = GetGlobalDataCountInLeaf(left_leaf); |
| 824 | data_size_t num_data_in_right_child = GetGlobalDataCountInLeaf(right_leaf); |
| 825 | // only have root |
| 826 | if (right_leaf < 0) { |
| 827 | smaller_leaf = -1; |
| 828 | } else if (num_data_in_left_child < num_data_in_right_child) { |
| 829 | smaller_leaf = left_leaf; |
| 830 | } else { |
| 831 | smaller_leaf = right_leaf; |
| 832 | } |
| 833 | |
| 834 | // Copy indices, gradients and hessians as early as possible |
| 835 | if (smaller_leaf >= 0 && num_dense_feature_groups_) { |
| 836 | // only need to initialize for smaller leaf |
| 837 | // Get leaf boundary |
| 838 | const data_size_t* indices = data_partition_->indices(); |
| 839 | data_size_t begin = data_partition_->leaf_begin(smaller_leaf); |
| 840 | data_size_t end = begin + data_partition_->leaf_count(smaller_leaf); |
| 841 | |
| 842 | // copy indices to the GPU: |
| 843 | #if GPU_DEBUG >= 2 |
| 844 | Log::Info("Copying indices, gradients and hessians to GPU..."); |
| 845 | printf("Indices size %d being copied (left = %d, right = %d)\n", end - begin, num_data_in_left_child, num_data_in_right_child); |
| 846 | #endif |
| 847 | indices_future_ = boost::compute::copy_async(indices + begin, indices + end, device_data_indices_->begin(), queue_); |
| 848 | |
| 849 | if (!is_constant_hessian_) { |
| 850 | #pragma omp parallel for schedule(static) |
| 851 | for (data_size_t i = begin; i < end; ++i) { |
| 852 | ordered_hessians_[i - begin] = hessians_[indices[i]]; |
| 853 | } |
| 854 | // copy ordered hessians to the GPU: |
| 855 | hessians_future_ = queue_.enqueue_write_buffer_async(device_hessians_, 0, (end - begin) * sizeof(score_t), ptr_pinned_hessians_); |
| 856 | } |
| 857 | |
| 858 | #pragma omp parallel for schedule(static) |
| 859 | for (data_size_t i = begin; i < end; ++i) { |
| 860 | ordered_gradients_[i - begin] = gradients_[indices[i]]; |
| 861 | } |
| 862 | // copy ordered gradients to the GPU: |
| 863 | gradients_future_ = queue_.enqueue_write_buffer_async(device_gradients_, 0, (end - begin) * sizeof(score_t), ptr_pinned_gradients_); |
| 864 | |
| 865 | #if GPU_DEBUG >= 2 |
| 866 | Log::Info("Gradients/hessians/indices copied to device with size %d", end - begin); |
| 867 | #endif |
| 868 | } |
| 869 | return SerialTreeLearner::BeforeFindBestSplit(tree, left_leaf, right_leaf); |
| 870 | } |
| 871 | |
| 872 | bool GPUTreeLearner::ConstructGPUHistogramsAsync( |
| 873 | const std::vector<int8_t>& is_feature_used, |
nothing calls this directly
no test coverage detected