From 56395d120b29a51a0980f64bd1e36285dbf11400 Mon Sep 17 00:00:00 2001 From: Philip Hyunsu Cho Date: Wed, 31 Aug 2022 12:42:08 -0700 Subject: [PATCH] Work around MSVC behavior wrt constexpr capture (#8211) * Work around MSVC behavior wrt constexpr capture * Fix lint --- src/tree/gpu_hist/histogram.cu | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/tree/gpu_hist/histogram.cu b/src/tree/gpu_hist/histogram.cu index 209363405..aa04eaf2c 100644 --- a/src/tree/gpu_hist/histogram.cu +++ b/src/tree/gpu_hist/histogram.cu @@ -275,7 +275,7 @@ void BuildGradientHistogram(EllpackDeviceAccessor const& matrix, constexpr int kItemsPerThread = 8; constexpr int kItemsPerTile = kBlockThreads * kItemsPerThread; - auto runit = [&](auto kernel) { + auto runit = [&, kMinItemsPerBlock = kItemsPerTile](auto kernel) { if (shared) { dh::safe_cuda(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 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 // There are fixed costs to launching each block, e.g. zeroing shared memory // 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()); // Average number of matrix elements processed by each group std::size_t items_per_group = d_ridx.size() * columns_per_group;