   for (auto i = 0; i < trees.tree_limit; i++) {
     TreeEnsemble tree;
     trees.get_tree(tree, i);
-    std::vector<gpu_treeshap::PathElement> tmp_path;
-    tmp_path.reserve(tree.max_depth);
-    tmp_path.emplace_back(0, -1, 0, -inf, inf, false, 1.0, 0.0f);
-    RecurseTree(0, tree, &tmp_path, &paths, &path_idx);
+    for (auto j = 0ull; j < tree.num_outputs; j++) {
+      std::vector<gpu_treeshap::PathElement> tmp_path;
+      tmp_path.reserve(tree.max_depth);
+      tmp_path.emplace_back(0, -1, 0, -inf, inf, false, 1.0, 0.0f);
+      RecurseTree(0, tree, &tmp_path, &paths, &path_idx, j, tree.num_outputs);
+    }
   }
   return paths;
 }

                         static_cast<float>(tree.thresholds[pos]), false,
                         left_zero_fraction, 0.0f);

  RecurseTree(left_child, tree, tmp_path, paths, path_idx);

  // Add left split to the path
  tmp_path->back() = gpu_treeshap::PathElement(
      0, tree.features[pos], 0, static_cast<float>(tree.thresholds[pos]), inf,
      false, 1.0 - left_zero_fraction, 0.0f);

  RecurseTree(tree.children_right[pos], tree, tmp_path, paths, path_idx);

  tmp_path->pop_back();
}

                               num_rows, num_features);
  }
};

// Helper function that accepts an index into a flat contiguous array and the
// dimensions of a tensor and returns the indices with respect to the tensor
template <typename T, size_t N>
__device__ void flat_idx_to_tensor_idx(T flat_idx, const T (&shape)[N],
                                       T (&out_idx)[N]) {
  T current_size = shape[0];
  for (auto i = 1ull; i < N; i++) {
    current_size *= shape[i];
  }
  for (auto i = 0ull; i < N; i++) {
    current_size /= shape[i];
    out_idx[i] = flat_idx / current_size;
    flat_idx -= current_size * out_idx[i];
  }
}

// Given a shape and coordinates into a tensor, return the index into the
// backing storage one-dimensional array
template <typename T, size_t N>
__device__ T tensor_idx_to_flat_idx(const T (&shape)[N],
                                    const T (&tensor_idx)[N]) {
  T current_size = shape[0];
  for (auto i = 1ull; i < N; i++) {
    current_size *= shape[i];
  }
  T idx = 0;
  for (auto i = 0ull; i < N; i++) {
    current_size /= shape[i];
    idx += tensor_idx[i] * current_size;
  }
  return idx;
}

inline void dense_tree_path_dependent_gpu(
    const TreeEnsemble &trees, const ExplanationDataset &data,
