Merge pull request #7 from dhuth/master

V0.2.1

1 files changed, 0 insertions, 3734 deletions

diff --git a/loop_cuda_rose.cc b/loop_cuda_rose.cc
deleted file mode 100644
index c5633ee..0000000
--- a/loop_cuda_rose.cc
+++ /dev/null
@@ -1,3734 +0,0 @@
-/*****************************************************************************
- Copyright (C) 2009 University of Utah
- All Rights Reserved.
-
- Purpose:
- Cudaize methods
-
- Notes:
-
- History:
- 1/7/10 Created by Gabe Rudy by migrating code from loop.cc
- 31/1/11 Modified by Protonu Basu
-*****************************************************************************/
-#define TRANSFORMATION_FILE_INFO Sg_File_Info::generateDefaultFileInfoForTransformationNode()
-#include <code_gen/CG_stringBuilder.h>
-#include <codegen.h>
-#include <code_gen/CG_utils.h>
-#include <code_gen/CG_outputRepr.h>
-#include "loop_cuda_rose.hh"
-#include "loop.hh"
-#include <math.h>
-//#include <useful.h>
-#include "omegatools.hh"
-#include "ir_cudarose.hh"
-#include "ir_rose.hh"
-#include "ir_rose_utils.hh"
-#include "chill_error.hh"
-#include <vector>
-#include "Outliner.hh"
-//#define DEBUG
-using namespace omega;
-using namespace SageBuilder;
-using namespace SageInterface;
-//using namespace Outliner;
-//using namespace ASTtools;
-char *k_cuda_texture_memory; //protonu--added to track texture memory type
-//extern char *omega::k_cuda_texture_memory; //protonu--added to track texture memory type
-extern char *omega::k_ocg_comment;
-
-static int cudaDebug;
-class CudaStaticInit {
-public:
-  CudaStaticInit() {
-    cudaDebug = 0; //Change this to 1 for debug
-  }
-};
-static CudaStaticInit junkInitInstance__;
-
-std::string& upcase(std::string& s) {
-  for (int i = 0; i < s.size(); i++)
-    s[i] = toupper(s[i]);
-  return s;
-}
-
-void printVs(const std::vector<std::string>& curOrder) {
-  if (!cudaDebug) return;
-  for (int i = 0; i < curOrder.size(); i++) {
-    if (i > 0)
-      printf(",");
-    printf("%s", curOrder[i].c_str());
-  }
-  printf("\n");
-}
-
-void printVS(const std::vector<std::string>& curOrder) {
-  if(!cudaDebug) return;
-  for (int i = 0; i < curOrder.size(); i++) {
-    if (i > 0)
-      printf(",");
-    printf("%s", curOrder[i].c_str());
-  }
-  printf("\n");
-}
-
-LoopCuda::~LoopCuda() {
-  const int m = stmt.size();
-  for (int i = 0; i < m; i++)
-    stmt[i].code->clear();
-}
-
-bool LoopCuda::symbolExists(std::string s) {
-  
-  if (body_symtab->find_variable(SgName(s.c_str()))
-      || parameter_symtab->find_variable(SgName(s.c_str())))
-    return true;
-  if (globals->lookup_variable_symbol(SgName(s.c_str())))
-    return true;
-  for (int i = 0; i < idxNames.size(); i++)
-    for (int j = 0; j < idxNames[i].size(); j++)
-      if (strcmp(idxNames[i][j].c_str(), s.c_str()) == 0)
-        return true;
-  return false;
-}
-
-void LoopCuda::addSync(int stmt_num, std::string idxName) {
-  //we store these and code-gen inserts sync to omega comments where stmt
-  //in loop that has idxName being generated
-  syncs.push_back(make_pair(stmt_num, idxName));
-}
-
-void LoopCuda::renameIndex(int stmt_num, std::string idx, std::string newName) {
-  int level = findCurLevel(stmt_num, idx);
-  if (idxNames.size() <= stmt_num || idxNames[stmt_num].size() < level)
-    throw std::runtime_error("Invalid statment number of index");
-  idxNames[stmt_num][level - 1] = newName.c_str();
-}
-
-enum Type {
-  Int
-};
-
-SgNode* wrapInIfFromMinBound(SgNode* then_part, SgForStatement* loop,
-                             SgScopeStatement* symtab, SgVariableSymbol* bound_sym) {
-  // CG_roseBuilder *ocg = new CG_roseBuilder(
-  
-  SgBinaryOp* test_expr = isSgBinaryOp(loop->get_test_expr());
-  SgExpression* upperBound;
-  SgExpression* conditional;
-  upperBound = test_expr->get_rhs_operand();
-  CG_outputRepr *ifstmt;
-  
-  SgCallExpression *call;
-  if (call = isSgCallExpression(upperBound))
-    if (isSgVarRefExp(call->get_function())->get_symbol()->get_name().getString()
-        == "__rose_lt") {
-      SgExprListExp* arg_list = call->get_args();
-      SgExpression *if_bound = *(arg_list->get_expressions().begin());
-      /*This relies on the minimum expression being the rhs operand of
-       * the min instruction.
-       */
-      SgIfStmt *ifstmt = buildIfStmt(
-        buildLessOrEqualOp(buildVarRefExp(bound_sym), if_bound),
-        isSgStatement(then_part), NULL);
-      return isSgNode(ifstmt);
-      
-    }
-  
-/*  if (isSgConditionalExp(upperBound)) {
-    conditional = isSgConditionalExp(upperBound)->get_conditional_exp();
-    
-    if (isSgBinaryOp(conditional)) {
-    SgBinaryOp* binop = isSgBinaryOp(conditional);
-    
-    if (isSgLessThanOp(binop) || isSgLessOrEqualOp(binop)) {
-    SgIfStmt *ifstmt = buildIfStmt(
-    buildLessOrEqualOp(buildVarRefExp(bound_sym),
-    test_expr), isSgStatement(then_part), NULL);
-    return isSgNode(ifstmt);
-    }
-    
-    }
-    
-    }
-*/
-  return then_part;
-}
-
-/**
- * This would be better if it was done by a CHiLL xformation instead of at codegen
- *
- * state:
- * for(...)
- *   for(...)
- *     cur_body
- *   stmt1
- *
- * stm1 is in-between two loops that are going to be reduced. The
- * solution is to put stmt1 at the end of cur_body but conditionally run
- * in on the last step of the for loop.
- *
- * A CHiLL command that would work better:
- *
- * for(...)
- *   stmt0
- *   for(for i=0; i<n; i++)
- *     cur_body
- *   stmt1
- * =>
- * for(...)
- *   for(for i=0; i<n; i++)
- *     if(i==0) stmt0
- *     cur_body
- *     if(i==n-1) stmt1
- */
-
-std::vector<SgForStatement*> findCommentedFors(const char* index, SgNode* tnl) {
-  std::vector<SgForStatement *> result;
-  bool next_loop_ok = false;
-  
-  if (isSgBasicBlock(tnl)) {
-    
-    SgStatementPtrList& list = isSgBasicBlock(tnl)->get_statements();
-    
-    for (SgStatementPtrList::iterator it = list.begin(); it != list.end();
-         it++) {
-      std::vector<SgForStatement*> t = findCommentedFors(index,
-                                                         isSgNode(*it));
-      std::copy(t.begin(), t.end(), back_inserter(result));
-    }
-  } else if (isSgForStatement(tnl)) {
-    
-    AstTextAttribute* att =
-      (AstTextAttribute*) (isSgNode(tnl)->getAttribute(
-                             "omega_comment"));
-    std::string comment = att->toString();
-    
-    if (comment.find("~cuda~") != std::string::npos
-        && comment.find("preferredIdx: ") != std::string::npos) {
-      std::string idx = comment.substr(
-        comment.find("preferredIdx: ") + 14, std::string::npos);
-      if (idx.find(" ") != std::string::npos)
-        idx = idx.substr(0, idx.find(" "));
-      if (strcmp(idx.c_str(), index) == 0)
-        next_loop_ok = true;
-    }
-    
-    if (next_loop_ok) {
-      //printf("found loop %s\n", static_cast<tree_for *>(tn)->index()->name());
-      result.push_back(isSgForStatement(tnl));
-    } else {
-      //printf("looking down for loop %s\n", static_cast<tree_for *>(tn)->index()->name());
-      std::vector<SgForStatement*> t = findCommentedFors(index,
-                                                         isSgForStatement(tnl)->get_loop_body());
-      std::copy(t.begin(), t.end(), back_inserter(result));
-    }
-    next_loop_ok = false;
-  } else if (isSgIfStmt(tnl)) {
-    //printf("looking down if\n");
-    SgIfStmt *tni = isSgIfStmt(tnl);
-    std::vector<SgForStatement*> t = findCommentedFors(index,
-                                                       tni->get_true_body());
-    std::copy(t.begin(), t.end(), back_inserter(result));
-  }
-  
-  return result;
-}
-
-SgNode* forReduce(SgForStatement* loop, SgVariableSymbol* reduceIndex,
-                  SgScopeStatement* body_syms) {
-  //We did the replacements all at once with recursiveFindPreferedIdxs
-  //replacements r;
-  //r.oldsyms.append(loop->index());
-  //r.newsyms.append(reduceIndex);
-  //tree_for* new_loop = (tree_for*)loop->clone_helper(&r, true);
-  SgForStatement* new_loop = loop;
-  
-  //return body one loops in
-  SgNode* tnl = loop_body_at_level(new_loop, 1);
-  //wrap in conditional if necessary
-  tnl = wrapInIfFromMinBound(tnl, new_loop, body_syms, reduceIndex);
-  return tnl;
-}
-
-void recursiveFindRefs(SgNode* code, std::set<const SgVariableSymbol *>& syms,
-                       SgFunctionDefinition* def) {
-  
-  SgStatement* s = isSgStatement(code);
-  // L = {symbols defined within 's'}, local variables declared within 's'
-  ASTtools::VarSymSet_t L;
-  ASTtools::collectDefdVarSyms(s, L);
-  //dump (L, "L = ");
-  
-  // U = {symbols used within 's'}
-  ASTtools::VarSymSet_t U;
-  ASTtools::collectRefdVarSyms(s, U);
-  //dump (U, "U = ");
-  
-  // U - L = {symbols used within 's' but not defined in 's'}
-  // variable references to non-local-declared variables
-  ASTtools::VarSymSet_t diff_U_L;
-  set_difference(U.begin(), U.end(), L.begin(), L.end(),
-                 inserter(diff_U_L, diff_U_L.begin()));
-  //dump (diff_U_L, "U - L = ");
-  
-  // Q = {symbols defined within the function surrounding 's' that are
-  // visible at 's'}, including function parameters
-  ASTtools::VarSymSet_t Q;
-  ASTtools::collectLocalVisibleVarSyms(def->get_declaration(), s, Q);
-//    dump (Q, "Q = ");
-  
-  // (U - L) \cap Q = {variables that need to be passed as parameters
-  // to the outlined function}
-  // a sub set of variables that are not globally visible (no need to pass at all)
-  // It excludes the variables with a scope between global and the enclosing function
-  set_intersection(diff_U_L.begin(), diff_U_L.end(), Q.begin(), Q.end(),
-                   inserter(syms, syms.begin()));
-  
-  /* std::vector<SgVariableSymbol *> scalars;
-  //SgNode  *tnl = static_cast<const omega::CG_roseRepr *>(repr)->GetCode();
-  SgStatement* stmt;
-  SgExpression* exp;
-  if (tnl != NULL) {
-  if(stmt = isSgStatement(tnl)){
-  if(isSgBasicBlock(stmt)){
-  SgStatementPtrList& stmts = isSgBasicBlock(stmt)->get_statements();
-  for(int i =0; i < stmts.size(); i++){
-  //omega::CG_roseRepr *r = new omega::CG_roseRepr(isSgNode(stmts[i]));
-  std::vector<SgVariableSymbol *> a = recursiveFindRefs(isSgNode(stmts[i]));
-  //delete r;
-  std::copy(a.begin(), a.end(), back_inserter(scalars));
-  }
-  
-  }
-  else if(isSgForStatement(stmt)){
-  
-  SgForStatement *tnf =  isSgForStatement(stmt);
-  //omega::CG_roseRepr *r = new omega::CG_roseRepr(isSgStatement(tnf->get_loop_body()));
-  std::vector<SgVariableSymbol *> a = recursiveFindRefs(isSgNode(tnf->get_loop_body()));
-  //delete r;
-  std::copy(a.begin(), a.end(), back_inserter(scalars));
-  }
-  else if(isSgFortranDo(stmt)){
-  SgFortranDo *tfortran =  isSgFortranDo(stmt);
-  omega::CG_roseRepr *r = new omega::CG_roseRepr(isSgStatement(tfortran->get_body()));
-  std::vector<SgVariableSymbol *> a = recursiveFindRefs(r);
-  delete r;
-  std::copy(a.begin(), a.end(), back_inserter(scalars));
-  }
-  
-  else if(isSgIfStmt(stmt) ){
-  SgIfStmt* tni = isSgIfStmt(stmt);
-  //omega::CG_roseRepr *r = new omega::CG_roseRepr(isSgNode(tni->get_conditional()));
-  std::vector<SgVariableSymbol *> a = recursiveFindRefs(isSgNode(tni->get_conditional()));
-  //delete r;
-  std::copy(a.begin(), a.end(), back_inserter(scalars));
-  //r = new omega::CG_roseRepr(isSgNode(tni->get_true_body()));
-  a = recursiveFindRefs(isSgNode(tni->get_true_body()));
-  //delete r;
-  std::copy(a.begin(), a.end(), back_inserter(scalars));
-  //r = new omega::CG_roseRepr(isSgNode(tni->get_false_body()));
-  a = recursiveFindRefs(isSgNode(tni->get_false_body()));
-  //delete r;
-  std::copy(a.begin(), a.end(), back_inserter(scalars));
-  }
-  else if(isSgExprStatement(stmt)) {
-  //omega::CG_roseRepr *r = new omega::CG_roseRepr(isSgExpression(isSgExprStatement(stmt)->get_expression()));
-  std::vector<SgVariableSymbol *> a = recursiveFindRefs(isSgNode(isSgExprStatement(stmt)->get_expression()));
-  //delete r;
-  std::copy(a.begin(), a.end(), back_inserter(scalars));
-  
-  }
-  }
-  }
-  else{
-  SgExpression* op = isSgExpression(tnl);
-  if(isSgVarRefExp(op)){
-  
-  scalars.push_back(isSgVarRefExp(op)->get_symbol());
-  
-  }
-  else if( isSgAssignOp(op)){
-  //omega::CG_roseRepr *r1 = new omega::CG_roseRepr(isSgAssignOp(op)->get_lhs_operand());
-  std::vector<SgVariableSymbol *> a1 = recursiveFindRefs(isSgNode(isSgAssignOp(op)->get_lhs_operand()));
-  //delete r1;
-  std::copy(a1.begin(), a1.end(), back_inserter(scalars));
-  //omega::CG_roseRepr *r2 = new omega::CG_roseRepr(isSgAssignOp(op)->get_rhs_operand());
-  std::vector<SgVariableSymbol *> a2 = recursiveFindRefs(isSgNode(isSgAssignOp(op)->get_rhs_operand()));
-  //delete r2;
-  std::copy(a2.begin(), a2.end(), back_inserter(scalars));
-  
-  }
-  else if(isSgBinaryOp(op)){
-  // omega::CG_roseRepr *r1 = new omega::CG_roseRepr(isSgBinaryOp(op)->get_lhs_operand());
-  std::vector<SgVariableSymbol *> a1 = recursiveFindRefs(isSgNode(isSgBinaryOp(op)->get_lhs_operand()));
-  //delete r1;
-  std::copy(a1.begin(), a1.end(), back_inserter(scalars));
-  //omega::CG_roseRepr *r2 = new omega::CG_roseRepr(isSgBinaryOp(op)->get_rhs_operand());
-  std::vector<SgVariableSymbol *> a2 = recursiveFindRefs((isSgBinaryOp(op)->get_rhs_operand()));
-  //delete r2;
-  std::copy(a2.begin(), a2.end(), back_inserter(scalars));
-  }
-  else if(isSgUnaryOp(op)){
-  //omega::CG_roseRepr *r1 = new omega::CG_roseRepr(isSgUnaryOp(op)->get_operand());
-  std::vector<SgVariableSymbol *> a1 = recursiveFindRefs(isSgNode(isSgUnaryOp(op)->get_operand()));
-  //delete r1;
-  std::copy(a1.begin(), a1.end(), back_inserter(scalars));
-  }
-  
-  }
-  return scalars;
-  
-  
-  */
-  
-}
-
-SgNode* recursiveFindReplacePreferedIdxs(SgNode* code, SgSymbolTable* body_syms,
-                                         SgSymbolTable* param_syms, SgScopeStatement* body,
-                                         std::map<std::string, SgVariableSymbol*>& loop_idxs,
-                                         SgGlobal* globalscope, bool sync = false) {
-  //tree_node_list* tnl = new tree_node_list;
-  //tree_node_list_iter tnli(code);
-  SgVariableSymbol* idxSym = 0;
-  std::vector<SgStatement*> r1;
-  std::vector<SgNode*> r2;
-  SgNode* tnli;
-  SgNode* tnli1;
-  SgNode* tnli2;
-  SgBasicBlock * clone;
-  
-  if (isSgForStatement(code)) {
-    AstTextAttribute* att =
-      (AstTextAttribute*) (isSgNode(code)->getAttribute(
-                             "omega_comment"));
-    
-    std::string comment;
-    if (att != NULL)
-      comment = att->toString();
-    
-    if (comment.find("~cuda~") != std::string::npos
-        && comment.find("preferredIdx: ") != std::string::npos) {
-      std::string idx = comment.substr(
-        comment.find("preferredIdx: ") + 14, std::string::npos);
-      if (idx.find(" ") != std::string::npos)
-        idx = idx.substr(0, idx.find(" "));
-      if (loop_idxs.find(idx) != loop_idxs.end())
-        idxSym = loop_idxs.find(idx)->second;
-      //Get the proc variable sybol for this preferred index
-      if (idxSym == 0) {
-        idxSym = body_syms->find_variable(idx.c_str());
-        if (!idxSym)
-          idxSym = param_syms->find_variable(idx.c_str());
-        //printf("idx not found: lookup %p\n", idxSym);
-        if (!idxSym) {
-          SgVariableDeclaration* defn = buildVariableDeclaration(
-            SgName((char*) idx.c_str()), buildIntType());
-          //idxSym = new var_sym(type_s32, (char*)idx.c_str());
-          SgInitializedNamePtrList& variables = defn->get_variables();
-          SgInitializedNamePtrList::const_iterator i =
-            variables.begin();
-          SgInitializedName* initializedName = *i;
-          SgVariableSymbol* vs = new SgVariableSymbol(
-            initializedName);
-          prependStatement(defn, body);
-          vs->set_parent(body_syms);
-          body_syms->insert(SgName((char*) idx.c_str()), vs);
-          idxSym = vs;
-          //printf("idx created and inserted\n");
-        }
-        //Now insert into our map for future
-	if (cudaDebug)
-	    std::cout << idx << "\n\n";
-        loop_idxs.insert(make_pair(idx, idxSym));
-      }
-      //See if we have a sync as well
-      if (comment.find("sync") != std::string::npos) {
-        //printf("Inserting sync after current block\n");
-        sync = true;
-      }
-      
-    }
-    if (idxSym) {
-      SgForInitStatement* list =
-        isSgForStatement(code)->get_for_init_stmt();
-      SgStatementPtrList& initStatements = list->get_init_stmt();
-      SgStatementPtrList::const_iterator j = initStatements.begin();
-      const SgVariableSymbol* index;
-      
-      if (SgExprStatement *expr = isSgExprStatement(*j))
-        if (SgAssignOp* op = isSgAssignOp(expr->get_expression()))
-          if (SgVarRefExp* var_ref = isSgVarRefExp(
-                op->get_lhs_operand()))
-            index = var_ref->get_symbol();
-      
-      std::vector<SgVarRefExp *> array = substitute(code, index, NULL,
-                                                    isSgNode(body_syms));
-      
-      for (int j = 0; j < array.size(); j++)
-        array[j]->set_symbol(idxSym);
-    }
-    
-    SgStatement* body_ = isSgStatement(
-      recursiveFindReplacePreferedIdxs(
-        isSgNode((isSgForStatement(code)->get_loop_body())),
-        body_syms, param_syms, body, loop_idxs, globalscope));
-    
-    omega::CG_roseRepr * tnl = new omega::CG_roseRepr(code);
-    omega::CG_outputRepr* block = tnl->clone();
-    tnli = static_cast<const omega::CG_roseRepr *>(block)->GetCode();
-    
-    isSgForStatement(tnli)->set_loop_body(body_);
-    body_->set_parent(tnli);
-    
-    if (idxSym) {
-      SgForInitStatement* list =
-        isSgForStatement(tnli)->get_for_init_stmt();
-      SgStatementPtrList& initStatements = list->get_init_stmt();
-      SgStatementPtrList::const_iterator j = initStatements.begin();
-      const SgVariableSymbol* index;
-      
-      if (SgExprStatement *expr = isSgExprStatement(*j))
-        if (SgAssignOp* op = isSgAssignOp(expr->get_expression()))
-          if (SgVarRefExp* var_ref = isSgVarRefExp(
-                op->get_lhs_operand()))
-            index = var_ref->get_symbol();
-      
-      std::vector<SgVarRefExp *> array = substitute(tnli, index, NULL,
-                                                    isSgNode(body_syms));
-      
-      for (int j = 0; j < array.size(); j++)
-        array[j]->set_symbol(idxSym);
-    }
-    //  std::cout << isSgNode(body_)->unparseToString() << "\n\n";
-    if (att != NULL)
-      tnli->setAttribute("omega_comment", att);
-    
-    if (sync) {
-      SgName name_syncthreads("__syncthreads");
-      SgFunctionSymbol * syncthreads_symbol =
-        globalscope->lookup_function_symbol(name_syncthreads);
-      
-      // Create a call to __syncthreads():
-      SgFunctionCallExp * syncthreads_call = buildFunctionCallExp(
-        syncthreads_symbol, buildExprListExp());
-      
-      SgExprStatement* stmt = buildExprStatement(syncthreads_call);
-      
-      /*    if (SgBasicBlock* bb = isSgBasicBlock(
-            isSgForStatement(code)->get_loop_body()))
-            appendStatement(isSgStatement(stmt), bb);
-            
-            else if (SgStatement* ss = isSgStatement(
-            isSgForStatement(code)->get_loop_body())) {
-            SgBasicBlock* bb2 = buildBasicBlock();
-            
-            isSgNode(ss)->set_parent(bb2);
-            appendStatement(ss, bb2);
-            
-            appendStatement(isSgStatement(stmt), bb2);
-            isSgNode(stmt)->set_parent(bb2);
-            isSgForStatement(code)->set_loop_body(bb2);
-            isSgNode(bb2)->set_parent(code);
-            }
-      */
-      
-      SgBasicBlock* bb2 = buildBasicBlock();
-      
-      bb2->append_statement(isSgStatement(tnli));
-      bb2->append_statement(stmt);
-      /* SgNode* parent = code->get_parent();
-         if(!isSgStatement(parent))
-         throw loop_error("Parent not a statement");
-         
-         if(isSgForStatement(parent)){
-         if(SgStatement *ss = isSgForStatement(isSgForStatement(parent)->get_loop_body())){
-         omega::CG_roseRepr * tnl = new omega::CG_roseRepr(ss);
-         omega::CG_outputRepr* block= tnl->clone();
-         
-         SgNode *new_ss = static_cast<const omega::CG_roseRepr *>(block)->GetCode();
-         SgBasicBlock* bb2 = buildBasicBlock();
-         
-         isSgNode(new_ss)->set_parent(bb2);
-         appendStatement(isSgStatement(new_ss), bb2);
-         appendStatement(isSgStatement(stmt), bb2);
-         isSgNode(stmt)->set_parent(bb2);
-         
-         isSgStatement(parent)->replace_statement_from_basicBlock(ss, isSgStatement(bb2));
-         
-         }else if(isSgBasicBlock(isSgForStatement(parent)->get_loop_body()))
-         isSgStatement(isSgForStatement(parent)->get_loop_body())->insert_statement(isSgStatement(code), stmt, false);
-         else
-         throw loop_error("parent statement type undefined!!");
-         
-         }
-         else if(isSgBasicBlock(parent))
-         isSgStatement(parent)->insert_statement(isSgStatement(code), stmt, false);
-         else
-         throw loop_error("parent statement type undefined!!");
-         
-         //tnl->print();
-         *
-         *
-         */
-      sync = true;
-      return isSgNode(bb2);
-      
-    } else
-      return tnli;
-  } else if (isSgIfStmt(code)) {
-    SgStatement* body_ = isSgStatement(
-      recursiveFindReplacePreferedIdxs(
-        isSgNode((isSgIfStmt(code)->get_true_body())),
-        body_syms, param_syms, body, loop_idxs, globalscope));
-    
-    omega::CG_roseRepr * tnl = new omega::CG_roseRepr(code);
-    omega::CG_outputRepr* block = tnl->clone();
-    tnli = static_cast<const omega::CG_roseRepr *>(block)->GetCode();
-    
-    isSgIfStmt(tnli)->set_true_body(body_);
-    
-    if ((isSgIfStmt(code)->get_false_body()))
-      isSgIfStmt(tnli)->set_false_body(
-        isSgStatement(
-          recursiveFindReplacePreferedIdxs(
-            isSgNode(
-              (isSgIfStmt(code)->get_false_body())),
-            body_syms, param_syms, body, loop_idxs,
-            globalscope)));
-    
-    return tnli;
-  } else if (isSgStatement(code) && !isSgBasicBlock(code)) {
-    omega::CG_roseRepr * tnl = new omega::CG_roseRepr(code);
-    omega::CG_outputRepr* block = tnl->clone();
-    tnli = static_cast<const omega::CG_roseRepr *>(block)->GetCode();
-    
-    return tnli;
-    
-  } else if (isSgBasicBlock(code)) {
-    SgStatementPtrList& tnl = isSgBasicBlock(code)->get_statements();
-    
-    SgStatementPtrList::iterator temp;
-    clone = buildBasicBlock();
-    bool sync_found = false;
-    for (SgStatementPtrList::const_iterator it = tnl.begin();
-         it != tnl.end(); it++) {
-      
-      if (isSgForStatement(*it)) {
-        AstTextAttribute* att =
-          (AstTextAttribute*) (isSgNode(*it)->getAttribute(
-                                 "omega_comment"));
-        
-        std::string comment;
-        if (att != NULL)
-          comment = att->toString();
-        
-        if (comment.find("~cuda~") != std::string::npos
-            && comment.find("preferredIdx: ")
-            != std::string::npos) {
-          std::string idx = comment.substr(
-            comment.find("preferredIdx: ") + 14,
-            std::string::npos);
-          if (idx.find(" ") != std::string::npos)
-            idx = idx.substr(0, idx.find(" "));
-          //printf("sym_tab preferred index: %s\n", idx.c_str());
-          if (loop_idxs.find(idx) != loop_idxs.end())
-            idxSym = loop_idxs.find(idx)->second;
-          //Get the proc variable sybol for this preferred index
-          if (idxSym == 0) {
-            idxSym = body_syms->find_variable(idx.c_str());
-            if (!idxSym)
-              idxSym = param_syms->find_variable(idx.c_str());
-            //printf("idx not found: lookup %p\n", idxSym);
-            if (!idxSym) {
-              SgVariableDeclaration* defn =
-                buildVariableDeclaration(
-                  SgName((char*) idx.c_str()),
-                  buildIntType());
-              //idxSym = new var_sym(type_s32, (char*)idx.c_str());
-              SgInitializedNamePtrList& variables =
-                defn->get_variables();
-              SgInitializedNamePtrList::const_iterator i =
-                variables.begin();
-              SgInitializedName* initializedName = *i;
-              SgVariableSymbol* vs = new SgVariableSymbol(
-                initializedName);
-              prependStatement(defn, body);
-              vs->set_parent(body_syms);
-              body_syms->insert(SgName((char*) idx.c_str()), vs);
-              //printf("idx created and inserted\n");
-              idxSym = vs;
-            }
-            //Now insert into our map for future
-	    if (cudaDebug)
-		std::cout << idx << "\n\n";
-            loop_idxs.insert(make_pair(idx, idxSym));
-            
-          }
-          //See if we have a sync as well
-          if (comment.find("sync") != std::string::npos) {
-            //printf("Inserting sync after current block\n");
-            sync = true;
-          }
-          
-        }
-        if (idxSym) {
-          SgForInitStatement* list =
-            isSgForStatement(*it)->get_for_init_stmt();
-          SgStatementPtrList& initStatements = list->get_init_stmt();
-          SgStatementPtrList::const_iterator j =
-            initStatements.begin();
-          const SgVariableSymbol* index;
-          
-          if (SgExprStatement *expr = isSgExprStatement(*j))
-            if (SgAssignOp* op = isSgAssignOp(
-                  expr->get_expression()))
-              if (SgVarRefExp* var_ref = isSgVarRefExp(
-                    op->get_lhs_operand()))
-                index = var_ref->get_symbol();
-          
-          std::vector<SgVarRefExp *> array = substitute(*it, index,
-                                                        NULL, isSgNode(body_syms));
-          
-          for (int j = 0; j < array.size(); j++)
-            array[j]->set_symbol(idxSym);
-          
-        }
-        
-        SgStatement* body_ =
-          isSgStatement(
-            recursiveFindReplacePreferedIdxs(
-              isSgNode(
-                (isSgForStatement(*it)->get_loop_body())),
-              body_syms, param_syms, body, loop_idxs,
-              globalscope));
-        
-        omega::CG_roseRepr * tnl = new omega::CG_roseRepr(*it);
-        omega::CG_outputRepr* block = tnl->clone();
-        tnli =
-          static_cast<const omega::CG_roseRepr *>(block)->GetCode();
-        
-        isSgForStatement(tnli)->set_loop_body(body_);
-        body_->set_parent(tnli);
-        if (idxSym) {
-          SgForInitStatement* list =
-            isSgForStatement(tnli)->get_for_init_stmt();
-          SgStatementPtrList& initStatements = list->get_init_stmt();
-          SgStatementPtrList::const_iterator j =
-            initStatements.begin();
-          const SgVariableSymbol* index;
-          
-          if (SgExprStatement *expr = isSgExprStatement(*j))
-            if (SgAssignOp* op = isSgAssignOp(
-                  expr->get_expression()))
-              if (SgVarRefExp* var_ref = isSgVarRefExp(
-                    op->get_lhs_operand()))
-                index = var_ref->get_symbol();
-          
-          std::vector<SgVarRefExp *> array = substitute(tnli, index,
-                                                        NULL, isSgNode(body_syms));
-          
-          for (int j = 0; j < array.size(); j++)
-            array[j]->set_symbol(idxSym);
-        }
-        idxSym = 0;
-        //  std::cout << isSgNode(body_)->unparseToString() << "\n\n";
-        if (att != NULL)
-          tnli->setAttribute("omega_comment", att);
-        clone->append_statement(isSgStatement(tnli));
-        if (sync) {
-          SgName name_syncthreads("__syncthreads");
-          SgFunctionSymbol * syncthreads_symbol =
-            globalscope->lookup_function_symbol(
-              name_syncthreads);
-          
-          // Create a call to __syncthreads():
-          SgFunctionCallExp * syncthreads_call = buildFunctionCallExp(
-            syncthreads_symbol, buildExprListExp());
-          
-          SgExprStatement* stmt = buildExprStatement(
-            syncthreads_call);
-          
-          /*    if (SgBasicBlock* bb = isSgBasicBlock(
-                isSgForStatement(code)->get_loop_body()))
-                appendStatement(isSgStatement(stmt), bb);
-                
-                else if (SgStatement* ss = isSgStatement(
-                isSgForStatement(code)->get_loop_body())) {
-                SgBasicBlock* bb2 = buildBasicBlock();
-                
-                isSgNode(ss)->set_parent(bb2);
-                appendStatement(ss, bb2);
-                
-                appendStatement(isSgStatement(stmt), bb2);
-                isSgNode(stmt)->set_parent(bb2);
-                isSgForStatement(code)->set_loop_body(bb2);
-                isSgNode(bb2)->set_parent(code);
-                }
-          */
-          
-          //SgBasicBlock* bb2 = buildBasicBlock();
-          clone->append_statement(stmt);
-          /* SgNode* parent = code->get_parent();
-             if(!isSgStatement(parent))
-             throw loop_error("Parent not a statement");
-             
-             if(isSgForStatement(parent)){
-             if(SgStatement *ss = isSgForStatement(isSgForStatement(parent)->get_loop_body())){
-             omega::CG_roseRepr * tnl = new omega::CG_roseRepr(ss);
-             omega::CG_outputRepr* block= tnl->clone();
-             
-             SgNode *new_ss = static_cast<const omega::CG_roseRepr *>(block)->GetCode();
-             SgBasicBlock* bb2 = buildBasicBlock();
-             
-             isSgNode(new_ss)->set_parent(bb2);
-             appendStatement(isSgStatement(new_ss), bb2);
-             appendStatement(isSgStatement(stmt), bb2);
-             isSgNode(stmt)->set_parent(bb2);
-             
-             isSgStatement(parent)->replace_statement_from_basicBlock(ss, isSgStatement(bb2));
-             
-             }else if(isSgBasicBlock(isSgForStatement(parent)->get_loop_body()))
-             isSgStatement(isSgForStatement(parent)->get_loop_body())->insert_statement(isSgStatement(code), stmt, false);
-             else
-             throw loop_error("parent statement type undefined!!");
-             
-             }
-             else if(isSgBasicBlock(parent))
-             isSgStatement(parent)->insert_statement(isSgStatement(code), stmt, false);
-             else
-             throw loop_error("parent statement type undefined!!");
-             
-             //tnl->print();
-             *
-             *
-             */
-          sync = true;
-          //    return isSgNode(bb2);
-          
-        }
-        
-        //  return tnli;
-      } else if (isSgIfStmt(*it)) {
-        SgStatement* body_ = isSgStatement(
-          recursiveFindReplacePreferedIdxs(
-            isSgNode((isSgIfStmt(*it)->get_true_body())),
-            body_syms, param_syms, body, loop_idxs,
-            globalscope));
-        
-        omega::CG_roseRepr * tnl = new omega::CG_roseRepr(*it);
-        omega::CG_outputRepr* block = tnl->clone();
-        tnli1 =
-          static_cast<const omega::CG_roseRepr *>(block)->GetCode();
-        
-        isSgIfStmt(tnli1)->set_true_body(body_);
-        
-        if ((isSgIfStmt(*it)->get_false_body()))
-          isSgIfStmt(tnli1)->set_false_body(
-            isSgStatement(
-              recursiveFindReplacePreferedIdxs(
-                isSgNode(
-                  (isSgIfStmt(*it)->get_false_body())),
-                body_syms, param_syms, body,
-                loop_idxs, globalscope)));
-        
-        clone->append_statement(isSgStatement(tnli1));
-        //return tnli;
-      } else if (isSgStatement(*it)) {
-        omega::CG_roseRepr * tnl = new omega::CG_roseRepr(*it);
-        omega::CG_outputRepr* block = tnl->clone();
-        tnli2 =
-          static_cast<const omega::CG_roseRepr *>(block)->GetCode();
-        
-        clone->append_statement(isSgStatement(tnli2));
-        //return tnli;
-        
-      }
-    }
-    
-    return isSgNode(clone);
-    
-  }
-  
-  /*    if (!isSgBasicBlock(
-        recursiveFindReplacePreferedIdxs(isSgNode(*it), body_syms,
-        param_syms, body, loop_idxs, globalscope))) {
-        SgStatement *to_push = isSgStatement(
-        recursiveFindReplacePreferedIdxs(isSgNode(*it),
-        body_syms, param_syms, body, loop_idxs,
-        globalscope, sync));
-        clone->append_statement(to_push);
-        
-        if ((sync_found) && isSgForStatement(to_push)) {
-        SgName name_syncthreads("__syncthreads");
-        SgFunctionSymbol * syncthreads_symbol =
-        globalscope->lookup_function_symbol(
-        name_syncthreads);
-        
-        // Create a call to __syncthreads():
-        SgFunctionCallExp * syncthreads_call = buildFunctionCallExp(
-        syncthreads_symbol, buildExprListExp());
-        
-        SgExprStatement* stmt = buildExprStatement(
-        syncthreads_call);
-        
-        clone->append_statement(isSgStatement(stmt));
-        }
-        //  std::cout<<isSgNode(*it)->unparseToString()<<"\n\n";
-        } else {
-        
-        SgStatementPtrList& tnl2 = isSgBasicBlock(
-        recursiveFindReplacePreferedIdxs(isSgNode(*it),
-        body_syms, param_syms, body, loop_idxs,
-        globalscope))->get_statements();
-        for (SgStatementPtrList::const_iterator it2 = tnl2.begin();
-        it2 != tnl2.end(); it2++) {
-        clone->append_statement(*it2);
-        
-        sync_found = true;
-        //  std::cout<<isSgNode(*it2)->unparseToString()<<"\n\n";
-        }
-        }
-        
-        }
-        return isSgNode(clone);
-        }
-  */
-//  return tnl;
-}
-
-// loop_vars -> array references
-// loop_idxs -> <idx_name,idx_sym> map for when we encounter a loop with a different preferredIndex
-// dim_vars -> out param, fills with <old,new> var_sym pair for 2D array dimentions (messy stuff)
-SgNode* swapVarReferences(SgNode* code,
-                          std::set<const SgVariableSymbol *>& syms, SgSymbolTable* param,
-                          SgSymbolTable* body, SgScopeStatement* body_stmt) {
-  //Iterate over every expression, looking up each variable and type
-  //reference used and possibly replacing it or adding it to our symbol
-  //table
-  //
-  //We use the built-in cloning helper methods to seriously help us with this!
-  
-  //Need to do a recursive mark
-  
-  std::set<const SgVariableSymbol *>::iterator myIterator;
-  for (myIterator = syms.begin(); myIterator != syms.end(); myIterator++) {
-    SgName var_name = (*myIterator)->get_name();
-    std::string x = var_name.getString();
-    
-    if ((param->find_variable(var_name) == NULL)
-        && (body->find_variable(var_name) == NULL)) {
-      SgInitializedName* decl = (*myIterator)->get_declaration();
-      
-      SgVariableSymbol* dvs = new SgVariableSymbol(decl);
-      SgVariableDeclaration* var_decl = buildVariableDeclaration(
-        dvs->get_name(), dvs->get_type());
-      
-      AstTextAttribute* att = (AstTextAttribute*) (isSgNode(
-                                                     decl->get_declaration())->getAttribute("__shared__"));
-      if (isSgNode(decl->get_declaration())->attributeExists(
-            "__shared__"))
-        var_decl->get_declarationModifier().get_storageModifier().setCudaShared();
-      
-      appendStatement(var_decl, body_stmt);
-      
-      dvs->set_parent(body);
-      body->insert(var_name, dvs);
-    }
-    
-    std::vector<SgVarRefExp *> array = substitute(code, *myIterator, NULL,
-                                                  isSgNode(body));
-    
-    SgVariableSymbol* var = (SgVariableSymbol*) (*myIterator);
-    for (int j = 0; j < array.size(); j++)
-      array[j]->set_symbol(var);
-  }
-  
-  return code;
-}
-
-bool LoopCuda::validIndexes(int stmt, const std::vector<std::string>& idxs) {
-  for (int i = 0; i < idxs.size(); i++) {
-    bool found = false;
-    for (int j = 0; j < idxNames[stmt].size(); j++) {
-      if (strcmp(idxNames[stmt][j].c_str(), idxs[i].c_str()) == 0) {
-        found = true;
-      }
-    }
-    if (!found) {
-      return false;
-    }
-  }
-  return true;
-}
-
-bool LoopCuda::cudaize_v2(std::string kernel_name,
-                          std::map<std::string, int> array_dims,
-                          std::vector<std::string> blockIdxs,
-                          std::vector<std::string> threadIdxs) {
-  CG_outputBuilder *ocg = ir->builder();
-  int stmt_num = 0;
-  if (cudaDebug) {
-    printf("cudaize_v2(%s, {", kernel_name.c_str());
-    //for(
-    printf("}, blocks={");
-    printVs(blockIdxs);
-    printf("}, thread={");
-    printVs(threadIdxs);
-    printf("})\n");
-  }
-  
-  this->array_dims = array_dims;
-  if (!validIndexes(stmt_num, blockIdxs)) {
-    throw std::runtime_error("One of the indexes in the block list was not "
-                             "found in the current set of indexes.");
-  }
-  if (!validIndexes(stmt_num, threadIdxs)) {
-    throw std::runtime_error(
-      "One of the indexes in the thread list was not "
-      "found in the current set of indexes.");
-  }
-  if (blockIdxs.size() == 0)
-    throw std::runtime_error("Cudaize: Need at least one block dimention");
-  int block_level = 0;
-  //Now, we will determine the actual size (if possible, otherwise
-  //complain) for the block dimentions and thread dimentions based on our
-  //indexes and the relations for our stmt;
-  for (int i = 0; i < blockIdxs.size(); i++) {
-    int level = findCurLevel(stmt_num, blockIdxs[i]);
-    int ub, lb;
-    CG_outputRepr* ubrepr = extractCudaUB(stmt_num, level, ub, lb);
-    if (lb != 0) {
-      //attempt to "normalize" the loop with an in-place tile and then re-check our bounds
-      if (cudaDebug)
-        printf(
-          "Cudaize: doing tile at level %d to try and normalize lower bounds\n",
-          level);
-      tile(stmt_num, level, 1, level, CountedTile);
-      idxNames[stmt_num].insert(idxNames[stmt_num].begin() + (level), ""); //TODO: possibly handle this for all sibling stmts
-      ubrepr = extractCudaUB(stmt_num, level, ub, lb);
-    }
-    if (lb != 0) {
-      char buf[1024];
-      sprintf(buf,
-              "Cudaize: Loop at level %d does not have 0 as it's lower bound",
-              level);
-      throw std::runtime_error(buf);
-    }
-    if (ub < 0) {
-      char buf[1024];
-      sprintf(buf,
-              "Cudaize: Loop at level %d does not have a hard upper bound",
-              level);
-      //Anand: Commenting out error indication for lack of constant upper bound
-      //throw std::runtime_error(buf);
-    }
-    if (cudaDebug)
-      printf("block idx %s level %d lb: %d ub %d\n", blockIdxs[i].c_str(),
-             level, lb, ub);
-    if (i == 0) {
-      block_level = level;
-      if (ubrepr == NULL) {
-        cu_bx = ub + 1;
-        cu_bx_repr = NULL;
-      } else {
-        cu_bx = 0;
-        cu_bx_repr = ocg->CreatePlus(ubrepr, ocg->CreateInt(1));
-      }
-      idxNames[stmt_num][level - 1] = "bx";
-    } else if (i == 1) {
-      if (ubrepr == NULL) {
-        cu_by = ub + 1;
-        cu_by_repr = NULL;
-      } else {
-        cu_by = 0;
-        cu_by_repr = ocg->CreatePlus(ubrepr, ocg->CreateInt(1));
-      }
-      idxNames[stmt_num][level - 1] = "by";
-    }
-  }
-  if (!cu_by && !cu_by_repr)
-    block_level = 0;
-  int thread_level1 = 0;
-  int thread_level2 = 0;
-  for (int i = 0; i < threadIdxs.size(); i++) {
-    int level = findCurLevel(stmt_num, threadIdxs[i]);
-    int ub, lb;
-    CG_outputRepr* ubrepr = extractCudaUB(stmt_num, level, ub, lb);
-    if (lb != 0) {
-      //attempt to "normalize" the loop with an in-place tile and then re-check our bounds
-      if (cudaDebug)
-        printf(
-          "Cudaize: doing tile at level %d to try and normalize lower bounds\n",
-          level);
-      tile(stmt_num, level, 1, level, CountedTile);
-      idxNames[stmt_num].insert(idxNames[stmt_num].begin() + (level), "");
-      ubrepr = extractCudaUB(stmt_num, level, ub, lb);
-    }
-    if (lb != 0) {
-      char buf[1024];
-      sprintf(buf,
-              "Cudaize: Loop at level %d does not have 0 as it's lower bound",
-              level);
-      throw std::runtime_error(buf);
-    }
-    if (ub < 0) {
-      char buf[1024];
-      sprintf(buf,
-              "Cudaize: Loop at level %d does not have a hard upper bound",
-              level);
-      //Anand: Commenting out error indication for lack of constant upper bound
-      //throw std::runtime_error(buf);
-    }
-    
-    if (cudaDebug)
-      printf("thread idx %s level %d lb: %d ub %d\n",
-             threadIdxs[i].c_str(), level, lb, ub);
-    if (i == 0) {
-      thread_level1 = level;
-      if (ubrepr == NULL) {
-        cu_tx = ub + 1;
-        cu_tx_repr = NULL;
-      } else {
-        cu_tx = 0;
-        cu_tx_repr = ocg->CreatePlus(ubrepr, ocg->CreateInt(1));
-      }
-      idxNames[stmt_num][level - 1] = "tx";
-    } else if (i == 1) {
-      thread_level2 = level;
-      if (ubrepr == NULL) {
-        cu_ty = ub + 1;
-        cu_ty_repr = NULL;
-      } else {
-        cu_ty = 0;
-        cu_ty_repr = ocg->CreatePlus(ubrepr, ocg->CreateInt(1));
-      }
-      idxNames[stmt_num][level - 1] = "ty";
-    } else if (i == 2) {
-      if (ubrepr == NULL) {
-        cu_tz = ub + 1;
-        cu_tz_repr = NULL;
-      } else {
-        cu_tz = 0;
-        cu_tz_repr = ocg->CreatePlus(ubrepr, ocg->CreateInt(1));
-      }
-      idxNames[stmt_num][level - 1] = "tz";
-    }
-  }
-  if (!cu_ty && !cu_ty_repr)
-    thread_level1 = 0;
-  if (!cu_tz && !cu_tz_repr)
-    thread_level2 = 0;
-  
-  //Make changes to nonsplitlevels
-  const int m = stmt.size();
-  for (int i = 0; i < m; i++) {
-    if (block_level) {
-      //stmt[i].nonSplitLevels.append((block_level)*2);
-      stmt_nonSplitLevels[i].push_back((block_level) * 2);
-    }
-    if (thread_level1) {
-      //stmt[i].nonSplitLevels.append((thread_level1)*2);
-      stmt_nonSplitLevels[i].push_back((thread_level1) * 2);
-    }
-    if (thread_level2) {
-      //stmt[i].nonSplitLevels.append((thread_level1)*2);
-      stmt_nonSplitLevels[i].push_back((thread_level1) * 2);
-    }
-  }
-  
-  if (cudaDebug) {
-    printf("Codegen: current names: ");
-    printVS(idxNames[stmt_num]);
-  }
-  //Set codegen flag
-  code_gen_flags |= GenCudaizeV2;
-  
-  //Save array dimention sizes
-  this->array_dims = array_dims;
-  cu_kernel_name = kernel_name.c_str();
-  
-}
-
-/*
- * setupConstantVar
- * handles constant variable declaration
- * and adds a global constant variable
- * parameters:
- *   constant - the constant_memory_mapping object for this loop
- *   arr_def  - the VarDefs object for the mapped variable
- *   globals  - Rose Global variables
- *   i        - an index to keep new variable names unique
- *   symtab   - global symbol table
- */
-static void setupConstantVar(constant_memory_mapping* constant, VarDefs* arr_def, SgGlobal* globals, int i, SgSymbolTable* symtab) {
-  char* buf1 = new char[32];
-  snprintf(buf1, 32, "cs%dRef", i+1);
-  arr_def->secondName = buf1;
-  
-  char buf2[64];
-  snprintf(buf2, 64, "__device__ __constant__ float");
-  
-  SgVariableDeclaration* consvar_decl = buildVariableDeclaration(
-                     SgName(std::string(buf1)), buildArrayType(
-                                 buildOpaqueType(SgName(buf2),globals),
-                                 arr_def->size_expr));
-  SgInitializedNamePtrList& variables = consvar_decl->get_variables();
-  SgInitializedNamePtrList::const_iterator j = variables.begin();
-  SgInitializedName* initializedName = *j;
-  SgVariableSymbol* consvar_sym = new SgVariableSymbol(initializedName);
-  prependStatement(consvar_decl, globals);
-
-  consvar_sym->set_parent(symtab);
-  symtab->insert(SgName(std::string(buf1)), consvar_sym);
-  
-  constant->set_mapped_symbol(arr_def->original_name.c_str(), consvar_sym);
-  constant->set_vardef(arr_def->original_name.c_str(), arr_def);
-}
-
-/*
- * cudaBindConstantVar
- * allocs a variable to constant memory
- *   constant  - the constant mapping object
- *   arr_def   - the VarDefs abject
- *   globals   - global symbol table
- *   stmt_list - the GPU functions' statement list
- */
-static void cudaBindConstantVar(constant_memory_mapping* constant, VarDefs* arr_def, SgGlobal* globals, SgStatementPtrList* stmt_list) {
-  SgName cudaMemcpyToSymbol_name("cudaMemcpyToSymbol");
-  SgFunctionDeclaration* cudaMemcpyToSymbol_decl = buildNondefiningFunctionDeclaration(
-    cudaMemcpyToSymbol_name, buildVoidType(), buildFunctionParameterList(), globals);
-  SgExprListExp* args = buildExprListExp();
-  args->append_expression(buildCastExp(constant->get_mapped_symbol_exp(arr_def->original_name.c_str()),
-                                       buildPointerType(buildVoidType())));
-  args->append_expression(buildVarRefExp(arr_def->in_data));
-  args->append_expression(arr_def->size_expr);
-  stmt_list->push_back(buildExprStatement(
-    buildFunctionCallExp(buildFunctionRefExp(cudaMemcpyToSymbol_decl), args)));
-}
-
-static void consmapArrayRefs(constant_memory_mapping* constant, std::vector<IR_ArrayRef*>* refs, SgGlobal* globals, IR_Code* ir, CG_roseBuilder* ocg) {
-  // if constant mapping is not being used, ignore this function
-  if(constant == NULL) return;
-  for(int i = 0; i < refs->size(); i++) {
-    IR_ArrayRef* aref = (*refs)[i];
-    if(constant->is_array_mapped(aref->name().c_str())) {
-      // get array reference dimensions
-      int dims = aref->symbol()->n_dim();
-      if(dims > 2) {
-        printf(" \n CHiLL does not handle constant memory mapping for more than 2D arrays.\n");
-        return;
-      }
-      
-      SgExpression* varexp = constant->get_mapped_symbol_exp(aref->name().c_str());
-      SgExpression* index_exp;
-      // build index expression
-      if(dims == 1) {
-        index_exp = static_cast<omega::CG_roseRepr*>(aref->index(0)->clone())->GetExpression();
-      }
-      if(dims == 2) {
-        VarDefs* arr_def = constant->get_vardef(aref->name().c_str());
-        CG_outputRepr* i0 = aref->index(0)->clone();
-        CG_outputRepr* i1 = aref->index(1)->clone();
-        CG_outputRepr* sz = new CG_roseRepr(buildIntVal(arr_def->size_multi_dim[0]));
-        CG_outputRepr* exp = ocg->CreatePlus(ocg->CreateTimes(sz->clone(), i0), i1);
-        index_exp = static_cast<omega::CG_roseRepr*>(exp->clone())->GetExpression();
-      }
-      ir->ReplaceExpression(aref, new CG_roseRepr(buildPntrArrRefExp(varexp, index_exp)));
-    }
-  }
-}
-
-/*
- * setupTexmappingVar
- * handles texture variable declaration
- * and adds a global texture object 
- * parameters:
- *   texture    - the texture_memory_mapping object
- *   arr_def    - the VarDefs object for the mapped variable
- *   globals    - Rose Global variables
- *   i          - an index to keep the new variable names unique
- *   devptr_sym - the devptr that the original variable is associated with
- *   symtab     - GPU function symbol table
- */
-static void setupTexmappingVar(texture_memory_mapping* texture, VarDefs* arr_def, SgGlobal* globals, int i, SgVariableSymbol* devptr_sym, SgSymbolTable* symtab) {  
-  char* buf1 = new char[32];
-  snprintf(buf1, 32, "tex%dRef", i+1);
-  arr_def->secondName = buf1;
-  
-  char buf2[64];
-  // single-dimensional 
-  snprintf(buf2, 64, "texture<float, %d, cudaReadModeElementType>", 1);
-  // multi-dimensional
-  // snprintf(buf2, 64, "texture<float, %d, cudaReadModeElemetType>", (int)(arr_def->size_multi_dim.size())); //*/
-  
-  SgVariableDeclaration* texvar_decl = buildVariableDeclaration(SgName(std::string(buf1)), buildOpaqueType(buf2, globals));
-  
-  SgInitializedNamePtrList& variables = texvar_decl->get_variables();
-  SgInitializedNamePtrList::const_iterator j = variables.begin();
-  SgInitializedName* initializedName = *j;
-  SgVariableSymbol* texvar_sym = new SgVariableSymbol(initializedName);
-  prependStatement(texvar_decl, globals);
-  
-  texvar_sym->set_parent(symtab);
-  symtab->insert(SgName(buf1), texvar_sym);
-  
-  texture->set_mapped_symbol(arr_def->original_name.c_str(), texvar_sym);
-  texture->set_devptr_symbol(arr_def->original_name.c_str(), devptr_sym);
-  texture->set_vardef(arr_def->original_name.c_str(), arr_def);
-}
-
-
-/*
- * One dimensional version of cudaBindTexture
- * see cudaBindTexture for details
- */
-static SgFunctionCallExp* cudaBindTexture1D(texture_memory_mapping* texture, VarDefs* arr_def, SgGlobal* globals) {
-  SgName cudaBindTexture_name("cudaBindTexture");
-  SgFunctionDeclaration* cudaBindTexture_decl = buildNondefiningFunctionDeclaration(
-      cudaBindTexture_name, buildVoidType(), buildFunctionParameterList(), globals);
-  
-  SgExprListExp* args = buildExprListExp();
-  args->append_expression(buildIntVal(0));
-  args->append_expression(texture->get_mapped_symbol_exp(arr_def->original_name.c_str()));
-  args->append_expression(texture->get_devptr_symbol_exp(arr_def->original_name.c_str()));
-  args->append_expression(arr_def->size_expr);
-  return buildFunctionCallExp(buildFunctionRefExp(cudaBindTexture_decl), args);
-}
-
-/*
- * Two dimensional version of cudaBindTexture
- * see cudaBindTexture for details
- */
-//static SgFunctionCallExp* cudaBindTexture2D(texture_memory_mapping* texture, VarDefs* arr_def, SgGlobal* globals) {
-//  SgName cudaBindTexture_name("cudaBindTexture2D");
-//  SgFunctionDeclaration* cudaBindTexture_decl = buildNondefiningFunctionDeclaration(
-//      cudaBindTexture_name, buildVoidType(), buildFunctionParameterList(), globals);
-//  
-//  SgExprListExp* args = buildExprListExp();
-//  args->append_expression(buildIntVal(0));
-//  args->append_expression(texture->get_tex_mapped_symbol_exp(arr_def->original_name.c_str()));
-//  args->append_expression(texture->get_devptr_symbol_exp(arr_def->original_name.c_str()));
-//  args->append_expression(buildIntVal(texture->get_dim_length(arr_def->original_name.c_str(), 0)));
-//  args->append_expression(buildIntVal(texture->get_dim_length(arr_def->original_name.c_str(), 1)));
-//  args->append_expression(arr_def->size_expr);
-//  return buildFunctionCallExp(buildFunctionRefExp(cudaBindTexture_decl), args);
-//}
-
-/*
- * cudaBindTexture
- * binds a variable to a texture
- * parameters:
- *    texture   - the texture mapping object
- *    arr_def   - the VarDefs object
- *    globals   - global symbol table
- *    stmt_list - the GPU functions' statement list
- * notes:
- *    only supports binding 1D textures, may need to consider cudaBindTexture2D for 2D textures
- */
-static void cudaBindTexture(texture_memory_mapping* texture, VarDefs* arr_def, SgGlobal* globals, SgStatementPtrList* stmt_list) {
-  //int dims = (int)(arr_def->size_multi_dim.size());
-  //int dims = texture->get_dims(arr_def->original_name.c_str());
-  //if(dims == 1)
-    stmt_list->push_back(
-      buildExprStatement(cudaBindTexture1D(texture, arr_def, globals)));
-  //if(dims == 2)
-  //  stmt_list->push_back(
-  //    buildExprStatement(cudaBindTexture2D(texture, arr_def, globals)));
-}
-
-/*
- * texmapArrayRefs
- * maps array reference expresions of texture mapped variables to the tex1D function
- * parameters:
- *    texture - the texture mapping object
- *    refs    - a list of all array read operations
- *    globals - global symbol table
- *    ir      - handles IR_Code operations
- *    ocg     - handles CG_roseBuilder operations
-**/
-static void texmapArrayRefs(texture_memory_mapping* texture, std::vector<IR_ArrayRef*>* refs, SgGlobal* globals, IR_Code* ir, CG_roseBuilder *ocg) {
-  // if texture mapping is not being used, ignore this function
-  if(texture == NULL) return;
-  for(int i = 0; i < refs->size(); i++) {
-    IR_ArrayRef* aref = (*refs)[i];
-    if(texture->is_array_mapped(aref->name().c_str())) {
-      
-      // get array dimensions
-      VarDefs* arr_def = texture->get_vardef(aref->name().c_str());
-      int dims = aref->symbol()->n_dim();
-      if(dims > 2) {
-        printf(" \n CHiLL does not handle texture mapping for more than 2D arrays.\n");
-        // TODO throw some sort of error. or handle in texture_copy function
-        return;
-      }
-      
-      // build texture lookup function declaration
-      char texNDfetch_strName[16];
-      sprintf(texNDfetch_strName, "tex%dDfetch", 1); // for now, only support tex1Dfetch
-      //sprintf(texNDfetch_strName, "tex%dDfetch", dims);
-      SgFunctionDeclaration* fetch_decl = buildNondefiningFunctionDeclaration(
-          SgName(texNDfetch_strName), buildFloatType(), buildFunctionParameterList(), globals);
-      
-      // build args
-      SgExprListExp* args = buildExprListExp();
-      args->append_expression(texture->get_mapped_symbol_exp(aref->name().c_str()));
-      
-      // set indexing args
-      //for(int i = 0; i < dims; i++) {
-      //  args->append_expression((static_cast<omega::CG_roseRepr*>(aref->index(i)->clone()))->GetExpression());
-      //}
-      if(dims == 1) {
-        args->append_expression(static_cast<omega::CG_roseRepr*>(aref->index(0)->clone())->GetExpression());
-      }
-      else if(dims == 2) {
-        CG_outputRepr* i0 = aref->index(0)->clone();
-        CG_outputRepr* i1 = aref->index(1)->clone();
-        CG_outputRepr* sz = new CG_roseRepr(buildIntVal(arr_def->size_multi_dim[0]));
-        CG_outputRepr* expr = ocg->CreatePlus(ocg->CreateTimes(sz->clone(), i0), i1);
-        args->append_expression(static_cast<omega::CG_roseRepr*>(expr->clone())->GetExpression());
-      }
-      
-      // build function call and replace original array ref
-      SgFunctionCallExp* fetch_call = buildFunctionCallExp(buildFunctionRefExp(fetch_decl), args);
-      ir->ReplaceExpression(aref, new CG_roseRepr(fetch_call));
-    }
-  }
-}
-
-SgNode* LoopCuda::cudaize_codegen_v2() {
-    if(cudaDebug)
-	printf("cudaize codegen V2\n");
-  CG_roseBuilder *ocg = dynamic_cast<CG_roseBuilder*>(ir->builder());
-  if (!ocg)
-    return false;
-
- //protonu--adding an annote to track texture memory type
-  //ANNOTE(k_cuda_texture_memory, "cuda texture memory", TRUE);
-  //ANNOTE(k_cuda_constant_memory, "cuda constant memory", TRUE);
-  int tex_mem_on = 0;
-  int cons_mem_on = 0;
-
-
-  
-  CG_outputRepr* repr;
-  std::vector<VarDefs> arrayVars;
-  std::vector<VarDefs> localScopedVars;
-  
-  std::vector<IR_ArrayRef *> ro_refs;
-  std::vector<IR_ArrayRef *> wo_refs;
-  std::set<std::string> uniqueRefs;
-  std::set<std::string> uniqueWoRefs;
-  std::set<const SgVariableSymbol *> syms;
-  std::set<const SgVariableSymbol *> psyms;
-  std::set<const SgVariableSymbol *> pdSyms;
-  SgStatementPtrList* replacement_list = new SgStatementPtrList;
-  
-  for (int j = 0; j < stmt.size(); j++) {
-    std::vector<IR_ArrayRef *> refs = ir->FindArrayRef(stmt[j].code);
-    for (int i = 0; i < refs.size(); i++) {
-      //printf("ref %s wo %d\n", static_cast<const char*>(refs[i]->name()), refs[i]->is_write());
-      SgVariableSymbol* var = body_symtab->find_variable(
-        SgName((char*) refs[i]->name().c_str()));
-      SgVariableSymbol* var2 = parameter_symtab->find_variable(
-        SgName((char*) refs[i]->name().c_str()));
-      
-      //If the array is not a parameter, then it's a local array and we
-      //want to recreate it as a stack variable in the kernel as opposed to
-      //passing it in.
-      if (var != NULL) {
-        //anand-- needs modification, if variable is parameter it wont be part of the
-        // block's symbol table but the functiond definition's symbol table
-        
-        continue;
-      }
-      if (uniqueRefs.find(refs[i]->name()) == uniqueRefs.end()) {
-        
-        uniqueRefs.insert(refs[i]->name());
-        if (refs[i]->is_write()) {
-          uniqueWoRefs.insert(refs[i]->name());
-          wo_refs.push_back(refs[i]);
-        } else
-          ro_refs.push_back(refs[i]);
-      }
-      if (refs[i]->is_write()
-          && uniqueWoRefs.find(refs[i]->name())
-          == uniqueWoRefs.end()) {
-        uniqueWoRefs.insert(refs[i]->name());
-        wo_refs.push_back(refs[i]);
-        //printf("adding %s to wo\n", static_cast<const char*>(refs[i]->name()));
-      }
-      pdSyms.insert((const SgVariableSymbol*) var2);
-    }
-  }
-  
-  if (cudaDebug) {
-      printf("reading from array ");
-      for (int i = 0; i < ro_refs.size(); i++)
-	  printf("'%s' ", ro_refs[i]->name().c_str());
-      printf("and writing to array ");
-      for (int i = 0; i < wo_refs.size(); i++)
-	  printf("'%s' ", wo_refs[i]->name().c_str());
-      printf("\n");
-  }  
-  const char* gridName = "dimGrid";
-  const char* blockName = "dimBlock";
-  
-  //TODO: Could allow for array_dims_vars to be a mapping from array
-  //references to to variable names that define their length.
-  SgVariableSymbol* dim1 = 0;
-  SgVariableSymbol* dim2 = 0;
-  
-  for (int i = 0; i < wo_refs.size(); i++) {
-    //TODO: Currently assume all arrays are floats of one or two dimentions
-    SgVariableSymbol* outArray = 0;
-    std::string name = wo_refs[i]->name();
-    outArray = body_symtab->find_variable(SgName((char*) name.c_str()));
-    int size_n_d;
-    if (outArray == NULL)
-      outArray = parameter_symtab->find_variable(
-        SgName((char*) name.c_str()));
-    
-    VarDefs v;
-    v.size_multi_dim = std::vector<int>();
-    char buf[32];
-    snprintf(buf, 32, "devO%dPtr", i + 1);
-    v.name = buf;
-    if (isSgPointerType(outArray->get_type())) {
-      if (isSgArrayType(
-            isSgNode(
-              isSgPointerType(outArray->get_type())->get_base_type()))) {
-        //  v.type = ((array_type *)(((ptr_type *)(outArray->type()))->ref_type()))->elem_type();
-        SgType* t =
-          isSgPointerType(outArray->get_type())->get_base_type();
-        /*   SgExprListExp* dimList = t->get_dim_info();
-             SgExpressionPtrList::iterator j= dimList->get_expressions().begin();
-             SgExpression* expr=NULL;
-             for (; j != dimList->get_expressions().end(); j++)
-             expr = *j;
-        */
-        while (isSgArrayType(t))
-          t = isSgArrayType(t)->get_base_type();
-        
-        if (!isSgType(t)) {
-          char buf[1024];
-          sprintf(buf, "CudaizeCodeGen: Array type undetected!");
-          throw std::runtime_error(buf);
-          
-        }
-        
-        v.type = t;
-      } else
-        v.type = isSgPointerType(outArray->get_type())->get_base_type();
-    } else if (isSgArrayType(outArray->get_type())) {
-      if (isSgArrayType(
-            isSgNode(
-              isSgArrayType(outArray->get_type())->get_base_type()))) {
-        //  v.type = ((array_type *)(((ptr_type *)(outArray->type()))->ref_type()))->elem_type();
-        SgType* t =
-          isSgArrayType(outArray->get_type())->get_base_type();
-        /*   SgExprListExp* dimList = t->get_dim_info();
-             SgExpressionPtrList::iterator j= dimList->get_expressions().begin();
-             SgExpression* expr=NULL;
-             for (; j != dimList->get_expressions().end(); j++)
-             expr = *j;
-        */
-        while (isSgArrayType(t))
-          t = isSgArrayType(t)->get_base_type();
-        
-        if (!isSgType(t)) {
-          char buf[1024];
-          sprintf(buf, "CudaizeCodeGen: Array type undetected!");
-          throw std::runtime_error(buf);
-          
-        }
-        
-        v.type = t;
-      } else
-        v.type = isSgArrayType(outArray->get_type())->get_base_type();
-    } else
-      v.type = buildFloatType();
-    v.tex_mapped = false;
-    v.cons_mapped = false;
-    v.original_name = wo_refs[i]->name();
-    //Size of the array = dim1 * dim2 * num bytes of our array type
-    
-    //If our input array is 2D (non-linearized), we want the actual
-    //dimentions of the array
-    CG_outputRepr* size;
-    //Lookup in array_dims
-    std::map<std::string, int>::iterator it = array_dims.find(name.c_str());
-    if (isSgPointerType(outArray->get_type())
-        && isSgArrayType(
-          isSgNode(
-            isSgPointerType(outArray->get_type())->get_base_type()))) {
-      SgType* t = isSgPointerType(outArray->get_type())->get_base_type();
-      /*   SgExprListExp* dimList = t->get_dim_info();
-           SgExpressionPtrList::iterator j= dimList->get_expressions().begin();
-           SgExpression* expr=NULL;
-           for (; j != dimList->get_expressions().end(); j++)
-           expr = *j;
-      */
-      if (isSgIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedLongVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedLongVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgLongIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgLongLongIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgLongIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedLongLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedLongLongIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgAddOp(isSgArrayType(t)->get_index())) {
-        SgAddOp *op_add = isSgAddOp(isSgArrayType(t)->get_index());
-        
-        SgExpression *lhs = op_add->get_lhs_operand();
-        SgExpression *rhs = op_add->get_rhs_operand();
-        
-        if (isSgIntVal(lhs))
-          size_n_d = (int) isSgIntVal(lhs)->get_value() + (int) (isSgIntVal(rhs)->get_value());
-        else if (isSgUnsignedIntVal(lhs))
-          size_n_d = (int) isSgUnsignedIntVal(lhs)->get_value()
-            + (int) isSgUnsignedIntVal(rhs)->get_value();
-        else if (isSgUnsignedLongVal(lhs))
-          size_n_d = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongIntVal(lhs))
-          size_n_d = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongLongIntVal(lhs))
-          size_n_d = (int) (isSgLongLongIntVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongIntVal(lhs))
-          size_n_d = (int) (isSgLongIntVal(lhs)->get_value()
-                            + isSgLongIntVal(rhs)->get_value());
-        else if (isSgUnsignedLongLongIntVal(lhs))
-          size_n_d =
-            (int) (isSgUnsignedLongLongIntVal(lhs)->get_value()
-                   + isSgUnsignedLongLongIntVal(rhs)->get_value());
-        
-      }
-      t = isSgArrayType(t)->get_base_type();
-      while (isSgArrayType(t)) {
-        int dim;
-        if (isSgIntVal(isSgArrayType(t)->get_index()))
-          dim =
-            (int) (isSgIntVal(isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedLongVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedLongVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedLongLongIntVal(
-                   isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedLongLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgAddOp(isSgArrayType(t)->get_index())) {
-          SgAddOp *op_add = isSgAddOp(isSgArrayType(t)->get_index());
-          
-          SgExpression *lhs = op_add->get_lhs_operand();
-          SgExpression *rhs = op_add->get_rhs_operand();
-          
-          if (isSgIntVal(lhs))
-            dim = (int) isSgIntVal(lhs)->get_value()
-              + (int) (isSgIntVal(rhs)->get_value());
-          else if (isSgUnsignedIntVal(lhs))
-            dim = (int) isSgUnsignedIntVal(lhs)->get_value()
-              + (int) isSgUnsignedIntVal(rhs)->get_value();
-          else if (isSgUnsignedLongVal(lhs))
-            dim = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongIntVal(lhs))
-            dim = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongLongIntVal(lhs))
-            dim = (int) (isSgLongLongIntVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongIntVal(lhs))
-            dim = (int) (isSgLongIntVal(lhs)->get_value()
-                         + isSgLongIntVal(rhs)->get_value());
-          else if (isSgUnsignedLongLongIntVal(lhs))
-            dim =
-              (int) (isSgUnsignedLongLongIntVal(lhs)->get_value()
-                     + isSgUnsignedLongLongIntVal(rhs)->get_value());
-          
-        }
-        size_n_d *= dim;
-        v.size_multi_dim.push_back(dim);
-        t = isSgArrayType(t)->get_base_type();
-      }
-      //v.size_2d = (int) (isSgIntVal(t->get_index())->get_value());
-      
-      if (cudaDebug)
-	  printf("Detected Multi-dimensional array sized of %d for %s\n",
-		 size_n_d, (char*) wo_refs[i]->name().c_str());
-      size = ocg->CreateInt(size_n_d);
-    } else if (isSgArrayType(outArray->get_type())
-               && isSgArrayType(
-                 isSgNode(
-                   isSgArrayType(outArray->get_type())->get_base_type()))) {
-      SgType* t = outArray->get_type();
-      /*   SgExprListExp* dimList = t->get_dim_info();
-           SgExpressionPtrList::iterator j= dimList->get_expressions().begin();
-           SgExpression* expr=NULL;
-           for (; j != dimList->get_expressions().end(); j++)
-           expr = *j;
-      */
-      
-      if (isSgIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedLongVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedLongVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgLongIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgLongLongIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgLongIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedLongLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedLongLongIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgAddOp(isSgArrayType(t)->get_index())) {
-        SgAddOp *op_add = isSgAddOp(isSgArrayType(t)->get_index());
-        
-        SgExpression *lhs = op_add->get_lhs_operand();
-        SgExpression *rhs = op_add->get_rhs_operand();
-        
-        if (isSgIntVal(lhs))
-          size_n_d = (int) isSgIntVal(lhs)->get_value() + (int) (isSgIntVal(rhs)->get_value());
-        else if (isSgUnsignedIntVal(lhs))
-          size_n_d = (int) isSgUnsignedIntVal(lhs)->get_value()
-            + (int) isSgUnsignedIntVal(rhs)->get_value();
-        else if (isSgUnsignedLongVal(lhs))
-          size_n_d = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongIntVal(lhs))
-          size_n_d = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongLongIntVal(lhs))
-          size_n_d = (int) (isSgLongLongIntVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongIntVal(lhs))
-          size_n_d = (int) (isSgLongIntVal(lhs)->get_value()
-                            + isSgLongIntVal(rhs)->get_value());
-        else if (isSgUnsignedLongLongIntVal(lhs))
-          size_n_d =
-            (int) (isSgUnsignedLongLongIntVal(lhs)->get_value()
-                   + isSgUnsignedLongLongIntVal(rhs)->get_value());
-        
-      }
-      t = isSgArrayType(t)->get_base_type();
-      while (isSgArrayType(t)) {
-        int dim;
-        if (isSgIntVal(isSgArrayType(t)->get_index()))
-          dim =
-            (int) (isSgIntVal(isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedLongVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedLongVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedLongLongIntVal(
-                   isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedLongLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgAddOp(isSgArrayType(t)->get_index())) {
-          SgAddOp *op_add = isSgAddOp(isSgArrayType(t)->get_index());
-          
-          SgExpression *lhs = op_add->get_lhs_operand();
-          SgExpression *rhs = op_add->get_rhs_operand();
-          
-          if (isSgIntVal(lhs))
-            dim = (int) isSgIntVal(lhs)->get_value()
-              + (int) (isSgIntVal(rhs)->get_value());
-          else if (isSgUnsignedIntVal(lhs))
-            dim = (int) isSgUnsignedIntVal(lhs)->get_value()
-              + (int) isSgUnsignedIntVal(rhs)->get_value();
-          else if (isSgUnsignedLongVal(lhs))
-            dim = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongIntVal(lhs))
-            dim = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongLongIntVal(lhs))
-            dim = (int) (isSgLongLongIntVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongIntVal(lhs))
-            dim = (int) (isSgLongIntVal(lhs)->get_value()
-                         + isSgLongIntVal(rhs)->get_value());
-          else if (isSgUnsignedLongLongIntVal(lhs))
-            dim =
-              (int) (isSgUnsignedLongLongIntVal(lhs)->get_value()
-                     + isSgUnsignedLongLongIntVal(rhs)->get_value());
-          
-        }
-        size_n_d *= dim;
-        v.size_multi_dim.push_back(dim);
-        t = isSgArrayType(t)->get_base_type();
-      }
-      
-      //v.size_2d = (int) (isSgIntVal(t->get_index())->get_value());
-      
-      if (cudaDebug)
-	  printf("Detected Multi-Dimensional array sized of %d for %s\n",
-		 size_n_d, (char*) wo_refs[i]->name().c_str());
-      size = ocg->CreateInt(size_n_d);
-    } else if (it != array_dims.end()) {
-      int ref_size = it->second;
-      //size =
-      //        ocg->CreateInt(
-      //                isSgIntVal(
-      //                        isSgArrayType(outArray->get_type())->get_index())->get_value());
-      //v.size_2d = isSgArrayType(outArray->get_type())->get_rank();
-      //v.var_ref_size = ref_size;
-      size = ocg->CreateInt(ref_size);
-      
-    } else {
-      if (dim1) {
-        size = ocg->CreateTimes(
-          new CG_roseRepr(isSgExpression(buildVarRefExp(dim1))),
-          new CG_roseRepr(isSgExpression(buildVarRefExp(dim2))));
-      } else {
-        char buf[1024];
-        sprintf(buf,
-                "CudaizeCodeGen: Array reference %s does not have a "
-                "detectable size or specififed dimentions",
-                name.c_str());
-        throw std::runtime_error(buf);
-      }
-    }
-    
-    v.size_expr =
-      static_cast<CG_roseRepr*>(ocg->CreateTimes(size,
-                                                 new omega::CG_roseRepr(
-                                                   isSgExpression(buildSizeOfOp(v.type)))))->GetExpression();
-    
-    v.in_data = 0;
-    v.out_data = outArray;
-    //Check for in ro_refs and remove it at this point
-    std::vector<IR_ArrayRef *>::iterator it_;
-    for (it_ = ro_refs.begin(); it_ != ro_refs.end(); it_++) {
-      if ((*it_)->name() == wo_refs[i]->name()) {
-        break;
-      }
-    }
-    if (it_ != ro_refs.end()) {
-      v.in_data = outArray;
-      ro_refs.erase(it_);
-    }
-    
-    arrayVars.push_back(v);
-    
-  }
-  
-  //protonu-- assuming that all texture mapped memories were originally read only mems
-  //there should be safety checks for that, will implement those later
-  
-  for (int i = 0; i < ro_refs.size(); i++) {
-    SgVariableSymbol* inArray = 0;
-    std::string name = ro_refs[i]->name();
-    inArray = body_symtab->find_variable(SgName((char*) name.c_str()));
-    if (inArray == NULL)
-      inArray = parameter_symtab->find_variable(
-        SgName((char*) name.c_str()));
-    
-    VarDefs v;
-    v.size_multi_dim = std::vector<int>();
-    char buf[32];
-    snprintf(buf, 32, "devI%dPtr", i + 1);
-    v.name = buf;
-    int size_n_d;
-    if (isSgPointerType(inArray->get_type())) {
-      if (isSgArrayType(
-            isSgNode(
-              isSgPointerType(inArray->get_type())->get_base_type()))) {
-        
-        SgType* t =
-          isSgPointerType(inArray->get_type())->get_base_type();
-        
-        while (isSgArrayType(t))
-          t = isSgArrayType(t)->get_base_type();
-        
-        if (!isSgType(t)) {
-          char buf[1024];
-          sprintf(buf, "CudaizeCodeGen: Array type undetected!");
-          throw std::runtime_error(buf);
-          
-        }
-        v.type = t;
-      } else
-        v.type = isSgPointerType(inArray->get_type())->get_base_type();
-    } else if (isSgArrayType(inArray->get_type())) {
-      if (isSgArrayType(
-            isSgNode(
-              isSgArrayType(inArray->get_type())->get_base_type()))) {
-        
-        SgType* t = inArray->get_type();
-        while (isSgArrayType(t))
-          t = isSgArrayType(t)->get_base_type();
-        
-        if (!isSgType(t)) {
-          char buf[1024];
-          sprintf(buf, "CudaizeCodeGen: Array type undetected!");
-          throw std::runtime_error(buf);
-          
-        }
-        v.type = t;
-      } else
-        v.type = isSgArrayType(inArray->get_type())->get_base_type();
-    }
-    
-    else
-      v.type = buildFloatType();
-    
-    v.tex_mapped = false;
-    v.cons_mapped = false;
-    v.original_name = ro_refs[i]->name();
-    
-    //derick -- adding texture and constant mapping
-    if ( texture != NULL)
-      v.tex_mapped = (texture->is_array_mapped(name.c_str()))? true:false; //protonu-track tex mapped vars
-    if (v.tex_mapped){
-      printf("this variable  %s is mapped to texture memory", name.c_str());
-    }
-    //derick -- this is commented out until constant memory is implemeted
-    if ( constant_mem != NULL)
-      v.cons_mapped = (constant_mem->is_array_mapped(name.c_str()))? true:false; //protonu-track tex mapped vars
-    if (v.cons_mapped){
-      printf("this variable  %s is mapped to constant memory", name.c_str());
-    }
-    
-    //Size of the array = dim1 * dim2 * num bytes of our array type
-    //If our input array is 2D (non-linearized), we want the actual
-    //dimentions of the array (as it might be less than cu_n
-    CG_outputRepr* size;
-    //Lookup in array_dims
-    std::map<std::string, int>::iterator it = array_dims.find(name.c_str());
-    if (isSgPointerType(inArray->get_type())
-        && isSgArrayType(
-          isSgPointerType(inArray->get_type())->get_base_type())) {
-      //SgArrayType* t = isSgArrayType(isSgArrayType(inArray->get_type())->get_base_type());
-      //v.size_2d = t->get_rank();
-      SgType* t = isSgPointerType(inArray->get_type())->get_base_type();
-      /*   SgExprListExp* dimList = t->get_dim_info();
-           SgExpressionPtrList::iterator j= dimList->get_expressions().begin();
-           SgExpression* expr=NULL;
-           for (; j != dimList->get_expressions().end(); j++)
-           expr = *j;
-      */
-      //v.size_2d = 1;
-      if (isSgIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedLongVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedLongVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgLongIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgLongLongIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgLongIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedLongLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedLongLongIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgAddOp(isSgArrayType(t)->get_index())) {
-        SgAddOp *op_add = isSgAddOp(isSgArrayType(t)->get_index());
-        
-        SgExpression *lhs = op_add->get_lhs_operand();
-        SgExpression *rhs = op_add->get_rhs_operand();
-        
-        if (isSgIntVal(lhs))
-          size_n_d = (int) isSgIntVal(lhs)->get_value() + (int) (isSgIntVal(rhs)->get_value());
-        else if (isSgUnsignedIntVal(lhs))
-          size_n_d = (int) isSgUnsignedIntVal(lhs)->get_value()
-            + (int) isSgUnsignedIntVal(rhs)->get_value();
-        else if (isSgUnsignedLongVal(lhs))
-          size_n_d = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongIntVal(lhs))
-          size_n_d = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongLongIntVal(lhs))
-          size_n_d = (int) (isSgLongLongIntVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongIntVal(lhs))
-          size_n_d = (int) (isSgLongIntVal(lhs)->get_value()
-                            + isSgLongIntVal(rhs)->get_value());
-        else if (isSgUnsignedLongLongIntVal(lhs))
-          size_n_d =
-            (int) (isSgUnsignedLongLongIntVal(lhs)->get_value()
-                   + isSgUnsignedLongLongIntVal(rhs)->get_value());
-        
-      }
-      t = isSgArrayType(t)->get_base_type();
-      while (isSgArrayType(t)) {
-        int dim;
-        if (isSgIntVal(isSgArrayType(t)->get_index()))
-          dim =
-            (int) (isSgIntVal(isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedLongVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedLongVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedLongLongIntVal(
-                   isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedLongLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgAddOp(isSgArrayType(t)->get_index())) {
-          SgAddOp *op_add = isSgAddOp(isSgArrayType(t)->get_index());
-          
-          SgExpression *lhs = op_add->get_lhs_operand();
-          SgExpression *rhs = op_add->get_rhs_operand();
-          
-          if (isSgIntVal(lhs))
-            dim = (int) isSgIntVal(lhs)->get_value()
-              + (int) (isSgIntVal(rhs)->get_value());
-          else if (isSgUnsignedIntVal(lhs))
-            dim = (int) isSgUnsignedIntVal(lhs)->get_value()
-              + (int) isSgUnsignedIntVal(rhs)->get_value();
-          else if (isSgUnsignedLongVal(lhs))
-            dim = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongIntVal(lhs))
-            dim = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongLongIntVal(lhs))
-            dim = (int) (isSgLongLongIntVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongIntVal(lhs))
-            dim = (int) (isSgLongIntVal(lhs)->get_value()
-                         + isSgLongIntVal(rhs)->get_value());
-          else if (isSgUnsignedLongLongIntVal(lhs))
-            dim =
-              (int) (isSgUnsignedLongLongIntVal(lhs)->get_value()
-                     + isSgUnsignedLongLongIntVal(rhs)->get_value());
-          
-        }
-        size_n_d *= dim;
-        v.size_multi_dim.push_back(dim);
-        t = isSgArrayType(t)->get_base_type();
-      }
-      if (cudaDebug)
-	  printf("Detected Multi-dimensional array sized of %d for %s\n",
-		 size_n_d, (char*) ro_refs[i]->name().c_str());
-      size = ocg->CreateInt(size_n_d);
-    } else if (isSgArrayType(inArray->get_type())
-               && isSgArrayType(
-                 isSgArrayType(inArray->get_type())->get_base_type())) {
-      //SgArrayType* t = isSgArrayType(isSgArrayType(inArray->get_type())->get_base_type());
-      //v.size_2d = t->get_rank();
-      SgType* t = inArray->get_type();
-      /*   SgExprListExp* dimList = t->get_dim_info();
-           SgExpressionPtrList::iterator j= dimList->get_expressions().begin();
-           SgExpression* expr=NULL;
-           for (; j != dimList->get_expressions().end(); j++)
-           expr = *j;
-      */
-      
-      if (isSgIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedLongVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedLongVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgLongIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgLongLongIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d =
-          (int) (isSgLongIntVal(isSgArrayType(t)->get_index())->get_value());
-      else if (isSgUnsignedLongLongIntVal(isSgArrayType(t)->get_index()))
-        size_n_d = (int) (isSgUnsignedLongLongIntVal(
-                            isSgArrayType(t)->get_index())->get_value());
-      else if (isSgAddOp(isSgArrayType(t)->get_index())) {
-        SgAddOp *op_add = isSgAddOp(isSgArrayType(t)->get_index());
-        
-        SgExpression *lhs = op_add->get_lhs_operand();
-        SgExpression *rhs = op_add->get_rhs_operand();
-        
-        if (isSgIntVal(lhs))
-          size_n_d = (int) isSgIntVal(lhs)->get_value() + (int) (isSgIntVal(rhs)->get_value());
-        else if (isSgUnsignedIntVal(lhs))
-          size_n_d = (int) isSgUnsignedIntVal(lhs)->get_value()
-            + (int) isSgUnsignedIntVal(rhs)->get_value();
-        else if (isSgUnsignedLongVal(lhs))
-          size_n_d = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongIntVal(lhs))
-          size_n_d = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongLongIntVal(lhs))
-          size_n_d = (int) (isSgLongLongIntVal(lhs)->get_value()
-                            + isSgUnsignedLongVal(rhs)->get_value());
-        else if (isSgLongIntVal(lhs))
-          size_n_d = (int) (isSgLongIntVal(lhs)->get_value()
-                            + isSgLongIntVal(rhs)->get_value());
-        else if (isSgUnsignedLongLongIntVal(lhs))
-          size_n_d =
-            (int) (isSgUnsignedLongLongIntVal(lhs)->get_value()
-                   + isSgUnsignedLongLongIntVal(rhs)->get_value());
-        
-      }
-      t = isSgArrayType(t)->get_base_type();
-      while (isSgArrayType(t)) {
-        int dim;
-        if (isSgIntVal(isSgArrayType(t)->get_index()))
-          dim =
-            (int) (isSgIntVal(isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedLongVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedLongVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgLongIntVal(isSgArrayType(t)->get_index()))
-          dim = (int) (isSgLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgUnsignedLongLongIntVal(
-                   isSgArrayType(t)->get_index()))
-          dim = (int) (isSgUnsignedLongLongIntVal(
-                         isSgArrayType(t)->get_index())->get_value());
-        else if (isSgAddOp(isSgArrayType(t)->get_index())) {
-          SgAddOp *op_add = isSgAddOp(isSgArrayType(t)->get_index());
-          
-          SgExpression *lhs = op_add->get_lhs_operand();
-          SgExpression *rhs = op_add->get_rhs_operand();
-          
-          if (isSgIntVal(lhs))
-            dim = (int) isSgIntVal(lhs)->get_value()
-              + (int) (isSgIntVal(rhs)->get_value());
-          else if (isSgUnsignedIntVal(lhs))
-            dim = (int) isSgUnsignedIntVal(lhs)->get_value()
-              + (int) isSgUnsignedIntVal(rhs)->get_value();
-          else if (isSgUnsignedLongVal(lhs))
-            dim = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongIntVal(lhs))
-            dim = (int) (isSgUnsignedLongVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongLongIntVal(lhs))
-            dim = (int) (isSgLongLongIntVal(lhs)->get_value()
-                         + isSgUnsignedLongVal(rhs)->get_value());
-          else if (isSgLongIntVal(lhs))
-            dim = (int) (isSgLongIntVal(lhs)->get_value()
-                         + isSgLongIntVal(rhs)->get_value());
-          else if (isSgUnsignedLongLongIntVal(lhs))
-            dim =
-              (int) (isSgUnsignedLongLongIntVal(lhs)->get_value()
-                     + isSgUnsignedLongLongIntVal(rhs)->get_value());
-          
-        }
-        size_n_d *= dim;
-        v.size_multi_dim.push_back(dim);
-        t = isSgArrayType(t)->get_base_type();
-      }
-      if (cudaDebug)
-	  printf("Detected Multi-Dimensional array sized of %d for %s\n",
-		 size_n_d, (char*) ro_refs[i]->name().c_str());
-      size = ocg->CreateInt(size_n_d);
-    }
-    
-    else if (it != array_dims.end()) {
-      int ref_size = it->second;
-      //                v.var_ref_size = ref_size;
-      size = ocg->CreateInt(ref_size);
-    } else {
-      if (dim1) {
-        size = ocg->CreateTimes(
-          new CG_roseRepr(isSgExpression(buildVarRefExp(dim1))),
-          new CG_roseRepr(isSgExpression(buildVarRefExp(dim2))));
-      } else {
-        char buf[1024];
-        sprintf(buf,
-                "CudaizeCodeGen: Array reference %s does not have a "
-                "detectable size or specififed dimentions",
-                name.c_str());
-        throw std::runtime_error(buf);
-      }
-    }
-    v.size_expr =
-      static_cast<CG_roseRepr*>(ocg->CreateTimes(size,
-                                                 new omega::CG_roseRepr(
-                                                   isSgExpression(buildSizeOfOp(v.type)))))->GetExpression();
-    
-    v.in_data = inArray;
-    v.out_data = 0;
-    arrayVars.push_back(v);
-  }
-  
-  if (arrayVars.size() < 2) {
-    fprintf(stderr,
-            "cudaize error: Did not find two arrays being accessed\n");
-    return false;
-  }
-
-  //protonu--debugging tool--the printf statement
-  //tex_mem_on signals use of tex mem
-  /* derick -- texmapping near malloc mcopy
-  for(int i=0; i<arrayVars.size(); i++)
-  {
-	  //printf("var name %s, tex_mem used %s\n", arrayVars[i].name.c_str(), (arrayVars[i].tex_mapped)?"true":"false");
-	  if (arrayVars[i].tex_mapped  ) tex_mem_on ++;
-	  //if (arrayVars[i].cons_mapped  ) cons_mem_on ++;
-  }
-  */
-  
-  //Add our mallocs (and input array memcpys)
-  for (int i = 0; i < arrayVars.size(); i++) {
-    if(arrayVars[i].cons_mapped) {
-      setupConstantVar(constant_mem, &arrayVars[i], globals, i, symtab);
-      SgStatementPtrList *tnl = new SgStatementPtrList;
-      cudaBindConstantVar(constant_mem, &arrayVars[i], globals, tnl);
-      setup_code = ocg->StmtListAppend(setup_code, new CG_roseRepr(tnl));
-    } 
-    else {
-      SgVariableDeclaration* defn = buildVariableDeclaration(
-        SgName(arrayVars[i].name.c_str()),
-        buildPointerType(arrayVars[i].type));
-      SgInitializedNamePtrList& variables = defn->get_variables();
-      SgInitializedNamePtrList::const_iterator j = variables.begin();
-      SgInitializedName* initializedName = *j;
-      SgVariableSymbol* dvs = new SgVariableSymbol(initializedName);
-      prependStatement(defn, func_body);
-      
-      dvs->set_parent(body_symtab);
-      body_symtab->insert(SgName(arrayVars[i].name.c_str()), dvs);
-    
-//    SgVariableSymbol* dvs = body_symtab->find_variable(SgName(arrayVars[i].name.c_str()));
-    
-    //  if(dvs == NULL)
-    //      dvs =  parameter_symtab->find_variable(SgName(arrayVars[i].name.c_str()));
-    
-    //cudaMalloc args
-    // SgBasicBlock* block = buildBasicBlock();
-      SgName name_cuda_malloc("cudaMalloc");
-      SgFunctionDeclaration * decl_cuda_malloc =
-        buildNondefiningFunctionDeclaration(name_cuda_malloc,
-                                            buildVoidType(), buildFunctionParameterList(), globals);
-      
-      SgName name_cuda_copy("cudaMemcpy");
-      SgFunctionDeclaration * decl_cuda_copy =
-        buildNondefiningFunctionDeclaration(name_cuda_copy,
-                                            buildVoidType(), buildFunctionParameterList(), globals);
-      
-      SgExprListExp* args = buildExprListExp();
-      args->append_expression(
-        buildCastExp(buildAddressOfOp(buildVarRefExp(dvs)),
-                     buildPointerType(buildPointerType(buildVoidType()))));
-      args->append_expression(arrayVars[i].size_expr);
-    
-//    decl_cuda_malloc->get_parameterList()->append_arg
-      SgFunctionCallExp *the_call = buildFunctionCallExp(
-        buildFunctionRefExp(decl_cuda_malloc), args);
-      
-      SgExprStatement* stmt = buildExprStatement(the_call);
-    
-    //  (*replacement_list).push_back (stmt);
-    
-      SgStatementPtrList* tnl = new SgStatementPtrList;
-      (*tnl).push_back(stmt);
-      setup_code = ocg->StmtListAppend(setup_code, new CG_roseRepr(tnl));
-      if (arrayVars[i].in_data) {
-      
-        SgExprListExp * cuda_copy_in_args = buildExprListExp();
-        cuda_copy_in_args->append_expression(
-          isSgExpression(buildVarRefExp(dvs)));
-        cuda_copy_in_args->append_expression(
-          isSgExpression(buildVarRefExp(arrayVars[i].in_data)));
-        CG_roseRepr* size_exp = new CG_roseRepr(arrayVars[i].size_expr);
-        cuda_copy_in_args->append_expression(
-          static_cast<CG_roseRepr*>(size_exp->clone())->GetExpression());
-        cuda_copy_in_args->append_expression(
-          buildOpaqueVarRefExp("cudaMemcpyHostToDevice", globals));
-      
-//                                      cuda_copy_in_args->append_expression(
-//                                              new SgVarRefExp(sourceLocation, )
-//                                      );
-        SgFunctionCallExp * cuda_copy_in_func_call = buildFunctionCallExp(
-          buildFunctionRefExp(decl_cuda_copy), cuda_copy_in_args);
-      
-        SgExprStatement* stmt = buildExprStatement(cuda_copy_in_func_call);
-      
-        SgStatementPtrList *tnl = new SgStatementPtrList;
-        (*tnl).push_back(stmt);
-        setup_code = ocg->StmtListAppend(setup_code, new CG_roseRepr(tnl));
-      
-        if(arrayVars[i].tex_mapped) {
-          setupTexmappingVar(texture, &arrayVars[i], globals, i, dvs, symtab);
-          SgStatementPtrList *tnl = new SgStatementPtrList;
-          cudaBindTexture(texture, &arrayVars[i], globals, tnl);
-          setup_code = ocg->StmtListAppend(setup_code, new CG_roseRepr(tnl));
-        }
-      }
-    }
-  }
-  
-  //Build dimGrid dim3 variables based on loop dimentions and ti/tj
-  char blockD1[120];
-  char blockD2[120];
-  if (dim1) {
-    snprintf(blockD1, 120, "%s/%d",
-             dim1->get_declaration()->get_name().getString().c_str(), cu_tx);
-    snprintf(blockD2, 120, "%s/%d",
-             dim2->get_declaration()->get_name().getString().c_str(), cu_ty);
-  } else {
-    snprintf(blockD1, 120, "%d", cu_bx);
-    snprintf(blockD2, 120, "%d", cu_by);
-    //snprintf(blockD1, 120, "%d/%d", cu_nx, cu_tx);
-    //snprintf(blockD2, 120, "%d/%d", cu_ny, cu_ty);
-  }
-  
-  SgInitializedName* arg1 = buildInitializedName("i", buildIntType());
-  SgInitializedName* arg2 = buildInitializedName("j", buildIntType());
-  SgInitializedName* arg3 = buildInitializedName("k", buildIntType());
-  SgName type_name("dim3");
-  //SgClassSymbol * type_symbol = globalScope->lookup_class_symbol(type_name);
-  
-  //ROSE_ASSERT(type_symbol != NULL);
-  
-  //SgClassDeclaration * dim3classdecl = isSgClassDeclaration(
-  //        type_symbol->get_declaration());
-  
-  SgFunctionDeclaration * funcdecl = buildNondefiningFunctionDeclaration(
-    SgName("dim3"), buildOpaqueType("dim3", globalScope),
-    //isSgType(dim3classdecl->get_type()),
-    buildFunctionParameterList(arg1, arg2, arg3), globalScope);
-  
-  if (cu_bx && cu_by)
-    repr = ocg->CreateDim3((const char*) gridName, ocg->CreateInt(cu_bx),
-                           ocg->CreateInt(cu_by));
-  else if (cu_bx_repr && cu_by_repr)
-    repr = ocg->CreateDim3((const char*) gridName, cu_bx_repr, cu_by_repr);
-  else if (cu_bx_repr)
-    repr = ocg->CreateDim3((const char*) gridName, cu_bx_repr,
-                           ocg->CreateInt(1));
-  setup_code = ocg->StmtListAppend(setup_code, repr);
-  //SgStatementPtrList* dimList = static_cast<CG_roseRepr *>(repr)->GetList();
-  
-  //for(SgStatementPtrList::iterator it = (*dimList).begin(); it != (*dimList).end(); it++)
-  //    (*replacement_list).push_back (*it);
-  
-  //  repr = ocg->CreateDim3((const char*)blockName, cu_tx,cu_ty);
-  
-  if (cu_tz > 1 || cu_tz_repr) {
-    
-    if (cu_tx && cu_ty && cu_tz)
-      repr = ocg->CreateDim3((char*) blockName, ocg->CreateInt(cu_tx),
-                             ocg->CreateInt(cu_ty), ocg->CreateInt(cu_tz));
-    else if (cu_tx_repr && cu_ty_repr && cu_tz_repr)
-      repr = ocg->CreateDim3((char*) blockName, cu_tx_repr, cu_ty_repr,
-                             cu_tz_repr);
-    // SgStatementPtrList* dimList = static_cast<CG_roseRepr *>(repr)->GetList();
-    
-    // for(SgStatementPtrList::iterator it = (*dimList).begin(); it != (*dimList).end(); it++)
-    //    (*replacement_list).push_back (*it);
-    
-  } else {
-    if (cu_tx && cu_ty)
-      repr = ocg->CreateDim3((char*) blockName, ocg->CreateInt(cu_tx),
-                             ocg->CreateInt(cu_ty));
-    else if (cu_tx_repr && cu_ty_repr)
-      repr = ocg->CreateDim3((char*) blockName, cu_tx_repr, cu_ty_repr);
-    //SgStatementPtrList* dimList = static_cast<CG_roseRepr *>(repr)->GetList();
-    
-    //for(SgStatementPtrList::iterator it = (*dimList).begin(); it != (*dimList).end(); it++)
-    //   (*replacement_list).push_back (*it);
-    
-  }
-  
-  setup_code = ocg->StmtListAppend(setup_code, repr);
-  
-  SgCudaKernelExecConfig* config = new SgCudaKernelExecConfig(
-    buildVarRefExp(gridName), buildVarRefExp(blockName), NULL, NULL);
-  //SgCudaKernelExecConfig* config = new SgCudaKernelExecConfig(buildIntVal(cu_bx), , NULL, NULL);
-  SgExprListExp* iml = new SgExprListExp();
-  SgCastExp* dim_s;
-  
-  //Creating Kernel function
-  SgBasicBlock* bb = new SgBasicBlock(TRANSFORMATION_FILE_INFO);
-  SgFunctionDefinition* kernel_defn = new SgFunctionDefinition(
-    TRANSFORMATION_FILE_INFO, bb);
-  SgFunctionDeclaration* kernel_decl_ = new SgFunctionDeclaration(
-    TRANSFORMATION_FILE_INFO, SgName((char*)cu_kernel_name.c_str()),buildFunctionType(buildVoidType(), buildFunctionParameterList()), kernel_defn);
-  SgFunctionDeclaration* kernel_decl = new SgFunctionDeclaration(
-    TRANSFORMATION_FILE_INFO, SgName((char*)cu_kernel_name.c_str()),buildFunctionType(buildVoidType(), buildFunctionParameterList()), kernel_defn);
-  
-  //((kernel_decl->get_declarationModifier()).get_storageModifier()).setStatic();
-  
-  kernel_decl->set_definingDeclaration(kernel_decl);
-  kernel_defn->set_parent(kernel_decl);
-  bb->set_parent(kernel_defn);
-  bb->set_endOfConstruct(TRANSFORMATION_FILE_INFO);
-  bb->get_endOfConstruct()->set_parent(bb);
-  
-  //SgFunctionSymbol* functionSymbol = new SgFunctionSymbol(kernel_decl_);
-  //globals->insert_symbol(SgName((char*) cu_kernel_name.c_str()),
-  //            functionSymbol);
-  SgFunctionSymbol* functionSymbol2 = new SgFunctionSymbol(kernel_decl);
-  
-  globals->insert_symbol(SgName((char*) cu_kernel_name.c_str()),
-                         functionSymbol2);
-  
-  kernel_decl_->set_parent(globals);
-  
-  kernel_decl_->set_scope(globals);
-  
-  kernel_decl_->setForward();
-  
-  globals->prepend_declaration(kernel_decl_);
-  
-  kernel_decl->set_endOfConstruct(TRANSFORMATION_FILE_INFO);
-  kernel_decl->get_endOfConstruct()->set_parent(kernel_decl);
-  
-  kernel_decl->set_parent(globals);
-  kernel_decl->set_scope(globals);
-  
-  kernel_decl->get_definition()->set_endOfConstruct(TRANSFORMATION_FILE_INFO);
-  kernel_decl->get_definition()->get_endOfConstruct()->set_parent(
-    kernel_decl->get_definition());
-  
-  globals->append_statement(kernel_decl);
-  
-  //printf("%s %s\n", static_cast<const char*>(cu_kernel_name), dims);
-  //--derick - kernel function parameters  
-  for (int i = 0; i < arrayVars.size(); i++)
-    //Throw in a type cast if our kernel takes 2D array notation
-    //like (float(*) [1024])
-  {
-    //protonu--throwing in another hack to stop the caller from passing tex mapped
-    //vars to the kernel.
-    if (arrayVars[i].tex_mapped == true || arrayVars[i].cons_mapped)
-      continue;
-    if (!(arrayVars[i].size_multi_dim.empty())) {
-      //snprintf(dims,120,"(float(*) [%d])%s", arrayVars[i].size_2d,
-      //         const_cast<char*>(arrayVars[i].name.c_str()));
-      
-      SgType* t = arrayVars[i].type;
-      for (int k = arrayVars[i].size_multi_dim.size() - 1; k >= 0; k--) {
-        t = buildArrayType(t,
-                           buildIntVal(arrayVars[i].size_multi_dim[k]));
-      }
-      SgVariableSymbol* temp = body_symtab->find_variable(
-        SgName((char*) arrayVars[i].name.c_str()));
-      if (temp == NULL)
-        temp = parameter_symtab->find_variable(
-          SgName((char*) arrayVars[i].name.c_str()));
-      
-      dim_s = buildCastExp(buildVarRefExp(temp), buildPointerType(t),
-                           SgCastExp::e_C_style_cast);
-      
-      //printf("%d %s\n", i, dims);
-      iml->append_expression(dim_s);
-      
-      SgInitializedName* id = buildInitializedName(
-        (char*) arrayVars[i].original_name.c_str(),
-        buildPointerType(t));
-      kernel_decl->get_parameterList()->append_arg(id);
-      kernel_decl_->get_parameterList()->append_arg(id);
-      id->set_file_info(TRANSFORMATION_FILE_INFO);
-      
-      // DQ (9/8/2007): We now test this, so it has to be set explicitly.
-      id->set_scope(kernel_decl->get_definition());
-      
-      // DQ (9/8/2007): Need to add variable symbol to global scope!
-      //printf ("Fixing up the symbol table in scope = %p = %s for SgInitializedName = %p = %s \n",globalScope,globalScope->class_name().c_str(),var1_init_name,var1_init_name->get_name().str());
-      SgVariableSymbol *var_symbol = new SgVariableSymbol(id);
-      kernel_decl->get_definition()->insert_symbol(id->get_name(),
-                                                   var_symbol);
-      
-      // if(kernel_decl->get_definition()->get_symbol_table()->find((const) id) == NULL)
-      
-    } else {
-      //printf("%d %s\n", i, static_cast<const char*>(arrayVars[i].name));
-      SgVariableSymbol* temp = body_symtab->find_variable(
-        SgName((char*) arrayVars[i].name.c_str()));
-      if (temp == NULL)
-        temp = parameter_symtab->find_variable(
-          SgName((char*) arrayVars[i].name.c_str()));
-      iml->append_expression(buildVarRefExp(temp));
-      SgInitializedName* id = buildInitializedName(
-        (char*) arrayVars[i].original_name.c_str(),
-        buildPointerType(arrayVars[i].type));
-      kernel_decl->get_parameterList()->append_arg(id);
-      kernel_decl_->get_parameterList()->append_arg(id);
-      id->set_file_info(TRANSFORMATION_FILE_INFO);
-      
-      // DQ (9/8/2007): We now test this, so it has to be set explicitly.
-      id->set_scope(kernel_decl->get_definition());
-      
-      // DQ (9/8/2007): Need to add variable symbol to global scope!
-      //printf ("Fixing up the symbol table in scope = %p = %s for SgInitializedName = %p = %s \n"$
-      SgVariableSymbol *var_symbol = new SgVariableSymbol(id);
-      kernel_decl->get_definition()->insert_symbol(id->get_name(),
-                                                   var_symbol);
-      
-    }
-    
-  }
-  if (dim1) {
-    iml->append_expression(buildVarRefExp(dim1));
-    SgInitializedName* id = buildInitializedName(
-      dim1->get_name().getString().c_str(), dim1->get_type());
-    kernel_decl->get_parameterList()->append_arg(id);
-    
-    iml->append_expression(buildVarRefExp(dim2));
-    SgInitializedName* id2 = buildInitializedName(
-      dim2->get_name().getString().c_str(), dim2->get_type());
-    
-    kernel_decl->get_parameterList()->append_arg(id);
-    kernel_decl_->get_parameterList()->append_arg(id);
-  }
-  
-  kernel_decl->get_functionModifier().setCudaKernel();
-  kernel_decl_->get_functionModifier().setCudaKernel();
-  SgCudaKernelCallExp * cuda_call_site = new SgCudaKernelCallExp(
-    TRANSFORMATION_FILE_INFO, buildFunctionRefExp(kernel_decl), iml,config);
-  
-  //  SgStatementPtrList *tnl2 = new SgStatementPtrList;
-  
-  (*replacement_list).push_back(buildExprStatement(cuda_call_site));
-  
-  setup_code = ocg->StmtListAppend(setup_code,
-                                   new CG_roseRepr(replacement_list));
-  
-  //cuda free variables
-  for (int i = 0; i < arrayVars.size(); i++) {
-    if (arrayVars[i].out_data) {
-      
-      SgName name_cuda_copy("cudaMemcpy");
-      SgFunctionDeclaration * decl_cuda_copyout =
-        buildNondefiningFunctionDeclaration(name_cuda_copy,
-                                            buildVoidType(), buildFunctionParameterList(),
-                                            globals);
-      
-      SgExprListExp* args = buildExprListExp();
-      SgExprListExp * cuda_copy_out_args = buildExprListExp();
-      cuda_copy_out_args->append_expression(
-        isSgExpression(buildVarRefExp(arrayVars[i].out_data)));
-      cuda_copy_out_args->append_expression(
-        isSgExpression(buildVarRefExp(arrayVars[i].name)));
-      CG_roseRepr* size_exp = new CG_roseRepr(arrayVars[i].size_expr);
-      cuda_copy_out_args->append_expression(
-        static_cast<CG_roseRepr*>(size_exp->clone())->GetExpression());
-      cuda_copy_out_args->append_expression(
-        buildOpaqueVarRefExp("cudaMemcpyDeviceToHost", globals));
-      
-//                                      cuda_copy_in_args->append_expression(
-//                                              new SgVarRefExp(sourceLocation, )
-//                                      );
-      SgFunctionCallExp * cuda_copy_out_func_call = buildFunctionCallExp(
-        buildFunctionRefExp(decl_cuda_copyout), cuda_copy_out_args);
-      
-      SgFunctionCallExp *the_call = buildFunctionCallExp(
-        buildFunctionRefExp(decl_cuda_copyout), cuda_copy_out_args);
-      
-      SgExprStatement* stmt = buildExprStatement(the_call);
-      
-      SgStatementPtrList* tnl3 = new SgStatementPtrList;
-      
-      (*tnl3).push_back(stmt);
-      
-      //   tree_node_list* tnl = new tree_node_list;
-      //   tnl->append(new tree_instr(the_call));
-      setup_code = ocg->StmtListAppend(setup_code, new CG_roseRepr(tnl3));
-      
-    }
-    if(!arrayVars[i].cons_mapped) {
-      SgName name_cuda_free("cudaFree");
-      SgFunctionDeclaration * decl_cuda_free =
-        buildNondefiningFunctionDeclaration(name_cuda_free,
-                                            buildVoidType(), buildFunctionParameterList(), globals);
-      
-      SgExprListExp* args3 = buildExprListExp();
-      
-      SgVariableSymbol* tmp = body_symtab->find_variable(
-        SgName(arrayVars[i].name.c_str()));
-      if (tmp == NULL)
-        tmp = parameter_symtab->find_variable(
-          SgName(arrayVars[i].name.c_str()));
-      
-      args3->append_expression(buildVarRefExp(tmp));
-      
-      SgFunctionCallExp *the_call2 = buildFunctionCallExp(
-        buildFunctionRefExp(decl_cuda_free), args3);
-      
-      SgExprStatement* stmt2 = buildExprStatement(the_call2);
-      
-      SgStatementPtrList* tnl4 = new SgStatementPtrList;
-      
-      (*tnl4).push_back(stmt2);
-      //(*replacement_list).push_back (stmt2);
-    
-      setup_code = ocg->StmtListAppend(setup_code, new CG_roseRepr(tnl4));
-    }
-  }
-  
-  // ---------------
-  // BUILD THE KERNEL
-  // ---------------
-  
-  //Extract out kernel body
-  SgNode* code = getCode();
-  //Create kernel function body
-  //Add Params
-  std::map<std::string, SgVariableSymbol*> loop_vars;
-  //In-Out arrays
-  for (int i = 0; i < arrayVars.size(); i++) {
-    /*   if(arrayVars[i].in_data)
-         fptr = arrayVars[i].in_data->type()->clone();
-         else
-         fptr = arrayVars[i].out_data->type()->clone();
-    */
-    
-    // fptr = new_proc_syms->install_type(fptr);
-    std::string name =
-      arrayVars[i].in_data ?
-      arrayVars[i].in_data->get_declaration()->get_name().getString() :
-      arrayVars[i].out_data->get_declaration()->get_name().getString();
-    //SgVariableSymbol* sym = new var_sym(fptr, arrayVars[i].in_data ? arrayVars[i].in_data->name() : arrayVars[i].out_data->name());
-    
-    SgVariableSymbol* sym =
-      kernel_decl->get_definition()->get_symbol_table()->find_variable(
-        (const char*) name.c_str());
-    /* SgVariableDeclaration*  defn = buildVariableDeclaration(SgName(name.c_str()), buildFloatType());
-       SgInitializedNamePtrList& variables = defn->get_variables();
-       SgInitializedNamePtrList::const_iterator i = variables.begin();
-       SgInitializedName* initializedName = *i;
-       SgVariableSymbol* sym = new SgVariableSymbol(initializedName);
-       prependStatement(defn, isSgScopeStatement(root_));
-       
-       vs->set_parent(symtab2_);
-       symtab2_->insert(SgName(_s.c_str()), vs);
-    */
-    
-    if (sym != NULL)
-      loop_vars.insert(
-        std::pair<std::string, SgVariableSymbol*>(std::string(name),
-                                                  sym));
-  }
-  
-  //Figure out which loop variables will be our thread and block dimention variables
-  std::vector<SgVariableSymbol *> loop_syms;
-  //Get our indexes
-  std::vector<const char*> indexes; // = get_loop_indexes(code,cu_num_reduce);
-  int threadsPos = 0;
-  
-  CG_outputRepr *body = NULL;
-  SgFunctionDefinition* func_d = func_definition;
-  //std::vector<SgVariableSymbol *> symbols =  recursiveFindRefs(code);
-  
-  SgName name_sync("__syncthreads");
-  SgFunctionDeclaration * decl_sync = buildNondefiningFunctionDeclaration(
-    name_sync, buildVoidType(), buildFunctionParameterList(),
-    globalScope);
-  
-  recursiveFindRefs(code, syms, func_d);
-  
-  //SgFunctionDeclaration* func = Outliner::generateFunction (code, (char*)cu_kernel_name.c_str(), syms, pdSyms, psyms, NULL, globalScope);
-  
-  if (cu_bx > 1 || cu_bx_repr) {
-    indexes.push_back("bx");
-    SgName type_name("blockIdx.x");
-    SgClassSymbol * type_symbol = globalScope->lookup_class_symbol(
-      type_name);
-    SgVariableDeclaration * var_decl = buildVariableDeclaration("bx",
-                                                                buildIntType(), NULL,
-                                                                isSgScopeStatement(kernel_decl->get_definition()->get_body()));
-    SgStatementPtrList *tnl = new SgStatementPtrList;
-    // (*tnl).push_back(isSgStatement(var_decl));
-    appendStatement(var_decl, kernel_decl->get_definition()->get_body());
-    
-    SgVariableSymbol* bx =
-      kernel_decl->get_definition()->get_body()->lookup_variable_symbol(
-        SgName("bx"));
-    SgStatement* assign = isSgStatement(
-      buildAssignStatement(buildVarRefExp(bx),
-                           buildOpaqueVarRefExp("blockIdx.x",
-                                                kernel_decl->get_definition()->get_body())));
-    (*tnl).push_back(assign);
-    // body = ocg->StmtListAppend(body,
-    //                                  new CG_roseRepr(tnl));
-    appendStatement(assign, kernel_decl->get_definition()->get_body());
-    
-  }
-  if (cu_by > 1 || cu_by_repr) {
-    indexes.push_back("by");
-    SgName type_name("blockIdx.y");
-    SgClassSymbol * type_symbol = globalScope->lookup_class_symbol(
-      type_name);
-    SgVariableDeclaration * var_decl = buildVariableDeclaration("by",
-                                                                buildIntType(), NULL,
-                                                                isSgScopeStatement(kernel_decl->get_definition()->get_body()));
-    // SgStatementPtrList *tnl = new SgStatementPtrList;
-    // (*tnl).push_back(isSgStatement(var_decl));
-    appendStatement(var_decl, kernel_decl->get_definition()->get_body());
-    
-    SgVariableSymbol* by =
-      kernel_decl->get_definition()->get_body()->lookup_variable_symbol(
-        SgName("by"));
-    SgStatement* assign = isSgStatement(
-      buildAssignStatement(buildVarRefExp(by),
-                           buildOpaqueVarRefExp("blockIdx.y",
-                                                kernel_decl->get_definition()->get_body())));
-    //(*tnl).push_back(assign);
-    // body = ocg->StmtListAppend(body,
-    //                                  new CG_roseRepr(tnl));
-    appendStatement(assign, kernel_decl->get_definition()->get_body());
-    
-  }
-  if (cu_tx_repr || cu_tx > 1) {
-    threadsPos = indexes.size();
-    indexes.push_back("tx");
-    SgName type_name("threadIdx.x");
-    SgClassSymbol * type_symbol = globalScope->lookup_class_symbol(
-      type_name);
-    SgVariableDeclaration * var_decl = buildVariableDeclaration("tx",
-                                                                buildIntType(), NULL,
-                                                                isSgScopeStatement(kernel_decl->get_definition()->get_body()));
-    //  SgStatementPtrList *tnl = new SgStatementPtrList;
-    //  (*tnl).push_back(isSgStatement(var_decl));
-    appendStatement(var_decl, kernel_decl->get_definition()->get_body());
-    
-    SgVariableSymbol* tx =
-      kernel_decl->get_definition()->get_body()->lookup_variable_symbol(
-        SgName("tx"));
-    SgStatement* assign = isSgStatement(
-      buildAssignStatement(buildVarRefExp(tx),
-                           buildOpaqueVarRefExp("threadIdx.x",
-                                                kernel_decl->get_definition()->get_body())));
-    //(*tnl).push_back(assign);
-    // body = ocg->StmtListAppend(body,
-    //                                  new CG_roseRepr(tnl));
-    appendStatement(assign, kernel_decl->get_definition()->get_body());
-    
-  }
-  if (cu_ty_repr || cu_ty > 1) {
-    indexes.push_back("ty");
-    SgName type_name("threadIdx.y");
-    SgClassSymbol * type_symbol = globalScope->lookup_class_symbol(
-      type_name);
-    SgVariableDeclaration * var_decl = buildVariableDeclaration("ty",
-                                                                buildIntType(), NULL,
-                                                                isSgScopeStatement(kernel_decl->get_definition()->get_body()));
-    appendStatement(var_decl, kernel_decl->get_definition()->get_body());
-    
-    // SgStatementPtrList *tnl = new SgStatementPtrList;
-    // (*tnl).push_back(isSgStatement(var_decl));
-    SgVariableSymbol* ty =
-      kernel_decl->get_definition()->get_body()->lookup_variable_symbol(
-        SgName("ty"));
-    SgStatement* assign = isSgStatement(
-      buildAssignStatement(buildVarRefExp(ty),
-                           buildOpaqueVarRefExp("threadIdx.y",
-                                                kernel_decl->get_definition()->get_body())));
-    // (*tnl).push_back(assign);
-    //  body = ocg->StmtListAppend(body,
-    //                                   new CG_roseRepr(tnl));
-    appendStatement(assign, kernel_decl->get_definition()->get_body());
-    
-  }
-  if (cu_tz_repr || cu_tz > 1) {
-    indexes.push_back("tz");
-    SgName type_name("threadIdx.z");
-    SgClassSymbol * type_symbol = globalScope->lookup_class_symbol(
-      type_name);
-    SgVariableDeclaration * var_decl = buildVariableDeclaration("tz",
-                                                                buildIntType(), NULL,
-                                                                isSgScopeStatement(kernel_decl->get_definition()->get_body()));
-    //   SgStatementPtrList *tnl = new SgStatementPtrList;
-    //   (*tnl).push_back(isSgStatement(var_decl));
-    appendStatement(var_decl, kernel_decl->get_definition()->get_body());
-    
-    SgVariableSymbol* tz =
-      kernel_decl->get_definition()->get_body()->lookup_variable_symbol(
-        "tz");
-    SgStatement* assign = isSgStatement(
-      buildAssignStatement(buildVarRefExp(tz),
-                           buildOpaqueVarRefExp("threadIdx.z",
-                                                kernel_decl->get_definition()->get_body())));
-    //    (*tnl).push_back(assign);
-    //     body = ocg->StmtListAppend(body,
-    //                                      new CG_roseRepr(tnl));
-    appendStatement(assign, kernel_decl->get_definition()->get_body());
-    
-  }
-  
-  std::map<std::string, SgVariableSymbol*> loop_idxs; //map from idx names to their new syms
-  
-  SgNode* swapped_ = swapVarReferences(code, syms,
-                                       kernel_decl->get_definition()->get_symbol_table(),
-                                       kernel_decl->get_definition()->get_body()->get_symbol_table(),
-                                       kernel_decl->get_definition()->get_body());
-  
-  //std::cout << swapped_->unparseToString() << std::endl << std::endl;
-  
-  SgNode *swapped = recursiveFindReplacePreferedIdxs(swapped_,
-                                                     kernel_decl->get_definition()->get_body()->get_symbol_table(),
-                                                     kernel_decl->get_definition()->get_symbol_table(),
-                                                     kernel_decl->get_definition()->get_body(), loop_idxs, globalScope); //in-place swapping
-  //swapped->print();
-  
-  if (!isSgBasicBlock(swapped)) {
-    appendStatement(isSgStatement(swapped),
-                    kernel_decl->get_definition()->get_body());
-    swapped->set_parent(
-      isSgNode(kernel_decl->get_definition()->get_body()));
-  } else {
-    
-    for (SgStatementPtrList::iterator it =
-           isSgBasicBlock(swapped)->get_statements().begin();
-         it != isSgBasicBlock(swapped)->get_statements().end(); it++) {
-      appendStatement(*it, kernel_decl->get_definition()->get_body());
-      (*it)->set_parent(
-        isSgNode(kernel_decl->get_definition()->get_body()));
-      
-    }
-    
-  }
-  
-  for (int i = 0; i < indexes.size(); i++) {
-    std::vector<SgForStatement*> tfs = findCommentedFors(indexes[i],
-                                                         swapped);
-    for (int k = 0; k < tfs.size(); k++) {
-      //printf("replacing %p tfs for index %s\n", tfs[k], indexes[i]);
-      SgNode* newBlock = forReduce(tfs[k], loop_idxs[indexes[i]],
-                                   kernel_decl->get_definition());
-      //newBlock->print();
-      swap_node_for_node_list(tfs[k], newBlock);
-      //printf("AFTER SWAP\n");        newBlock->print();
-    }
-  }
-
-  //--derick replace array refs of texture mapped vars here
-  body = new CG_roseRepr(kernel_decl->get_definition()->get_body());
-  std::vector<IR_ArrayRef*> refs = ir->FindArrayRef(body);
-  texmapArrayRefs(texture, &refs, globals, ir, ocg);
-  // do the same for constant mapped vars
-  consmapArrayRefs(constant_mem, &refs, globals, ir, ocg);
-  
-  return swapped;
-}
-
-//Order taking out dummy variables
-std::vector<std::string> cleanOrder(std::vector<std::string> idxNames) {
-  std::vector<std::string> results;
-  for (int j = 0; j < idxNames.size(); j++) {
-    if (idxNames[j].length() != 0)
-      results.push_back(idxNames[j]);
-  }
-  return results;
-}
-
-//First non-dummy level in ascending order
-int LoopCuda::nonDummyLevel(int stmt, int level) {
-  //level comes in 1-basd and should leave 1-based
-  for (int j = level - 1; j < idxNames[stmt].size(); j++) {
-    if (idxNames[stmt][j].length() != 0) {
-      //printf("found non dummy level of %d with idx: %s when searching for %d\n", j+1, (const char*) idxNames[stmt][j], level);
-      return j + 1;
-    }
-  }
-  char buf[128];
-  sprintf(buf, "%d", level);
-  throw std::runtime_error(
-    std::string("Unable to find a non-dummy level starting from ")
-    + std::string(buf));
-}
-
-int LoopCuda::findCurLevel(int stmt, std::string idx) {
-  for (int j = 0; j < idxNames[stmt].size(); j++) {
-    if (strcmp(idxNames[stmt][j].c_str(), idx.c_str()) == 0)
-      return j + 1;
-  }
-  throw std::runtime_error(
-    std::string("Unable to find index ") + idx
-    + std::string(" in current list of indexes"));
-}
-
-void LoopCuda::permute_cuda(int stmt,
-                            const std::vector<std::string>& curOrder) {
-  //printf("curOrder: ");
-  //printVs(curOrder);
-  //printf("idxNames: ");
-  //printVS(idxNames[stmt]);
-  std::vector<std::string> cIdxNames = cleanOrder(idxNames[stmt]);
-  bool same = true;
-  std::vector<int> pi;
-  for (int i = 0; i < curOrder.size(); i++) {
-    bool found = false;
-    for (int j = 0; j < cIdxNames.size(); j++) {
-      if (strcmp(cIdxNames[j].c_str(), curOrder[i].c_str()) == 0) {
-        pi.push_back(j + 1);
-        found = true;
-        if (j != i)
-          same = false;
-      }
-    }
-    if (!found) {
-      throw std::runtime_error(
-        "One of the indexes in the permute order were not "
-        "found in the current set of indexes.");
-    }
-  }
-  for (int i = curOrder.size(); i < cIdxNames.size(); i++) {
-    pi.push_back(i);
-  }
-  if (same)
-    return;
-  permute(stmt, pi);
-  //Set old indexe names as new
-  for (int i = 0; i < curOrder.size(); i++) {
-    idxNames[stmt][i] = curOrder[i].c_str(); //what about sibling stmts?
-  }
-}
-
-bool LoopCuda::permute(int stmt_num, const std::vector<int> &pi) {
-// check for sanity of parameters
-  if (stmt_num >= stmt.size() || stmt_num < 0)
-    throw std::invalid_argument("invalid statement " + to_string(stmt_num));
-  const int n = stmt[stmt_num].xform.n_out();
-  if (pi.size() > (n - 1) / 2)
-    throw std::invalid_argument(
-      "iteration space dimensionality does not match permute dimensionality");
-  int first_level = 0;
-  int last_level = 0;
-  for (int i = 0; i < pi.size(); i++) {
-    if (pi[i] > (n - 1) / 2 || pi[i] <= 0)
-      throw std::invalid_argument(
-        "invalid loop level " + to_string(pi[i])
-        + " in permuation");
-    
-    if (pi[i] != i + 1) {
-      if (first_level == 0)
-        first_level = i + 1;
-      last_level = i + 1;
-    }
-  }
-  if (first_level == 0)
-    return true;
-  
-  std::vector<int> lex = getLexicalOrder(stmt_num);
-  std::set<int> active = getStatements(lex, 2 * first_level - 2);
-  Loop::permute(active, pi);
-}
-
-void LoopCuda::tile_cuda(int stmt, int level, int outer_level) {
-  tile_cuda(stmt, level, 1, outer_level, "", "", CountedTile);
-}
-void LoopCuda::tile_cuda(int level, int tile_size, int outer_level,
-                         std::string idxName, std::string ctrlName, TilingMethodType method) {
-  tile_cuda(0, level, tile_size, outer_level, idxName, ctrlName, method);
-}
-
-void LoopCuda::tile_cuda(int stmt, int level, int tile_size, int outer_level,
-                         std::string idxName, std::string ctrlName, TilingMethodType method) {
-  //Do regular tile but then update the index and control loop variable
-  //names as well as the idxName to reflect the current state of things.
-  //printf("tile(%d,%d,%d,%d)\n", stmt, level, tile_size, outer_level);
-  //printf("idxNames before: ");
-  //printVS(idxNames[stmt]);
-  
-  tile(stmt, level, tile_size, outer_level, method);
-  
-  if (idxName.size())
-    idxNames[stmt][level - 1] = idxName.c_str();
-  if (tile_size == 1) {
-    //potentially rearrange loops
-    if (outer_level < level) {
-      std::string tmp = idxNames[stmt][level - 1];
-      for (int i = level - 1; i > outer_level - 1; i--) {
-        if (i - 1 >= 0)
-          idxNames[stmt][i] = idxNames[stmt][i - 1];
-      }
-      idxNames[stmt][outer_level - 1] = tmp;
-    }
-    //TODO: even with a tile size of one, you need a insert (of a dummy loop)
-    idxNames[stmt].insert(idxNames[stmt].begin() + (level), "");
-  } else {
-    if (!ctrlName.size())
-      throw std::runtime_error("No ctrl loop name for tile");
-    //insert
-    idxNames[stmt].insert(idxNames[stmt].begin() + (outer_level - 1),
-                          ctrlName.c_str());
-  }
-  
-  //printf("idxNames after: ");
-  //printVS(idxNames[stmt]);
-}
-
-bool LoopCuda::datacopy_privatized_cuda(int stmt_num, int level,
-                                        const std::string &array_name,
-                                        const std::vector<int> &privatized_levels, bool allow_extra_read,
-                                        int fastest_changing_dimension, int padding_stride,
-                                        int padding_alignment, bool cuda_shared) {
-  int old_stmts = stmt.size();
-  //  printf("before datacopy_privatized:\n");
-  printIS();
-  //datacopy_privatized(stmt_num, level, array_name, privatized_levels, allow_extra_read, fastest_changing_dimension, padding_stride, padding_alignment, cuda_shared);
-  if (cuda_shared)
-    datacopy_privatized(stmt_num, level, array_name, privatized_levels,
-                        allow_extra_read, fastest_changing_dimension, padding_stride,
-                        padding_alignment, 1);
-  else
-    datacopy_privatized(stmt_num, level, array_name, privatized_levels,
-                        allow_extra_read, fastest_changing_dimension, padding_stride,
-                        padding_alignment, 0);
-  //  printf("after datacopy_privatized:\n");
-  printIS();
-  
-  //Adjust idxNames to reflect updated state
-  std::vector<std::string> cIdxNames = cleanOrder(idxNames[stmt_num]);
-  int new_stmts = stmt.size();
-  for (int i = old_stmts; i < new_stmts; i++) {
-    //printf("fixing up statement %d\n", i);
-    std::vector<std::string> idxs;
-    
-    //protonu-making sure the vector of nonSplitLevels grows along with
-    //the statement structure
-    stmt_nonSplitLevels.push_back(std::vector<int>());
-    
-    //Indexes up to level will be the same
-    for (int j = 0; j < level - 1; j++)
-      idxs.push_back(cIdxNames[j]);
-    
-    //Expect privatized_levels to match
-    for (int j = 0; j < privatized_levels.size(); j++)
-      idxs.push_back(cIdxNames[privatized_levels[j] - 1]);//level is one-based
-    
-    //all further levels should match order they are in originally
-    if (privatized_levels.size()) {
-      int last_privatized = privatized_levels.back();
-      int top_level = last_privatized
-        + (stmt[i].IS.n_set() - idxs.size());
-      //printf("last privatized_levels: %d top_level: %d\n", last_privatized, top_level);
-      for (int j = last_privatized; j < top_level; j++) {
-        idxs.push_back(cIdxNames[j]);
-        //printf("pushing back: %s\n", (const char*)cIdxNames[j]);
-      }
-    }
-    idxNames.push_back(idxs);
-  }
-}
-
-bool LoopCuda::datacopy_cuda(int stmt_num, int level,
-                             const std::string &array_name, 
-                             const std::vector<std::string> new_idxs,
-                             bool allow_extra_read, int fastest_changing_dimension,
-                             int padding_stride, int padding_alignment, bool cuda_shared) {
-  
-  int old_stmts = stmt.size();
-  //datacopy(stmt_num,level,array_name,allow_extra_read,fastest_changing_dimension,padding_stride,padding_alignment,cuda_shared);
-  //  printf("before datacopy:\n");
-  //  printIS();
-  if (cuda_shared)
-    datacopy(stmt_num, level, array_name, allow_extra_read,
-             fastest_changing_dimension, padding_stride, padding_alignment,
-             1);
-  else
-    datacopy(stmt_num, level, array_name, allow_extra_read,
-             fastest_changing_dimension, padding_stride, padding_alignment,
-             0);
-  //  printf("after datacopy:\n");
-  printIS();
-
-  //Adjust idxNames to reflect updated state
-  std::vector<std::string> cIdxNames = cleanOrder(idxNames[stmt_num]);
-  int new_stmts = stmt.size();
-  for (int i = old_stmts; i < new_stmts; i++) {
-    //printf("fixing up statement %d\n", i);
-    std::vector<std::string> idxs;
-    
-    //protonu-making sure the vector of nonSplitLevels grows along with
-    //the statement structure
-    stmt_nonSplitLevels.push_back(std::vector<int>());
-    
-    //Indexes up to level will be the same
-    for (int j = 0; j < level - 1; j++)
-      idxs.push_back(cIdxNames[j]);
-    
-    //all further levels should get names from new_idxs
-    int top_level = stmt[i].IS.n_set();
-    //printf("top_level: %d level: %d\n", top_level, level);
-    if (new_idxs.size() < top_level - level + 1)
-      throw std::runtime_error(
-        "Need more new index names for new datacopy loop levels");
-    
-    for (int j = level - 1; j < top_level; j++) {
-      idxs.push_back(new_idxs[j - level + 1].c_str());
-      //printf("pushing back: %s\n", new_idxs[j-level+1].c_str());
-    }
-    idxNames.push_back(idxs);
-  }
-}
-
-bool LoopCuda::unroll_cuda(int stmt_num, int level, int unroll_amount) {
-  int old_stmts = stmt.size();
-  //bool b= unroll(stmt_num, , unroll_amount);
-  
-  int dim = 2 * level - 1;
-  std::vector<int> lex = getLexicalOrder(stmt_num);
-  std::set<int> same_loop = getStatements(lex, dim - 1);
-  
-  level = nonDummyLevel(stmt_num, level);
-  //printf("unrolling %d at level %d\n", stmt_num,level);
-  
-  //protonu--using the new version of unroll, which returns
-  //a set of ints instead of a bool. To keep Gabe's logic
-  //I'll check the size of the set, if it's 0 return true
-  //bool b= unroll(stmt_num, level, unroll_amount);
-  std::set<int> b_set = unroll(stmt_num, level, unroll_amount, idxNames);
-  bool b = false;
-  if (b_set.size() == 0)
-    b = true;
-  //end--protonu
-  
-  //Adjust idxNames to reflect updated state
-  std::vector<std::string> cIdxNames = cleanOrder(idxNames[stmt_num]);
-  std::vector<std::string> origSource = idxNames[stmt_num];
-  ;
-  //Drop index names at level
-  if (unroll_amount == 0) {
-    //For all statements that were in this unroll together, drop index name for unrolled level
-    idxNames[stmt_num][level - 1] = "";
-    for (std::set<int>::iterator i = same_loop.begin();
-         i != same_loop.end(); i++) {
-      //printf("in same loop as %d is %d\n", stmt_num, (*i));
-      //idxNames[(*i)][level-1] = "";
-      idxNames[(*i)] = idxNames[stmt_num];
-    }
-  }
-  
-  lex = getLexicalOrder(stmt_num);
-  same_loop = getStatements(lex, dim - 1);
-  
-  bool same_as_source = false;
-  int new_stmts = stmt.size();
-  for (int i = old_stmts; i < new_stmts; i++) {
-    //Check whether we had a sync for the statement we are unrolling, if
-    //so, propogate that to newly created statements so that if they are
-    //in a different loop structure, they will also get a syncthreads
-    int size = syncs.size();
-    for (int j = 0; j < size; j++) {
-      if (syncs[j].first == stmt_num)
-        syncs.push_back(make_pair(i, syncs[j].second));
-    }
-    
-    //protonu-making sure the vector of nonSplitLevels grows along with
-    //the statement structure
-    stmt_nonSplitLevels.push_back(std::vector<int>());
-    
-    //We expect that new statements have a constant for the variable in
-    //stmt[i].IS at level (as seen with print_with_subs), otherwise there
-    //will be a for loop at level and idxNames should match stmt's
-    //idxNames pre-unrolled
-    Relation IS = stmt[i].IS;
-    //Ok, if you know how the hell to get anything out of a Relation, you
-    //should probably be able to do this more elegantly. But for now, I'm
-    //hacking it.
-    std::string s = IS.print_with_subs_to_string();
-    //s looks looks like
-    //{[_t49,8,_t51,_t52,128]: 0 <= _t52 <= 3 && 0 <= _t51 <= 15 && 0 <= _t49 && 64_t49+16_t52+_t51 <= 128}
-    //where level == 5, you see a integer in the input set
-    
-    //If that's not an integer and this is the first new statement, then
-    //we think codegen will have a loop at that level. It's not perfect,
-    //not sure if it can be determined without round-tripping to codegen.
-    int sIdx = 0;
-    int eIdx = 0;
-    for (int j = 0; j < level - 1; j++) {
-      sIdx = s.find(",", sIdx + 1);
-      if (sIdx < 0)
-        break;
-    }
-    if (sIdx > 0) {
-      eIdx = s.find("]");
-      int tmp = s.find(",", sIdx + 1);
-      if (tmp > 0 && tmp < eIdx)
-        eIdx = tmp; //", before ]"
-      if (eIdx > 0) {
-        sIdx++;
-        std::string var = s.substr(sIdx, eIdx - sIdx);
-        //printf("%s\n", s.c_str());
-        //printf("set var for stmt %d at level %d is %s\n", i, level, var.c_str());
-        if (atoi(var.c_str()) == 0 && i == old_stmts) {
-          //TODO:Maybe do see if this new statement would be in the same
-          //group as the original and if it would, don't say
-          //same_as_source
-          if (same_loop.find(i) == same_loop.end()) {
-            printf(
-              "stmt %d level %d, newly created unroll statement should have same level indexes as source\n",
-              i, level);
-            same_as_source = true;
-          }
-        }
-      }
-    }
-    
-    //printf("fixing up statement %d n_set %d with %d levels\n", i, stmt[i].IS.n_set(), level-1);
-    if (same_as_source)
-      idxNames.push_back(origSource);
-    else
-      idxNames.push_back(idxNames[stmt_num]);
-  }
-  
-  return b;
-}
-
-void LoopCuda::copy_to_texture(const char *array_name) {
-  //protonu--placeholder for now
-  //set the bool for using cuda memory as true
-  //in a vector of strings, put the names of arrays to tex mapped
-  if (!texture)
-    texture = new texture_memory_mapping(true, array_name);
-  else
-    texture->add(array_name);
-  
-}
-
-//void LoopCuda::copy_to_texture_2d(const char *array_name, int width, int height) {
-//  if (!texture)
-//    texture = new texture_memory_mapping(true, array_name, width, height);
-//  else
-//    texture->add(array_name, width, height);
-//}
-
-void LoopCuda::copy_to_constant(const char *array_name) {
-  if(!constant_mem)
-    constant_mem = new constant_memory_mapping(true, array_name);
-  else
-    constant_mem->add(array_name);
-}
-
-//protonu--moving this from Loop
-SgNode* LoopCuda::codegen() {
-  if (code_gen_flags & GenCudaizeV2)
-    return cudaize_codegen_v2();
-  //Do other flagged codegen methods, return plain vanilla generated code
-  return getCode();
-}
-
-//These three are in Omega code_gen.cc and are used as a massive hack to
-//get out some info from MMGenerateCode. Yea for nasty side-effects.
-namespace omega {
-  extern int checkLoopLevel;
-  extern int stmtForLoopCheck;
-  extern int upperBoundForLevel;
-  extern int lowerBoundForLevel;
-}
-
-CG_outputRepr* LoopCuda::extractCudaUB(int stmt_num, int level,
-                                       int &outUpperBound, int &outLowerBound) {
-  // check for sanity of parameters
-  const int m = stmt.size();
-  if (stmt_num >= m || stmt_num < 0)
-    throw std::invalid_argument("invalid statement " + to_string(stmt_num));
-  const int n = stmt[stmt_num].xform.n_out();
-  if (level > (n - 1) / 2 || level <= 0)
-    throw std::invalid_argument("invalid loop level " + to_string(level));
-  
-  int dim = 2 * level - 1;
-  
-  std::vector<int> lex = getLexicalOrder(stmt_num);
-  std::set<int> same_loop = getStatements(lex, dim - 1);
-  
-  // extract the intersection of the iteration space to be considered
-  Relation hull;
-  {
-    hull = Relation::True(n);
-    for (std::set<int>::iterator i = same_loop.begin();
-         i != same_loop.end(); i++) {
-      Relation r = getNewIS(*i);
-      for (int j = dim + 2; j <= r.n_set(); j++)
-        r = Project(r, r.set_var(j));
-      hull = Intersection(hull, r);
-      hull.simplify(2, 4);
-    }
-    
-    for (int i = 2; i <= dim + 1; i += 2) {
-      //std::string name = std::string("_t") + to_string(t_counter++);
-      std::string name = std::string("_t")
-        + to_string(tmp_loop_var_name_counter++);
-      hull.name_set_var(i, name);
-    }
-    hull.setup_names();
-  }
-  
-  // extract the exact loop bound of the dimension to be unrolled
-  if (is_single_iteration(hull, dim)) {
-    throw std::runtime_error(
-      "No loop availabe at level to extract upper bound.");
-  }
-  Relation bound = get_loop_bound(hull, dim);
-  if (!bound.has_single_conjunct() || !bound.is_satisfiable()
-      || bound.is_tautology())
-    throw loop_error(
-      "loop error: unable to extract loop bound for cudaize");
-  
-  // extract the loop stride
-  EQ_Handle stride_eq;
-  /*int stride = 1;
-    {
-    bool simple_stride = true;
-    int strides = countStrides(bound.query_DNF()->single_conjunct(),
-    bound.set_var(dim + 1), stride_eq, simple_stride);
-    if (strides > 1)
-    throw loop_error("loop error: too many strides");
-    else if (strides == 1) {
-    int sign = stride_eq.get_coef(bound.set_var(dim + 1));
-    //      assert(sign == 1 || sign == -1);
-    Constr_Vars_Iter it(stride_eq, true);
-    stride = abs((*it).coef / sign);
-    }
-    }
-  */
-  int stride = 1;
-  {
-    
-    coef_t stride;
-    std::pair<EQ_Handle, Variable_ID> result = find_simplest_stride(bound,
-                                                                    bound.set_var(dim + 1));
-    if (result.second == NULL)
-      stride = 1;
-    else
-      stride = abs(result.first.get_coef(result.second))
-        / gcd(abs(result.first.get_coef(result.second)),
-              abs(result.first.get_coef(bound.set_var(dim + 1))));
-    
-    if (stride > 1)
-      throw loop_error("loop error: too many strides");
-    /*else if (stride == 1) {
-      int sign = result.first.get_coef(bound.set_var(dim+1));
-      assert(sign == 1 || sign == -1);
-      } */
-  }
-  
-  if (stride != 1) {
-    char buf[1024];
-    sprintf(buf, "Cudaize: Loop at level %d has non-one stride of %d",
-            level, stride);
-    throw std::runtime_error(buf);
-  }
-  
-  //Use code generation system to build tell us our bound information. We
-  //need a hard upper bound a 0 lower bound.
-  
-  checkLoopLevel = level * 2;
-  stmtForLoopCheck = stmt_num;
-  upperBoundForLevel = -1;
-  lowerBoundForLevel = -1;
-  printCode(1, false);
-  checkLoopLevel = 0;
-  
-  outUpperBound = upperBoundForLevel;
-  outLowerBound = lowerBoundForLevel;
-  
-  if (outUpperBound == -1) {
-    
-    CG_result* temp = last_compute_cgr_;
-    
-    while (temp) {
-      CG_loop * loop;
-      if (loop = dynamic_cast<CG_loop*>(temp)) {
-        if (loop->level_ == 2 * level) {
-          Relation bound = copy(loop->bounds_);
-          Variable_ID v = bound.set_var(2 * level);
-          for (GEQ_Iterator e(
-                 const_cast<Relation &>(bound).single_conjunct()->GEQs());
-               e; e++) {
-            if ((*e).get_coef(v) < 0
-                && (*e).is_const_except_for_global(v))
-              return output_upper_bound_repr(ir->builder(), *e, v,
-                                             bound,
-                                             std::vector<std::pair<CG_outputRepr *, int> >(
-                                               bound.n_set(),
-                                               std::make_pair(
-                                                 static_cast<CG_outputRepr *>(NULL),
-                                                 0)));
-          }
-        }
-        if (loop->level_ > 2 * level)
-          break;
-        else
-          temp = loop->body_;
-      } else
-        break;
-    }
-  }
-  
-  return NULL;
-}
-
-void LoopCuda::printCode(int effort, bool actuallyPrint) const {
-  const int m = stmt.size();
-  if (m == 0)
-    return;
-  const int n = stmt[0].xform.n_out();
-  
-  /*or (int i = 0; i < m; i++) {
-    IS[i + 1] = stmt[i].IS;
-    xform[i + 1] = stmt[i].xform;
-    
-    //nonSplitLevels[i+1] = stmt[i].nonSplitLevels;
-    }
-  */
-  
-  // invalidate saved codegen computation
-  if (last_compute_cgr_ != NULL) {
-    delete last_compute_cgr_;
-    last_compute_cgr_ = NULL;
-  }
-  
-  if (last_compute_cg_ != NULL) {
-    delete last_compute_cg_;
-    last_compute_cg_ = NULL;
-  }
-  
-  //Relation known = Extend_Set(copy(this->known), n - this->known.n_set());
-  /*CG_stringBuilder *ocg = new CG_stringBuilder();
-    Tuple<CG_outputRepr *> nameInfo;
-    for (int i = 1; i <= m; i++)
-    nameInfo.append(new CG_stringRepr("s" + to_string(i)));
-  */
-  
-  // -- replacing MMGenerateCode
-  // -- formally CG_outputRepr* repr = MMGenerateCode(ocg, xform, IS, nameInfo, known, nonSplitLevels, syncs, idxTupleNames, effort);
-  // -- in the future, these if statements need to be cleaned up.
-  // -- something like check_lastComputeCG might be a decent protected member function
-  // -- and/or something that returns a std::vector<CG_outputRepr*> that also checks last_compute_cg_
-  //if (last_compute_cg_ == NULL) {
-  std::vector<Relation> IS(m);
-  std::vector<Relation> xforms(m);
-  std::vector<std::vector<int> > nonSplitLevels(m);
-  
-  /*    std::vector < std::vector <std::string> > idxTupleNames;
-        if (useIdxNames) {
-        for (int i = 0; i < idxNames.size(); i++) {
-        Tuple<std::string> idxs;
-        for (int j = 0; j < idxNames[i].size(); j++)
-        idxs.append(idxNames[i][j]);
-        idxTupleNames.append(idxs);
-        }
-        }
-  */
-  for (int i = 0; i < m; i++) {
-    IS[i] = stmt[i].IS;
-    xforms[i] = stmt[i].xform;
-    nonSplitLevels[i] = stmt_nonSplitLevels[i];
-  }
-  Relation known = Extend_Set(copy(this->known), n - this->known.n_set());
-  
-  last_compute_cg_ = new CodeGen(xforms, IS, known, nonSplitLevels, idxNames,
-                                 syncs);
-
-  delete last_compute_cgr_;  // this was just done  above? 
-  last_compute_cgr_ = NULL;
-  //}
-  
-  if (last_compute_cgr_ == NULL || last_compute_effort_ != effort) {
-    delete last_compute_cgr_;
-    last_compute_cgr_ = last_compute_cg_->buildAST(effort);
-    last_compute_effort_ = effort;
-  }
-  
-  //std::vector<CG_outputRepr *> stmts(m);
-  //for (int i = 0; i < m; i++)
-  //    stmts[i] = stmt[i].code;
-  //CG_outputRepr* repr = last_compute_cgr_->printRepr(ocg, stmts);
-  // -- end replacing MMGenerateCode
-  std::string repr = last_compute_cgr_->printString();
-  
-  if (actuallyPrint)
-    std::cout << repr << std::endl;
-  //std::cout << static_cast<CG_stringRepr*>(repr)->GetString();
-  /*
-    for (int i = 1; i <= m; i++)
-    delete nameInfo[i];
-  */
-  
-  //delete ocg;
-}
-
-void LoopCuda::printRuntimeInfo() const {
-  for (int i = 0; i < stmt.size(); i++) {
-    Relation IS = stmt[i].IS;
-    Relation xform = stmt[i].xform;
-    printf("stmt[%d]\n", i);
-    printf("IS\n");
-    IS.print_with_subs();
-    
-    printf("xform[%d]\n", i);
-    xform.print_with_subs();
-    
-  }
-}
-
-void LoopCuda::printIndexes() const {
-  for (int i = 0; i < stmt.size(); i++) {
-    printf("stmt %d nset %d ", i, stmt[i].IS.n_set());
-    
-    for (int j = 0; j < idxNames[i].size(); j++) {
-      if (j > 0)
-        printf(",");
-      printf("%s", idxNames[i][j].c_str());
-    }
-    printf("\n");
-  }
-}
-
-SgNode* LoopCuda::getCode(int effort) const {
-  const int m = stmt.size();
-  if (m == 0)
-    return new SgNode;
-  const int n = stmt[0].xform.n_out();
-  /*
-    Tuple<CG_outputRepr *> ni(m);
-    Tuple < Relation > IS(m);
-    Tuple < Relation > xform(m);
-    vector < vector <int> > nonSplitLevels(m);
-    for (int i = 0; i < m; i++) {
-    ni[i + 1] = stmt[i].code;
-    IS[i + 1] = stmt[i].IS;
-    xform[i + 1] = stmt[i].xform;
-    nonSplitLevels[i + 1] = stmt_nonSplitLevels[i];
-    
-    //nonSplitLevels[i+1] = stmt[i].nonSplitLevels;
-    }
-  */
-  //Relation known = Extend_Set(copy(this->known), n - this->known.n_set());
-//#ifdef DEBUG
-//#endif
-  //std::cout << GetString(MMGenerateCode(new CG_stringBuilder(), xform, IS, ni, known,
-  //                nonSplitLevels, syncs, idxTupleNames, effort));
-  if (last_compute_cgr_ != NULL) {
-    delete last_compute_cgr_;
-    last_compute_cgr_ = NULL;
-  }
-  
-  if (last_compute_cg_ != NULL) {
-    delete last_compute_cg_;
-    last_compute_cg_ = NULL;
-  }
-  
-  CG_outputBuilder *ocg = ir->builder();
-  // -- replacing MMGenerateCode
-  // -- formally CG_outputRepr* repr = MMGenerateCode(ocg, xform, IS, nameInfo, known, nonSplitLevels, syncs, idxTupleNames, effort);
-  // -- in the future, these if statements need to be cleaned up.
-  // -- something like check_lastComputeCG might be a decent protected member function
-  // -- and/or something that returns a std::vector<CG_outputRepr*> that also checks last_compute_cg_
-  //if (last_compute_cg_ == NULL) {
-  std::vector<Relation> IS(m);
-  std::vector<Relation> xforms(m);
-  std::vector<std::vector<int> > nonSplitLevels(m);
-  for (int i = 0; i < m; i++) {
-    IS[i] = stmt[i].IS;
-    xforms[i] = stmt[i].xform;
-    nonSplitLevels[i] = stmt_nonSplitLevels[i];
-  }
-  
-  /*std::vector < std::vector<std::string> > idxTupleNames;
-    if (useIdxNames) {
-    for (int i = 0; i < idxNames.size(); i++) {
-    std::vector<std::string> idxs;
-    for (int j = 0; j < idxNames[i].size(); j++)
-    idxs.push_back(idxNames[i][j]);
-    idxTupleNames.push_back(idxs);
-    }
-    }
-  */
-  Relation known = Extend_Set(copy(this->known), n - this->known.n_set());
-  
-  last_compute_cg_ = new CodeGen(xforms, IS, known, nonSplitLevels, idxNames,
-                                 syncs);
-  delete last_compute_cgr_;
-  last_compute_cgr_ = NULL;
-  //}
-  
-  if (last_compute_cgr_ == NULL || last_compute_effort_ != effort) {
-    delete last_compute_cgr_;
-    last_compute_cgr_ = last_compute_cg_->buildAST(effort);
-    last_compute_effort_ = effort;
-  }
-  
-  std::vector<CG_outputRepr *> stmts(m);
-  for (int i = 0; i < m; i++)
-    stmts[i] = stmt[i].code;
-  CG_outputRepr* repr = last_compute_cgr_->printRepr(ocg, stmts);
-  // -- end replacing MMGenerateCode
-  
-  //CG_outputRepr *overflow_initialization = ocg->CreateStmtList();
-  CG_outputRepr *overflow_initialization = ocg->StmtListAppend(NULL, NULL);
-  for (std::map<int, std::vector<Free_Var_Decl *> >::const_iterator i =
-         overflow.begin(); i != overflow.end(); i++)
-    for (std::vector<Free_Var_Decl *>::const_iterator j = i->second.begin();
-         j != i->second.end(); j++)
-      //overflow_initialization = ocg->StmtListAppend(overflow_initialization, ocg->CreateStmtList(ocg->CreateAssignment(0, ocg->CreateIdent((*j)->base_name()), ocg->CreateInt(0))));
-      overflow_initialization = ocg->StmtListAppend(
-        overflow_initialization,
-        ocg->StmtListAppend(
-          ocg->CreateAssignment(0,
-                                ocg->CreateIdent((*j)->base_name()),
-                                ocg->CreateInt(0)), NULL));
-  
-  repr = ocg->StmtListAppend(overflow_initialization, repr);
-  SgNode *tnl = static_cast<CG_roseRepr *>(repr)->GetCode();
-  SgStatementPtrList *list = static_cast<CG_roseRepr *>(repr)->GetList();
-  
-  if (tnl != NULL)
-    return tnl;
-  else if (tnl == NULL && list != NULL) {
-    SgBasicBlock* bb2 = buildBasicBlock();
-    
-    for (SgStatementPtrList::iterator it = (*list).begin();
-         it != (*list).end(); it++)
-      bb2->append_statement(*it);
-    
-    tnl = isSgNode(bb2);
-  } else
-    throw loop_error("codegen failed");
-  
-  delete repr;
-  /*
-    for (int i = 1; i <= m; i++)
-    delete ni[i];
-  */
-  return tnl;
-  
-}
-
-//protonu--adding constructors for the new derived class
-LoopCuda::LoopCuda() :
-  Loop(), code_gen_flags(GenInit) {
-}
-
-LoopCuda::LoopCuda(IR_Control *irc, int loop_num) :
-  Loop(irc) {
-  setup_code = NULL;
-  teardown_code = NULL;
-  code_gen_flags = 0;
-  cu_bx = cu_by = cu_tx = cu_ty = cu_tz = 1;
-  cu_bx_repr = NULL;
-  cu_tx_repr = NULL;
-  cu_by_repr = NULL;
-  cu_ty_repr = NULL;
-  cu_tz_repr = NULL;
-  
-  cu_num_reduce = 0;
-  cu_mode = GlobalMem;
-  texture = NULL;
-  constant_mem = NULL;
-  
-  int m = stmt.size();
-  //printf("\n the size of stmt(initially) is: %d\n", stmt.size());
-  for (int i = 0; i < m; i++)
-    stmt_nonSplitLevels.push_back(std::vector<int>());
-  
-  globals = ((IR_cudaroseCode *) ir)->gsym_;
-  globalScope = ((IR_cudaroseCode *) ir)->first_scope;
-  parameter_symtab = ((IR_cudaroseCode *) ir)->parameter;
-  body_symtab = ((IR_cudaroseCode *) ir)->body;
-  func_body = ((IR_cudaroseCode *) ir)->defn;
-  func_definition = ((IR_cudaroseCode *) ir)->func_defn;
-  std::vector<SgForStatement *> tf = ((IR_cudaroseCode *) ir)->get_loops();
-  
-  symtab = tf[loop_num]->get_symbol_table();
-  
-  std::vector<SgForStatement *> deepest = find_deepest_loops(
-    isSgNode(tf[loop_num]));
-  
-  for (int i = 0; i < deepest.size(); i++) {
-    SgVariableSymbol* vs;
-    SgForInitStatement* list = deepest[i]->get_for_init_stmt();
-    SgStatementPtrList& initStatements = list->get_init_stmt();
-    SgStatementPtrList::const_iterator j = initStatements.begin();
-    if (SgExprStatement *expr = isSgExprStatement(*j))
-      if (SgAssignOp* op = isSgAssignOp(expr->get_expression()))
-        if (SgVarRefExp* var_ref = isSgVarRefExp(op->get_lhs_operand()))
-          vs = var_ref->get_symbol();
-    
-    index.push_back(vs->get_name().getString().c_str()); //reflects original code index names
-  }
-  
-  for (int i = 0; i < stmt.size(); i++)
-    idxNames.push_back(index); //refects prefered index names (used as handles in cudaize v2)
-  useIdxNames = false;
-  
-}
-
-void LoopCuda::printIS() {
- if (!cudaDebug) return;
-  int k = stmt.size();
-  for (int i = 0; i < k; i++) {
-    printf(" printing statement:%d\n", i);
-    stmt[i].IS.print();
-  }
-}
-