 }
 
 __global__ void ClearBuffersKernel(
-    BitField result_buffer_self, BitField result_buffer_input, BitField feature_buffer) {
+    BitField result_buffer_output, BitField result_buffer_input) {
   auto tid = blockIdx.x * blockDim.x + threadIdx.x;
-  if (tid < result_buffer_self.Size()) {
-    result_buffer_self.Clear(tid);
+  if (tid < result_buffer_output.Size()) {
+    result_buffer_output.Clear(tid);
   }
   if (tid < result_buffer_input.Size()) {
     result_buffer_input.Clear(tid);

  if (!has_constraint_ || nid == 0) {
    return feature_list;
  }
  auto selected = this->QueryNode(nid);
  CHECK_EQ(input_buffer_bits_.Size(), output_buffer_bits_.Size());
  int constexpr kBlockThreads = 256;
  const int n_grids = static_cast<int>(
      dh::DivRoundUp(output_buffer_bits_.Size(), kBlockThreads));

  QueryFeatureListKernel<<<n_grids, kBlockThreads>>>
      (feature_list,
       selected,
       input_buffer_bits_,
       output_buffer_bits_);

  thrust::counting_iterator<int32_t> begin(0);
  thrust::counting_iterator<int32_t> end(result_buffer_.size());

  const int n_grids = static_cast<int>(
      dh::DivRoundUp(input_buffer_bits_.Size(), kBlockThreads));
  ClearBuffersKernel<<<n_grids, kBlockThreads>>>(
      output_buffer_bits_, input_buffer_bits_);
}

common::Span<int32_t> FeatureInteractionConstraint::QueryNode(int32_t node_id) {
