diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/chillASTs.cc | 474 | ||||
| -rwxr-xr-x | src/ir_clang.cc | 8 | ||||
| -rw-r--r-- | src/printer/cfamily.cpp | 353 | ||||
| -rw-r--r-- | src/printer/dump.cpp | 8 | ||||
| -rw-r--r-- | src/printer/generic.cpp | 3 |
5 files changed, 257 insertions, 589 deletions
diff --git a/src/chillASTs.cc b/src/chillASTs.cc index 69fe3ac..bee56df 100644 --- a/src/chillASTs.cc +++ b/src/chillASTs.cc @@ -5,6 +5,7 @@ #include <stack> #include "chillAST.h" #include "printer/dump.h" +#include "printer/cfamily.h" using namespace std; @@ -435,8 +436,6 @@ chillAST_VarDecl *chillAST_RecordDecl::findSubpart(const char *nam) { } fprintf(stderr, "chillAST_RecordDecl::findSubpart() couldn't find member NAMED %s in ", nam); print(); - printf("\n\n"); - fflush(stdout); return NULL; } @@ -455,33 +454,6 @@ chillAST_VarDecl *chillAST_RecordDecl::findSubpartByType(const char *typ) { return NULL; } - -void chillAST_RecordDecl::print(int indent, FILE *fp) { - CHILL_DEBUG_PRINT("chillAST_RecordDecl::print()\n"); - if (isUnnamed) return; - - chillindent(indent, fp); - if (isStruct) { - fprintf(fp, "struct "); - if (strncmp("unnamed", name, 7)) fprintf(fp, "%s\n", name); - - chillindent(indent, fp); - fprintf(fp, "{\n"); - for (int i = 0; i < subparts.size(); i++) { - subparts[i]->print(indent + 1, fp); - fprintf(fp, ";\n"); - } - fprintf(fp, "} "); - fprintf(fp, - "\n"); // ?? need semicolon when defining struct. can't have it when part of a typedef. - } else { - CHILL_ERROR("/* UNKNOWN RECORDDECL print() */ "); - exit(-1); - } - fflush(fp); -} - - chillAST_SymbolTable *chillAST_RecordDecl::addVariableToSymbolTable(chillAST_VarDecl *vd) { // for now, just bail. or do we want the struct to have an actual symbol table? //fprintf(stderr, "chillAST_RecordDecl::addVariableToSymbolTable() ignoring struct member %s vardecl\n", vd->varname); @@ -505,20 +477,6 @@ void chillAST_RecordDecl::printStructure(int indent, FILE *fp) { fflush(fp); } -void chillAST_FunctionDecl::printParameterTypes( FILE *fp ) { // also prints names - //fprintf(stderr, "\n\n%s chillAST_FunctionDecl::printParameterTypes()\n", functionName); - fprintf(fp, "( "); - int numparameters = parameters->size(); - for (int i=0; i<numparameters; i++) { - if (i!=0) fprintf(fp, ", "); - chillAST_VarDecl *p = (*parameters)[i]; - p->print(0, fp); // note: no indent, as this is in the function parens - } - fprintf(fp, " )"); // end of input parameters - -} - - chillAST_FunctionDecl::chillAST_FunctionDecl(const char *rt, const char *fname, void *unique) { CHILL_DEBUG_PRINT("chillAST_FunctionDecl::chillAST_FunctionDecl with unique %p\n", unique); if (fname) @@ -576,52 +534,6 @@ void chillAST_FunctionDecl::addChild(chillAST_Node *node) { node->setParent(this); // this, or body?? } -void chillAST_FunctionDecl::print(int indent, FILE *fp) { - //fprintf(fp, "\n// functiondecl %p \n", this); - //chillindent(indent, fp); - //fprintf(fp, "//(functiondecl) %d parameters\n", numparameters); - - fprintf(fp, "\n"); - chillindent(indent, fp); - - if (externfunc) fprintf(fp, "extern "); - - if (function_type == CHILLAST_FUNCTION_GPU) fprintf(fp, "__global__ "); - fprintf(fp, "%s %s", returnType, functionName); - printParameterTypes(fp); - - - - // non-parameter variables (now must have explicit vardecl in the body) - //int numvars = symbol_table.size(); - //for (int i=0; i<numvars; i++) { - // symbol_table[i]->print(1,fp); - // fprintf(fp, ";\n"); - //} - - // now the body - if (!(externfunc || forwarddecl)) { - if (body) { - fprintf(fp, "\n{\n"); - //chillindent(indent+1, fp); fprintf(fp, "//body\n"); fflush(fp); - body->print(indent + 1, fp); - fprintf(fp, "\n"); - //chillindent(indent+1, fp); fprintf(fp, "//END body\n"); fflush(fp); - - // tidy up - chillindent(indent, fp); - fprintf(fp, "}\n"); - } // if body - else { - fprintf(fp, "{}\n"); // should never happen, but not external and no body - } - } else { // extern func or forward decl. just end forward declaration - fprintf(fp, "; // fwd decl\n"); - } - - fflush(fp); -} - void chillAST_FunctionDecl::gatherVarDecls(vector<chillAST_VarDecl *> &decls) { //fprintf(stderr, "chillAST_FunctionDecl::gatherVarDecls()\n"); //if (0 < children.size()) fprintf(stderr, "functiondecl has %d children\n", children.size()); @@ -1373,8 +1285,6 @@ void chillAST_ForStmt::loseLoopWithLoopVar(char *var) { //fprintf(stderr, "loop condition RHS is ternary\nCondition RHS"); C->print(); - printf("\n"); - fflush(stdout); chillAST_Node *l = C->lhs; if (l->isParenExpr()) l = ((chillAST_ParenExpr *) l)->subexpr; chillAST_Node *r = C->rhs; @@ -1468,27 +1378,6 @@ chillAST_IntegerLiteral *chillAST_BinaryOperator::evalAsIntegerLiteral() { return new chillAST_IntegerLiteral(evalAsInt()); // ?? } -void chillAST_BinaryOperator::print(int indent, FILE *fp) { - chillindent(indent, fp); - bool needparens = getPrec()<lhs->getPrec(); - - if (needparens) fprintf(fp, "("); - if (lhs) lhs->print(0, fp); - else fprintf(fp, "(NULL)"); - if (needparens) fprintf(fp, ")"); - - fprintf(fp, " %s ", op); - - needparens = getPrec()<=rhs->getPrec(); - - if (needparens) fprintf(fp, "("); - if (rhs) rhs->print(0, fp); - else fprintf(fp, "(NULL)"); - if (needparens) fprintf(fp, ")"); - fflush(fp); -} - - char *chillAST_BinaryOperator::stringRep(int indent) { std::string s = string(lhs->stringRep()) + " " + op + " " + string(lhs->stringRep()); return strdup(s.c_str()); @@ -1706,18 +1595,6 @@ chillAST_TernaryOperator::chillAST_TernaryOperator(const char *oper, chillAST_No if (rhs) rhs->setParent(this); } -void chillAST_TernaryOperator::print(int indent, FILE *fp) { - chillindent(indent, fp); - fprintf(fp, "("); - condition->print(0, fp); - fprintf(fp, "%s", op); - lhs->print(0, fp); - fprintf(fp, ":"); - rhs->print(0, fp); - fprintf(fp, ")"); - fflush(fp); -} - void chillAST_TernaryOperator::replaceChild(chillAST_Node *old, chillAST_Node *newchild) { //fprintf(stderr, "\nbinop::replaceChild( old 0x%x, new ) lhs 0x%x rhd 0x%x\n", old, lhs, rhs); @@ -2239,8 +2116,6 @@ chillAST_Node *chillAST_MemberExpr::multibase2() { /*fprintf(stderr, "ME MB2\n" chillAST_VarDecl *chillAST_MemberExpr::getUnderlyingVarDecl() { fprintf(stderr, "chillAST_MemberExpr:getUnderlyingVarDecl()\n"); print(); - printf("\n"); - fflush(stdout); exit(-1); // find the member with the correct name @@ -2266,12 +2141,8 @@ chillAST_VarDecl *chillAST_MemberExpr::multibase() { fprintf(stderr, "chillAST_MemberExpr::multibase() vardecl is not a struct??\n"); fprintf(stderr, "vd "); vd->print(); - printf("\n"); - fflush(stdout); fprintf(stderr, "vd "); vd->dump(); - printf("\n"); - fflush(stdout); exit(-1); } @@ -2326,14 +2197,6 @@ chillAST_DeclRefExpr *buildDeclRefExpr(chillAST_VarDecl *vd) { return dre; } -void chillAST_DeclRefExpr::print(int indent, FILE *fp) { - chillindent(indent, fp); - //fprintf(fp, "%s %s", declarationType, declarationName); // this is printing float *A - fprintf(fp, "%s", declarationName); // this is printing A - fflush(fp); -} - - char *chillAST_DeclRefExpr::stringRep(int indent) { return strdup(declarationName); } @@ -2625,44 +2488,6 @@ chillAST_FloatingLiteral::chillAST_FloatingLiteral(chillAST_FloatingLiteral *old filename = NULL; } - -void chillAST_FloatingLiteral::print(int indent, FILE *fp) { - chillindent(indent, fp); - //fprintf(fp, "%f", value); - // attempt to be more like rose output - char output[1024]; // warning, hardcoded - - if (allthedigits != NULL) { - strcpy(output, allthedigits); // if they have specified 100 digits of pi, give 'em 100 digits - //fprintf(stderr, "floatingliteral allthedigits = '%s'\n", allthedigits); - } else { - sprintf(output, "%f", value); - - // next part to avoid printing 123.4560000000000000000000000000 - char *dot = index(output, '.'); - if (dot) { - char *end = output + strlen(output); - char *onechar; - char *lastnonzero = dot; - for (onechar = output; onechar < end; onechar++) { - if (*onechar != '0') lastnonzero = onechar; - } - - if (lastnonzero == dot) - lastnonzero[2] = '\0'; // may be after end of string, but that should be OK - else lastnonzero[1] = '\0'; // may be after end of string, but that should be OK - } - } - if (precision == 1) { - int len = strlen(output); - output[len] = 'f'; // explicit single precision - output[len + 1] = '\0'; - } - - fprintf(fp, "%s", output); - fflush(fp); -} - chillAST_Node *chillAST_FloatingLiteral::constantFold() { return this; }; // NOOP chillAST_Node *chillAST_FloatingLiteral::clone() { @@ -2689,20 +2514,6 @@ void chillAST_UnaryOperator::gatherArrayRefs(std::vector<chillAST_ArraySubscript subexpr->gatherArrayRefs(refs, isAssignmentOp()); // } - -void chillAST_UnaryOperator::print(int indent, FILE *fp) { - bool needparens = false; - if (subexpr->isNotLeaf()) needparens = true; // may get more complicated - - chillindent(indent, fp); // will this ever be invoked? - if (prefix) fprintf(fp, "%s", op); - if (needparens) fprintf(fp, "("); - subexpr->print(0, fp); - if (needparens) fprintf(fp, ")"); - if (!prefix) fprintf(fp, "%s", op); - fflush(fp); -} - void chillAST_UnaryOperator::gatherVarLHSUsage(vector<chillAST_VarDecl *> &decls) { if ((!strcmp("++", op)) || (!strcmp("--", op))) { subexpr->gatherVarUsage(decls); // do all unary modify the subexpr? (no, - ) @@ -2886,30 +2697,6 @@ void chillAST_CStyleCastExpr::replaceVarDecls(chillAST_VarDecl *olddecl, chillAS subexpr->replaceVarDecls(olddecl, newdecl); } -void chillAST_CStyleCastExpr::print(int indent, FILE *fp) { - //fprintf(stderr, "CStyleCastExpr::print()\n"); - chillindent(indent, fp); - - // special cases? should probably walk the AST and change the literal itself - if (!strcmp("float", towhat) && subexpr->isIntegerLiteral()) { // (float) 3 => 3.0f - subexpr->print(0, fp); - fprintf(fp, ".0f"); - } else if (!strcmp("double", towhat) && subexpr->isIntegerLiteral()) { // (double) 3 => 3.0 - subexpr->print(0, fp); - fprintf(fp, ".0"); - } else if (!strcmp("float", towhat) && subexpr->isFloatingLiteral()) { // (float) 3.0 => 3.0f - subexpr->print(0, fp); - fprintf(fp, "f"); - } else { // general case - fprintf(fp, "(%s) ", towhat); - if (subexpr->getPrec()<getPrec()) fprintf(fp,"("); - if (subexpr->isVarDecl()) fprintf(fp, "%s", ((chillAST_VarDecl *) subexpr)->varname); - else subexpr->print(indent, fp); - if (subexpr->getPrec()<getPrec()) fprintf(fp,")"); - } - fflush(fp); -}; - class chillAST_Node *chillAST_CStyleCastExpr::constantFold() { subexpr = subexpr->constantFold(); return this; @@ -2964,15 +2751,6 @@ chillAST_CStyleAddressOf::chillAST_CStyleAddressOf(chillAST_Node *sub) { //fprintf(stderr, "chillAST_CStyleCastExpr (%s) sub 0x%x\n", towhat, sub ); } -void chillAST_CStyleAddressOf::print(int indent, FILE *fp) { - //fprintf(stderr, "CStyleAddressOf::print()\n"); - chillindent(indent, fp); - fprintf(fp, "(&"); - subexpr->print(0, fp); - fprintf(fp, ")"); - fflush(fp); -}; - class chillAST_Node *chillAST_CStyleAddressOf::constantFold() { subexpr = subexpr->constantFold(); return this; @@ -3066,16 +2844,6 @@ chillAST_CudaMalloc::chillAST_CudaMalloc(chillAST_Node *devmemptr, chillAST_Node sizeinbytes = size; // probably a multiply like sizeof(int) * 1024 }; -void chillAST_CudaMalloc::print(int indent, FILE *fp) { - chillindent(indent, fp); - fprintf(fp, "cudaMalloc("); - devPtr->print(0, fp); - fprintf(fp, ","); - sizeinbytes->print(0, fp); - fprintf(fp, ")"); - fflush(fp); -}; - class chillAST_Node *chillAST_CudaMalloc::constantFold() { devPtr = devPtr->constantFold(); return this; @@ -3127,12 +2895,6 @@ chillAST_CudaFree::chillAST_CudaFree(chillAST_VarDecl *var) { variable = var; }; -void chillAST_CudaFree::print(int indent, FILE *fp) { - chillindent(indent, fp); - fprintf(fp, "cudaFree(%s)", variable->varname); - fflush(fp); -}; - class chillAST_Node *chillAST_CudaFree::constantFold() { return this; } @@ -3237,30 +2999,11 @@ void chillAST_CudaMemcpy::gatherVarUsage(vector<chillAST_VarDecl *> &decls) { chillAST_CudaSyncthreads::chillAST_CudaSyncthreads() { } -void chillAST_CudaSyncthreads::print(int indent, FILE *fp) { - chillindent(indent, fp); - fprintf(fp, "__syncthreads()"); - fflush(fp); -} - chillAST_ReturnStmt::chillAST_ReturnStmt(chillAST_Node *retval) { returnvalue = retval; if (returnvalue) returnvalue->setParent(this); } - -void chillAST_ReturnStmt::print(int indent, FILE *fp) { - chillindent(indent, fp); - if (returnvalue != NULL) { - fprintf(fp, "return("); - returnvalue->print(0, fp); - fprintf(fp, ")"); // parent will add ";\n" ?? - } else { - fprintf(fp, "return"); - } - fflush(fp); -} - class chillAST_Node *chillAST_ReturnStmt::constantFold() { if (returnvalue) returnvalue = returnvalue->constantFold(); return this; @@ -3319,55 +3062,6 @@ void chillAST_CallExpr::addArg(chillAST_Node *a) { numargs += 1; } - -void chillAST_CallExpr::print(int indent, FILE *fp) { - chillindent(indent, fp); - chillAST_FunctionDecl *FD = NULL; - chillAST_MacroDefinition *MD = NULL; - - if (callee->isDeclRefExpr()) { - chillAST_DeclRefExpr *DRE = (chillAST_DeclRefExpr *) callee; - //fprintf(stderr, "DRE decl is 0x%x\n", DRE->decl); - if (!DRE->decl) { - // a macro? - fprintf(fp, "%s ", DRE->declarationName); - return; // ?? - } - - //fprintf(stderr, "DRE decl of type %s\n", DRE->decl->getTypeString()); - if ((DRE->decl)->isFunctionDecl()) FD = (chillAST_FunctionDecl *) DRE->decl; - else { - fprintf(stderr, "chillAST_CallExpr::print() DRE decl of type %s\n", DRE->decl->getTypeString()); - exit(-1); - } - } else if (callee->isFunctionDecl()) FD = (chillAST_FunctionDecl *) callee; - else if (callee->isMacroDefinition()) { - MD = (chillAST_MacroDefinition *) callee; - fprintf(fp, "%s(", MD->macroName); - } else { - fprintf(stderr, "\nchillAST_CallExpr::print() callee of unhandled type %s\n", callee->getTypeString()); - callee->dump(); - exit(-1); - } - - if (FD) { - fprintf(fp, "%s", FD->functionName); - fflush(fp); - if (grid && block) { - fprintf(fp, "<<<%s,%s>>>(", grid->varname, block->varname); // a - } else fprintf(fp, "("); - } - - - //callee->print( indent, fp); - for (int i = 0; i < args.size(); i++) { - if (i != 0) fprintf(fp, ", "); - args[i]->print(0, fp); - } - fprintf(fp, ")"); //a - fflush(fp); -} - void chillAST_CallExpr::gatherArrayRefs(std::vector<chillAST_ArraySubscriptExpr *> &refs, bool writtento) { for (int i = 0; i < args.size(); i++) { args[i]->gatherArrayRefs(refs, writtento); @@ -3751,22 +3445,6 @@ chillAST_CompoundStmt::chillAST_CompoundStmt() { typedefTable = new chillAST_TypedefTable; }; - -void chillAST_CompoundStmt::print(int indent, FILE *fp) { - int numchildren = children.size(); - for (int i = 0; i < numchildren; i++) { - children[i]->print(indent, fp); - if (children[i]->getType() != CHILLAST_NODE_FORSTMT - && children[i]->getType() != CHILLAST_NODE_IFSTMT - && children[i]->getType() != CHILLAST_NODE_COMPOUNDSTMT - //&& children[i]->getType() != CHILLAST_NODE_VARDECL // vardecl does its own ";\n" - ) { - fprintf(fp, ";\n"); // probably wrong - } - } - fflush(fp); -} - void chillAST_CompoundStmt::replaceChild(chillAST_Node *old, chillAST_Node *newchild) { //fprintf(stderr, "chillAST_CompoundStmt::replaceChild( old %s, new %s)\n", old->getTypeString(), newchild->getTypeString() ); vector<chillAST_Node *> dupe = children; @@ -3979,15 +3657,6 @@ chillAST_ParenExpr::chillAST_ParenExpr(chillAST_Node *sub) { subexpr->setParent(this); } -void chillAST_ParenExpr::print(int indent, FILE *fp) { - //fprintf(stderr, "chillAST_ParenExpr::print()\n"); - chillindent(indent, fp); // hard to believe this will ever do anything - fprintf(fp, "("); - subexpr->print(0, fp); - fprintf(fp, ")"); - fflush(fp); -} - void chillAST_ParenExpr::gatherArrayRefs(std::vector<chillAST_ArraySubscriptExpr *> &refs, bool writtento) { subexpr->gatherArrayRefs(refs, writtento); } @@ -4265,61 +3934,6 @@ chillAST_Node *chillAST_IfStmt::clone() { return IS; } -void chillAST_IfStmt::print(int indent, FILE *fp) { - chillindent(indent, fp); - fprintf(fp, "if ("); - if (cond) cond->print(0, fp); - else fprintf(fp, "(NULL cond)"); - - bool needbracket = true; - if (thenpart) { - if (thenpart->isBinaryOperator()) needbracket = false; - if (thenpart->isCompoundStmt()) { // almost always true - chillAST_CompoundStmt *CS = (chillAST_CompoundStmt *) thenpart; - if (CS->children.size() == 1 && CS->children[0]->isBinaryOperator()) needbracket = false; - } - - if (needbracket) fprintf(fp, ") {\n"); - else fprintf(fp, ")\n"); - - thenpart->print(indent + 1, fp); // end of line - - if (needbracket) { - //fprintf(fp, "\n"); - chillindent(indent, fp); - fprintf(fp, "}\n"); - } - } else fprintf(fp, "(NULL thenpart)"); - - - needbracket = true; - if (elsepart) { - if (elsepart->isBinaryOperator()) needbracket = false; - if (elsepart->isCompoundStmt()) { // almost always true - chillAST_CompoundStmt *CS = (chillAST_CompoundStmt *) elsepart; - - if (CS->children.size() == 1 && CS->children[0]->isBinaryOperator()) needbracket = false; - - } - - fprintf(fp, "\n"); - chillindent(indent, fp); - - if (needbracket) fprintf(fp, "else {\n"); - else fprintf(fp, "else\n"); - - elsepart->print(indent + 1, fp); - - if (needbracket) { - fprintf(fp, "\n"); - chillindent(indent, fp); - fprintf(fp, "}\n"); - } - } - //else fprintf(fp, "else { /* NOTHING */ }"); -} - - bool chillAST_IfStmt::findLoopIndexesToReplace(chillAST_SymbolTable *symtab, bool forcesync) { thenpart->findLoopIndexesToReplace(symtab); elsepart->findLoopIndexesToReplace(symtab); @@ -4478,78 +4092,6 @@ chillAST_Preprocessing::chillAST_Preprocessing(CHILLAST_PREPROCESSING_POSITION p blurb = strdup(text); } -void chillAST_Preprocessing::print(int indent, FILE *fp) { // probably very wrong - if (position == CHILLAST_PREPROCESSING_LINEAFTER) { - fprintf(fp, "\n"); - chillindent(indent, fp); - } - if (position == CHILLAST_PREPROCESSING_LINEBEFORE) { // ??? - //fprintf(fp, "\n"); - chillindent(indent, fp); - } - - fprintf(fp, "%s", blurb); - - if (position == CHILLAST_PREPROCESSING_TOTHERIGHT) { - fprintf(fp, "\n"); - } - - - if (position == CHILLAST_PREPROCESSING_LINEBEFORE) { - } - - - -} - -//! I'm just a bit lazy to write ifs ... -const char* binaryPrec[] = { - " :: ", - " . -> ", - "", - " .* ->* ", - " * / % ", - " + - ", - " << >> ", - " < <= > >=", - " == != ", - " & ", - " ^ ", - " | ", - " && ", - " || ", - " = += -= *= /= %= <<= >>= &= ^= |= ", - " , " -}; - -bool opInSet(const char* set,char* op) { - string tmp = op; - tmp=" "+tmp+" "; - return strstr(set, tmp.c_str()) != NULL; -} - -int chillAST_BinaryOperator::getPrec() { - for (int i = 0; i< 16;++i) - if (opInSet(binaryPrec[i],op)) return INT8_MAX+i+1; - return INT8_MAX; -} - -const char* unaryPrec[] = { - "", - " -- ++ ", - " -- ++ + - ! ~ * & ", -}; - -int chillAST_UnaryOperator::getPrec() { - if (prefix) - for (int i = 2;i>=0;--i) - if (opInSet(unaryPrec[i],op)) return INT8_MAX+i+1; - else - for (int i = 0;i<3;--i) - if (opInSet(unaryPrec[i],op)) return INT8_MAX+i+1; - return INT8_MAX; -} - void chillAST_Node::addVariableToScope(chillAST_VarDecl *vd) { CHILL_DEBUG_PRINT("addVariableToScope( %s )\n", vd->varname); if (!symbolTable) return; @@ -4609,8 +4151,22 @@ void chillAST_Node::dump(int indent, FILE *fp) { if (fp == stderr) { chill::printer::Dump d; d.printErr("",this); + fprintf(stderr,"\n"); } else { chill::printer::Dump d; d.printOut("",this); + fprintf(stdout,"\n"); + } +} + +void chillAST_Node::print(int indent, FILE *fp) { + if (fp == stderr) { + chill::printer::CFamily d; + d.printErr("",this); + fprintf(stderr,"\n"); + } else { + chill::printer::CFamily d; + d.printOut("",this); + fprintf(stdout,"\n"); } }
\ No newline at end of file diff --git a/src/ir_clang.cc b/src/ir_clang.cc index 408148f..e0699a5 100755 --- a/src/ir_clang.cc +++ b/src/ir_clang.cc @@ -1636,14 +1636,6 @@ bool IR_chillArrayRef::operator!=(const IR_Ref &that) const { return !op; } -void IR_chillArrayRef::Dump() const { - //fprintf(stderr, "IR_chillArrayRef::Dump() this 0x%x chillASE 0x%x\n", this, chillASE); - chillASE->print(); - printf("\n"); - fflush(stdout); -} - - bool IR_chillArrayRef::operator==(const IR_Ref &that) const { //fprintf(stderr, "IR_xxxxArrayRef::operator==\n"); //printf("I am\n"); chillASE->print(); printf("\n"); diff --git a/src/printer/cfamily.cpp b/src/printer/cfamily.cpp index a15c106..e1e2b9a 100644 --- a/src/printer/cfamily.cpp +++ b/src/printer/cfamily.cpp @@ -3,123 +3,170 @@ // #include "printer/cfamily.h" +#include <iomanip> + using namespace std; using namespace chill::printer; +bool opInSet(const char *set, char *op) { + string tmp = op; + tmp = " " + tmp + " "; + return strstr(set, tmp.c_str()) != NULL; +} + bool ifSemicolonFree(CHILLAST_NODE_TYPE t) { return t == CHILLAST_NODE_FUNCTIONDECL || t == CHILLAST_NODE_IFSTMT || t == CHILLAST_NODE_FORSTMT || t == CHILLAST_NODE_MACRODEFINITION; } -void dumpVector(GenericPrinter *p, string ident, chillAST_NodeList *n, ostream &o) { - for (int i = 0; i < n->size(); ++i) - p->print("", (*n)[i], o); -} - -void dumpVector(GenericPrinter *p, string ident, chillAST_SymbolTable *n, ostream &o) { - for (int i = 0; i < n->size(); ++i) - p->print("", (*n)[i], o); +void CFamily::printS(std::string ident, chillAST_ArraySubscriptExpr *n, std::ostream &o) { + print(ident, n->base, o); + o << "["; + print(ident, n->index, o); + o << "]"; +} + +//! I'm just a bit lazy to write ifs ... +const char *binaryPrec[] = { + " :: ", + " . -> ", + "", + " .* ->* ", + " * / % ", + " + - ", + " << >> ", + " < <= > >=", + " == != ", + " & ", + " ^ ", + " | ", + " && ", + " || ", + " = += -= *= /= %= <<= >>= &= ^= |= ", + " , " +}; + +int CFamily::getPrecS(chillAST_BinaryOperator *n) { + for (int i = 0; i < 16; ++i) + if (opInSet(binaryPrec[i], n->op)) return defGetPrecS() + i + 1; + CHILL_ERROR("Unrecognized binary operator: %s\n", n->op); + return defGetPrecS(); +} + +void CFamily::printS(std::string ident, chillAST_BinaryOperator *n, std::ostream &o) { + int prec = getPrec(n); + if (n->lhs) printPrec(ident, n->lhs, o, prec); + else o << "(NULL)"; + o << " " << n->op << " "; + if (n->rhs) printPrec(ident, n->rhs, o, prec); + else o << "(NULL)"; } -void dumpVector(GenericPrinter *p, string ident, chillAST_TypedefTable *n, ostream &o) { - for (int i = 0; i < n->size(); ++i) - p->print("", (*n)[i], o); +int CFamily::getPrecS(chillAST_CallExpr *n) { + return defGetPrecS() + 2; } -void Dump::print(string ident, chillAST_Node *n, ostream &o) { - o << "(" << n->getTypeString() << " "; - if (n->getParameters()) { - o << "(Params: "; - dumpVector(this, ident, n->getParameters(), o); - o << ") "; +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 (!(DRE->decl)) { + o << DRE->declarationName; + return; + } + 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; + if (FD) { + o << FD->functionName; + if (n->grid && n->block) + o << "<<<" << n->grid->varname << "," << n->block->varname << ">>>"; + o << "("; } - if (n->getSymbolTable()) { - o << "(VarScope: "; - dumpVector(this, ident, n->getSymbolTable(), o); - o << ") "; + if (MD && n->args.size()) + o << "("; + for (int i = 0; i < n->args.size(); ++i) { + if (i != 0) o << ", "; + print(ident, n->args[i], o); } - if (n->getTypedefTable()) { - o << "(TypeDef: "; - dumpVector(this, ident, n->getTypedefTable(), o); - o << ") "; + if (FD || n->args.size()) + o << ")"; +} + +void CFamily::printS(std::string ident, chillAST_CompoundStmt *n, std::ostream &o) { + chillAST_NodeList *c = n->getChildren(); + string nid = ident + identSpace; + if (c->size() > 1) o << "{"; + for (int i = 0; i < c->size(); ++i) { + o << "\n" << nid; + print(nid, c->at(i), o); + if (!ifSemicolonFree(c->at(i)->getType())) o << ";"; } - o << ": "; - // Recurse - GenericPrinter::print(ident, n, o); - o << ") "; -} - -void Dump::printS(std::string ident, chillAST_ArraySubscriptExpr *n, std::ostream &o) { - print(ident,n->base,o); - o<<"["; - print(ident,n->index,o); - o<<"]"; -} - -void Dump::printS(std::string ident, chillAST_BinaryOperator *n, std::ostream &o) { - o << n->op << " "; - if (n->lhs) print(ident, n->lhs, o); - else o << "(NULL) "; - if (n->rhs) print(ident, n->rhs, o); - else o << "(NULL) "; + if (c->size() > 1) o << "\n" << ident << "}"; } -void Dump::printS(std::string ident, chillAST_CallExpr *n, std::ostream &o) { - if (n->callee) - print(ident, n->callee, o); +int CFamily::getPrecS(chillAST_CStyleAddressOf *n) { + return defGetPrecS() + 3; } -void Dump::printS(std::string ident, chillAST_CompoundStmt *n, std::ostream &o) { - dumpVector(this, ident, n->getChildren(), o); +void CFamily::printS(std::string ident, chillAST_CStyleAddressOf *n, std::ostream &o) { + int prec = getPrec(n); + printPrec(ident, n->subexpr, o, prec); } -void Dump::printS(std::string ident, chillAST_CStyleAddressOf *n, std::ostream &o) { - print(ident, n->subexpr, o); +int CFamily::getPrecS(chillAST_CStyleCastExpr *n) { + return defGetPrecS() + 3; } -void Dump::printS(std::string ident, chillAST_CStyleCastExpr *n, std::ostream &o) { - o << n->towhat << " "; - print(ident, n->subexpr, o); +void CFamily::printS(std::string ident, chillAST_CStyleCastExpr *n, std::ostream &o) { + o << "(" << n->towhat << ")"; + printPrec(ident, n->subexpr, o, getPrec(n)); } -void Dump::printS(std::string ident, chillAST_CudaFree *n, std::ostream &o) { - o << n->variable->varname << " "; +void CFamily::printS(std::string ident, chillAST_CudaFree *n, std::ostream &o) { + o << "cudaFree(" << n->variable->varname << ")"; } -void Dump::printS(std::string ident, chillAST_CudaKernelCall *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_CudaKernelCall *n, std::ostream &o) { CHILL_ERROR("Not implemented"); } -void Dump::printS(std::string ident, chillAST_CudaMalloc *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_CudaMalloc *n, std::ostream &o) { + o << "cudaMalloc("; print(ident, n->devPtr, o); + o << ", "; print(ident, n->sizeinbytes, o); + o << ")"; } -void Dump::printS(std::string ident, chillAST_CudaMemcpy *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 << ", " << n->cudaMemcpyKind << ")"; } -void Dump::printS(std::string ident, chillAST_CudaSyncthreads *n, std::ostream &o) {} +void CFamily::printS(std::string ident, chillAST_CudaSyncthreads *n, std::ostream &o) { + o << "__syncthreads()"; +} -void Dump::printS(std::string ident, chillAST_DeclRefExpr *n, std::ostream &o) { - chillAST_VarDecl *vd = n->getVarDecl(); - if (vd) - if (vd->isAParameter) o << "ParmVar "; else o << "Var "; - o << n->declarationName << " "; - chillAST_FunctionDecl *fd = n->getFunctionDecl(); - if (fd) dumpVector(this, ident, fd->getParameters(), o); +void CFamily::printS(std::string ident, chillAST_DeclRefExpr *n, std::ostream &o) { + o << n->declarationName; } -void Dump::printS(std::string ident, chillAST_FloatingLiteral *n, std::ostream &o) { - if (n->precision == 1) o << "float "; - else o << "double "; - o << n->value; +void CFamily::printS(std::string ident, chillAST_FloatingLiteral *n, std::ostream &o) { + // Althedigits contaminates the result + o << showpoint << n->value; + if (n->getPrecision() == 1) + o << "f"; } -void Dump::printS(std::string ident, chillAST_ForStmt *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_ForStmt *n, std::ostream &o) { if (n->metacomment) o << "// " << n->metacomment << "\n"; o << "for ("; @@ -128,10 +175,8 @@ void Dump::printS(std::string ident, chillAST_ForStmt *n, std::ostream &o) { print(ident, n->getCond(), o); o << ";"; print(ident, n->getInc(), o); - o << ")"; - if (n->getBody()->getType() == CHILLAST_NODE_COMPOUNDSTMT) { - if (n->getBody()->getChildren()->size() < 2) o << "\n" << ident << identSpace; - else o << " "; + o << ") "; + if (n->getBody()->isCompoundStmt()) { print(ident, n->getBody(), o); } else { CHILL_ERROR("Body of for loop not COMPOUNDSTMT\n"); @@ -139,50 +184,78 @@ void Dump::printS(std::string ident, chillAST_ForStmt *n, std::ostream &o) { } } -void Dump::printS(std::string ident, chillAST_Free *n, std::ostream &o) {} +void CFamily::printS(std::string ident, chillAST_Free *n, std::ostream &o) {} -void Dump::printS(std::string ident, chillAST_FunctionDecl *n, std::ostream &o) { - if (n->filename) o << n->filename << " "; - if (n->isFromSourceFile) o << "FromSourceFile" << " "; - o << n->returnType << " " << n->functionName << " "; - if (n->getBody()) print(ident, n->getBody(), o); +void CFamily::printS(std::string ident, chillAST_FunctionDecl *n, std::ostream &o) { + if (n->isExtern()) o << "extern "; + if (n->getFunctionType() == CHILLAST_FUNCTION_GPU) o << "__global__ "; + o << n->returnType << " " << n->functionName << "("; + + chillAST_SymbolTable *pars = n->getParameters(); + for (int i = 0; i < pars->size(); ++i) { + if (i != 0) + o << ", "; + print(ident, pars->at(i), o); + } + o << ")"; + if (!(n->isExtern() || n->isForward())) { + o << " "; + if (n->getBody()) + print(ident, n->getBody(), o); + else { + CHILL_ERROR("Non-extern or forward function decl doesn't have a body"); + o << "{}"; + } + } else { + o << ";"; + } } -void Dump::printS(std::string ident, chillAST_IfStmt *n, std::ostream &o) { - print(ident, n->cond, o); - print(ident, n->thenpart, o); - if (n->elsepart) +void CFamily::printS(std::string ident, chillAST_IfStmt *n, std::ostream &o) { + o << "if ("; + print(ident, n->getCond(), o); + o << ") "; + if (!n->getThen()) { + CHILL_ERROR("If statement is without then part!"); + exit(-1); + } + print(ident, n->getThen(), o); + if (!(n->getThen()->isCompoundStmt())) + CHILL_ERROR("Then part is not a CompoundStmt!\n"); + if (n->getElse()) { + o << "else "; print(ident, n->elsepart, o); + } } -void Dump::printS(std::string ident, chillAST_IntegerLiteral *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_IntegerLiteral *n, std::ostream &o) { o << n->value; } -void Dump::printS(std::string ident, chillAST_ImplicitCastExpr *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_ImplicitCastExpr *n, std::ostream &o) { print(ident, n->subexpr, o); } -void Dump::printS(std::string ident, chillAST_MacroDefinition *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_MacroDefinition *n, std::ostream &o) { o << "#define" << n->macroName << " "; int np = n->getParameters()->size(); if (np) { o << "(" << n->getParameters()->at(0)->varname; for (int i = 1; i < np; ++i) o << ", " << n->getParameters()->at(i)->varname; - o<<")"; + o << ")"; } // TODO newline for multiline macro print(ident, n->getBody(), o); } -void Dump::printS(std::string ident, chillAST_Malloc *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_Malloc *n, std::ostream &o) { o << "malloc("; print(ident, n->sizeexpr, o); o << ")"; } -void Dump::printS(std::string ident, chillAST_MemberExpr *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_MemberExpr *n, std::ostream &o) { int prec = getPrec(n); if (n->base) printPrec(ident, n->base, o, prec); else o << "(NULL)"; @@ -192,48 +265,67 @@ void Dump::printS(std::string ident, chillAST_MemberExpr *n, std::ostream &o) { else o << "(NULL)"; } -void Dump::printS(std::string ident, chillAST_NULL *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_NULL *n, std::ostream &o) { o << "/* (NULL statement) */"; } -void Dump::printS(std::string ident, chillAST_NoOp *n, std::ostream &o) {} +void CFamily::printS(std::string ident, chillAST_NoOp *n, std::ostream &o) {} -void Dump::printS(std::string ident, chillAST_ParenExpr *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_ParenExpr *n, std::ostream &o) { + o << "("; print(ident, n->subexpr, o); + o << ")"; } -void Dump::printS(std::string ident, chillAST_Preprocessing *n, std::ostream &o) {} +void CFamily::printS(std::string ident, chillAST_Preprocessing *n, std::ostream &o) { + CHILL_ERROR("Not implemented\n"); +} -void Dump::printS(std::string ident, chillAST_RecordDecl *n, std::ostream &o) { - // TODO access control - o << n->getName() << " "; - o << n->isAStruct() << " "; - o << n->isAUnion() << " "; +void CFamily::printS(std::string ident, chillAST_RecordDecl *n, std::ostream &o) { + if (n->isUnnamed) return; + if (n->isAStruct()) { + string nid = ident + identSpace; + o << "struct " << n->getName() << " {"; + chillAST_SymbolTable *sp = n->getSubparts(); + for (int i = 0; i < sp->size(); ++i) { + o << "\n" << nid; + print(nid, sp->at(i), o); + o << ";"; + } + o << "\n" << ident << "}"; + } else { + CHILL_ERROR("Encountered Unkown record type"); + exit(-1); + } } -void Dump::printS(std::string ident, chillAST_ReturnStmt *n, std::ostream &o) { - if (n->returnvalue) print(ident, n->returnvalue, o); +void CFamily::printS(std::string ident, chillAST_ReturnStmt *n, std::ostream &o) { + o << "return"; + if (n->returnvalue) { + o << " "; + print(ident, n->returnvalue, o); + } } -void Dump::printS(std::string ident, chillAST_Sizeof *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_Sizeof *n, std::ostream &o) { o << "sizeof(" << n->thing << ")"; } -void Dump::printS(std::string ident, chillAST_SourceFile *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_SourceFile *n, std::ostream &o) { o << "// this source is derived from CHILL AST originally from file '" << n->filename << "' as parsed by frontend compiler " << n->frontend << "\n\n"; int nchild = n->getChildren()->size(); for (int i = 0; i < nchild; ++i) { if (n->getChild(i)->isFromSourceFile) { o << ident; - print(indent, n->getChild(i), o); - if (!isSemiColonFree(n->getChild(i)->getType())) o << ";\n"; - else o<<"\n"; + print(ident, n->getChild(i), o); + if (!ifSemicolonFree(n->getChild(i)->getType())) o << ";\n"; + else o << "\n"; } } } -void Dump::printS(std::string ident, chillAST_TypedefDecl *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_TypedefDecl *n, std::ostream &o) { if (n->isAStruct()) o << "/* A typedef STRUCT */\n"; o << ident << "typedef "; @@ -246,20 +338,43 @@ void Dump::printS(std::string ident, chillAST_TypedefDecl *n, std::ostream &o) { o << n->newtype; } -void Dump::printS(std::string ident, chillAST_TernaryOperator *n, std::ostream &o) { - o << n->op << " "; - print(ident, n->condition, o); - print(ident, n->lhs, o); - print(ident, n->rhs, o); +int CFamily::getPrecS(chillAST_TernaryOperator *n) { + return defGetPrecS() + 15; } -void Dump::printS(std::string ident, chillAST_UnaryOperator *n, std::ostream &o) { - if (n->prefix) o << "prefix "; - else o << "postfix "; - print(ident, n->subexpr, o); +void CFamily::printS(std::string ident, chillAST_TernaryOperator *n, std::ostream &o) { + int prec = getPrec(n); + printPrec(ident, n->condition, o, prec); + o << "" << n->op << ""; + printPrec(ident, n->lhs, o, prec); + o << ":"; + printPrec(ident, n->rhs, o, prec); +} + +const char *unaryPrec[] = { + "", + " -- ++ ", + " -- ++ + - ! ~ * & ", +}; + +int CFamily::getPrecS(chillAST_UnaryOperator *n) { + if (n->prefix) { + for (int i = 2; i >= 0; --i) + if (opInSet(unaryPrec[i], n->op)) return defGetPrecS() + i + 1; + } else + for (int i = 1; i < 3; ++i) + if (opInSet(unaryPrec[i], n->op)) return defGetPrecS() + i + 1; + return defGetPrecS(); +} + +void CFamily::printS(std::string ident, chillAST_UnaryOperator *n, std::ostream &o) { + int prec = getPrec(n); + if (n->prefix) o << n->op; + printPrec(ident, n->subexpr, o, prec); + if (!n->prefix) o << n->op; } -void Dump::printS(std::string ident, chillAST_VarDecl *n, std::ostream &o) { +void CFamily::printS(std::string ident, chillAST_VarDecl *n, std::ostream &o) { if (n->isDevice) o << "__device__ "; if (n->isShared) o << "__shared__ "; if (n->isRestrict) o << "__restrict__ "; diff --git a/src/printer/dump.cpp b/src/printer/dump.cpp index bf3f6b8..66bcc7e 100644 --- a/src/printer/dump.cpp +++ b/src/printer/dump.cpp @@ -10,15 +10,17 @@ using namespace std; void dumpVector(GenericPrinter *p, string ident, chillAST_NodeList *n, ostream &o) { for (int i = 0; i < n->size(); ++i) - p->print("", (*n)[i], o); + p->print(ident, (*n)[i], o); } + void dumpVector(GenericPrinter *p, string ident, chillAST_SymbolTable *n, ostream &o) { for (int i = 0; i < n->size(); ++i) - p->print("", (*n)[i], o); + p->print(ident, (*n)[i], o); } + void dumpVector(GenericPrinter *p, string ident, chillAST_TypedefTable *n, ostream &o) { for (int i = 0; i < n->size(); ++i) - p->print("", (*n)[i], o); + p->print(ident, (*n)[i], o); } void Dump::print(string ident, chillAST_Node *n, ostream &o) { diff --git a/src/printer/generic.cpp b/src/printer/generic.cpp index 67c694f..79ee312 100644 --- a/src/printer/generic.cpp +++ b/src/printer/generic.cpp @@ -198,3 +198,6 @@ int GenericPrinter::getPrec(chillAST_Node *n) { } } +void GenericPrinter::errorPrintS(std::string ident, chillAST_Node *n, std::ostream &o) { + CHILL_ERROR("Unhandled case in printer: %s\n", n->getTypeString()); +} |