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

Work around MSVC behavior wrt constexpr capture (#8211)

Browse Source 
* Work around MSVC behavior wrt constexpr capture

* Fix lint
This commit is contained in:
Philip Hyunsu Cho 2022-08-31 12:42:08 -07:00 committed by GitHub
parent a868498c18
commit 56395d120b
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
1 changed files with 1 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
src/tree/gpu_hist/histogram.cu
View File

@ -275,7 +275,7 @@ void BuildGradientHistogram(EllpackDeviceAccessor const& matrix,
constexpr int kItemsPerThread = 8; constexpr int kItemsPerThread = 8;
constexpr int kItemsPerTile = kBlockThreads * kItemsPerThread; constexpr int kItemsPerTile = kBlockThreads * kItemsPerThread;
auto runit = [&](auto kernel) { auto runit = [&, kMinItemsPerBlock = kItemsPerTile](auto kernel) {
if (shared) { if (shared) {
dh::safe_cuda(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, dh::safe_cuda(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize,
max_shared_memory)); max_shared_memory));
@ -295,7 +295,6 @@ void BuildGradientHistogram(EllpackDeviceAccessor const& matrix,
// Otherwise launch blocks such that each block has a minimum amount of work to do // Otherwise launch blocks such that each block has a minimum amount of work to do
// There are fixed costs to launching each block, e.g. zeroing shared memory // There are fixed costs to launching each block, e.g. zeroing shared memory
// The below amount of minimum work was found by experimentation // The below amount of minimum work was found by experimentation
constexpr int kMinItemsPerBlock = kItemsPerTile;
int columns_per_group = common::DivRoundUp(matrix.row_stride, feature_groups.NumGroups()); int columns_per_group = common::DivRoundUp(matrix.row_stride, feature_groups.NumGroups());
// Average number of matrix elements processed by each group // Average number of matrix elements processed by each group
std::size_t items_per_group = d_ridx.size() * columns_per_group; std::size_t items_per_group = d_ridx.size() * columns_per_group;