/*
 * This file belongs to the Galois project, a C++ library for exploiting parallelism.
 * The code is being released under the terms of the 3-Clause BSD License (a
 * copy is located in LICENSE.txt at the top-level directory).
 *
 * Copyright (C) 2018, The University of Texas at Austin. All rights reserved.
 * UNIVERSITY EXPRESSLY DISCLAIMS ANY AND ALL WARRANTIES CONCERNING THIS
 * SOFTWARE AND DOCUMENTATION, INCLUDING ANY WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR ANY PARTICULAR PURPOSE, NON-INFRINGEMENT AND WARRANTIES OF
 * PERFORMANCE, AND ANY WARRANTY THAT MIGHT OTHERWISE ARISE FROM COURSE OF
 * DEALING OR USAGE OF TRADE.  NO WARRANTY IS EITHER EXPRESS OR IMPLIED WITH
 * RESPECT TO THE USE OF THE SOFTWARE OR DOCUMENTATION. Under no circumstances
 * shall University be liable for incidental, special, indirect, direct or
 * consequential damages or loss of profits, interruption of business, or
 * related expenses which may arise from use of Software or Documentation,
 * including but not limited to those resulting from defects in Software and/or
 * Documentation, or loss or inaccuracy of data of any kind.
 */

#include <sstream>
#include <fstream>
#include <climits>
#include <vector>
#include <unordered_set>
#include "clang/Frontend/FrontendPluginRegistry.h"
#include "clang/AST/AST.h"
#include "clang/AST/ASTConsumer.h"
#include "clang/AST/RecursiveASTVisitor.h"
#include "clang/Frontend/CompilerInstance.h"
#include "clang/Sema/Sema.h"
#include "llvm/Support/raw_ostream.h"

/*
 *  * Matchers
 *   */

#include "clang/ASTMatchers/ASTMatchers.h"
#include "clang/ASTMatchers/ASTMatchFinder.h"

/*
 * Rewriter
 */

#include "clang/Rewrite/Core/Rewriter.h"

using namespace clang;
using namespace clang::ast_matchers;
using namespace llvm;

#define VARIABLE_NAME_CHARACTERS                                               \
  "abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789_"
#define COMPUTE_FILENAME "._compute_kernels"
#define ORCHESTRATOR_FILENAME "._orchestrator_kernel"

// ASSUMPTIONS
#define __GALOIS_PREPROCESS_GLOBAL_VARIABLE_PREFIX__                           \
  "local_" // TODO: detect automatically
#define __GALOIS_BITSET_VARIABLE_PREFIX__ "bitset_"
#define __GALOIS_ACCUMULATOR_TYPE__ "galois::DGAccumulator"
#define __GALOIS_REDUCEMAX_TYPE__ "galois::DGReduceMax"
#define __GALOIS_REDUCEMIN_TYPE__ "galois::DGReduceMin"

namespace {

typedef std::tuple<std::vector<std::string>, std::vector<std::string>,
                   std::vector<std::string>>
    tuple3strings;

// finds and replaces all occurrences
// ensures progress even if replaceWith contains toReplace
void findAndReplace(std::string& str, const std::string& toReplace,
                    const std::string& replaceWith) {
  std::size_t pos = 0;
  while (1) {
    std::size_t found = str.find(toReplace, pos);
    if (found != std::string::npos) {
      str = str.replace(found, toReplace.length(), replaceWith);
      pos = found + replaceWith.length();
    } else {
      break;
    }
  }
}

class IrGLOperatorVisitor : public RecursiveASTVisitor<IrGLOperatorVisitor> {
private:
  // ASTContext* astContext;
  Rewriter&
      rewriter; // FIXME: is this necessary? can ASTContext give source code?
  bool isTopological;
  std::unordered_set<const Stmt*> skipStmts;
  std::unordered_set<const Decl*> skipDecls;
  std::unordered_set<const Stmt*> elseStmts;
  std::map<std::string, std::string> nodeMap;
  std::string funcName;
  std::stringstream declString, varDeclString, bodyString;
  std::map<std::string, std::string> accumulatorToTypeMap;
  std::map<std::string, std::string> reducemaxToTypeMap;
  std::map<std::string, std::string> reduceminToTypeMap;
  std::map<std::string, std::string> parameterToTypeMap;
  std::map<std::string, std::string> globalVariableToTypeMap;
  std::map<std::string, std::string> bitsetVariableToTypeMap;
  std::unordered_set<std::string>
      symbolTable; // FIXME: does not handle scoped variables (nesting with same
                   // variable name)
  std::ofstream Output;

  std::map<std::string, std::string>& SharedVariablesToTypeMap;
  std::map<std::string, tuple3strings>& KernelToArgumentsMap;

  unsigned
      conditional; // if (generated) code is enclosed within an if-condition
  bool conditional_pop; // if (generated) code is enclosed within an (generated)
                        // if-condition for pop

public:
  explicit IrGLOperatorVisitor(
      ASTContext* context, Rewriter& R, bool isTopo,
      std::map<std::string, std::string> accumToTypeMap,
      std::map<std::string, std::string> maxToTypeMap,
      std::map<std::string, std::string> minToTypeMap,
      std::map<std::string, std::string>& sharedVariables,
      std::map<std::string, tuple3strings>& kernels)
      : // astContext(context),
        rewriter(R), isTopological(isTopo),
        accumulatorToTypeMap(accumToTypeMap), reducemaxToTypeMap(maxToTypeMap),
        reduceminToTypeMap(minToTypeMap),
        SharedVariablesToTypeMap(sharedVariables),
        KernelToArgumentsMap(kernels), conditional(0), conditional_pop(false) {}

  virtual ~IrGLOperatorVisitor() {}

  std::string FormatCBlock(std::string text) {
    assert(text.find("graph->getData") == std::string::npos);

    // replace node variables: array of structures to structure of arrays
    // FIXME: does not handle scoped variables (nesting with same variable name)
    for (auto& nodeVar : nodeMap) {
      std::size_t pos = text.find(nodeVar.first);
      while (pos != std::string::npos) {
        std::size_t end = text.find(".", pos);
        text.erase(pos, end - pos + 1);
        end = text.find_first_not_of(VARIABLE_NAME_CHARACTERS, pos);
        if (end == std::string::npos)
          text.append("[" + nodeVar.second + "]");
        else
          text.insert(end, "[" + nodeVar.second + "]");
        text.insert(pos, "p_");
        pos = text.find(nodeVar.first);
      }
    }

    /*while (1) {
      std::size_t found = text.find("nout(");
      if (found != std::string::npos) {
        std::size_t begin = text.find(",", found);
        std::size_t end = text.find(",", begin+1);
        std::string var = text.substr(begin+1, end - begin - 1);
        std::string replace = "graph->getOutDegree(" + var + ")";
        end = text.find(")", end);
        text = text.replace(found, end - found + 1, replace);
      } else {
        break;
      }
    }*/

    // replace std::distance(graph->edge_begin(NODE), graph->edge_end(NODE))
    // with getOutDegree(NODE)
    while (1) {
      std::string distance = "std::distance(graph->edge_begin(";
      std::size_t found    = text.find(distance);
      if (found != std::string::npos) {
        std::size_t begin   = found + distance.length();
        std::size_t end     = text.find(")", begin);
        std::string var     = text.substr(begin, end - begin);
        std::string replace = "graph.getOutDegree(" + var + ")";
        assert(text.find("graph->edge_end(" + var + ")", end) !=
               std::string::npos);
        end  = text.find(")", end + 1); // graph->edge_end()
        end  = text.find(")", end + 1); // std::distance()
        text = text.replace(found, end - found + 1, replace);
      } else {
        break;
      }
    }

    // replace ATOMIC_VAR[NODE].exchange(VALUE) with
    // atomicExch(&ATOMIC_VAR[NODE], VALUE)
    while (1) {
      std::string exchange = ".exchange(";
      std::size_t found    = text.find(exchange);
      if (found != std::string::npos) {
        std::size_t begin   = found + exchange.length();
        std::size_t end     = text.find(")", begin);
        std::string value   = text.substr(begin, end - begin);
        begin               = text.rfind("p_", found - 1);
        std::string var     = text.substr(begin, found - begin);
        std::string replace = "atomicExch(&" + var + ", " + value + ")";
        text                = text.replace(begin, end - begin + 1, replace);
      } else {
        break;
      }
    }

    // use & (address-of) operator for atomic variables in atomicAdd() etc
    {
      std::size_t start = 0;
      while (1) {
        std::size_t found = text.find("galois::atomic", start);
        if (found != std::string::npos) {
          std::size_t replace = text.find("(", start);
          text.insert(replace + 1, "&");
          start = replace;
        } else {
          break;
        }
      }
      findAndReplace(text, "galois::atomic", "atomicTest");
    }
    {
      std::size_t start = 0;
      while (1) {
        std::size_t found = text.find("galois::min", start);
        if (found != std::string::npos) {
          std::size_t replace = text.find("(", start);
          text.insert(replace + 1, "&");
          start = replace;
        } else {
          break;
        }
      }
      findAndReplace(text, "galois::min", "atomicTestMin");
    }
    {
      std::size_t start = 0;
      while (1) {
        std::size_t found = text.find("galois::add", start);
        if (found != std::string::npos) {
          std::size_t replace = text.find("(", start);
          text.insert(replace + 1, "&");
          start = replace;
        } else {
          break;
        }
      }
      findAndReplace(text, "galois::add", "atomicTestAdd");
    }

    {
      std::size_t found = text.find("ctx.push(");
      if (found != std::string::npos) {
        std::size_t begin = text.rfind("(");
        std::size_t end   = text.find(")", begin);
        std::string var   = text.substr(begin + 1, end - begin - 1);
        text              = "WL.push(\"" + var + "\")";
      }
    }

    // replace GRAPH->getGID(VAR) with GRAPH.node_data[VAR]
    while (1) {
      std::size_t found = text.find("->getGID");
      if (found != std::string::npos) {
        std::size_t begin   = text.find("(", found);
        std::size_t end     = text.find(")", begin);
        std::string var     = text.substr(begin + 1, end - begin - 1);
        std::string replace = ".node_data[" + var + "]";
        text                = text.replace(found, end - found + 1, replace);
      } else {
        break;
      }
    }

    findAndReplace(text, " this->",
                   " "); // because IrGL is not aware of the "this" variable
    findAndReplace(text, "->getEdgeDst", ".getAbsDestination");
    findAndReplace(text, "->getEdgeData", ".getAbsWeight");
    findAndReplace(text, "std::fabs", "fabs");
    findAndReplace(text, "\\",
                   "\\\\"); // FIXME: should it be only within quoted strings?

    for (auto& it : accumulatorToTypeMap) {
      auto accumulatorName = it.first;
      if (text.find(accumulatorName) == 0) {
        std::size_t begin = text.find("=");
        std::size_t end   = text.find(";", begin);
        std::string value = text.substr(begin + 1, end - begin - 3);
        text              = accumulatorName + ".reduce(" + value + ")";
      }
    }

    for (auto& it : reducemaxToTypeMap) {
      auto reducemaxName = it.first;
      if (text.find(reducemaxName) == 0) {
        std::size_t begin = text.find("(");
        std::size_t end   = text.find(")", begin);
        std::string value = text.substr(begin + 1, end - begin - 1);
        text              = reducemaxName + ".reduce(" + value + ")";
      }
    }

    for (auto& it : reduceminToTypeMap) {
      auto reduceminName = it.first;
      if (text.find(reduceminName) == 0) {
        std::size_t begin = text.find("(");
        std::size_t end   = text.find(")", begin);
        std::string value = text.substr(begin + 1, end - begin - 1);
        text              = reduceminName + ".reduce(" + value + ")";
      }
    }

    return text;
  }

  void WriteCBlock(std::string text) {
    text              = FormatCBlock(text);
    std::size_t found = text.find("printf");
    if (found != std::string::npos) {
      bodyString << "CBlock(['" << text << "']),\n";
    } else {
      if (text.find("WL.push") != std::string::npos) {
        bodyString << text << ",\n";
      } else if (text.find("ReturnFromParallelFor") != std::string::npos) {
        bodyString << text << ",\n";
      } else {
        bodyString << "CBlock([\"" << text << "\"]),\n";
      }
    }
  }

  std::string FormatType(std::string text) {
    findAndReplace(text, "GNode", "index_type");
    findAndReplace(text, "_Bool", "bool");

    // remove reference (&)
    if (text.rfind("&") == (text.length() - 1)) {
      text = text.substr(0, text.length() - 1);
    }

    // get basic type
    if ((text.find("std::atomic") == 0) ||
        (text.find("struct std::atomic") == 0)) {
      std::size_t begin = text.find("<");
      std::size_t end   = text.find_last_of(">");
      text              = text.substr(begin + 1, end - begin - 1);
    }
    if (text.find("cll::opt") == 0) {
      std::size_t begin = text.find("<");
      std::size_t end   = text.find_last_of(">");
      text              = text.substr(begin + 1, end - begin - 1);
    }

    return text;
  }

  virtual bool TraverseDecl(Decl* D) {
    bool traverse = RecursiveASTVisitor<IrGLOperatorVisitor>::TraverseDecl(D);
    if (traverse && D && isa<CXXMethodDecl>(D)) {
      if (conditional_pop) {
        bodyString << "]),\n"; // end If "pop" (no edge-iterator)
        conditional_pop = false;
      }
      assert(conditional == 0);
      bodyString << "]),\n"; // end ForAll
      for (auto& it : accumulatorToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        auto br   = "cub::BlockReduce<" + type + ", TB_SIZE>";
        bodyString << "CBlock([\"" << name << ".thread_exit<" << br << " >("
                   << name << "_ts)\"], parse = False),\n";
      }
      for (auto& it : reducemaxToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        auto br   = "cub::BlockReduce<" + type + ", TB_SIZE>";
        bodyString << "CBlock([\"" << name << ".thread_exit<" << br << " >("
                   << name << "_ts)\"], parse = False),\n";
      }
      for (auto& it : reduceminToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        auto br   = "cub::BlockReduce<" + type + ", TB_SIZE>";
        bodyString << "CBlock([\"" << name << ".thread_exit<" << br << " >("
                   << name << "_ts)\"], parse = False),\n";
      }
      bodyString << "]),\n"; // end Kernel
      std::vector<std::string> globalArguments, arguments, bitsetArguments;
      for (auto& global : globalVariableToTypeMap) {
        declString << ", ('" << global.second << "', '" << global.first << "')";
        globalArguments.push_back(global.first);
      }
      for (auto& param : parameterToTypeMap) {
        declString << ", ('" << param.second << " *', 'p_" << param.first
                   << "')";
        arguments.push_back(param.first);
        // FIXME: assert type matches if it exists
        SharedVariablesToTypeMap[param.first] = param.second;
      }
      for (auto& bitset : bitsetVariableToTypeMap) {
        declString << ", ('DynamicBitset&', '" << bitset.first << "')";
        auto str = bitset.first;
        findAndReplace(str, __GALOIS_BITSET_VARIABLE_PREFIX__, "");
        bitsetArguments.push_back(str);
      }
      for (auto& it : accumulatorToTypeMap) {
        declString << ", ('HGAccumulator<" << it.second << ">', '" << it.first
                   << "')";
      }
      for (auto& it : reducemaxToTypeMap) {
        declString << ", ('HGReduceMax<" << it.second << ">', '" << it.first
                   << "')";
      }
      for (auto& it : reduceminToTypeMap) {
        declString << ", ('HGReduceMin<" << it.second << ">', '" << it.first
                   << "')";
      }
      KernelToArgumentsMap[funcName] =
          std::make_tuple(globalArguments, arguments, bitsetArguments);
      declString << "],\n[\n";
      Output.open(COMPUTE_FILENAME, std::ios::app);
      Output << declString.str();
      Output << varDeclString.str();
      Output << bodyString.str();
      Output.close();
    }
    return traverse;
  }

  virtual bool TraverseCXXForRangeStmt(CXXForRangeStmt* forStmt) {
    assert(forStmt);
    bool traverse =
        RecursiveASTVisitor<IrGLOperatorVisitor>::TraverseCXXForRangeStmt(
            forStmt);
    if (traverse) {
      bodyString << "]),\n"; // end ForAll
      bodyString << "),\n";  // end ClosureHint
      // ASSUMPTION: no code after edge-loop
      // FIXME: assert that there is no code after the edge-loop
    }
    return traverse;
  }

  virtual bool TraverseForStmt(ForStmt* forStmt) {
    assert(forStmt);
    bool traverse =
        RecursiveASTVisitor<IrGLOperatorVisitor>::TraverseForStmt(forStmt);
    if (traverse) {
      bodyString << "]),\n"; // end ForAll
      bodyString << "),\n";  // end ClosureHint
      // ASSUMPTION: no code after edge-loop
      // FIXME: assert that there is no code after the edge-loop
    }
    return traverse;
  }

  virtual bool TraverseIfStmt(IfStmt* ifStmt) {
    assert(ifStmt);
    const Stmt* elseStmt = ifStmt->getElse();
    if (elseStmt) {
      assert(elseStmts.find(elseStmt) == elseStmts.end());
      elseStmts.insert(elseStmt);
    }
    bool traverse =
        RecursiveASTVisitor<IrGLOperatorVisitor>::TraverseIfStmt(ifStmt);
    if (elseStmt) {
      assert(elseStmts.find(elseStmt) == elseStmts.end());
    }
    if (traverse) {
      if (conditional > 0) {
        bodyString << "]),\n"; // end If
        --conditional;
      }
    }
    return traverse;
  }

  virtual bool TraverseStmt(Stmt* S) {
    bool elseStmt = (elseStmts.find(S) != elseStmts.end());
    if (elseStmt) {
      bodyString << "],\n[\n"; // else of If
      elseStmts.erase(S);
    }
    bool traverse = RecursiveASTVisitor<IrGLOperatorVisitor>::TraverseStmt(S);
    return traverse;
  }

  virtual bool VisitCXXMethodDecl(CXXMethodDecl* func) {
    assert(declString.rdbuf()->in_avail() == 0);
    declString.str("");
    bodyString.str("");
    funcName = func->getParent()->getNameAsString();
    declString << "Kernel(\"" << funcName << "\", [G.param()";
    // ASSUMPTION: first parameter is GNode
    std::string vertexName = func->getParamDecl(0)->getNameAsString();
    if (isTopological) {
      declString << ", ('unsigned int', '__begin'), ('unsigned int', '__end')";
      for (auto& it : accumulatorToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        auto br   = "cub::BlockReduce<" + type + ", TB_SIZE>";
        bodyString << "CDecl([(\"__shared__ " << br << "::TempStorage\", \""
                   << name << "_ts\", \"\")]),\n";
        bodyString << "CBlock([\"" << name << ".thread_entry()\"]),\n";
      }
      for (auto& it : reducemaxToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        auto br   = "cub::BlockReduce<" + type + ", TB_SIZE>";
        bodyString << "CDecl([(\"__shared__ " << br << "::TempStorage\", \""
                   << name << "_ts\", \"\")]),\n";
        bodyString << "CBlock([\"" << name << ".thread_entry()\"]),\n";
      }
      for (auto& it : reduceminToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        auto br   = "cub::BlockReduce<" + type + ", TB_SIZE>";
        bodyString << "CDecl([(\"__shared__ " << br << "::TempStorage\", \""
                   << name << "_ts\", \"\")]),\n";
        bodyString << "CBlock([\"" << name << ".thread_entry()\"]),\n";
      }
      bodyString << "ForAll(\"" << vertexName
                 << "\", G.nodes(\"__begin\", \"__end\"),\n[\n";
      bodyString << "CDecl([(\"bool\", \"pop\", \" = " << vertexName
                 << " < __end\")]),\n"; // hack for IrGL compiler
      bodyString
          << "If(\"pop\", [\n"; // hack to handle nested parallelism (should be
                                // handled by IrGL compiler in the future)
      conditional_pop = true;
    } else {
      bodyString << "ForAll(\"wlvertex\", WL.items(),\n[\n";
      bodyString << "CDecl([(\"int\", \"" << vertexName << "\", \"\")]),\n";
      bodyString << "CDecl([(\"bool\", \"pop\", \"\")]),\n";
      bodyString << "WL.pop(\"pop\", \"wlvertex\", \"" << vertexName
                 << "\"),\n";
      bodyString
          << "If(\"pop\", [\n"; // hack to handle unsuccessful pop (should be
                                // handled by IrGL compiler in the future)
      conditional_pop = true;
    }
    symbolTable.insert(vertexName);
    return true;
  }

  virtual bool VisitCXXForRangeStmt(CXXForRangeStmt* forStmt) {
    assert(forStmt->getLoopVariable());
    skipDecls.insert(forStmt->getLoopVariable());
    DeclStmt* rangeStmt = forStmt->getRangeStmt();
    for (auto decl : rangeStmt->decls()) {
      if (VarDecl* varDecl = dyn_cast<VarDecl>(decl)) {
        const Expr* expr = varDecl->getAnyInitializer();
        skipStmts.insert(expr);
      }
    }
    std::string variableName = forStmt->getLoopVariable()->getNameAsString();
    symbolTable.insert(variableName);
    const Expr* expr       = forStmt->getRangeInit();
    std::string vertexName = rewriter.getRewrittenText(expr->getSourceRange());
    std::size_t found      = vertexName.find("graph->edges");
    assert(found != std::string::npos);
    std::size_t begin = vertexName.find("(", found);
    std::size_t end   = vertexName.find(",", begin);
    if (end == std::string::npos)
      end = vertexName.find(")", begin);
    vertexName = vertexName.substr(begin + 1, end - begin - 1);
    while (conditional > 0) {
      // ASSUMPTION: for loop is in If block, which does not have a
      // corresponding Else block
      bodyString << "], [ CBlock([\"pop = false\"]), ]),\n"; // end If
      --conditional;
    }
    bodyString << "]),\n"; // end If "pop"
    conditional_pop = false;
    bodyString
        << "UniformConditional(If(\"!pop\", [CBlock(\"continue\")]), "
           "uniform_only = False, _only_if_np = True),\n"; // hack to handle
                                                           // non-nested
                                                           // parallelism
    bodyString << "ClosureHint(\n";
    bodyString << "ForAll(\"" << variableName << "\", G.edges(\"" << vertexName
               << "\"),\n[\n";
    return true;
  }

  virtual bool VisitForStmt(ForStmt* forStmt) {
    skipStmts.insert(forStmt->getInit());
    skipStmts.insert(forStmt->getCond());
    skipStmts.insert(forStmt->getInc());
    std::string variableName, vertexName;
    DeclStmt* declStmt = dyn_cast<DeclStmt>(forStmt->getInit());
    for (auto decl : declStmt->decls()) {
      if (VarDecl* varDecl = dyn_cast<VarDecl>(decl)) {
        std::string text = rewriter.getRewrittenText(varDecl->getSourceRange());
        std::size_t found = text.find("graph->edge_begin");
        if (found != std::string::npos) {
          std::size_t begin = text.find("(", found);
          std::size_t end   = text.find(",", begin);
          if (end == std::string::npos)
            end = text.find(")", begin);
          vertexName   = text.substr(begin + 1, end - begin - 1);
          variableName = varDecl->getNameAsString();
          break;
        }
      }
    }
    symbolTable.insert(variableName);
    while (conditional > 0) {
      // ASSUMPTION: for loop is in If block, which does not have a
      // corresponding Else block
      bodyString << "], [ CBlock([\"pop = false\"]), ]),\n"; // end If
      --conditional;
    }
    bodyString << "]),\n"; // end If "pop"
    conditional_pop = false;
    bodyString
        << "UniformConditional(If(\"!pop\", [CBlock(\"continue\")]), "
           "uniform_only = False, _only_if_np = True),\n"; // hack to handle
                                                           // non-nested
                                                           // parallelism
    bodyString << "ClosureHint(\n";
    bodyString << "ForAll(\"" << variableName << "\", G.edges(\"" << vertexName
               << "\"),\n[\n";
    return true;
  }

  virtual bool VisitDeclStmt(DeclStmt* declStmt) {
    bool skipStmt = skipStmts.find(declStmt) != skipStmts.end();
    for (auto decl : declStmt->decls()) {
      if (VarDecl* varDecl = dyn_cast<VarDecl>(decl)) {
        const Expr* expr = varDecl->getAnyInitializer();
        skipStmts.insert(expr);
        if (!skipStmt && (skipDecls.find(varDecl) == skipDecls.end())) {
          symbolTable.insert(varDecl->getNameAsString());
          std::size_t pos = std::string::npos;
          std::string initText;
          if (expr != NULL) {
            initText = rewriter.getRewrittenText(expr->getSourceRange());
            pos      = initText.find("graph->getData");
          }
          if (pos == std::string::npos) {
            std::string decl = varDecl->getType().getAsString();
            if (conditional_pop) {
              varDeclString << "CDecl([(\"" << FormatType(decl) << "\", \""
                            << varDecl->getNameAsString() << "\", \"\")]),\n";
            } else {
              bodyString << "CDecl([(\"" << FormatType(decl) << "\", \""
                         << varDecl->getNameAsString() << "\", \"\")]),\n";
            }

            if (!initText.empty()) {
              WriteCBlock(varDecl->getNameAsString() + " = " + initText);
            }
          } else {
            // get the first argument to graph->getData()
            std::size_t begin = initText.find("(", pos);
            std::size_t end   = initText.find(",", pos);
            if (end == std::string::npos) {
              end = initText.find(")", pos);
            }
            std::string index = initText.substr(begin + 1, end - begin - 1);
            nodeMap[varDecl->getNameAsString()] = index;
          }
        }
      }
    }
    return true;
  }

  virtual bool VisitBinaryOperator(BinaryOperator* binaryOp) {
    skipStmts.insert(binaryOp->getLHS());
    skipStmts.insert(binaryOp->getRHS());
    if (skipStmts.find(binaryOp) == skipStmts.end()) {
      assert(binaryOp->isAssignmentOp());
      WriteCBlock(rewriter.getRewrittenText(binaryOp->getSourceRange()));
    }
    return true;
  }

  virtual bool VisitCallExpr(CallExpr* callExpr) {
    symbolTable.insert(
        callExpr->getCalleeDecl()->getAsFunction()->getNameAsString());
    if (skipStmts.find(callExpr) == skipStmts.end()) {
      CXXMemberCallExpr* cxxCallExpr = dyn_cast<CXXMemberCallExpr>(callExpr);
      CXXRecordDecl* record          = NULL;
      if (cxxCallExpr) {
        record = cxxCallExpr->getRecordDecl();
      }
      if (record &&
          (record->getNameAsString().compare("GReducible") == 0)) { // not used
        Expr* implicit = cxxCallExpr->getImplicitObjectArgument();
        if (CastExpr* cast = dyn_cast<CastExpr>(implicit)) {
          implicit = cast->getSubExpr();
        }
        SourceLocation start;
        if (MemberExpr* member = dyn_cast<MemberExpr>(implicit)) {
          start = member->getMemberLoc();
        } else {
          start = implicit->getLocStart();
        }
        SourceRange range(start, implicit->getLocEnd());
        std::string reduceVar = rewriter.getRewrittenText(range);

        assert(callExpr->getNumArgs() == 1);
        Expr* argument  = callExpr->getArg(0);
        std::string var = rewriter.getRewrittenText(argument->getSourceRange());

        WriteCBlock("p_" + reduceVar + "[vertex] = " + var);
        std::string type =
            argument->getType().getUnqualifiedType().getAsString();
        parameterToTypeMap[reduceVar] = FormatType(type);
      } else {
        if (record &&
            (record->getNameAsString().compare("DynamicBitSet") == 0)) {
          auto varName = rewriter.getRewrittenText(
              cxxCallExpr->getImplicitObjectArgument()->getSourceRange());
          if (bitsetVariableToTypeMap.find(varName) ==
              bitsetVariableToTypeMap.end()) {
            bitsetVariableToTypeMap[varName] = "DynamicBitset"; // type ignored
          }
        }
        WriteCBlock(rewriter.getRewrittenText(callExpr->getSourceRange()));
      }
    }
    for (unsigned i = 0; i < callExpr->getNumArgs(); ++i) {
      skipStmts.insert(callExpr->getArg(i));
    }
    return true;
  }

  virtual bool VisitIfStmt(IfStmt* ifStmt) {
    const Expr* expr = ifStmt->getCond();
    std::string text =
        FormatCBlock(rewriter.getRewrittenText(expr->getSourceRange()));
    bodyString << "If(\"" << text << "\",\n[\n";
    ++conditional;
    skipStmts.insert(expr);
    return true;
  }

  virtual bool VisitCastExpr(CastExpr* castExpr) {
    if (skipStmts.find(castExpr) != skipStmts.end()) {
      skipStmts.insert(castExpr->getSubExpr());
    }
    return true;
  }

  virtual bool VisitCXXConstructExpr(CXXConstructExpr* constructExpr) {
    if (skipStmts.find(constructExpr) != skipStmts.end()) {
      for (auto arg : constructExpr->arguments()) {
        skipStmts.insert(arg);
      }
    }
    return true;
  }

  virtual bool
  VisitMaterializeTemporaryExpr(MaterializeTemporaryExpr* tempExpr) {
    if (skipStmts.find(tempExpr) != skipStmts.end()) {
      skipStmts.insert(tempExpr->GetTemporaryExpr());
    }
    return true;
  }

  virtual bool VisitParenExpr(ParenExpr* parenExpr) {
    if (skipStmts.find(parenExpr) != skipStmts.end()) {
      skipStmts.insert(parenExpr->getSubExpr());
    }
    return true;
  }

  virtual bool
  VisitConditionalOperator(ConditionalOperator* conditionalOperator) {
    if (skipStmts.find(conditionalOperator) != skipStmts.end()) {
      skipStmts.insert(conditionalOperator->getCond());
      skipStmts.insert(conditionalOperator->getTrueExpr());
      skipStmts.insert(conditionalOperator->getFalseExpr());
    }
    return true;
  }

  virtual bool VisitMemberExpr(MemberExpr* memberExpr) {
    std::string objectName =
        rewriter.getRewrittenText(memberExpr->getBase()->getSourceRange());
    const std::string& varName = memberExpr->getMemberNameInfo().getAsString();
    std::string type           = memberExpr->getMemberDecl()
                           ->getType()
                           .getUnqualifiedType()
                           .getAsString();
    if (nodeMap.find(objectName) != nodeMap.end()) {
      parameterToTypeMap[varName] = FormatType(type);
    } else if ((symbolTable.find(varName) == symbolTable.end()) &&
               (accumulatorToTypeMap.find(varName) ==
                accumulatorToTypeMap.end()) &&
               (reducemaxToTypeMap.find(varName) == reducemaxToTypeMap.end()) &&
               (reduceminToTypeMap.find(varName) == reduceminToTypeMap.end()) &&
               (type.find("Graph") != 0)) {
      if (globalVariableToTypeMap.find(varName) ==
          globalVariableToTypeMap.end()) {
        assert(varName.find(__GALOIS_PREPROCESS_GLOBAL_VARIABLE_PREFIX__) == 0);
        globalVariableToTypeMap[varName] = FormatType(type);
      }
    }
    return true;
  }
};

class IrGLOrchestratorVisitor
    : public RecursiveASTVisitor<IrGLOrchestratorVisitor> {
private:
  ASTContext* astContext;
  Rewriter&
      rewriter; // FIXME: is this necessary? can ASTContext give source code?
  bool requiresWorklist; // shared across kernels
  std::ofstream Output;

  std::map<std::string, std::string> accumulatorToTypeMap;
  std::map<std::string, std::string> reducemaxToTypeMap;
  std::map<std::string, std::string> reduceminToTypeMap;
  std::map<std::string, std::string> SharedVariablesToTypeMap;
  std::map<std::string, tuple3strings> KernelToArgumentsMap;
  std::map<std::string, std::vector<std::string>> HostKernelsToArgumentsMap;

  std::string FileNamePath;

public:
  explicit IrGLOrchestratorVisitor(ASTContext* context, Rewriter& R)
      : astContext(context), rewriter(R), requiresWorklist(false) {}

  virtual ~IrGLOrchestratorVisitor() {}

  void GenerateIrGLASTToFile() const {
    std::string filename;
    std::size_t found = FileNamePath.rfind("/");
    if (found != std::string::npos)
      filename = FileNamePath.substr(found + 1, FileNamePath.length() - 1);
    else
      filename = FileNamePath;

    std::ofstream header;
    header.open(FileNamePath + "_cuda.h");
    header << "#pragma once\n\n";
    header << "#include \"galois/runtime/DataCommMode.h\"\n";
    header << "#include \"galois/cuda/HostDecls.h\"\n\n";
    if (requiresWorklist) {
      header << "\nstruct CUDA_Worklist {\n";
      header << "\tint *in_items;\n";
      header << "\tint num_in_items;\n";
      header << "\tint *out_items;\n";
      header << "\tint num_out_items;\n";
      header << "\tint max_size;\n";
      header << "};\n\n";
    }
    for (auto& var : SharedVariablesToTypeMap) {
      header << "void get_bitset_" << var.first
             << "_cuda(struct CUDA_Context* ctx, uint64_t* bitset_compute);\n";
      header << "void bitset_" << var.first
             << "_reset_cuda(struct CUDA_Context* ctx);\n";
      header << "void bitset_" << var.first
             << "_reset_cuda(struct CUDA_Context* ctx, size_t begin, size_t "
                "end);\n";
      header << var.second << " get_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned LID);\n";
      header << "void set_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned LID, " << var.second
             << " v);\n";
      header << "void add_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned LID, " << var.second
             << " v);\n";
      header << "bool min_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned LID, " << var.second
             << " v);\n";
      header << "void batch_get_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, uint8_t* v);\n";
      header << "void batch_get_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, size_t* v_size, DataCommMode* data_mode);\n";
      header << "void batch_get_mirror_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, uint8_t* v);\n";
      header << "void batch_get_mirror_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, size_t* v_size, DataCommMode* data_mode);\n";
      header << "void batch_get_reset_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, " << var.second << " i);\n";
      header << "void batch_get_reset_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, size_t* v_size, DataCommMode* data_mode, "
             << var.second << " i);\n";
      header << "void batch_set_mirror_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t *v, DataCommMode data_mode);\n";
      header << "void batch_set_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, DataCommMode data_mode);\n";
      header << "void batch_add_mirror_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, DataCommMode data_mode);\n";
      header << "void batch_add_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, DataCommMode data_mode);\n";
      header << "void batch_min_mirror_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, DataCommMode data_mode);\n";
      header << "void batch_min_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
             << "uint8_t* v, DataCommMode data_mode);\n";
      header << "void batch_reset_node_" << var.first
             << "_cuda(struct CUDA_Context* ctx, size_t begin, size_t end, "
             << var.second << " v);\n\n";
    }
    for (auto& kernel : HostKernelsToArgumentsMap) {
      header << "void " << kernel.first << "_cuda(";
      for (auto& argument : kernel.second) {
        header << argument << ", ";
      }
      header << "struct CUDA_Context* ctx);\n";
    }
    header.close();

    std::ofstream cuheader;
    cuheader.open(FileNamePath + "_cuda.cuh");
    cuheader << "#pragma once\n";
    cuheader << "#include <cuda.h>\n";
    cuheader << "#include <stdio.h>\n";
    cuheader << "#include <sys/types.h>\n";
    cuheader << "#include <unistd.h>\n";
    cuheader << "#include \"" << filename << "_cuda.h\"\n";
    cuheader << "#include \"galois/runtime/cuda/DeviceSync.h\"\n\n";
    cuheader << "struct CUDA_Context : public CUDA_Context_Common {\n";
    for (auto& var : SharedVariablesToTypeMap) {
      cuheader << "\tstruct CUDA_Context_Field<" << var.second << "> "
               << var.first << ";\n";
    }
    if (requiresWorklist) {
      cuheader << "\tWorklist2 in_wl;\n";
      cuheader << "\tWorklist2 out_wl;\n";
      cuheader << "\tstruct CUDA_Worklist *shared_wl;\n";
    }
    cuheader << "};\n\n";
    cuheader << "struct CUDA_Context* get_CUDA_context(int id) {\n";
    cuheader << "\tstruct CUDA_Context* ctx;\n";
    cuheader << "\tctx = (struct CUDA_Context* ) calloc(1, sizeof(struct "
                "CUDA_Context));\n";
    cuheader << "\tctx->id = id;\n";
    cuheader << "\treturn ctx;\n";
    cuheader << "}\n\n";
    cuheader
        << "bool init_CUDA_context(struct CUDA_Context* ctx, int device) {\n";
    cuheader << "\treturn init_CUDA_context_common(ctx, device);\n";
    cuheader << "}\n\n";
    cuheader << "void load_graph_CUDA(struct CUDA_Context* ctx, ";
    if (requiresWorklist) {
      cuheader << "struct CUDA_Worklist *wl, double wl_dup_factor, ";
    }
    cuheader << "MarshalGraph &g, unsigned num_hosts) {\n";
    cuheader << "\tsize_t mem_usage = mem_usage_CUDA_common(g, num_hosts);\n";
    for (auto& var : SharedVariablesToTypeMap) {
      cuheader << "\tmem_usage += mem_usage_CUDA_field(&ctx->" << var.first
               << ", g, num_hosts);\n";
    }
    cuheader << "\tprintf(\"[%d] Host memory for communication context: %3u "
                "MB\\n\", ctx->id, mem_usage/1048756);\n";
    cuheader << "\tload_graph_CUDA_common(ctx, g, num_hosts);\n";
    for (auto& var : SharedVariablesToTypeMap) {
      cuheader << "\tload_graph_CUDA_field(ctx, &ctx->" << var.first
               << ", num_hosts);\n";
    }
    if (requiresWorklist) {
      cuheader << "\twl->max_size = wl_dup_factor*g.nnodes;\n";
      cuheader << "\tctx->in_wl = Worklist2((size_t)wl->max_size);\n";
      cuheader << "\tctx->out_wl = Worklist2((size_t)wl->max_size);\n";
      cuheader << "\twl->num_in_items = -1;\n";
      cuheader << "\twl->num_out_items = -1;\n";
      cuheader << "\twl->in_items = ctx->in_wl.wl;\n";
      cuheader << "\twl->out_items = ctx->out_wl.wl;\n";
      cuheader << "\tctx->shared_wl = wl;\n";
      cuheader << "\tprintf(\"[%d] load_graph_GPU: worklist size %d\\n\", "
                  "ctx->id, (size_t)wl_dup_factor*g.nnodes);\n";
    }
    cuheader << "\treset_CUDA_context(ctx);\n";
    cuheader << "}\n\n";
    cuheader << "void reset_CUDA_context(struct CUDA_Context* ctx) {\n";
    for (auto& var : SharedVariablesToTypeMap) {
      cuheader << "\tctx->" << var.first
               << ".data.zero_gpu();\n"; // FIXME: should do this only for
                                         // std::atomic variables?
    }
    cuheader << "}\n\n";
    for (auto& var : SharedVariablesToTypeMap) {
      cuheader
          << "void get_bitset_" << var.first
          << "_cuda(struct CUDA_Context* ctx, uint64_t* bitset_compute) {\n";
      cuheader << "\tctx->" << var.first
               << ".is_updated.cpu_rd_ptr()->copy_to_cpu(bitset_compute);\n";
      cuheader << "}\n\n";
      cuheader << "void bitset_" << var.first
               << "_reset_cuda(struct CUDA_Context* ctx) {\n";
      cuheader << "\tctx->" << var.first
               << ".is_updated.cpu_rd_ptr()->reset();\n";
      cuheader << "}\n\n";
      cuheader << "void bitset_" << var.first
               << "_reset_cuda(struct CUDA_Context* ctx, size_t begin, size_t "
                  "end) {\n";
      cuheader << "\treset_bitset_field(&ctx->" << var.first
               << ", begin, end);\n";
      cuheader << "}\n\n";
      cuheader << var.second << " get_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned LID) {\n";
      cuheader << "\t" << var.second << " *" << var.first << " = ctx->"
               << var.first << ".data.cpu_rd_ptr();\n";
      cuheader << "\treturn " << var.first << "[LID];\n";
      cuheader << "}\n\n";
      cuheader << "void set_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned LID, " << var.second
               << " v) {\n";
      cuheader << "\t" << var.second << " *" << var.first << " = ctx->"
               << var.first << ".data.cpu_wr_ptr();\n";
      cuheader << "\t" << var.first << "[LID] = v;\n";
      cuheader << "}\n\n";
      cuheader << "void add_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned LID, " << var.second
               << " v) {\n";
      cuheader << "\t" << var.second << " *" << var.first << " = ctx->"
               << var.first << ".data.cpu_wr_ptr();\n";
      cuheader << "\t" << var.first << "[LID] += v;\n";
      cuheader << "}\n\n";
      cuheader << "bool min_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned LID, " << var.second
               << " v) {\n";
      cuheader << "\t" << var.second << " *" << var.first << " = ctx->"
               << var.first << ".data.cpu_wr_ptr();\n";
      cuheader << "\tif (" << var.first << "[LID] > v){\n";
      cuheader << "\t\t" << var.first << "[LID] = v;\n";
      cuheader << "\t\treturn true;\n";
      cuheader << "\t}\n";
      cuheader << "\treturn false;\n";
      cuheader << "}\n\n";
      cuheader << "void batch_get_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v) {\n";
      cuheader << "\tbatch_get_shared_field<" << var.second
               << ", sharedMaster, false>(ctx, &ctx->" << var.first
               << ", from_id, v);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_get_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, size_t* v_size, DataCommMode* data_mode) {\n";
      cuheader << "\tbatch_get_shared_field<" << var.second
               << ", sharedMaster, false>(ctx, &ctx->" << var.first
               << ", from_id, v, v_size, data_mode);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_get_mirror_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v) {\n";
      cuheader << "\tbatch_get_shared_field<" << var.second
               << ", sharedMirror, false>(ctx, &ctx->" << var.first
               << ", from_id, v);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_get_mirror_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, size_t* v_size, DataCommMode* data_mode) {\n";
      cuheader << "\tbatch_get_shared_field<" << var.second
               << ", sharedMirror, false>(ctx, &ctx->" << var.first
               << ", from_id, v, v_size, data_mode);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_get_reset_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, " << var.second << " i) {\n";
      cuheader << "\tbatch_get_shared_field<" << var.second
               << ", sharedMirror, true>(ctx, &ctx->" << var.first
               << ", from_id, v, i);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_get_reset_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, size_t* v_size, DataCommMode* data_mode, "
               << var.second << " i) {\n";
      cuheader
          << "\tbatch_get_shared_field<" << var.second
          << ", sharedMirror, true>(ctx, &ctx->" << var.first
          << ", from_id, v, v_size, data_mode, i);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_set_mirror_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, DataCommMode data_mode) {\n";
      cuheader << "\tbatch_set_shared_field<" << var.second
               << ", sharedMirror, setOp>(ctx, &ctx->" << var.first
               << ", from_id, v, data_mode);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_set_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, DataCommMode data_mode) {\n";
      cuheader << "\tbatch_set_shared_field<" << var.second
               << ", sharedMaster, setOp>(ctx, &ctx->" << var.first
               << ", from_id, v, data_mode);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_add_mirror_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, DataCommMode data_mode) {\n";
      cuheader << "\tbatch_set_shared_field<" << var.second
               << ", sharedMirror, addOp>(ctx, &ctx->" << var.first
               << ", from_id, v, data_mode);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_add_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, DataCommMode data_mode) {\n";
      cuheader << "\tbatch_set_shared_field<" << var.second
               << ", sharedMaster, addOp>(ctx, &ctx->" << var.first
               << ", from_id, v, data_mode);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_min_mirror_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, DataCommMode data_mode) {\n";
      cuheader << "\tbatch_set_shared_field<" << var.second
               << ", sharedMirror, minOp>(ctx, &ctx->" << var.first
               << ", from_id, v, data_mode);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_min_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, unsigned from_id, "
               << "uint8_t* v, DataCommMode data_mode) {\n";
      cuheader << "\tbatch_set_shared_field<" << var.second
               << ", sharedMaster, minOp>(ctx, &ctx->" << var.first
               << ", from_id, v, data_mode);\n";
      cuheader << "}\n\n";
      cuheader << "void batch_reset_node_" << var.first
               << "_cuda(struct CUDA_Context* ctx, size_t begin, size_t end, "
               << var.second << " v) {\n";
      cuheader << "\treset_data_field<" << var.second << ">(&ctx->" << var.first
               << ", begin, end, v);\n";
      cuheader << "}\n\n";
    }
    cuheader.close();

    std::ofstream IrGLAST;
    IrGLAST.open(FileNamePath + "_cuda.py");
    IrGLAST << "from gg.ast import *\n";
    IrGLAST << "from gg.lib.graph import Graph\n";
    IrGLAST << "from gg.lib.wl import Worklist\n";
    IrGLAST << "from gg.ast.params import GraphParam\n";
    IrGLAST << "import cgen\n";
    IrGLAST << "G = Graph(\"graph\")\n";
    IrGLAST << "WL = Worklist()\n";
    IrGLAST << "ast = Module([\n";
    IrGLAST << "CBlock([cgen.Include(\"kernels/reduce.cuh\", "
            << "system = False)], parse = False),\n";
    IrGLAST << "CBlock([cgen.Include(\"" << filename << "_cuda.cuh\", "
            << "system = False)], parse = False),\n";
    for (auto& var : SharedVariablesToTypeMap) {
      std::string upper = var.first;
      std::transform(upper.begin(), upper.end(), upper.begin(), ::toupper);
    }

    std::ifstream compute(COMPUTE_FILENAME);
    IrGLAST << compute.rdbuf();
    compute.close();
    remove(COMPUTE_FILENAME);

    std::ifstream orchestrator(ORCHESTRATOR_FILENAME);
    IrGLAST << orchestrator.rdbuf();
    orchestrator.close();
    remove(ORCHESTRATOR_FILENAME);

    IrGLAST << "])\n"; // end Module
    IrGLAST.close();

    llvm::errs() << "IrGL file and headers generated:\n";
    llvm::errs() << FileNamePath << "_cuda.py\n";
    llvm::errs() << FileNamePath << "_cuda.cuh\n";
    llvm::errs() << FileNamePath << "_cuda.h\n";
  }

  std::string FormatType(std::string text) {
    findAndReplace(text, "GNode", "index_type");
    findAndReplace(text, "_Bool", "bool");

    // get basic type
    if (text.find("std::atomic") == 0) {
      std::size_t begin = text.find("<");
      std::size_t end   = text.find_last_of(">");
      text              = text.substr(begin + 1, end - begin - 1);
    }
    if (text.find("cll::opt") == 0) {
      std::size_t begin = text.find("<");
      std::size_t end   = text.find_last_of(">");
      text              = text.substr(begin + 1, end - begin - 1);
    }

    return text;
  }

  virtual bool TraverseDecl(Decl* D) {
    bool traverse =
        RecursiveASTVisitor<IrGLOrchestratorVisitor>::TraverseDecl(D);
    return traverse;
  }

  virtual bool VisitCallExpr(CallExpr* callExpr) {
    std::string text   = rewriter.getRewrittenText(callExpr->getSourceRange());
    bool isTopological = true;
    std::size_t begin  = text.find("galois::do_all");
    if (begin != 0) {
      begin         = text.find("galois::for_each");
      isTopological = false;
    }
    if (begin == 0) {
      // generate kernel for operator
      assert(callExpr->getNumArgs() > 2);
      // ASSUMPTION: Argument 1 (0-indexed) is operator
      auto arg_type = callExpr->getArg(1)->getType();
      assert(arg_type->isStructureType());
      // TODO: use anonymous function instead of class-operator()
      RecordDecl* record       = arg_type->getAsStructureType()->getDecl();
      CXXRecordDecl* cxxrecord = dyn_cast<CXXRecordDecl>(record);
      // find if there are any accumulators being used
      for (auto decl : cxxrecord->decls()) {
        auto field = dyn_cast<FieldDecl>(decl);
        if (field) {
          auto templateType = field->getType()
                                  .getNonReferenceType()
                                  .getSingleStepDesugaredType(*astContext)
                                  .getTypePtr()
                                  ->getAs<TemplateSpecializationType>();
          if (templateType) {
            std::string type =
                templateType->getCanonicalTypeInternal().getAsString();
            if (type.find(__GALOIS_ACCUMULATOR_TYPE__) != std::string::npos) {
              auto accumulatorName = field->getNameAsString();
              auto argType = templateType->getArg(0).getAsType().getAsString();
              accumulatorToTypeMap[accumulatorName] = argType;
            } else if (type.find(__GALOIS_REDUCEMAX_TYPE__) !=
                       std::string::npos) {
              auto reducemaxName = field->getNameAsString();
              auto argType = templateType->getArg(0).getAsType().getAsString();
              reducemaxToTypeMap[reducemaxName] = argType;
            } else if (type.find(__GALOIS_REDUCEMIN_TYPE__) !=
                       std::string::npos) {
              auto reduceminName = field->getNameAsString();
              auto argType = templateType->getArg(0).getAsType().getAsString();
              reduceminToTypeMap[reduceminName] = argType;
            }
          }
        }
      }
      for (auto method : cxxrecord->methods()) {
        // FIXME: handle overloading/polymorphism
        if (method->getNameAsString().compare("operator()") == 0) {
          // FIXME: should the arguments to the call be parsed?
          // llvm::errs() << "\nVertex Operator:\n" <<
          //  rewriter.getRewrittenText(method->getSourceRange()) << "\n";
          // method->dump();
          IrGLOperatorVisitor operatorVisitor(
              astContext, rewriter, isTopological, accumulatorToTypeMap,
              reducemaxToTypeMap, reduceminToTypeMap, SharedVariablesToTypeMap,
              KernelToArgumentsMap);
          operatorVisitor.TraverseDecl(method);
          break;
        }
      }

      std::string kernelName = cxxrecord->getNameAsString();

      // kernel signature
      Output.open(ORCHESTRATOR_FILENAME, std::ios::app);
      std::stringstream Output_topo_all, Output_topo_all_body;
      std::stringstream Output_topo_masters, Output_topo_masters_body;
      std::stringstream Output_topo_nodeswithedges,
          Output_topo_nodeswithedges_body;
      std::stringstream Output_topo_common, Output_topo_common_body;
      Output << "Kernel(\"" << kernelName << "_cuda\", [";
      if (isTopological) {
        Output_topo_all << "Kernel(\"" << kernelName << "_allNodes_cuda\", [";
        Output_topo_all_body << "CBlock([\"" << kernelName << "_cuda(";
        Output_topo_all_body << "0, ctx->gg.nnodes, ";
        Output_topo_masters << "Kernel(\"" << kernelName
                            << "_masterNodes_cuda\", [";
        Output_topo_masters_body << "CBlock([\"" << kernelName << "_cuda(";
        Output_topo_masters_body
            << "ctx->beginMaster, ctx->beginMaster + ctx->numOwned, ";
        Output_topo_nodeswithedges << "Kernel(\"" << kernelName
                                   << "_nodesWithEdges_cuda\", [";
        Output_topo_nodeswithedges_body << "CBlock([\"" << kernelName
                                        << "_cuda(";
        Output_topo_nodeswithedges_body << "0, ctx->numNodesWithEdges, ";
      }
      std::vector<std::string> arguments, arguments_topo_common;
      if (isTopological) {
        Output << "('unsigned int ', '__begin'), ('unsigned int ', '__end'), ";
        arguments.push_back("unsigned int __begin");
        arguments.push_back("unsigned int __end");
      }
      for (auto& it : accumulatorToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        Output << "('" << type << " &', '" << name << "'), ";
        arguments.push_back(type + " & " + name);
        if (isTopological) {
          Output_topo_common << "('" << type << " &', '" << name << "'), ";
          Output_topo_common_body << name << ", ";
          arguments_topo_common.push_back(type + " & " + name);
        }
      }
      for (auto& it : reducemaxToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        Output << "('" << type << " &', '" << name << "'), ";
        arguments.push_back(type + " & " + name);
        if (isTopological) {
          Output_topo_common << "('" << type << " &', '" << name << "'), ";
          Output_topo_common_body << name << ", ";
          arguments_topo_common.push_back(type + " & " + name);
        }
      }
      for (auto& it : reduceminToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        Output << "('" << type << " &', '" << name << "'), ";
        arguments.push_back(type + " & " + name);
        if (isTopological) {
          Output_topo_common << "('" << type << " &', '" << name << "'), ";
          Output_topo_common_body << name << ", ";
          arguments_topo_common.push_back(type + " & " + name);
        }
      }
      for (auto field : cxxrecord->fields()) {
        std::string name = field->getNameAsString();
        if (name.find(__GALOIS_PREPROCESS_GLOBAL_VARIABLE_PREFIX__) == 0) {
          std::string type = FormatType(field->getType().getAsString());
          Output << "('" << type << "', '" << name << "'), ";
          arguments.push_back(type + " " + name);
          if (isTopological) {
            Output_topo_common << "('" << type << "', '" << name << "'), ";
            Output_topo_common_body << name << ", ";
            arguments_topo_common.push_back(type + " " + name);
          }
        }
      }
      HostKernelsToArgumentsMap[kernelName] = arguments;
      Output << "('struct CUDA_Context* ', 'ctx')],\n[\n";
      if (isTopological) {
        HostKernelsToArgumentsMap[kernelName + "_allNodes"] =
            arguments_topo_common;
        HostKernelsToArgumentsMap[kernelName + "_masterNodes"] =
            arguments_topo_common;
        HostKernelsToArgumentsMap[kernelName + "_nodesWithEdges"] =
            arguments_topo_common;
        Output_topo_common << "('struct CUDA_Context* ', 'ctx')],\n[\n";
        Output_topo_common_body << "ctx)\"]),\n";
      }

      // initialize blocks and threads
      Output << "CDecl([(\"dim3\", \"blocks\", \"\")]),\n";
      Output << "CDecl([(\"dim3\", \"threads\", \"\")]),\n";
      Output << "CBlock([\"kernel_sizing(blocks, threads)\"]),\n";

      if (!isTopological) {
        requiresWorklist = true;
        Output << "CBlock([\"ctx->in_wl.update_gpu(ctx->shared_wl->num_in_"
                  "items)\"]),\n";
        Output << "CBlock([\"ctx->out_wl.will_write()\"]),\n";
        Output << "CBlock([\"ctx->out_wl.reset()\"]),\n";
      }

      for (auto& it : accumulatorToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        Output << "CDecl([(\"Shared<" << type << ">\", \"" << name
               << "val\", \" = Shared<" << type << ">(1)\")]),\n";
        Output << "CDecl([(\"HGAccumulator<" << type << ">\", \"_" << name
               << "\", \"\")]),\n";
        Output << "CBlock([\"*(" << name << "val.cpu_wr_ptr()) = 0\"]),\n";
        Output << "CBlock([\"_" << name << ".rv = " << name
               << "val.gpu_wr_ptr()\"]),\n";
      }
      for (auto& it : reducemaxToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        Output << "CDecl([(\"Shared<" << type << ">\", \"" << name
               << "val\", \" = Shared<" << type << ">(1)\")]),\n";
        Output << "CDecl([(\"HGReduceMax<" << type << ">\", \"_" << name
               << "\", \"\")]),\n";
        Output << "CBlock([\"*(" << name << "val.cpu_wr_ptr()) = 0\"]),\n";
        Output << "CBlock([\"_" << name << ".rv = " << name
               << "val.gpu_wr_ptr()\"]),\n";
      }
      for (auto& it : reduceminToTypeMap) {
        auto name = it.first;
        auto type = it.second;
        Output << "CDecl([(\"Shared<" << type << ">\", \"" << name
               << "val\", \" = Shared<" << type << ">(1)\")]),\n";
        Output << "CDecl([(\"HGReduceMin<" << type << ">\", \"_" << name
               << "\", \"\")]),\n";
        Output << "CBlock([\"*(" << name << "val.cpu_wr_ptr()) = 0\"]),\n";
        Output << "CBlock([\"_" << name << ".rv = " << name
               << "val.gpu_wr_ptr()\"]),\n";
      }

      // generate call to kernel
      Output << "Invoke(\"" << kernelName << "\", ";
      Output << "(\"ctx->gg\"";
      if (isTopological) {
        Output << ", \"__begin\", \"__end\"";
      }
      auto& karguments = KernelToArgumentsMap[kernelName];
      for (auto& argument : std::get<0>(karguments)) {
        Output << ", \"" << argument << "\"";
      }
      for (auto& argument : std::get<1>(karguments)) {
        Output << ", \"ctx->" << argument << ".data.gpu_wr_ptr()\"";
      }
      if (!isTopological) {
        Output << ", \"ctx->in_wl\"";
        Output << ", \"ctx->out_wl\"";
      }
      for (auto& argument : std::get<2>(karguments)) {
        Output << ", \"*(ctx->" << argument << ".is_updated.gpu_rd_ptr())\"";
      }
      for (auto& it : accumulatorToTypeMap) {
        Output << ", \"_" << it.first << "\"";
      }
      for (auto& it : reducemaxToTypeMap) {
        Output << ", \"_" << it.first << "\"";
      }
      for (auto& it : reduceminToTypeMap) {
        Output << ", \"_" << it.first << "\"";
      }
      Output << ")";
      Output << "),\n";
      Output << "CBlock([\"check_cuda_kernel\"], parse = False),\n";

      for (auto& it : accumulatorToTypeMap) {
        auto name = it.first;
        Output << "CBlock([\"" << name << " = *(" << name
               << "val.cpu_rd_ptr())\"]),\n";
      }
      for (auto& it : reducemaxToTypeMap) {
        auto name = it.first;
        Output << "CBlock([\"" << name << " = *(" << name
               << "val.cpu_rd_ptr())\"]),\n";
      }
      for (auto& it : reduceminToTypeMap) {
        auto name = it.first;
        Output << "CBlock([\"" << name << " = *(" << name
               << "val.cpu_rd_ptr())\"]),\n";
      }

      if (!isTopological) {
        Output << "CBlock([\"ctx->out_wl.update_cpu()\"]),\n";
        Output << "CBlock([\"ctx->shared_wl->num_out_items = "
                  "ctx->out_wl.nitems()\"]),\n";
      }

      Output << "], host = True),\n"; // end Kernel

      if (isTopological) {
        Output << Output_topo_all.str();
        Output << Output_topo_common.str();
        Output << Output_topo_all_body.str();
        Output << Output_topo_common_body.str();
        Output << "], host = True),\n"; // end Kernel
        Output << Output_topo_masters.str();
        Output << Output_topo_common.str();
        Output << Output_topo_masters_body.str();
        Output << Output_topo_common_body.str();
        Output << "], host = True),\n"; // end Kernel
        Output << Output_topo_nodeswithedges.str();
        Output << Output_topo_common.str();
        Output << Output_topo_nodeswithedges_body.str();
        Output << Output_topo_common_body.str();
        Output << "], host = True),\n"; // end Kernel
      }

      Output.close();
      accumulatorToTypeMap.clear();
      reducemaxToTypeMap.clear();
      reduceminToTypeMap.clear();

      std::string fileName =
          rewriter.getSourceMgr().getFilename(cxxrecord->getLocStart()).str();
      std::size_t found = fileName.rfind(".");
      assert(found != std::string::npos);
      FileNamePath = fileName.substr(0, found);
    }
    return true;
  }
};

class FunctionDeclHandler : public MatchFinder::MatchCallback {
private:
  // ASTContext *astContext;
  // Rewriter &rewriter;
  IrGLOrchestratorVisitor orchestratorVisitor;

public:
  FunctionDeclHandler(ASTContext* context, Rewriter& R)
      : // astContext(context),
        // rewriter(R),
        orchestratorVisitor(context, R) {}
  const IrGLOrchestratorVisitor& getOrchestrator() {
    return orchestratorVisitor;
  }
  virtual void run(const MatchFinder::MatchResult& Results) {
    const CXXMethodDecl* decl =
        Results.Nodes.getNodeAs<clang::CXXMethodDecl>("graphAlgorithm");
    if (decl) {
      // llvm::errs() << "\nGraph Operator:\n" <<
      //  rewriter.getRewrittenText(decl->getSourceRange()) << "\n";
      // decl->dump();
      orchestratorVisitor.TraverseDecl(const_cast<CXXMethodDecl*>(decl));
    }
  }
};

class IrGLFunctionsConsumer : public ASTConsumer {
private:
  // Rewriter &rewriter;
  MatchFinder Matchers;
  FunctionDeclHandler functionDeclHandler;

public:
  IrGLFunctionsConsumer(CompilerInstance& CI, Rewriter& R)
      : // rewriter(R),
        functionDeclHandler(&(CI.getASTContext()), R) {
    Matchers.addMatcher(functionDecl(hasName("go"),
                                     // hasOverloadedOperatorName("()"),
                                     hasAnyParameter(hasName("_graph")))
                            .bind("graphAlgorithm"),
                        &functionDeclHandler);
  }

  virtual void HandleTranslationUnit(ASTContext& Context) {
    Matchers.matchAST(Context);
    functionDeclHandler.getOrchestrator().GenerateIrGLASTToFile();
  }
};

class IrGLFunctionsAction : public PluginASTAction {
private:
  Rewriter rewriter;

protected:
  std::unique_ptr<ASTConsumer> CreateASTConsumer(CompilerInstance& CI,
                                                 llvm::StringRef) override {
    rewriter.setSourceMgr(CI.getSourceManager(), CI.getLangOpts());
    return llvm::make_unique<IrGLFunctionsConsumer>(CI, rewriter);
  }

  bool ParseArgs(const CompilerInstance& CI,
                 const std::vector<std::string>& args) override {
    return true;
  }
};

} // namespace

static FrontendPluginRegistry::Add<IrGLFunctionsAction>
    X("irgl", "translate LLVM AST to IrGL AST");
