Cache GPU histogram kernel configuration. (dmlc#10538)

src/tree/gpu_hist/histogram.cu

@@ -248,13 +248,13 @@ __global__ void __launch_bounds__(kBlockThreads)
     agent.BuildHistogramWithGlobal();
   }
 }
-
 namespace {
 constexpr std::int32_t kBlockThreads = 1024;
 constexpr std::int32_t kItemsPerThread = 8;
 constexpr std::int32_t ItemsPerTile() { return kBlockThreads * kItemsPerThread; }
 }  // namespace
 
+
 // Use auto deduction guide to workaround compiler error.
 template <auto Global = SharedMemHistKernel<false, kBlockThreads, kItemsPerThread>,
           auto Shared = SharedMemHistKernel<true, kBlockThreads, kItemsPerThread>>