SYCLops fails to generate a some SYCL code.

[Ralender]

I was experimenting with SYCLops to see what is capable of and it did pretty good but I found a case in which it failed.

here is how to reproduce: (test_sycl.cpp is below)

> ./build/bin/clang++ -fsycl-device-only -fno-sycl-instrument-device-code -fno-sycl-dead-args-optimization -O2 -mllvm -disable-loop-idiom-memset -mllvm -sycl-opt=false -fno-legacy-pass-manager -fno-unroll-loops -fno-vectorize -ffp-contract=off -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -Xsycl-target-frontend -fno-exceptions -std=c++20 -S -emit-llvm
test_sycl.cpp -o out.ll
> ./build/bin/syclops out.ll -emit-mlir -o out.mlir
Error: Dominator tree node contains more than 3 children - Control flow contains unexpected jumps

The clang command line was took from the test configuration with -Xsycl-target-frontend -fno-exceptions -std=c++20 added to compile the specific sycl example

the sycl code sample is inspired from https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/test/vitis/simple_tests/evaluation_tree.cpp

#include <cassert>
#include <optional>
#include <sycl/sycl.hpp>
#include <type_traits>
#include <variant>

namespace detail {

template <typename>
struct variant_trait {};

template <typename... Tys>
struct variant_trait<std::variant<Tys...>> {
  using indexes = std::make_index_sequence<sizeof...(Tys)>;
};

template <typename ret_type, typename Func, typename Var>
[[noreturn]] inline ret_type
visit_single_impl(Func &&, std::integer_sequence<size_t>, Var &&) {
  __builtin_unreachable();
}

template <typename ret_type, typename Func, typename Var, auto First,
          auto... Idx>
inline ret_type visit_single_impl(Func &&f,
                                  std::integer_sequence<size_t, First, Idx...>,
                                  Var &&var) {
  if (var.index() == First)
    return std::forward<Func>(f)(std::get<First>(var));
  return visit_single_impl<ret_type>(std::forward<Func>(f),
                                     std::integer_sequence<size_t, Idx...>{},
                                     std::forward<Var>(var));
}

template <typename Func, typename Var>
decltype(auto) visit_single(Func &&f, Var &&var) {
  using ret_type =
      std::invoke_result_t<Func, decltype(std::get<0>(std::declval<Var>()))>;
  return visit_single_impl<ret_type>(
      std::forward<Func>(f),
      typename variant_trait<
          std::remove_cv_t<std::remove_reference_t<Var>>>::indexes{},
      std::forward<Var>(var));
}
} // namespace detail

template <typename Func, typename Var, typename... Rest>
auto dev_visit(Func &&f, Var &&var, Rest &&...rest) {
  if constexpr (sizeof...(Rest) == 0)
    return detail::visit_single(std::forward<Func>(f), std::forward<Var>(var));
  else
    return detail::visit_single(
        [&](auto &&First) {
          return dev_visit(
              [&](auto &&...Others) {
                std::forward<Func>(f)(
                    std::forward<decltype(First)>(First),
                    std::forward<decltype(Others)>(Others)...);
              },
              std::forward<Rest>(rest)...);
        },
        std::forward<Var>(var));
}

template <typename T, typename IntTy = std::ptrdiff_t>
class rel_ptr {
  IntTy offset;

  char *get_this() const {
    return reinterpret_cast<char *>(const_cast<rel_ptr *>(this));
  }
  IntTy get_offset(T *t) const {
    return (reinterpret_cast<char *>(t) - get_this());
  }

public:
  T *get() const { return (T *)(get_this() + offset); }
  T &operator*() const { return *get(); }
  T *operator->() const { return get(); }
  T &operator[](std::size_t e) const { return get()[e]; }
  explicit operator bool() const { return get(); }
  bool operator!() const { return !get(); }

  rel_ptr() : rel_ptr(nullptr) {}
  explicit rel_ptr(T *t) : offset(get_offset(t)) {}
  rel_ptr(const rel_ptr &other) : rel_ptr(other.get()) {}
  rel_ptr &operator=(const rel_ptr &other) {
    offset = get_offset(other.get());
    return *this;
  }

  friend rel_ptr operator+(const rel_ptr &ptr, IntTy off) {
    return rel_ptr(ptr.get() + off);
  }
  friend rel_ptr operator-(const rel_ptr &ptr, IntTy off) {
    return rel_ptr(ptr.get() - off);
  }
  IntTy operator-(const rel_ptr &other) const { return get() - other.get(); }

  bool operator==(const rel_ptr &other) const { return get() == other.get(); }
  friend bool operator==(const rel_ptr &p1, T *p2) { return p1.get() == p2; }

  std::weak_ordering operator<=>(const rel_ptr &other) {
    return get() <=> other.get();
  }
  std::weak_ordering operator<=>(T *p2) { return get() <=> p2; }

  rel_ptr &operator++() {
    *this = (*this) + 1;
    return *this;
  }
  rel_ptr operator++(int) {
    rel_ptr tmp = *this;
    ++*this;
    return tmp;
  }
  rel_ptr &operator--() {
    *this = (*this) - 1;
    return *this;
  }
  rel_ptr operator--(int) {
    rel_ptr tmp = *this;
    --*this;
    return tmp;
  }
};

struct Constant {
  int d;
  int compute() { return d; }
};

struct node;

struct BinOp {
  rel_ptr<node> elem[2];
};
struct AddOp : BinOp {
  int compute();
};
struct MulOp : BinOp {
  int compute();
};

struct node {
  std::variant<Constant, AddOp, MulOp> n;
  int cache;
  node(auto v) : n(v) {}
  void compute() {
    cache = dev_visit([&](auto &e) { return e.compute(); }, n);
  }
  int get_value() { return cache; }
};

int AddOp::compute() {
  return BinOp::elem[0]->get_value() + BinOp::elem[1]->get_value();
}

int MulOp::compute() {
  return BinOp::elem[0]->get_value() * BinOp::elem[1]->get_value();
}

int main() {
  // Computation graph for (1 + 2) * 3
  std::array<node, 5> data{
      {Constant{1}, Constant{2}, AddOp{rel_ptr{&data[0]}, rel_ptr{&data[1]}},
       Constant{3}, MulOp{rel_ptr{&data[2]}, rel_ptr{&data[3]}}}};
  sycl::buffer<node> In{data.data(), {data.size()}};
  sycl::buffer<int> Out{1};

  sycl::queue Queue;

  Queue.submit([&](sycl::handler &cgh) {
    auto AIn = In.get_access<sycl::access::mode::read_write>(cgh);
    auto AOut = Out.get_access<sycl::access::mode::write>(cgh);
    int root_node = data.size() - 1;
    cgh.single_task<class Kernel>([=] {
      for (int i = 0; i < AIn.size(); i++)
        AIn[i].compute();
      AOut[0] = AIn[root_node].get_value();
    });
  });
  {
    auto AOut = Out.get_access<sycl::access::mode::read>();
    std::cout << AOut[0] << std::endl;
    assert(AOut[0] == 9);
  }
  return 0;
}