[REVIEW] ENH Decision Tree new backend computeSplit*Kernel histogram calculation optimization
#3674
Conversation
* using atomics to calculate PDFs and then using blockScan to get required CDFs that was originally issueing too many atomicAdds to shared memory
* was earlier hard-coded
* dynamically assigning based on occupancy while ceil-ing it to minimum 4 blocks
* pruning unnecessary comments and code
* improving doxygen comments
* adding some explanatory comments
…stogram-calculation-optimization-for-computesplitclassificationkernel
* shift the blockscan code to a reusable device function
* change appropriately in `computeSplitClassificationKernel`
…ogram-calculation-optimization-computeSplitRegressionKernel
venkywonka
added
improvement
| typedef cub::BlockScan<DataT, TPB> BlockScan; | ||
| __shared__ typename BlockScan::TempStorage temp_storage; | ||
|
|
||
| for (IdxT tix = threadIdx.x; tix < max(TPB, nbins); tix += blockDim.x) { |
There was a problem hiding this comment.
If nbins > TPB, then the resulting scan will be incorrect (because the total sum of the previous iteration is not being carried forward for the next iteration).
There was a problem hiding this comment.
IOW, InclusiveSum provides an option to also get the total sum with this function: https://nvlabs.github.io/cub/classcub_1_1_block_scan.html#a99222ab9b122e6df879ee04b4e8244da
There was a problem hiding this comment.
😅 Thank you for that, have rectified it 👍🏻
| // locations | ||
| offset_cdf += nbins; | ||
| //convert pdf to cdf | ||
| pdf_to_cdf<int, IdxT, TPB>(pdf_shist + offset_pdf, cdf_shist + offset_cdf, |
There was a problem hiding this comment.
why not just compute this cdf using the cdf from the above and its total sum? That way, we could avoid an extra block-scan operation.
There was a problem hiding this comment.
I did think of that initially, but leveraging total-sum from previous sumscan did not strike me! I have added it thejaswi 🙏🏻
…ogram-calculation-optimization-computeSplitRegressionKernel
* incorporating block_aggregate in inclusive sumscan * using total sum instead of doing a right-to-left sumscan
venkywonka
changed the title
venkywonka
changed the title
computeSplit*Kernel histogram calculation optimization
…ogram-calculation-optimization-computeSplitRegressionKernel
Codecov Report
@@ Coverage Diff @@
## branch-0.19 #3674 +/- ##
===============================================
+ Coverage 80.70% 82.92% +2.21%
===============================================
Files 227 227
Lines 17615 17591 -24
===============================================
+ Hits 14217 14587 +370
+ Misses 3398 3004 -394
Flags with carried forward coverage won't be shown. Click here to find out more.
Continue to review full report at Codecov.
|
|
@venkywonka please resolve conflicts |
…stogram-calculation-optimization-computeSplitRegressionKernel
venkywonka
force-pushed
the
enh-ext-histogram-calculation-optimization-computeSplitRegressionKernel
branch
from
607a926
to
6e39fef
Compare
|
@gpucibot merge |
This is a follow-up of PR #3616 and should be merged after that.
This PR introduces:
pdf_to_cdfconversion using inclusive-sumscan into a device function so that it can be reused by both theML::DecisionTree::computeSplitClassificationKernelandML::DecisionTree::computeSplitRegressionKernelML::DecisionTree::computeSplitRegressionKernel. These histograms are used for node-splitting in decision trees for the task of regression.computeSplitClassificationKernelhistogram calculation and occupancy optimization #3616