diff options
author | Tuowen Zhao <ztuowen@gmail.com> | 2016-10-01 11:29:57 -0600 |
---|---|---|
committer | Tuowen Zhao <ztuowen@gmail.com> | 2016-10-01 11:29:57 -0600 |
commit | f7c9c6e95f5a7333db5cef9d92c6c3438adbd6ee (patch) | |
tree | fb1f71863457aae8be68724d18d752437d6c19d1 /src | |
parent | 2028f3ddc680b3af6476ace8840d6bcc85b88d73 (diff) | |
download | chill-f7c9c6e95f5a7333db5cef9d92c6c3438adbd6ee.tar.gz chill-f7c9c6e95f5a7333db5cef9d92c6c3438adbd6ee.tar.bz2 chill-f7c9c6e95f5a7333db5cef9d92c6c3438adbd6ee.zip |
fixes
Diffstat (limited to 'src')
-rw-r--r-- | src/ast/node.cpp | 76 | ||||
-rw-r--r-- | src/chillASTs.cc | 268 | ||||
-rw-r--r-- | src/printer/cfamily.cpp | 34 | ||||
-rw-r--r-- | src/printer/dump.cpp | 12 |
4 files changed, 127 insertions, 263 deletions
diff --git a/src/ast/node.cpp b/src/ast/node.cpp index 3d28d97..dc0a913 100644 --- a/src/ast/node.cpp +++ b/src/ast/node.cpp @@ -9,11 +9,59 @@ using namespace std; +int chillAST_Node::chill_scalar_counter = 0; +int chillAST_Node::chill_array_counter = 1; + +const char *ChillAST_Node_Names[] = { + "Unknown AST node type", + "SourceFile", + "TypedefDecl", + "VarDecl", + // "ParmVarDecl", not used any more + "FunctionDecl", + "RecordDecl", + "MacroDefinition", + "CompoundStmt", + "ForStmt", + "TernaryOperator", + "BinaryOperator", + "UnaryOperator", + "ArraySubscriptExpr", + "MemberExpr", + "DeclRefExpr", + "IntegerLiteral", + "FloatingLiteral", + "ImplicitCastExpr", // not sure we need this + "ReturnStmt", + "CallExpr", + "DeclStmt", + "ParenExpr", + "CStyleCastExpr", + "CStyleAddressOf", + "IfStmt", + "SizeOf", + "Malloc", + "Free", + "NoOp", +// CUDA specific + "CudaMalloc", + "CudaFree", + "CudaMemcpy", + "CudaKernelCall", + "CudaSyncthreads", + "fake1", + "fake2", + "fake3" +}; + +const char* chillAST_Node::getTypeString() { + return ChillAST_Node_Names[getType()]; +}; + void chillAST_Node::fixChildInfo() {} void chillAST_Node::addChild(chillAST_Node *c) { c->parent = this; - // check to see if it's already there for (int i = 0; i < children.size(); i++) { if (c == children[i]) { CHILL_ERROR("addchild ALREADY THERE\n"); @@ -47,27 +95,13 @@ int chillAST_Node::findChild(chillAST_Node *c) { } void chillAST_Node::replaceChild(chillAST_Node *old, chillAST_Node *newchild) { - CHILL_DEBUG_PRINT("(%s) forgot to implement replaceChild() ... using generic\n", getTypeString()); - CHILL_DEBUG_PRINT("%d children\n", children.size()); - for (int i = 0; i < children.size(); i++) { - if (children[i] == old) { - children[i] = newchild; - newchild->setParent(this); + for (int i = 0; i < getNumChildren(); i++) { + if (getChild(i) == old) { + setChild(i,newchild); return; } } - CHILL_ERROR("%s %p generic replaceChild called with oldchild that was not a child\n", - getTypeString(), this); - CHILL_DEBUG_BEGIN - fprintf(stderr, "printing\n"); - print(); - fprintf(stderr, "\nchild: "); - if (!old) fprintf(stderr, "oldchild NULL!\n"); - old->print(); - fprintf(stderr, "\nnew: "); - newchild->print(); - fprintf(stderr, "\n"); - CHILL_DEBUG_END + CHILL_ERROR("Replace Child Falure, oldchild not a child\n"); exit(-1); }; @@ -78,9 +112,7 @@ void chillAST_Node::loseLoopWithLoopVar(char *var) { } } -//! recursive walk parent links, looking for loops, and grabbing the declRefExpr in the loop init and cond. -void chillAST_Node::gatherLoopIndeces( - std::vector<chillAST_VarDecl *> &indeces) { +void chillAST_Node::gatherLoopIndeces(std::vector<chillAST_VarDecl *> &indeces) { // you can quit when you get to certain nodes CHILL_DEBUG_PRINT("%s::gatherLoopIndeces()\n", getTypeString()); diff --git a/src/chillASTs.cc b/src/chillASTs.cc index 70c37e8..bea2360 100644 --- a/src/chillASTs.cc +++ b/src/chillASTs.cc @@ -9,52 +9,6 @@ using namespace std; -int chillAST_Node::chill_scalar_counter = 0; -int chillAST_Node::chill_array_counter = 1; - - -const char *ChillAST_Node_Names[] = { - "Unknown AST node type", - "SourceFile", - "TypedefDecl", - "VarDecl", - // "ParmVarDecl", not used any more - "FunctionDecl", - "RecordDecl", - "MacroDefinition", - "CompoundStmt", - "ForStmt", - "TernaryOperator", - "BinaryOperator", - "UnaryOperator", - "ArraySubscriptExpr", - "MemberExpr", - "DeclRefExpr", - "IntegerLiteral", - "FloatingLiteral", - "ImplicitCastExpr", // not sure we need this - "ReturnStmt", - "CallExpr", - "DeclStmt", - "ParenExpr", - "CStyleCastExpr", - "CStyleAddressOf", - "IfStmt", - "SizeOf", - "Malloc", - "Free", - "NoOp", -// CUDA specific - "CudaMalloc", - "CudaFree", - "CudaMemcpy", - "CudaKernelCall", - "CudaSyncthreads", - "fake1", - "fake2", - "fake3" -}; - bool streq(const char *a, const char *b) { return !strcmp(a, b); } @@ -625,15 +579,6 @@ void chillAST_MacroDefinition::setBody(chillAST_Node *bod) { bod->setParent(this); // well, ... } -void chillAST_MacroDefinition::insertChild(int i, chillAST_Node *node) { - body->insertChild(i, node); -} - -void chillAST_MacroDefinition::addChild(chillAST_Node *node) { - body->addChild(node); - node->setParent(this); // this, or body?? -} - chillAST_ForStmt::chillAST_ForStmt() { children.push_back(NULL); // init children.push_back(NULL); // cond @@ -646,7 +591,7 @@ chillAST_ForStmt::chillAST_ForStmt() { chillAST_ForStmt::chillAST_ForStmt(chillAST_Node *ini, chillAST_Node *con, chillAST_Node *inc, chillAST_Node *bod) - :chillAST_ForStmt() { + : chillAST_ForStmt() { setInit(ini); setCond(con); setInc(inc); @@ -763,7 +708,8 @@ void chillAST_ForStmt::printControl(int in, FILE *fp) { } chillAST_Node *chillAST_ForStmt::clone() { - chillAST_ForStmt *fs = new chillAST_ForStmt(getInit()->clone(), getCond()->clone(), getInc()->clone(), getBody()->clone()); + chillAST_ForStmt *fs = new chillAST_ForStmt(getInit()->clone(), getCond()->clone(), getInc()->clone(), + getBody()->clone()); fs->isFromSourceFile = isFromSourceFile; if (filename) fs->filename = strdup(filename); fs->setParent(getParent()); @@ -1113,7 +1059,7 @@ chillAST_BinaryOperator::chillAST_BinaryOperator() { } chillAST_BinaryOperator::chillAST_BinaryOperator(chillAST_Node *l, const char *oper, chillAST_Node *r) - :chillAST_BinaryOperator() { + : chillAST_BinaryOperator() { //fprintf(stderr, "chillAST_BinaryOperator::chillAST_BinaryOperator( l %p %s r %p, parent %p) this %p\n", l, oper, r, par, this); CHILL_DEBUG_PRINT("( l %s r )\n", oper); @@ -1259,7 +1205,7 @@ chillAST_TernaryOperator::chillAST_TernaryOperator() { } chillAST_TernaryOperator::chillAST_TernaryOperator(const char *oper, chillAST_Node *c, chillAST_Node *l, - chillAST_Node *r):chillAST_TernaryOperator() { + chillAST_Node *r) : chillAST_TernaryOperator() { if (op) op = strdup(oper); setCond(c); setLHS(l); @@ -2074,7 +2020,8 @@ chillAST_UnaryOperator::chillAST_UnaryOperator() { children.push_back(NULL); } -chillAST_UnaryOperator::chillAST_UnaryOperator(const char *oper, bool pre, chillAST_Node *sub):chillAST_UnaryOperator() { +chillAST_UnaryOperator::chillAST_UnaryOperator(const char *oper, bool pre, chillAST_Node *sub) + : chillAST_UnaryOperator() { op = strdup(oper); prefix = pre; setSubExpr(sub); @@ -2144,7 +2091,7 @@ chillAST_ImplicitCastExpr::chillAST_ImplicitCastExpr() { children.push_back(NULL); } -chillAST_ImplicitCastExpr::chillAST_ImplicitCastExpr(chillAST_Node *sub):chillAST_ImplicitCastExpr() { +chillAST_ImplicitCastExpr::chillAST_ImplicitCastExpr(chillAST_Node *sub) : chillAST_ImplicitCastExpr() { setSubExpr(sub); } @@ -2159,7 +2106,7 @@ chillAST_CStyleCastExpr::chillAST_CStyleCastExpr() { children.push_back(NULL); } -chillAST_CStyleCastExpr::chillAST_CStyleCastExpr(const char *to, chillAST_Node *sub):chillAST_CStyleCastExpr() { +chillAST_CStyleCastExpr::chillAST_CStyleCastExpr(const char *to, chillAST_Node *sub) : chillAST_CStyleCastExpr() { towhat = strdup(to); setSubExpr(sub); } @@ -2175,7 +2122,8 @@ class chillAST_Node *chillAST_CStyleCastExpr::clone() { chillAST_CStyleAddressOf::chillAST_CStyleAddressOf() { children.push_back(NULL); } -chillAST_CStyleAddressOf::chillAST_CStyleAddressOf(chillAST_Node *sub):chillAST_CStyleAddressOf() { + +chillAST_CStyleAddressOf::chillAST_CStyleAddressOf(chillAST_Node *sub) : chillAST_CStyleAddressOf() { setSubExpr(sub); } @@ -2191,7 +2139,7 @@ chillAST_Malloc::chillAST_Malloc() { children.push_back(NULL); } -chillAST_Malloc::chillAST_Malloc(chillAST_Node *numthings):chillAST_Malloc() { +chillAST_Malloc::chillAST_Malloc(chillAST_Node *numthings) : chillAST_Malloc() { setSize(numthings); isFromSourceFile = true; // default filename = NULL; @@ -2210,7 +2158,7 @@ chillAST_CudaMalloc::chillAST_CudaMalloc() { children.push_back(NULL); } -chillAST_CudaMalloc::chillAST_CudaMalloc(chillAST_Node *devmemptr, chillAST_Node *size):chillAST_CudaMalloc() { +chillAST_CudaMalloc::chillAST_CudaMalloc(chillAST_Node *devmemptr, chillAST_Node *size) : chillAST_CudaMalloc() { setDevPtr(devmemptr); setSize(size); }; @@ -2224,71 +2172,45 @@ class chillAST_Node *chillAST_CudaMalloc::clone() { } void chillAST_CudaMalloc::gatherArrayRefs(std::vector<chillAST_ArraySubscriptExpr *> &refs, bool w) { - chillAST_Node::gatherArrayRefs(refs,false); + chillAST_Node::gatherArrayRefs(refs, false); } void chillAST_CudaMalloc::gatherScalarRefs(std::vector<chillAST_DeclRefExpr *> &refs, bool writtento) { - chillAST_Node::gatherScalarRefs(refs,false); + chillAST_Node::gatherScalarRefs(refs, false); } -chillAST_CudaFree::chillAST_CudaFree(chillAST_VarDecl *var) { - variable = var; -}; - -class chillAST_Node *chillAST_CudaFree::constantFold() { - return this; +chillAST_CudaFree::chillAST_CudaFree() { + children.push_back(NULL); } +chillAST_CudaFree::chillAST_CudaFree(chillAST_Node *var) : chillAST_CudaFree() { + setParent(var); +}; + class chillAST_Node *chillAST_CudaFree::clone() { - chillAST_CudaFree *CF = new chillAST_CudaFree(variable); + chillAST_CudaFree *CF = new chillAST_CudaFree(getPointer()->clone()); CF->setParent(getParent()); CF->isFromSourceFile = isFromSourceFile; if (filename) CF->filename = strdup(filename); return CF; } -void chillAST_CudaFree::gatherArrayRefs(std::vector<chillAST_ArraySubscriptExpr *> &refs, bool w) {} - -void chillAST_CudaFree::gatherScalarRefs(std::vector<chillAST_DeclRefExpr *> &refs, bool writtento) {} - -void chillAST_CudaFree::gatherVarDecls(vector<chillAST_VarDecl *> &decls) { - variable->gatherVarDecls(decls); -} - - -void chillAST_CudaFree::gatherScalarVarDecls(vector<chillAST_VarDecl *> &decls) { - variable->gatherScalarVarDecls(decls); +chillAST_CudaMemcpy::chillAST_CudaMemcpy() { + addChild(NULL); + addChild(NULL); + addChild(NULL); } - -void chillAST_CudaFree::gatherArrayVarDecls(vector<chillAST_VarDecl *> &decls) { - variable->gatherArrayVarDecls(decls); -} - - -void chillAST_CudaFree::gatherVarUsage(vector<chillAST_VarDecl *> &decls) { - variable->gatherVarUsage(decls); -} - - -chillAST_CudaMemcpy::chillAST_CudaMemcpy(chillAST_VarDecl *d, chillAST_VarDecl *s, chillAST_Node *siz, char *kind) { - dest = d; - src = s; - //fprintf(stderr, "chillAST_CudaMemcpy::chillAST_CudaMemcpy( dest %s, src %s, ...)\n", d->varname, s->varname ); - size = siz; +chillAST_CudaMemcpy::chillAST_CudaMemcpy(chillAST_Node *d, chillAST_Node *s, chillAST_Node *siz, char *kind) { + setDest(d); + setSrc(s); + setSize(siz); cudaMemcpyKind = kind; }; -class chillAST_Node *chillAST_CudaMemcpy::constantFold() { - dest = (chillAST_VarDecl *) dest->constantFold(); - src = (chillAST_VarDecl *) src->constantFold(); - size = size->constantFold(); - return this; -} - class chillAST_Node *chillAST_CudaMemcpy::clone() { - chillAST_CudaMemcpy *CMCPY = new chillAST_CudaMemcpy((chillAST_VarDecl *) (dest->clone()), - (chillAST_VarDecl *) (src->clone()), size->clone(), + chillAST_CudaMemcpy *CMCPY = new chillAST_CudaMemcpy(getDest()->clone(), + getSrc()->clone(), getSize()->clone(), strdup(cudaMemcpyKind)); CMCPY->setParent(getParent()); CMCPY->isFromSourceFile = isFromSourceFile; @@ -2297,52 +2219,21 @@ class chillAST_Node *chillAST_CudaMemcpy::clone() { } void chillAST_CudaMemcpy::gatherArrayRefs(std::vector<chillAST_ArraySubscriptExpr *> &refs, bool w) { - dest->gatherArrayRefs(refs, false); - src->gatherArrayRefs(refs, false); - size->gatherArrayRefs(refs, false); + chillAST_Node::gatherArrayRefs(refs, false); } void chillAST_CudaMemcpy::gatherScalarRefs(std::vector<chillAST_DeclRefExpr *> &refs, bool writtento) { - dest->gatherScalarRefs(refs, false); - src->gatherScalarRefs(refs, false); - size->gatherScalarRefs(refs, false); -} - -void chillAST_CudaMemcpy::gatherVarDecls(vector<chillAST_VarDecl *> &decls) { - dest->gatherVarDecls(decls); - src->gatherVarDecls(decls); - size->gatherVarDecls(decls); -} - - -void chillAST_CudaMemcpy::gatherScalarVarDecls(vector<chillAST_VarDecl *> &decls) { - dest->gatherScalarVarDecls(decls); - src->gatherScalarVarDecls(decls); - size->gatherScalarVarDecls(decls); -} - - -void chillAST_CudaMemcpy::gatherArrayVarDecls(vector<chillAST_VarDecl *> &decls) { - dest->gatherArrayVarDecls(decls); - src->gatherArrayVarDecls(decls); - size->gatherArrayVarDecls(decls); -} - - -void chillAST_CudaMemcpy::gatherVarUsage(vector<chillAST_VarDecl *> &decls) { - dest->gatherVarUsage(decls); - src->gatherVarUsage(decls); - size->gatherVarUsage(decls); + chillAST_Node::gatherScalarRefs(refs, false); } - chillAST_CudaSyncthreads::chillAST_CudaSyncthreads() { } chillAST_ReturnStmt::chillAST_ReturnStmt() { children.push_back(NULL); } -chillAST_ReturnStmt::chillAST_ReturnStmt(chillAST_Node *retval):chillAST_ReturnStmt() { + +chillAST_ReturnStmt::chillAST_ReturnStmt(chillAST_Node *retval) : chillAST_ReturnStmt() { setRetVal(retval); } @@ -2356,87 +2247,21 @@ class chillAST_Node *chillAST_ReturnStmt::clone() { return RS; } -chillAST_CallExpr::chillAST_CallExpr(chillAST_Node *c) { //, int numofargs, chillAST_Node **theargs ) { - - //fprintf(stderr, "chillAST_CallExpr::chillAST_CallExpr callee type %s\n", c->getTypeString()); - callee = c; - //callee->setParent( this ); // ?? - numargs = 0; - grid = block = NULL; -} - - -void chillAST_CallExpr::addArg(chillAST_Node *a) { - args.push_back(a); - a->setParent(this); - numargs += 1; -} - -void chillAST_CallExpr::gatherArrayRefs(std::vector<chillAST_ArraySubscriptExpr *> &refs, bool writtento) { - for (int i = 0; i < args.size(); i++) { - args[i]->gatherArrayRefs(refs, writtento); - } -} - -void chillAST_CallExpr::gatherScalarRefs(std::vector<chillAST_DeclRefExpr *> &refs, bool writtento) { - for (int i = 0; i < args.size(); i++) { - args[i]->gatherScalarRefs(refs, writtento); - } -} - - -void chillAST_CallExpr::gatherVarDecls(vector<chillAST_VarDecl *> &decls) { - for (int i = 0; i < args.size(); i++) { - args[i]->gatherVarDecls(decls); - } +chillAST_CallExpr::chillAST_CallExpr() { + addChild(NULL); } - -void chillAST_CallExpr::gatherScalarVarDecls(vector<chillAST_VarDecl *> &decls) { - for (int i = 0; i < args.size(); i++) { - args[i]->gatherScalarVarDecls(decls); - } -} - - -void chillAST_CallExpr::gatherArrayVarDecls(vector<chillAST_VarDecl *> &decls) { - for (int i = 0; i < args.size(); i++) { - args[i]->gatherArrayVarDecls(decls); - } -} - - -void chillAST_CallExpr::gatherDeclRefExprs(vector<chillAST_DeclRefExpr *> &refs) { - for (int i = 0; i < args.size(); i++) { - args[i]->gatherDeclRefExprs(refs); - } -} - -void chillAST_CallExpr::replaceVarDecls(chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl) { - for (int i = 0; i < args.size(); i++) args[i]->replaceVarDecls(olddecl, newdecl); -} - -void chillAST_CallExpr::gatherVarUsage(vector<chillAST_VarDecl *> &decls) { - for (int i = 0; i < args.size(); i++) { - args[i]->gatherVarUsage(decls); - } -} - - -chillAST_Node *chillAST_CallExpr::constantFold() { - numargs = args.size(); // wrong place for this - for (int i = 0; i < numargs; i++) { - args[i] = args[i]->constantFold(); - } - return this; +chillAST_CallExpr::chillAST_CallExpr(chillAST_Node *c) : chillAST_CallExpr() { + setCallee(c); + grid = block = NULL; } chillAST_Node *chillAST_CallExpr::clone() { //fprintf(stderr, "chillAST_CallExpr::clone()\n"); //print(0, stderr); fprintf(stderr, "\n"); - chillAST_CallExpr *CE = new chillAST_CallExpr(callee->clone()); - for (int i = 0; i < args.size(); i++) CE->addArg(args[i]->clone()); + chillAST_CallExpr *CE = new chillAST_CallExpr(getCallee()->clone()); + for (int i = 1; i < getNumChildren(); i++) CE->addArg(getChild(i)->clone()); CE->isFromSourceFile = isFromSourceFile; if (filename) CE->filename = strdup(filename); return CE; @@ -2964,6 +2789,7 @@ bool chillAST_CompoundStmt::findLoopIndexesToReplace(chillAST_SymbolTable *symta chillAST_ParenExpr::chillAST_ParenExpr() { children.push_back(NULL); } + chillAST_ParenExpr::chillAST_ParenExpr(chillAST_Node *sub) { setSubExpr(sub); } @@ -3082,18 +2908,18 @@ chillAST_IfStmt::chillAST_IfStmt() { children.push_back(NULL); } -chillAST_IfStmt::chillAST_IfStmt(chillAST_Node *c, chillAST_Node *t, chillAST_Node *e):chillAST_IfStmt() { +chillAST_IfStmt::chillAST_IfStmt(chillAST_Node *c, chillAST_Node *t, chillAST_Node *e) : chillAST_IfStmt() { setCond(c); setThen(t); setElse(e); } void chillAST_IfStmt::gatherArrayRefs(std::vector<chillAST_ArraySubscriptExpr *> &refs, bool writtento) { - chillAST_Node::gatherArrayRefs(refs,0); + chillAST_Node::gatherArrayRefs(refs, 0); } void chillAST_IfStmt::gatherScalarRefs(std::vector<chillAST_DeclRefExpr *> &refs, bool writtento) { - chillAST_Node::gatherScalarRefs(refs,0); + chillAST_Node::gatherScalarRefs(refs, 0); } chillAST_Node *chillAST_IfStmt::clone() { diff --git a/src/printer/cfamily.cpp b/src/printer/cfamily.cpp index 3427c80..d8f9057 100644 --- a/src/printer/cfamily.cpp +++ b/src/printer/cfamily.cpp @@ -70,8 +70,8 @@ int CFamily::getPrecS(chillAST_CallExpr *n) { void CFamily::printS(std::string ident, chillAST_CallExpr *n, std::ostream &o) { chillAST_FunctionDecl *FD = NULL; chillAST_MacroDefinition *MD = NULL; - if (n->callee->isDeclRefExpr()) { - chillAST_DeclRefExpr *DRE = (chillAST_DeclRefExpr *) (n->callee); + if (n->getCallee()->isDeclRefExpr()) { + chillAST_DeclRefExpr *DRE = (chillAST_DeclRefExpr *) (n->getCallee()); if (!(DRE->decl)) { o << DRE->declarationName; return; @@ -79,23 +79,23 @@ void CFamily::printS(std::string ident, chillAST_CallExpr *n, std::ostream &o) { if (DRE->decl->isFunctionDecl()) FD = (chillAST_FunctionDecl *) (DRE->decl); else CHILL_ERROR("Function DRE of type %s\n", DRE->decl->getTypeString()); - } else if (n->callee->isFunctionDecl()) - FD = (chillAST_FunctionDecl *) n->callee; - else if (n->callee->isMacroDefinition()) - MD = (chillAST_MacroDefinition *) n->callee; + } else if (n->getCallee()->isFunctionDecl()) + FD = (chillAST_FunctionDecl *) n->getCallee(); + else if (n->getCallee()->isMacroDefinition()) + MD = (chillAST_MacroDefinition *) n->getCallee(); if (FD) { o << FD->functionName; if (n->grid && n->block) o << "<<<" << n->grid->varname << "," << n->block->varname << ">>>"; o << "("; } - if (MD && n->args.size()) + if (MD && n->getNumChildren()-1) o << "("; - for (int i = 0; i < n->args.size(); ++i) { + for (int i = 1; i < n->getNumChildren(); ++i) { if (i != 0) o << ", "; - print(ident, n->args[i], o); + print(ident, n->getChild(i), o); } - if (FD || n->args.size()) + if (FD || n->getNumChildren()-1) o << ")"; } @@ -130,7 +130,9 @@ void CFamily::printS(std::string ident, chillAST_CStyleCastExpr *n, std::ostream } void CFamily::printS(std::string ident, chillAST_CudaFree *n, std::ostream &o) { - o << "cudaFree(" << n->variable->varname << ")"; + o << "cudaFree("; + print(ident, n->getPointer(), o); + o << ")"; } void CFamily::printS(std::string ident, chillAST_CudaKernelCall *n, std::ostream &o) { @@ -146,8 +148,12 @@ void CFamily::printS(std::string ident, chillAST_CudaMalloc *n, std::ostream &o) } void CFamily::printS(std::string ident, chillAST_CudaMemcpy *n, std::ostream &o) { - o << "cudaMemcpy(" << n->dest->varname << ", " << n->src->varname << ", "; - print(ident, n->size, o); + o << "cudaMemcpy("; + print(ident, n->getDest(), o); + o << ", "; + print(ident, n->getSrc(), o); + o << ", "; + print(ident, n->getSize(), o); o << ", " << n->cudaMemcpyKind << ")"; } @@ -161,7 +167,7 @@ void CFamily::printS(std::string ident, chillAST_DeclRefExpr *n, std::ostream &o void CFamily::printS(std::string ident, chillAST_FloatingLiteral *n, std::ostream &o) { if (n->allthedigits) - o<<n->allthedigits; + o << n->allthedigits; else { o << showpoint << n->value; if (n->getPrecision() == 1) diff --git a/src/printer/dump.cpp b/src/printer/dump.cpp index 0e66f29..f3fe2e5 100644 --- a/src/printer/dump.cpp +++ b/src/printer/dump.cpp @@ -70,8 +70,8 @@ void Dump::printS(std::string ident, chillAST_BinaryOperator *n, std::ostream &o } void Dump::printS(std::string ident, chillAST_CallExpr *n, std::ostream &o) { - if (n->callee) - print(ident, n->callee, o); + if (n->getCallee()) + print(ident, n->getCallee(), o); } void Dump::printS(std::string ident, chillAST_CompoundStmt *n, std::ostream &o) { @@ -88,7 +88,7 @@ void Dump::printS(std::string ident, chillAST_CStyleCastExpr *n, std::ostream &o } void Dump::printS(std::string ident, chillAST_CudaFree *n, std::ostream &o) { - o << n->variable->varname << " "; + print(ident, n->getPointer(), o); } void Dump::printS(std::string ident, chillAST_CudaKernelCall *n, std::ostream &o) { @@ -102,9 +102,9 @@ void Dump::printS(std::string ident, chillAST_CudaMalloc *n, std::ostream &o) { void Dump::printS(std::string ident, chillAST_CudaMemcpy *n, std::ostream &o) { o << n->cudaMemcpyKind << " "; - print(ident, n->dest, o); - print(ident, n->src, o); - print(ident, n->size, o); + print(ident, n->getDest(), o); + print(ident, n->getSrc(), o); + print(ident, n->getSize(), o); } void Dump::printS(std::string ident, chillAST_CudaSyncthreads *n, std::ostream &o) {} |