     if (leaf_num_data == num_data) {
       if (use_all_features) {
         if (!is_constant_hessian)
-          histogram64<<<4*num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
+          histogram64<<<num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
                   arg3, arg4, arg5,
                   arg6, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
         else
-          histogram64<<<4*num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
+          histogram64<<<num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
                   arg3, arg4, arg5,
                   arg6_const, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
       } else {
         if (!is_constant_hessian)
-          histogram64_fulldata<<<4*num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
+          histogram64_fulldata<<<num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
                   arg3, arg4, arg5,
                   arg6, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
         else
-          histogram64_fulldata<<<4*num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
+          histogram64_fulldata<<<num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
                   arg3, arg4, arg5,
                   arg6_const, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
       }

    if (leaf_num_data == num_data) {
      if (use_all_features) {
        if (!is_constant_hessian)
          histogram16<<<16*num_workgroups, 16, 0, stream>>>(arg0, arg1, arg2,
                  arg3, arg4, arg5,
                  arg6, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
        else
           histogram16<<<16*num_workgroups, 16, 0, stream>>>(arg0, arg1, arg2,
                  arg3, arg4, arg5,
                  arg6_const, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
      } else {
        if (!is_constant_hessian)
           histogram16_fulldata<<<16*num_workgroups, 16, 0, stream>>>(arg0, arg1, arg2,
                  arg3, arg4, arg5,
                  arg6, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
        else
           histogram16_fulldata<<<16*num_workgroups, 16, 0, stream>>>(arg0, arg1, arg2,
                  arg3, arg4, arg5,
                  arg6_const, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
      }

      if (use_all_features) {
        // seems all features is always enabled, so this should be the same as fulldata
        if (!is_constant_hessian)
          histogram64<<<num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
                  arg3, arg4, arg5,
                  arg6, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
        else
          histogram64<<<num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
                  arg3, arg4, arg5,
                  arg6_const, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
      } else {
        if (!is_constant_hessian)
          histogram64<<<num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
                  arg3, arg4, arg5,
                  arg6, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
        else
          histogram64<<<num_workgroups, 64, 0, stream>>>(arg0, arg1, arg2,
                  arg3, arg4, arg5,
                  arg6_const, arg7, arg8, static_cast<acc_type*>(arg9), exp_workgroups_per_feature);
      }
