summaryrefslogtreecommitdiff
path: root/include
diff options
context:
space:
mode:
authorTuowen Zhao <ztuowen@gmail.com>2016-09-21 22:35:47 -0600
committerTuowen Zhao <ztuowen@gmail.com>2016-09-21 22:35:47 -0600
commitab016596602a4c6bdc27adf01c308b325af221f0 (patch)
tree4e86bfcf1f38fb00cc58082d540dc3570e0f126b /include
parent6983c09937baac3ffb7d3a45c3c5009c0eba7e6c (diff)
downloadchill-ab016596602a4c6bdc27adf01c308b325af221f0.tar.gz
chill-ab016596602a4c6bdc27adf01c308b325af221f0.tar.bz2
chill-ab016596602a4c6bdc27adf01c308b325af221f0.zip
something that only builds ...
Diffstat (limited to 'include')
-rw-r--r--include/chill_ast.hh2365
-rw-r--r--include/chill_error.hh4
-rw-r--r--include/dep.hh32
-rwxr-xr-xinclude/ir_clang.hh479
-rw-r--r--include/ir_code.hh282
-rw-r--r--include/ir_enums.hh24
-rw-r--r--include/ir_rose.hh267
-rw-r--r--include/ir_rose_utils.hh19
-rw-r--r--include/irtools.hh66
-rw-r--r--include/loop.hh171
-rw-r--r--include/omegatools.hh120
-rw-r--r--include/stencil.hh55
12 files changed, 3290 insertions, 594 deletions
diff --git a/include/chill_ast.hh b/include/chill_ast.hh
new file mode 100644
index 0000000..0803254
--- /dev/null
+++ b/include/chill_ast.hh
@@ -0,0 +1,2365 @@
+
+
+#ifndef _CHILL_AST_H_
+#define _CHILL_AST_H_
+
+
+#define CHILL_INDENT_AMOUNT 2
+
+#include <iostream>
+#include <stdio.h>
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+#include <vector> // std::vector
+
+#include <ir_enums.hh> // for IR_CONDITION_*
+
+char *parseUnderlyingType( char *sometype );
+char *parseArrayParts( char *sometype );
+bool isRestrict( const char *sometype );
+char *splitTypeInfo( char *underlyingtype );
+char *ulhack( char *brackets ); // change "1024UL" to "1024"
+char *restricthack( char *typeinfo ); // remove __restrict__ , MODIFIES the argument!
+
+
+enum CHILL_ASTNODE_TYPE {
+ CHILLAST_NODETYPE_UNKNOWN=0,
+ CHILLAST_NODETYPE_SOURCEFILE,
+ CHILLAST_NODETYPE_TYPEDEFDECL,
+ CHILLAST_NODETYPE_VARDECL,
+ // CHILLAST_NODETYPE_PARMVARDECL, not used any more
+ CHILLAST_NODETYPE_FUNCTIONDECL,
+ CHILLAST_NODETYPE_RECORDDECL, // struct or union (or class)
+ CHILLAST_NODETYPE_MACRODEFINITION,
+ CHILLAST_NODETYPE_COMPOUNDSTMT,
+ CHILLAST_NODETYPE_LOOP, // AKA ForStmt
+ CHILLAST_NODETYPE_TERNARYOPERATOR,
+ CHILLAST_NODETYPE_BINARYOPERATOR,
+ CHILLAST_NODETYPE_UNARYOPERATOR,
+ CHILLAST_NODETYPE_ARRAYSUBSCRIPTEXPR,
+ CHILLAST_NODETYPE_MEMBEREXPR, // structs/unions
+ CHILLAST_NODETYPE_DECLREFEXPR,
+ CHILLAST_NODETYPE_INTEGERLITERAL,
+ CHILLAST_NODETYPE_FLOATINGLITERAL,
+ CHILLAST_NODETYPE_IMPLICITCASTEXPR,
+ CHILLAST_NODETYPE_RETURNSTMT,
+ CHILLAST_NODETYPE_CALLEXPR,
+ CHILLAST_NODETYPE_DECLSTMT,
+ CHILLAST_NODETYPE_PARENEXPR,
+ CHILLAST_NODETYPE_CSTYLECASTEXPR,
+ CHILLAST_NODETYPE_CSTYLEADDRESSOF,
+ CHILLAST_NODETYPE_IFSTMT,
+ CHILLAST_NODETYPE_SIZEOF,
+ CHILLAST_NODETYPE_MALLOC,
+ CHILLAST_NODETYPE_FREE,
+ CHILLAST_NODETYPE_PREPROCESSING, // comments, #define, #include, whatever else works
+ CHILLAST_NODETYPE_NOOP, // NO OP
+ // CUDA specific
+ CHILLAST_NODETYPE_CUDAMALLOC,
+ CHILLAST_NODETYPE_CUDAFREE,
+ CHILLAST_NODETYPE_CUDAMEMCPY,
+ CHILLAST_NODETYPE_CUDAKERNELCALL,
+ CHILLAST_NODETYPE_CUDASYNCTHREADS,
+ CHILLAST_NODETYPE_NULL // explicit non-statement
+ // TODO
+
+} ;
+
+#define CHILLAST_NODETYPE_FORSTMT CHILLAST_NODETYPE_LOOP
+#define CHILLAST_NODETYPE_TRANSLATIONUNIT CHILLAST_NODETYPE_SOURCEFILE
+
+enum CHILL_FUNCTION_TYPE {
+ CHILL_FUNCTION_CPU = 0,
+ CHILL_FUNCTION_GPU
+};
+
+enum CHILL_MEMBER_EXP_TYPE {
+ CHILL_MEMBER_EXP_DOT = 0,
+ CHILL_MEMBER_EXP_ARROW
+};
+
+enum CHILL_PREPROCESSING_TYPE {
+ CHILL_PREPROCESSING_TYPEUNKNOWN = 0,
+ CHILL_PREPROCESSING_COMMENT,
+ CHILL_PREPROCESSING_POUNDDEFINE,
+ CHILL_PREPROCESSING_POUNDINCLUDE,
+ CHILL_PREPROCESSING_PRAGMA // unused so far
+};
+
+enum CHILL_PREPROCESSING_POSITION { // when tied to another statement
+ CHILL_PREPROCESSING_POSITIONUNKNOWN = 0,
+ CHILL_PREPROCESSING_LINEBEFORE, // previous line
+ CHILL_PREPROCESSING_LINEAFTER, // next line
+ CHILL_PREPROCESSING_TOTHERIGHT, // for this kind of comment, on same line
+ CHILL_PREPROCESSING_IMMEDIATELYBEFORE // on same line
+};
+
+
+
+
+
+extern const char* Chill_AST_Node_Names[]; // WARNING MUST BE KEPT IN SYNC WITH BELOW LIST
+
+
+// fwd declarations
+class chillAST_node; // the generic node. specific types derive from this
+class chillAST_NULL; // empty
+class chillAST_SourceFile; // ast for an entire source file (translationunit)
+
+class chillAST_TypedefDecl;
+class chillAST_VarDecl;
+//class chillAST_ParmVarDecl;
+class chillAST_FunctionDecl;
+class chillAST_RecordDecl; // structs and unions (and classes?)
+class chillAST_MacroDefinition;
+class chillAST_CompoundStmt; // just a bunch of other statements
+class chillAST_ForStmt; // AKA a LOOP
+class chillAST_TernaryOperator;
+class chillAST_BinaryOperator;
+class chillAST_ArraySubscriptExpr;
+class chillAST_MemberExpr;
+class chillAST_DeclRefExpr;
+class chillAST_IntegerLiteral;
+class chillAST_FloatingLiteral;
+class chillAST_UnaryOperator;
+class chillAST_ImplicitCastExpr;
+class chillAST_CStyleCastExpr;
+class chillAST_CStyleAddressOf;
+class chillAST_ReturnStmt;
+class chillAST_CallExpr;
+class chillAST_ParenExpr;
+class chillAST_Sizeof;
+class chillAST_Malloc;
+class chillAST_Free;
+class chillAST_NoOp;
+class chillAST_CudaMalloc;
+class chillAST_CudaFree;
+class chillAST_CudaMemcpy;
+class chillAST_CudaKernelCall;
+class chillAST_CudaSyncthreads;
+class chillAST_Preprocessing;
+
+typedef std::vector<chillAST_VarDecl *> chillAST_SymbolTable; // typedef
+typedef std::vector<chillAST_TypedefDecl *> chillAST_TypedefTable; // typedef
+
+bool symbolTableHasVariableNamed( chillAST_SymbolTable *table, const char *name ); // fwd decl
+chillAST_VarDecl *symbolTableFindVariableNamed( chillAST_SymbolTable *table, const char *name ); // fwd decl TODO too many similar named functions
+
+void printSymbolTable( chillAST_SymbolTable *st ); // fwd decl
+void printSymbolTableMoreInfo( chillAST_SymbolTable *st ); // fwd decl
+
+
+chillAST_node *lessthanmacro( chillAST_node *left, chillAST_node *right); // fwd declaration
+chillAST_SymbolTable *addSymbolToTable( chillAST_SymbolTable *st, chillAST_VarDecl *vd ); // fwd decl
+chillAST_TypedefTable *addTypedefToTable( chillAST_TypedefTable *tt, chillAST_TypedefDecl *td ); // fwd decl
+
+
+bool streq( const char *a, const char *b); // fwd decl
+void chillindent( int i, FILE *fp ); // fwd declaration
+void insertNewDeclAtLocationOfOldIfNeeded( chillAST_VarDecl *newdecl, chillAST_VarDecl *olddecl);
+
+chillAST_DeclRefExpr *buildDeclRefExpr( chillAST_VarDecl *);
+
+
+
+
+
+
+// an actual chill ast.
+// nodes based on clang AST which are in turn based on C++
+
+class chillAST_node { // generic node. a tree of these is the AST. this is virtual (can't instantiate)
+public:
+
+ static int chill_scalar_counter; // for manufactured scalars
+ static int chill_array_counter ; // for manufactured arrays
+ static int chill_pointer_counter ; // for manufactured arrays
+
+
+ CHILL_ASTNODE_TYPE asttype;
+
+ bool isSourceFile() { return (asttype == CHILLAST_NODETYPE_SOURCEFILE); };
+ bool isTypeDefDecl() { return (asttype == CHILLAST_NODETYPE_TYPEDEFDECL); };
+ bool isVarDecl() { return (asttype == CHILLAST_NODETYPE_VARDECL); };
+ bool isFunctionDecl() { return (asttype == CHILLAST_NODETYPE_FUNCTIONDECL); };
+ bool isRecordDecl() { return (asttype == CHILLAST_NODETYPE_RECORDDECL); };
+ bool isMacroDefinition() { return (asttype == CHILLAST_NODETYPE_MACRODEFINITION); };
+ bool isCompoundStmt() { return (asttype == CHILLAST_NODETYPE_COMPOUNDSTMT); };
+ bool isLoop() { return (asttype == CHILLAST_NODETYPE_LOOP); }; // AKA ForStmt
+ bool isForStmt() { return (asttype == CHILLAST_NODETYPE_LOOP); }; // AKA Loop
+ bool isIfStmt() { return (asttype == CHILLAST_NODETYPE_IFSTMT); };
+ bool isTernaryOperator() { return (asttype == CHILLAST_NODETYPE_TERNARYOPERATOR);};
+ bool isBinaryOperator() { return (asttype == CHILLAST_NODETYPE_BINARYOPERATOR); };
+ bool isUnaryOperator() { return (asttype == CHILLAST_NODETYPE_UNARYOPERATOR); };
+ bool isArraySubscriptExpr() { return (asttype == CHILLAST_NODETYPE_ARRAYSUBSCRIPTEXPR); };
+ bool isMemberExpr() { return (asttype == CHILLAST_NODETYPE_MEMBEREXPR); };
+ bool isDeclRefExpr() { return (asttype == CHILLAST_NODETYPE_DECLREFEXPR); };
+ bool isIntegerLiteral() { return (asttype == CHILLAST_NODETYPE_INTEGERLITERAL); };
+ bool isFloatingLiteral() { return (asttype == CHILLAST_NODETYPE_FLOATINGLITERAL); };
+ bool isImplicitCastExpr() { return (asttype == CHILLAST_NODETYPE_IMPLICITCASTEXPR); };
+ bool isReturnStmt() { return (asttype == CHILLAST_NODETYPE_RETURNSTMT); };
+ bool isCallExpr() { return (asttype == CHILLAST_NODETYPE_CALLEXPR); };
+ bool isParenExpr() { return (asttype == CHILLAST_NODETYPE_PARENEXPR); };
+ bool isSizeof() { return (asttype == CHILLAST_NODETYPE_SIZEOF); };
+ bool isMalloc() { return (asttype == CHILLAST_NODETYPE_MALLOC); };
+ bool isFree() { return (asttype == CHILLAST_NODETYPE_FREE); };
+ bool isPreprocessing() { return (asttype == CHILLAST_NODETYPE_PREPROCESSING); };
+ bool isNoOp() { return (asttype == CHILLAST_NODETYPE_NOOP); };
+ bool isNull() { return (asttype == CHILLAST_NODETYPE_NULL); };
+ bool isCStyleCastExpr() { return (asttype == CHILLAST_NODETYPE_CSTYLECASTEXPR); };
+ bool isCStyleAddressOf() { return (asttype == CHILLAST_NODETYPE_CSTYLEADDRESSOF); };
+ bool isCudaMalloc() { return (asttype == CHILLAST_NODETYPE_CUDAMALLOC); };
+ bool isCudaFree() { return (asttype == CHILLAST_NODETYPE_CUDAFREE); };
+ bool isCudaMemcpy() { return (asttype == CHILLAST_NODETYPE_CUDAMEMCPY); };
+ bool isCudaKERNELCALL() { return (asttype == CHILLAST_NODETYPE_CUDAKERNELCALL); };
+ bool isCudaSYNCTHREADS() { return (asttype == CHILLAST_NODETYPE_CUDASYNCTHREADS); };
+
+ bool isDeclStmt() { return (asttype == CHILLAST_NODETYPE_DECLSTMT); }; // doesn't exist
+
+ bool isConstant() { return (asttype == CHILLAST_NODETYPE_INTEGERLITERAL) || (asttype == CHILLAST_NODETYPE_FLOATINGLITERAL); }
+
+
+ virtual bool isAssignmentOp() { return false; };
+ virtual bool isComparisonOp() { return false; };
+ virtual bool isNotLeaf() { return false; };
+ virtual bool isLeaf() { return true; };
+ virtual bool isParmVarDecl() { return false; };
+
+ virtual bool isPlusOp() { return false; };
+ virtual bool isMinusOp() { return false; };
+ virtual bool isPlusMinusOp() { return false; };
+ virtual bool isMultDivOp() { return false; };
+
+ virtual bool isAStruct() { return false; };
+ virtual bool isAUnion() { return false; };
+
+ virtual bool hasSymbolTable() { return false; } ; // most nodes do NOT have a symbol table
+ virtual bool hasTypedefTable() { return false; } ; // most nodes do NOT have a typedef table
+ virtual chillAST_SymbolTable *getSymbolTable() { return NULL; } // most nodes do NOT have a symbol table
+
+ virtual chillAST_VarDecl *findVariableNamed( const char *name ); // recursive
+
+ chillAST_RecordDecl *findRecordDeclNamed( const char *name ); // recursive
+
+ // void addDecl( chillAST_VarDecl *vd); // recursive, adds to first symbol table it can find
+
+ // TODO decide how to hide some data
+ chillAST_node *parent;
+ bool isFromSourceFile; // false = #included
+ char *filename; // file this node is from
+
+ void segfault() { fprintf(stderr, "segfaulting on purpose\n"); int *i=0; int j = i[0]; }; // seg fault
+ int getNumChildren() { return children.size(); };
+ std::vector<chillAST_node*> children;
+ std::vector<chillAST_node*> getChildren() { return children; } ; // not usually useful
+ void setChildren( std::vector<chillAST_node*>&c ) { children = c; } ; // does not set parent. probably should
+ chillAST_node *getChild( int which) { return children[which]; };
+ void setChild( int which, chillAST_node *n ) { children[which] = n; children[which]->parent = this; } ;
+
+ char *metacomment; // for compiler internals, formerly a comment
+ void setMetaComment( char *c ) { metacomment = strdup(c); };
+
+ std::vector<chillAST_Preprocessing*> preprocessinginfo;
+
+ virtual void addChild( chillAST_node* c) {
+ //if (c->isFunctionDecl()) fprintf(stderr, "addchild FunctionDecl\n");
+ c->parent = this;
+ // check to see if it's already there
+ for (int i=0; i<children.size(); i++) {
+ if (c == children[i]) {
+ //fprintf(stderr, "addchild ALREADY THERE\n");
+ return; // already there
+ }
+ }
+ children.push_back(c);
+ } ; // not usually useful
+
+ virtual void insertChild(int i, chillAST_node* node) {
+ //fprintf(stderr, "%s inserting child of type %s at location %d\n", getTypeString(), node->getTypeString(), i);
+ node->parent = this;
+ children.insert( children.begin()+i, node );
+ };
+
+ virtual void removeChild(int i) {
+ children.erase( children.begin()+i );
+ };
+
+ int findChild( chillAST_node *c ) {
+ for (int i=0; i<children.size(); i++) {
+ if (children[i] == c) return i;
+ }
+ return -1;
+ }
+
+ virtual void replaceChild( chillAST_node *old, chillAST_node *newchild ) {
+ fprintf(stderr,"(%s) forgot to implement replaceChild() ... using generic\n" ,Chill_AST_Node_Names[asttype]);
+ fprintf(stderr, "%d children\n", children.size());
+ for (int i=0; i<children.size(); i++) {
+ if (children[i] == old) {
+ children[i] = newchild;
+ newchild->setParent( this );
+ return;
+ }
+ }
+ fprintf(stderr, "%s %p generic replaceChild called with oldchild that was not a child\n",
+ getTypeString(), this) ;
+ 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");
+ segfault(); // make easier for gdb
+ };
+
+ virtual void loseLoopWithLoopVar( char *var ) {
+ // walk tree. If a loop has this loop variable, replace the loop with the loop body,
+ // removing the loop. The loop will be spread across a bunch of cores that will each
+ // calculate their own loop variable.
+
+ // things that can not have loops as substatements can have a null version of this method
+ // things that have more complicated sets of "children" will have specialized versions
+
+ // this is the generic version of the method. It just recurses among its children.
+ // ForStmt is the only one that can actually remove itself. When it does, it will
+ // potentially change the children vector, which is not the simple array it might appear.
+ // so you have to make a copy of the vector to traverse
+
+ std::vector<chillAST_node*> dupe = children; // simple enough?
+ //fprintf(stderr, "node XXX has %d children\n", dupe.size());
+ //fprintf(stderr, "generic node %s has %d children\n", getTypeString(), dupe.size());
+ for (int i=0; i<dupe.size(); i++) { // recurse on all children
+ dupe[i]->loseLoopWithLoopVar( var );
+ }
+ }
+
+ virtual int evalAsInt() {
+ fprintf(stderr,"(%s) can't be evaluated as an integer??\n", Chill_AST_Node_Names[asttype]);
+ print(); fprintf(stderr, "\n");
+ segfault();
+ }
+
+ virtual const char* getUnderlyingType() {
+ fprintf(stderr,"(%s) forgot to implement getUnderlyingType()\n", Chill_AST_Node_Names[asttype]);
+ dump();
+ print();
+ fprintf(stderr, "\n\n");
+ segfault();
+ };
+
+ virtual chillAST_VarDecl* getUnderlyingVarDecl() {
+ fprintf(stderr,"(%s) forgot to implement getUnderlyingVarDecl()\n", Chill_AST_Node_Names[asttype]);
+ dump();
+ print();
+ fprintf(stderr, "\n\n");
+ segfault();
+ };
+
+
+ virtual chillAST_node *findref(){// find the SINGLE constant or data reference at this node or below
+ fprintf(stderr,"(%s) forgot to implement findref()\n" ,Chill_AST_Node_Names[asttype]);
+ dump();
+ print();
+ fprintf(stderr, "\n\n");
+ segfault();
+ };
+
+ virtual void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento ) {
+ fprintf(stderr,"(%s) forgot to implement gatherArrayRefs()\n" ,Chill_AST_Node_Names[asttype]);
+ dump();
+ print();
+ fprintf(stderr, "\n\n");
+ };
+
+ // TODO we MIGHT want the VarDecl // NOTHING IMPLEMENTS THIS? ???
+ virtual void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) {
+ fprintf(stderr,"(%s) forgot to implement gatherScalarRefs()\n" ,Chill_AST_Node_Names[asttype]);
+ dump();
+ print();
+ fprintf(stderr, "\n\n");
+ };
+
+ virtual void gatherLoopIndeces( std::vector<chillAST_VarDecl*> &indeces ) { // recursive walk parent links, looking for loops, and grabbing the declRefExpr in the loop init and cond.
+ // you can quit when you get to certain nodes
+
+ //fprintf(stderr, "%s::gatherLoopIndeces()\n", getTypeString());
+
+ if (isSourceFile() || isFunctionDecl() ) return; // end of the line
+
+ // just for debugging
+ //if (parent) {
+ // fprintf(stderr, "%s has parent of type %s\n", getTypeString(), parent->getTypeString());
+ //}
+ //else fprintf(stderr, "this %s %p has no parent???\n", getTypeString(), this);
+
+
+ if (!parent) return; // should not happen, but be careful
+
+ // for most nodes, this just recurses upwards
+ //fprintf(stderr, "%s::gatherLoopIndeces() %p recursing up\n", this);
+ parent->gatherLoopIndeces( indeces );
+ }
+
+
+ chillAST_ForStmt* findContainingLoop() { // recursive walk parent links, looking for loops
+ //fprintf(stderr, "%s::findContainingLoop() ", getTypeString());
+ //if (parent) fprintf(stderr, "parents is a %s\n", parent->getTypeString());
+ //else fprintf(stderr, "no parent\n");
+ // do not check SELF type, as we may want to find the loop containing a loop
+ if (!parent) return NULL;
+ if (parent->isForStmt()) return (chillAST_ForStmt*)parent;
+ return parent->findContainingLoop(); // recurse upwards
+ }
+
+ chillAST_node* findContainingNonLoop() { // recursive walk parent links, avoiding loops
+ fprintf(stderr, "%s::findContainingNonLoop() ", getTypeString());
+ //if (parent) fprintf(stderr, "parent is a %s\n", parent->getTypeString());
+ //else fprintf(stderr, "no parent\n");
+ // do not check SELF type, as we may want to find the loop containing a loop
+ if (!parent) return NULL;
+ if (parent->isCompoundStmt() && parent->getParent()->isForStmt()) return parent->getParent()->findContainingNonLoop(); // keep recursing
+ if (parent->isForStmt()) return parent->findContainingNonLoop(); // keep recursing
+ return (chillAST_node*)parent; // return non-loop
+ }
+
+ // TODO gather loop init and cond (and if cond) like gatherloopindeces
+
+ virtual void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs ){ // both scalar and arrays
+ fprintf(stderr,"(%s) forgot to implement gatherDeclRefExpr()\n" ,Chill_AST_Node_Names[asttype]);
+ };
+
+
+
+ virtual void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls ) {
+ fprintf(stderr,"(%s) forgot to implement gatherVarUsage()\n" ,Chill_AST_Node_Names[asttype]);
+ };
+
+ virtual void gatherVarLHSUsage( std::vector<chillAST_VarDecl*> &decls ) {
+ fprintf(stderr,"(%s) forgot to implement gatherVarLHSUsage()\n" ,Chill_AST_Node_Names[asttype]);
+ };
+
+
+ virtual void gatherVarDecls( std::vector<chillAST_VarDecl*> &decls ) { // ACTUAL Declaration
+ fprintf(stderr,"(%s) forgot to implement gatherVarDecls()\n" ,Chill_AST_Node_Names[asttype]);
+ };
+
+
+ virtual void gatherVarDeclsMore( std::vector<chillAST_VarDecl*> &decls ) { // even if the decl itself is not in the ast.
+ fprintf(stderr,"(%s) forgot to implement gatherVarDeclsMore()\n" ,Chill_AST_Node_Names[asttype]);
+ };
+
+ virtual void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls ) { // ACTUAL Declaration
+ fprintf(stderr,"(%s) forgot to implement gatherScalarVarDecls()\n" ,Chill_AST_Node_Names[asttype]);
+ };
+
+ virtual void gatherArrayVarDecls( std::vector<chillAST_VarDecl*> &decls ) { // ACTUAL Declaration
+ fprintf(stderr,"(%s) forgot to implement gatherArrayVarDecls()\n" ,Chill_AST_Node_Names[asttype]);
+ };
+
+ virtual chillAST_VarDecl *findArrayDecl( const char *name ) { // scoping TODO
+ if (!hasSymbolTable()) return parent->findArrayDecl( name ); // most things
+ else
+ fprintf(stderr,"(%s) forgot to implement gatherArrayVarDecls()\n" ,Chill_AST_Node_Names[asttype]);
+ }
+
+
+ virtual void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl) {
+ fprintf(stderr,"(%s) forgot to implement replaceVarDecls()\n" ,Chill_AST_Node_Names[asttype]);
+ };
+
+ virtual bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ) {
+ // this just looks for ForStmts with preferred index metacomment attached
+ fprintf(stderr,"(%s) forgot to implement findLoopIndexesToReplace()\n" ,Chill_AST_Node_Names[asttype]);
+ return false;
+ }
+
+
+ virtual chillAST_node* constantFold() { // hacky. TODO. make nice
+ fprintf(stderr,"(%s) forgot to implement constantFold()\n" ,Chill_AST_Node_Names[asttype]);
+ exit(-1); ;
+ };
+
+ virtual chillAST_node* clone() { // makes a deep COPY (?)
+ fprintf(stderr,"(%s) forgot to implement clone()\n" ,Chill_AST_Node_Names[asttype]);
+ exit(-1); ;
+ };
+ virtual void dump( int indent=0, FILE *fp = stderr ) {
+ fflush(fp);
+ fprintf(fp,"(%s) forgot to implement dump()\n" ,Chill_AST_Node_Names[asttype]); };// print ast
+
+ virtual void print( int indent=0, FILE *fp = stderr ) {
+ fflush(fp);
+ //fprintf(stderr, "generic chillAST_node::print() called!\n");
+ //fprintf(stderr, "asttype is %d\n", asttype);
+ fprintf(fp, "\n");
+ chillindent(indent, fp);
+ fprintf(fp,"(%s) forgot to implement print()\n" ,Chill_AST_Node_Names[asttype]);
+ };// print CODE
+
+ virtual void printName( int indent=0, FILE *fp = stderr ) {
+ fflush(fp);
+ //fprintf(stderr, "generic chillAST_node::printName() called!\n");
+ //fprintf(stderr, "asttype is %d\n", asttype);
+ fprintf(fp, "\n");
+ chillindent(indent, fp);
+ fprintf(fp,"(%s) forgot to implement printName()\n" ,Chill_AST_Node_Names[asttype]);
+ };// print CODE
+
+ virtual char *stringRep(int indent=0 ) { // the ast's print version
+ fflush(stdout);
+ // chillindent(indent, fp); TODO
+ fprintf(stderr,"(%s) forgot to implement stringRep()\n" ,Chill_AST_Node_Names[asttype]);
+ segfault();
+ }
+
+
+ virtual void printonly( int indent=0, FILE *fp = stderr ) { print( indent, fp); };
+
+ //virtual void printString( std::string &s ) {
+ // fprintf(stderr,"(%s) forgot to implement printString()\n" ,Chill_AST_Node_Names[asttype]);
+ //}
+
+
+ virtual void get_top_level_loops( std::vector<chillAST_ForStmt *> &loops) {
+ int n = children.size();
+ //fprintf(stderr, "get_top_level_loops of a %s with %d children\n", getTypeString(), n);
+ for (int i=0; i<n; i++) {
+ //fprintf(stderr, "child %d is a %s\n", i, children[i]->getTypeString());
+ if (children[i]->isForStmt()) {
+ loops.push_back( ((chillAST_ForStmt *)(children[i])) );
+ }
+ }
+ //fprintf(stderr, "found %d top level loops\n", loops.size());
+ }
+
+
+ virtual void repairParentChild() { // for nodes where all subnodes are children
+ int n = children.size();
+ for (int i=0; i<n; i++) {
+ if (children[i]->parent != this) {
+ fprintf(stderr, "fixing child %s that didn't know its parent\n", children[i]->getTypeString());
+ children[i]->parent = this;
+ }
+ }
+ }
+
+
+
+ virtual void get_deep_loops( std::vector<chillAST_ForStmt *> &loops) { // this is probably broken - returns ALL loops under it
+ int n = children.size();
+ //fprintf(stderr, "get_deep_loops of a %s with %d children\n", getTypeString(), n);
+ for (int i=0; i<n; i++) {
+ //fprintf(stderr, "child %d is a %s\n", i, children[i]->getTypeString());
+ children[i]->get_deep_loops( loops );
+ }
+ //fprintf(stderr, "found %d deep loops\n", loops.size());
+ }
+
+
+ // generic for chillAST_node with children
+ virtual void find_deepest_loops( std::vector<chillAST_ForStmt *> &loops) { // returns DEEPEST nesting of loops
+ std::vector<chillAST_ForStmt *>deepest; // deepest below here
+
+ int n = children.size();
+ //fprintf(stderr, "find_deepest_loops of a %s with %d children\n", getTypeString(), n);
+ for (int i=0; i<n; i++) {
+ std::vector<chillAST_ForStmt *> subloops; // loops below here among a child of mine
+
+ //fprintf(stderr, "child %d is a %s\n", i, children[i]->getTypeString());
+ children[i]->find_deepest_loops( subloops );
+
+ if (subloops.size() > deepest.size()) {
+ deepest = subloops;
+ }
+ }
+
+ // append deepest we see at this level to loops
+ for ( int i=0; i<deepest.size(); i++) {
+ loops.push_back( deepest[i] );
+ }
+
+ //fprintf(stderr, "found %d deep loops\n", loops.size());
+
+ }
+
+
+
+
+ const char *getTypeString() { return Chill_AST_Node_Names[asttype]; } ;
+ int getType() { return asttype; };
+ void setParent( chillAST_node *p) { parent = p; } ;
+ chillAST_node *getParent() { return parent; } ;
+
+ chillAST_SourceFile *getSourceFile() {
+ if (isSourceFile()) return ((chillAST_SourceFile *)this);
+ if (parent != NULL) return parent->getSourceFile();
+ fprintf(stderr, "UHOH, getSourceFile() called on node %p %s that does not have a parent and is not a source file\n", this, this->getTypeString());
+ this->print(); printf("\n\n"); fflush(stdout);
+ exit(-1);
+ }
+
+ virtual chillAST_node *findDatatype( char *t ) {
+ fprintf(stderr, "%s looking for datatype %s\n", getTypeString(), t);
+ if (parent != NULL) return parent->findDatatype(t); // most nodes do this
+ return NULL;
+ }
+
+
+ virtual chillAST_SymbolTable *addVariableToSymbolTable( chillAST_VarDecl *vd ) {
+ if (!parent) {
+ fprintf(stderr, "%s with no parent addVariableToSymbolTable()\n", getTypeString());
+ exit(-1);
+ }
+ //fprintf(stderr, "%s::addVariableToSymbolTable() (default) headed up\n", getTypeString());
+ return parent->addVariableToSymbolTable( vd ); // default, defer to parent
+ }
+
+ virtual void addTypedefToTypedefTable( chillAST_TypedefDecl *tdd ) {
+ parent->addTypedefToTypedefTable( tdd ); // default, defer to parent
+ }
+
+ void walk_parents() {
+ fprintf(stderr, "wp: (%s) ", getTypeString());
+ print(); printf("\n"); fflush(stdout);
+ if (isSourceFile()) { fprintf(stderr, "(top sourcefile)\n\n"); return;}
+
+ if (parent) parent->walk_parents();
+ else fprintf(stderr, "UHOH, %s has no parent??\n", getTypeString());
+ return;
+ }
+
+ virtual chillAST_node *getEnclosingStatement( int level = 0 );
+ virtual chillAST_VarDecl *multibase() {
+ fprintf(stderr,"(%s) forgot to implement multibase()\n", Chill_AST_Node_Names[asttype]);
+ exit(-1);
+ }
+ virtual chillAST_node *multibase2() {
+ fprintf(stderr,"(%s) forgot to implement multibase2()\n", Chill_AST_Node_Names[asttype]);
+ exit(-1);
+ }
+
+
+ virtual void gatherStatements( std::vector<chillAST_node*> &statements ) {
+ fprintf(stderr,"(%s) forgot to implement gatherStatements()\n" ,Chill_AST_Node_Names[asttype]);
+ dump();fflush(stdout);
+ print();
+ fprintf(stderr, "\n\n");
+ }
+
+
+ virtual bool isSameAs( chillAST_node *other ){ // for tree comparison
+ fprintf(stderr,"(%s) forgot to implement isSameAs()\n" ,Chill_AST_Node_Names[asttype]);
+ dump(); fflush(stdout);
+ print();
+ fprintf(stderr, "\n\n"); }
+
+ void printPreprocBEFORE( int indent, FILE *fp );
+ void printPreprocAFTER( int indent, FILE *fp );
+
+
+
+};
+
+class chillAST_NULL: public chillAST_node { // NOOP?
+public:
+ chillAST_NULL(chillAST_node *p = NULL) { parent = p; asttype = CHILLAST_NODETYPE_NULL; };
+ void print( int indent=0, FILE *fp = stderr ) {
+ chillindent( indent, fp );
+ fprintf(fp, "/* (NULL statement); */ ");
+ fflush(fp);
+ }
+ void dump( int indent=0, FILE *fp = stderr ) {
+ chillindent( indent, fp );
+ fprintf(fp, "(NULL statement) "); fflush(fp);
+ }
+};
+
+
+class chillAST_Preprocessing: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ CHILL_PREPROCESSING_POSITION position;
+ CHILL_PREPROCESSING_TYPE pptype;
+ char *blurb;
+
+ // constructors
+ chillAST_Preprocessing(); // not sure what this is good for
+ chillAST_Preprocessing( CHILL_PREPROCESSING_POSITION pos, CHILL_PREPROCESSING_TYPE t, char *text );
+
+ // other methods particular to this type of node
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ //void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+};
+
+
+//typedef is a keyword in the C and C++ programming languages. The purpose of typedef is to assign alternative names to existing types, most often those whose standard declaration is cumbersome, potentially confusing, or likely to vary from one implementation to another.
+class chillAST_TypedefDecl: public chillAST_node {
+private:
+ bool isStruct;
+ bool isUnion;
+ char *structname; // get rid of this?
+
+public:
+ char *newtype; // the new type name ??
+ char *underlyingtype; // float, int, "struct bubba" ?
+ char *arraypart; // string like "[1234][56]" ??
+
+ chillAST_RecordDecl *rd; // if it's a struct, point to the recorddecl ??
+ // TODO what if "typedef int[10] tenints; " ??
+ void setStructInfo( chillAST_RecordDecl *arrdee ) { rd = arrdee; };
+ chillAST_RecordDecl * getStructDef();
+
+
+ bool isAStruct() { return isStruct; };
+ bool isAUnion() { return isUnion; };
+ void setStruct(bool tf) { isStruct = tf; fprintf(stderr, "%s isStruct %d\n", structname, isStruct); };
+ void setUnion( bool tf) { isUnion = tf; };
+ void setStructName( const char *newname) { structname = strdup(newname); };
+ char *getStructName( ) { return structname; };
+ bool nameis( const char *n ) { return !strcmp(n, structname); };
+
+ // special for struct/unions rethink TODO
+ std::vector<chillAST_VarDecl *> subparts;
+ chillAST_VarDecl *findSubpart( const char *name );
+
+ //TODO hide data, set/get type and alias
+ chillAST_TypedefDecl();
+ chillAST_TypedefDecl(char *t, char *nt, chillAST_node *p);
+ chillAST_TypedefDecl(char *t, char *nt, char *a, chillAST_node *par);
+
+ const char* getUnderlyingType() { fprintf(stderr, "TypedefDecl getUnderLyingType()\n"); return underlyingtype; };
+ //virtual chillAST_VarDecl* getUnderlyingVarDecl() { return this; }; // ??
+
+ void dump( int indent=0, FILE *fp = stderr ) {
+ fprintf(fp, "(TypedefDecl %s %s %s)\n", underlyingtype, newtype, arraypart); };
+ void print( int indent=0, FILE *fp = stderr ) ;
+ //void printString( std::string &s );
+
+};
+
+
+class chillAST_VarDecl: public chillAST_node {
+public:
+ char *vartype; // should probably be an enum, except it's used for unnamed structs too
+
+ chillAST_RecordDecl *vardef;// the thing that says what the struct looks like
+ chillAST_TypedefDecl *typedefinition; // NULL for float, int, etc.
+ chillAST_RecordDecl * getStructDef(); // TODO make vardef private?
+
+ //bool insideAStruct; // this variable is itself part of a struct
+
+ char *underlyingtype;
+ char *varname;
+ char *arraypart; // [ 12 ] [ 34 ] if that is how it was defined
+ char *arraypointerpart;
+ char *arraysetpart;
+ void splitarraypart();
+
+ int numdimensions;
+ int *arraysizes; // TODO
+ bool knownArraySizes; // if this float *a or float a[128] ? true means we know ALL dimensions
+ int cudamallocsize; // usually specified in lua/python transformation file
+
+ bool isRestrict; // C++ __restrict__
+ bool isShared; // CUDA __shared__
+ bool isDevice; // CUDA __device__
+ bool isStruct;
+
+ int isAParameter;
+ bool byreference;
+ void setByReference( bool tf ) { byreference = tf; fprintf(stderr, "byref %d\n", tf); };
+
+ bool isABuiltin; // if variable is builtin, we don't need to declare it
+ void *uniquePtr; // DO NOT REFERENCE THROUGH THIS! just used to differentiate declarations
+ bool isArray() { return (numdimensions != 0); };
+ bool isAStruct() { return (isStruct || (typedefinition && typedefinition->isAStruct())); }
+ void setStruct( bool b ) {isStruct = b;/*fprintf(stderr,"vardecl %s IS A STRUCT\n",varname);*/ };
+ bool isPointer() { return isArray() && !knownArraySizes; } //
+
+ bool knowAllDimensions() { return knownArraySizes; } ;
+
+ chillAST_node *init;
+ void setInit( chillAST_node *i ) { init = i; i->setParent(this); };
+ bool hasInit() { return init != NULL; };
+ chillAST_node *getInit() { return init; };
+
+ chillAST_VarDecl();
+ chillAST_VarDecl( const char *t, const char *n, const char *a, chillAST_node *p);
+ chillAST_VarDecl( const char *t, const char *n, const char *a, void *ptr, chillAST_node *p);
+ chillAST_VarDecl( chillAST_TypedefDecl *tdd, const char *n, const char *arraypart, chillAST_node *par);
+ chillAST_VarDecl( chillAST_RecordDecl *astruct, const char *n, const char *arraypart, chillAST_node *par);
+
+ void dump( int indent=0, FILE *fp = stderr );
+ void print( int indent=0, FILE *fp = stderr );
+ void printName( int indent=0, FILE *fp = stderr );
+ bool isParmVarDecl() { return( isAParameter == 1 ); };
+ bool isBuiltin() { return( isABuiltin == 1 ); }; // designate variable as a builtin
+ void setLocation( void *ptr ) { uniquePtr = ptr; } ;
+
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls ) {}; // does nothing
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs ) {}; // does nothing
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl) {};
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+ const char* getUnderlyingType() { /* fprintf(stderr, "VarDecl getUnderLyingType()\n"); */return underlyingtype; };
+ virtual chillAST_VarDecl* getUnderlyingVarDecl() { return this; };
+
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+
+};
+
+
+class chillAST_DeclRefExpr: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ char *declarationType;
+ char *declarationName;
+ chillAST_node *decl; // the declaration of this variable or function ... uhoh
+ //char *functionparameters; // TODO probably should split this node into 2 types, one for variables, one for functions
+
+ // constructors
+ chillAST_DeclRefExpr();
+ chillAST_DeclRefExpr( const char *variablename, chillAST_node *p);
+ chillAST_DeclRefExpr( const char *vartype, const char *variablename, chillAST_node *p);
+ chillAST_DeclRefExpr( const char *vartype, const char *variablename, chillAST_node *dec, chillAST_node *p);
+ chillAST_DeclRefExpr( chillAST_VarDecl *vd, chillAST_node *p=NULL);
+ chillAST_DeclRefExpr( chillAST_FunctionDecl *fd, chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+ bool operator!=( chillAST_DeclRefExpr &other ) { return decl != other.decl ; };
+ bool operator==( chillAST_DeclRefExpr &other ) { return decl == other.decl ; }; // EXACT SAME VARECL BY ADDRESS
+
+ chillAST_node *getDecl() { return decl; };
+
+ chillAST_VarDecl *getVarDecl() {
+ if (!decl) return NULL; // should never happen
+ if (decl->isVarDecl()) return (chillAST_VarDecl *)decl;
+ return NULL;
+ };
+
+ chillAST_FunctionDecl *getFunctionDecl() {
+ if (!decl) return NULL; // should never happen
+ if (decl->isFunctionDecl()) return (chillAST_FunctionDecl *)decl;
+ return NULL;
+ };
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE
+ void dump( int indent=0, FILE *fp = stderr ); // print ast
+ char *stringRep(int indent=0 );
+
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento ) {}; // do nothing
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento );
+
+ // this is the AST node where these 2 differ
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls ) {}; // does nothing, to get the cvardecl using this method, the actual vardecl must be in the AST
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ); // returns the decl this declrefexpr references, even if the decl is not in the AST
+
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+ chillAST_node *findref(){return this;}// find the SINGLE constant or data reference at this node or below
+
+ const char* getUnderlyingType() {fprintf(stderr, "DeclRefExpr getUnderLyingType()\n"); return decl->getUnderlyingType();};
+
+ virtual chillAST_VarDecl* getUnderlyingVarDecl() { return decl->getUnderlyingVarDecl(); } // functions?? TODO
+
+ chillAST_VarDecl* multibase();
+ chillAST_node *multibase2() { return (chillAST_node *)multibase(); }
+};
+
+
+
+
+
+class chillAST_CompoundStmt: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_SymbolTable *symbol_table; // symbols defined inside this compound statement
+ chillAST_TypedefTable *typedef_table;
+
+ bool hasSymbolTable() { return true; } ;
+ bool hasTypeDefTable() { return true; } ;
+ chillAST_node *findDatatype( char *t ) {
+ fprintf(stderr, "chillAST_CompoundStmt::findDatatype( %s )\n", t);
+ if (typedef_table) {
+ for (int i=0; i< typedef_table->size(); i++) {
+ chillAST_TypedefDecl *tdd = (*typedef_table)[i];
+ if (tdd->nameis( t )) return tdd;
+ }
+ }
+ if (parent) return parent->findDatatype(t);
+ return NULL; // should not happen
+ }
+
+ chillAST_SymbolTable *getSymbolTable() { return symbol_table; }
+
+ chillAST_SymbolTable* addVariableToSymbolTable( chillAST_VarDecl *vd ) { // chillAST_CompoundStmt method
+ //fprintf(stderr, "\nchillAST_CompoundStmt addVariableToSymbolTable( %s )\n", vd->varname);
+ symbol_table = addSymbolToTable( symbol_table, vd );
+ //printSymbolTable( symbol_table );
+ return symbol_table;
+ }
+
+ void addTypedefToTypedefTable( chillAST_TypedefDecl *tdd ) {
+ typedef_table = addTypedefToTable( typedef_table, tdd );
+ }
+
+ // constructors
+ chillAST_CompoundStmt(); // never has any args ???
+
+ // other methods particular to this type of node
+
+
+ // required methods
+ void replaceChild( chillAST_node *old, chillAST_node *newchild );
+ void dump( int indent=0, FILE *fp = stderr );
+ void print( int indent=0, FILE *fp = stderr );
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false );
+ void loseLoopWithLoopVar( char *var ); // special case this for not for debugging
+
+ void gatherStatements( std::vector<chillAST_node*> &statements );
+};
+
+
+
+
+class chillAST_RecordDecl: public chillAST_node { // declaration of the shape of a struct or union
+private:
+ char *name; // could be NULL? for unnamed structs?
+ char *originalname;
+ bool isStruct;
+
+ bool isUnion;
+ std::vector<chillAST_VarDecl *> subparts;
+
+public:
+ chillAST_RecordDecl();
+ chillAST_RecordDecl( const char *nam, chillAST_node *p );
+ chillAST_RecordDecl( const char *nam, const char *orig, chillAST_node *p );
+
+ void setName( const char *newname) { name = strdup(newname); };
+ char *getName( ) { return name; };
+
+ bool isAUnion() { return isUnion; };
+ bool isAStruct() { return isStruct; };
+ bool isUnnamed;
+ void setUnnamed( bool b ) { isUnnamed = b; };
+
+
+ void setStruct(bool tf) { isStruct = tf; };
+ //fprintf(stderr, "%s isStruct %d\n", structname, isStruct); };
+ void setUnion( bool tf) { isUnion = tf; };
+
+ chillAST_SymbolTable *addVariableToSymbolTable( chillAST_VarDecl *vd ); // RecordDecl does NOTHING
+
+ int numSubparts() { return subparts.size(); };
+ void addSubpart( chillAST_VarDecl *s ) { subparts.push_back(s); };
+ chillAST_VarDecl *findSubpart( const char *name );
+ chillAST_VarDecl *findSubpartByType( const char *typ );
+
+ void dump( int indent=0, FILE *fp = stderr );
+ void print( int indent=0, FILE *fp = stderr ) ;
+ void printStructure( int indent=0, FILE *fp = stderr ) ;
+};
+
+
+
+
+class chillAST_FunctionDecl: public chillAST_node {
+private:
+ chillAST_CompoundStmt *body; // always a compound statement?
+ CHILL_FUNCTION_TYPE function_type; // CHILL_FUNCTION_CPU or CHILL_FUNCTION_GPU
+ bool externfunc; // function is external
+ bool builtin; // function is a builtin
+ bool forwarddecl;
+
+public:
+ char *returnType;
+ char *functionName;
+
+ // parameters
+ int numParameters() { return parameters.size(); } ;
+ chillAST_SymbolTable parameters;
+
+ // this is probably a mistake, but symbol_table here is pointing to BODY'S symbol table
+ //chillAST_SymbolTable *symbol_table; // symbols defined inside this function. REALLY the body's symbol table?
+
+ chillAST_TypedefTable *typedef_table; // function typedef table
+
+
+ bool hasSymbolTable() { return true; } ; // COULD HAVE
+ bool hasTypeDefTable(){ return true; } ; // COULD HAVE
+
+
+ //char *parametertypes; // a single string??
+ void printParameterTypes( FILE *fp );
+ void setName( char *n ) { functionName = strdup(n); /* probable memory leak */ };
+
+ void setBuiltin() { builtin = true; } ; // designate function as a builtin
+ bool isBuiltin() { return builtin; } ; // report whether is a builtin
+
+ void setExtern() { externfunc = true; }; // designate function as external
+ bool isExtern() { return externfunc; }; // report whether function is external
+
+ void setForward() { forwarddecl = true; }; // designate function as fwd declaration
+ bool isForward() { return forwarddecl; }; // report whether function is external
+
+ bool isFunctionCPU() { return( function_type == CHILL_FUNCTION_CPU ); };
+ bool isFunctionGPU() { return( function_type == CHILL_FUNCTION_GPU ); };
+ void setFunctionCPU() { function_type = CHILL_FUNCTION_CPU; };
+ void setFunctionGPU() { function_type = CHILL_FUNCTION_GPU; };
+
+ void *uniquePtr; // DO NOT REFERENCE THROUGH THIS! USED AS A UNIQUE ID
+
+
+
+
+ chillAST_FunctionDecl(); // { asttype = CHILLAST_NODETYPE_FUNCTIONDECL; numparameters = 0;};
+ chillAST_FunctionDecl(const char *rt, const char *fname, chillAST_node *p=NULL ) ;
+ chillAST_FunctionDecl(const char *rt, const char *fname, chillAST_node *p, void *unique ) ;
+
+ void addParameter( chillAST_VarDecl *p);
+ chillAST_VarDecl *hasParameterNamed( const char *name );
+ chillAST_VarDecl *findParameterNamed( const char *name ) { return hasParameterNamed( name ); };
+
+ void addDecl( chillAST_VarDecl *vd); // just adds to symbol table?? TODO
+
+ chillAST_VarDecl *funcHasVariableNamed( const char *name ); // functiondecl::hasVariableNamed
+ //chillAST_VarDecl *findVariableNamed( const char *name ) { return hasVariableNamed( name ); };
+
+ void addChild(chillAST_node* node); // special because inserts into BODY
+ void insertChild(int i, chillAST_node* node); // special because inserts into BODY
+
+ void setBody( chillAST_node * bod );
+ chillAST_CompoundStmt *getBody() { return( body); }
+
+ void print( int indent=0, FILE *fp = stderr ); // in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // in chill_ast.cc
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ chillAST_VarDecl *findArrayDecl( const char *name );
+ //void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ //void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento )
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void cleanUpVarDecls();
+
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false );
+
+ chillAST_node* constantFold();
+
+ chillAST_node *findDatatype( char *t ) {
+ fprintf(stderr, "%s looking for datatype %s\n", getTypeString(), t);
+ if (!typedef_table) { // not here
+ if (parent) return parent->findDatatype(t); // not here, check parents
+ else return NULL; // not defined here and no parent
+ }
+
+ //fprintf(stderr, "%d typedefs\n", typedef_table->size());
+ for (int i=0; i< typedef_table->size(); i++) {
+ chillAST_TypedefDecl *tdd = (*typedef_table)[i];
+ if ( tdd->nameis( t )) return tdd;
+ }
+ if (parent) return parent->findDatatype(t);
+ return NULL; // should not happen
+ }
+
+ chillAST_SymbolTable *getParameterSymbolTable() { return &parameters; }
+ chillAST_SymbolTable *getSymbolTable() { return body->getSymbolTable(); } //symbol_table; } //
+ void setSymbolTable( chillAST_SymbolTable *tab ) {
+ // no longer keeping a local ?? symbol_table = tab;
+ if (!body) { // can never happen now
+ body = new chillAST_CompoundStmt();
+ } // only if func is empty!
+ body->symbol_table = tab;
+ }
+
+ chillAST_SymbolTable* addVariableToSymbolTable( chillAST_VarDecl *vd ) { // chillAST_FunctionDecl method
+ //fprintf(stderr, "\nchillAST_FunctionDecl addVariableToSymbolTable( %s )\n", vd->varname);
+
+ // this is all dealing with the body's symbol table
+ // the function has a symbol table called "parameters" but that is a special case
+
+ addSymbolToTable( getSymbolTable(), vd );
+ if (!vd->parent) {
+ //fprintf(stderr, "setting parent of vardecl to be the function whose symbol table it is going into\n"); // ??
+ vd->setParent( this );
+ insertChild(0,vd);
+ }
+ //printSymbolTable( getSymbolTable() );
+ return getSymbolTable();
+ }
+
+
+ void addTypedefToTypedefTable( chillAST_TypedefDecl *tdd ) {
+ typedef_table = addTypedefToTable( typedef_table, tdd );
+ }
+
+ void replaceChild( chillAST_node *old, chillAST_node *newchild ) {
+ body->replaceChild( old, newchild );
+ }
+}; // end FunctionDecl
+
+
+
+
+class chillAST_SourceFile: public chillAST_node {
+public:
+
+ // constructors
+ chillAST_SourceFile(); // defined in chill_ast.cc
+ chillAST_SourceFile(const char *filename ); // defined in chill_ast.cc
+
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void printToFile( char *filename = NULL );
+
+ char *SourceFileName; // where this originated
+ char *FileToWrite;
+ char *frontend;
+
+ void setFileToWrite( char *f ) { FileToWrite = strdup( f ); };
+
+ void setFrontend( const char *compiler ) { if (frontend) free(frontend); frontend = strdup(compiler); }
+ // get, set filename ?
+
+ chillAST_SymbolTable *global_symbol_table; // (global) symbols defined inside this source file
+ chillAST_TypedefTable *global_typedef_table; // source file
+ chillAST_VarDecl *findVariableNamed( const char *name ); // looks in global_symbol_table;
+
+ bool hasSymbolTable() { return true; } ; // "has" vs "can have" TODO
+ bool hasTypeDefTable() { return true; } ;
+
+ chillAST_SymbolTable* addVariableToSymbolTable( chillAST_VarDecl *vd ) { // chillAST_SourceFile method
+ fprintf(stderr, "\nchillAST_SourceFile addVariableToSymbolTable( %s )\n", vd->varname);
+ global_symbol_table = addSymbolToTable( global_symbol_table, vd );
+ //addChild( vd ); // ??
+ //printSymbolTable( global_symbol_table );
+ return global_symbol_table;
+ }
+
+ void addTypedefToTypedefTable( chillAST_TypedefDecl *tdd ) {
+ //fprintf(stderr, "SOURCEFILE adding typedef %s to typedeftable\n", tdd->getStructName());
+ global_typedef_table = addTypedefToTable( global_typedef_table, tdd );
+ //fprintf(stderr, "now global typedef table has %d entries\n", global_typedef_table->size());
+ }
+
+ chillAST_node *findDatatype( char *t ) {
+ fprintf(stderr, "%s looking for datatype %s\n", getTypeString(), t);
+ fprintf(stderr, "%d global typedefs\n", global_typedef_table->size());
+ for (int i=0; i< global_typedef_table->size(); i++) {
+
+ chillAST_TypedefDecl *tdd = (*global_typedef_table)[i];
+ //fprintf(stderr, "comparing to %s\n", tdd->getStructName());
+ if (tdd->nameis( t )) {
+ //fprintf(stderr, "found it\n");
+ return (chillAST_node *)tdd;
+ }
+ }
+ return NULL;
+ }
+
+ std::vector< chillAST_FunctionDecl *> functions; // at top level, or anywhere?
+ std::vector< chillAST_MacroDefinition *> macrodefinitions;
+
+ chillAST_MacroDefinition* findMacro( const char *name ); // TODO ignores arguments
+ chillAST_FunctionDecl *findFunction( const char *name ); // TODO ignores arguments
+ chillAST_node *findCall( const char *name );
+ void addMacro(chillAST_MacroDefinition* md) {
+ macrodefinitions.push_back(md);
+ //fprintf(stderr, "addMacro(), now %d macros\n", macrodefinitions.size());
+ }
+ void addFunc(chillAST_FunctionDecl* fd) {
+ //fprintf(stderr, "chillAST_SourceFile::addFunc( %s %p)\n", fd->functionName, fd);
+
+ bool already = false;
+ for (int i=0; i<functions.size(); i++) {
+ //fprintf(stderr, "function %d is %s %p\n", i, functions[i]->functionName, functions[i]);
+ if (functions[i] == fd) {
+ //fprintf(stderr, "function %s was already in source functions\n", fd->functionName);
+ already = true;
+ }
+ }
+ if (!already) functions.push_back(fd);
+
+ // PROBABLY fd was created with sourcefile as its parent. Don't add it twice
+ addChild( (chillAST_node *)fd); }
+
+};
+
+
+/*
+ class chillAST_VarDecl: public chillAST_node { // now a SINGLE DECL. multiples in
+ public:
+ int howmany; // usually 1 but sometimes multiple declarations in a decl;
+ std::vector<class chillAST_SingleVarDecl*> decls;
+
+ chillAST_VarDecl();
+ chillAST_VarDecl( char *t, char *n, char *a);
+ void addDecl( char *t, char *n, char *a);
+
+ void dump( int indent=0, FILE *fp = stderr );
+ void print( int indent=0, FILE *fp = stderr );
+ };
+*/
+
+
+/*
+class chillAST_ParmVarDecl: public chillAST_node { // no longer used?
+public:
+ char *vartype; // should probably be an enum
+ char *varname;
+ char *arraypart;
+ int numdimensions; // TODO
+ int *arraysizes; // TODO
+ // hasDefaultArg
+ // getDefaultArg
+
+ chillAST_ParmVarDecl();
+ chillAST_ParmVarDecl( const char *type, const char *name, const char *ap, chillAST_node *p );
+
+ void dump( int indent=0, FILE *fp = stderr ) {
+ fprintf(fp, "(2VarDecl'%s' '%s' '%s')", vartype, varname, arraypart); };
+ void print( int indent=0, FILE *fp = stderr );
+};
+*/
+
+
+
+class chillAST_MacroDefinition: public chillAST_node {
+private:
+ chillAST_node *body; // rhs always a compound statement?
+ chillAST_SymbolTable *symbol_table;
+public:
+ char *macroName;
+ char *rhsString;
+
+ // parameters - these will be odd, in that they HAVE NO TYPE
+ int numParameters() { return parameters.size(); } ;
+ std::vector<chillAST_VarDecl *>parameters;
+
+ void setName( char *n ) { macroName = strdup(n); /* probable memory leak */ };
+ void setRhsString( char *n ) { rhsString = strdup(n); /* probable memory leak */ };
+ char *getRhsString() { return rhsString; }
+
+ chillAST_MacroDefinition();
+ chillAST_MacroDefinition( const char *name, chillAST_node *par);
+ chillAST_MacroDefinition( const char *name, const char *rhs, chillAST_node *par);
+
+ void addParameter( chillAST_VarDecl *p); // parameters have no TYPE ??
+ chillAST_VarDecl *hasParameterNamed( const char *name );
+ chillAST_VarDecl *findParameterNamed( const char *name ) { return hasParameterNamed( name ); };
+ void addChild(chillAST_node* node); // special because inserts into BODY
+ void insertChild(int i, chillAST_node* node); // special because inserts into BODY
+
+ void setBody( chillAST_node * bod );
+ chillAST_node *getBody() { return( body); }
+
+ void print( int indent=0, FILE *fp = stderr ); // in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // in chill_ast.cc
+
+ bool hasSymbolTable() { return true; } ;
+
+ //const std::vector<chillAST_VarDecl *> getSymbolTable() { return symbol_table; }
+ chillAST_SymbolTable *getSymbolTable() { return symbol_table; }
+ chillAST_SymbolTable* addVariableToSymbolTable( chillAST_VarDecl *vd ) { // chillAST_MacroDefinition method ??
+ //fprintf(stderr, "\nchillAST_MacroDefinition addVariableToSymbolTable( %s )\n", vd->varname);
+ symbol_table = addSymbolToTable( symbol_table, vd );
+ //printSymbolTable( symbol_table );
+ return symbol_table;
+ }
+
+
+ chillAST_node* clone();
+
+ // none of these make sense for macros
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls ){};
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls ){};
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls ){};
+ chillAST_VarDecl *findArrayDecl( const char *name ){};
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls ){};
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs ){};
+ void cleanUpVarDecls();
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){};
+ chillAST_node* constantFold(){};
+};
+
+
+
+
+
+
+class chillAST_ForStmt: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *init;
+ chillAST_node *cond;
+ chillAST_node *incr;
+ chillAST_node *body; // always a compoundstmt?
+ IR_CONDITION_TYPE conditionoperator; // from ir_code.hh
+
+ chillAST_SymbolTable *symbol_table; // symbols defined inside this forstmt (in init but not body?) body is compound stmt
+ bool hasSymbolTable() { return true; } ;
+
+ // constructors
+ chillAST_ForStmt();
+ chillAST_ForStmt( chillAST_node *ini, chillAST_node *con, chillAST_node *inc, chillAST_node *bod, chillAST_node *p);
+
+ // other methods particular to this type of node
+ void addSyncs();
+ void removeSyncComment();
+ chillAST_node *getInit() { return init; };
+ chillAST_node *getCond() { return cond; };
+ chillAST_node *getInc() { return incr; };
+ chillAST_node *getBody() { //fprintf(stderr, "chillAST_ForStmt::getBody(), returning a chillAST_node of type %s\n", body->getTypeString());
+ return body; };
+ void setBody( chillAST_node *b ) { body = b; b->parent = this; };
+
+ bool isNotLeaf() { return true; };
+ bool isLeaf() { return false; };
+
+
+ // required methods that I can't seem to get to inherit
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void printControl( int indent=0, FILE *fp = stderr ); // print just for ( ... ) but not body
+
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl); // will get called on inner loops
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false );
+
+ void gatherLoopIndeces( std::vector<chillAST_VarDecl*> &indeces );
+ void gatherLoopVars( std::vector<std::string> &loopvars ); // gather as strings ??
+
+ void get_deep_loops( std::vector<chillAST_ForStmt *> &loops) { // chillAST_ForStmt version
+ // ADD MYSELF!
+ loops.push_back( this );
+
+ int n = body->children.size();
+ //fprintf(stderr, "get_deep_loops of a %s with %d children\n", getTypeString(), n);
+ for (int i=0; i<n; i++) {
+ //fprintf(stderr, "child %d is a %s\n", i, body->children[i]->getTypeString());
+ body->children[i]->get_deep_loops( loops );
+ }
+ //fprintf(stderr, "found %d deep loops\n", loops.size());
+ }
+
+
+ void find_deepest_loops( std::vector<chillAST_ForStmt *> &loops) {
+ std::vector<chillAST_ForStmt *> b; // deepest loops below me
+
+ int n = body->children.size();
+ for (int i=0; i<n; i++) {
+ std::vector<chillAST_ForStmt *> l; // deepest loops below one child
+ body->children[i]->find_deepest_loops( l );
+ if ( l.size() > b.size() ) { // a deeper nesting than we've seen
+ b = l;
+ }
+ }
+
+ loops.push_back( this ); // add myself
+ for (int i=0; i<b.size(); i++) loops.push_back(b[i]);
+ }
+
+
+ void loseLoopWithLoopVar( char *var ); // chillAST_ForStmt
+ void replaceChild( chillAST_node *old, chillAST_node *newchild ) ;
+
+ chillAST_SymbolTable* addVariableToSymbolTable( chillAST_VarDecl *vd ) { // chillAST_ForStmt method
+ //fprintf(stderr, "\nchillAST_ForStmt addVariableToSymbolTable( %s )\n", vd->varname);
+ symbol_table = addSymbolToTable( symbol_table, vd );
+ //printSymbolTable( symbol_table );
+ return symbol_table;
+ }
+
+ void gatherStatements( std::vector<chillAST_node*> &statements );
+ bool lowerBound( int &l );
+ bool upperBound( int &u );
+
+};
+
+
+
+class chillAST_TernaryOperator: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ char *op; // TODO need enum so far, only "?" conditional operator
+ chillAST_node *condition;
+ chillAST_node *lhs; // keep name from binary
+ chillAST_node *rhs;
+
+
+ // constructors
+ chillAST_TernaryOperator();
+ chillAST_TernaryOperator(const char *op, chillAST_node *cond, chillAST_node *lhs, chillAST_node *rhs, chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+ bool isNotLeaf() { return true; };
+ bool isLeaf() { return false; };
+
+
+ char *getOp() { return op; }; // dangerous. could get changed!
+ chillAST_node *getCond() { return condition; };
+ chillAST_node *getRHS() { return rhs; };
+ chillAST_node *getLHS() { return lhs; };
+
+ void setCond( chillAST_node *newc ) { condition = newc; newc->setParent( this ); }
+ void setLHS( chillAST_node *newlhs ) { lhs = newlhs; newlhs->setParent( this ); }
+ void setRHS( chillAST_node *newrhs ) { rhs = newrhs; newrhs->setParent( this ); }
+
+
+
+
+ // required methods that I can't seem to get to inherit
+ void dump( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void printonly( int indent=0, FILE *fp = stderr );
+
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void replaceChild( chillAST_node *old, chillAST_node *newchild ) ;
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void gatherVarLHSUsage( std::vector<chillAST_VarDecl*> &decls );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+ void loseLoopWithLoopVar( char *var ){}; // ternop can't have loop as child?
+};
+
+
+
+class chillAST_BinaryOperator: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ char *op; // TODO need enum
+ chillAST_node *lhs;
+ chillAST_node *rhs;
+
+
+ // constructors
+ chillAST_BinaryOperator();
+ chillAST_BinaryOperator(chillAST_node *lhs, const char *op, chillAST_node *rhs, chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+ int evalAsInt();
+ chillAST_IntegerLiteral *evalAsIntegerLiteral();
+
+ bool isNotLeaf() { return true; };
+ bool isLeaf() { return false; };
+
+ chillAST_node *getRHS() { return rhs; };
+ chillAST_node *getLHS() { return lhs; };
+ void setLHS( chillAST_node *newlhs ) { lhs = newlhs; newlhs->setParent( this ); }
+ void setRHS( chillAST_node *newrhs ) { rhs = newrhs; newrhs->setParent( this ); }
+
+ char *getOp() { return op; }; // dangerous. could get changed!
+ bool isAugmentedAssignmentOp() {
+ return
+ (!strcmp(op, "*=")) || // BO_MulAssign,
+ (!strcmp(op, "/=")) || // BO_DivAssign
+ (!strcmp(op, "%=")) || // BO_RemAssign
+ (!strcmp(op, "+=")) || // BO_AddAssign
+ (!strcmp(op, "-=")) || // BO_SubAssign
+
+ (!strcmp(op, "<<=")) || // BO_ShlAssign
+ (!strcmp(op, ">>=")) || // BO_ShrAssign
+ (!strcmp(op, "&&=")) || // BO_AndAssign
+ (!strcmp(op, "||=")) || // BO_OrAssign
+
+ (!strcmp(op, "^=")) // BO_XorAssign
+ ;
+ }
+ bool isAssignmentOp() {
+ return( (!strcmp(op, "=")) || // BO_Assign,
+ isAugmentedAssignmentOp() );
+ };
+ bool isComparisonOp() { return (!strcmp(op,"<")) || (!strcmp(op,">")) || (!strcmp(op,"<=")) || (!strcmp(op,">=")); };
+
+ bool isPlusOp() { return (!strcmp(op,"+")); };
+ bool isMinusOp() { return (!strcmp(op,"-")); };
+ bool isPlusMinusOp() { return (!strcmp(op,"+")) || (!strcmp(op,"-")); };
+ bool isMultDivOp() { return (!strcmp(op,"*")) || (!strcmp(op,"/")); };
+
+ bool isStructOp() { return (!strcmp(op,".")) || (!strcmp(op,"->")); };
+
+
+ // required methods that I can't seem to get to inherit
+ void dump( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void printonly( int indent=0, FILE *fp = stderr );
+ char *stringRep(int indent=0 );
+
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void replaceChild( chillAST_node *old, chillAST_node *newchild ) ;
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento ); // chillAST_BinaryOperator
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void gatherVarLHSUsage( std::vector<chillAST_VarDecl*> &decls );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+ void loseLoopWithLoopVar( char *var ){}; // binop can't have loop as child?
+
+ void gatherStatements( std::vector<chillAST_node*> &statements ); //
+
+ bool isSameAs( chillAST_node *other );
+
+};
+
+
+
+
+
+
+
+
+class chillAST_ArraySubscriptExpr: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *base; // always a decl ref expr? No, for multidimensional array, is another ASE
+ chillAST_node *index;
+ bool imwrittento;
+ bool imreadfrom; // WARNING: ONLY used when both writtento and readfrom are true x += 1 and so on
+ chillAST_VarDecl *basedecl; // the vardecl that this refers to
+ void *uniquePtr; // DO NOT REFERENCE THROUGH THIS!
+
+ // constructors
+ chillAST_ArraySubscriptExpr();
+ chillAST_ArraySubscriptExpr( chillAST_node *bas, chillAST_node *indx, chillAST_node *p, void *unique);
+ chillAST_ArraySubscriptExpr( chillAST_node *bas, chillAST_node *indx, bool writtento, chillAST_node *p, void *unique);
+
+ chillAST_ArraySubscriptExpr( chillAST_VarDecl *v, std::vector<chillAST_node *> indeces, chillAST_node *p);
+
+ // other methods particular to this type of node
+ bool operator!=( const chillAST_ArraySubscriptExpr& ) ;
+ bool operator==( const chillAST_ArraySubscriptExpr& ) ;
+ chillAST_VarDecl *multibase(); // method for finding the basedecl
+ chillAST_node *multibase2() { return base->multibase2(); }
+
+ chillAST_node *getIndex(int dim);
+ void gatherIndeces( std::vector< chillAST_node * > &ind );
+
+ void replaceChild( chillAST_node *old, chillAST_node *newchild ); // will examine index
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void printonly( int indent=0, FILE *fp = stderr );
+ void print( int indent=0, FILE *fp = stderr ) const; // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ char *stringRep(int indent=0 );
+
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ chillAST_node *findref(){return this;}// find the SINGLE constant or data reference at this node or below
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+ const char* getUnderlyingType() {
+ //fprintf(stderr, "ASE getUnderlyingType() base of type %s\n", base->getTypeString()); base->print(); printf("\n"); fflush(stderr);
+ return base->getUnderlyingType(); };
+
+ virtual chillAST_VarDecl* getUnderlyingVarDecl() { return base->getUnderlyingVarDecl(); };
+
+};
+
+
+
+class chillAST_MemberExpr: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *base; // always a decl ref expr? No, can be Array Subscript Expr
+ char *member;
+ char *printstring;
+
+ chillAST_VarDecl *basedecl; // the vardecl that this refers to
+ void *uniquePtr; // DO NOT REFERENCE THROUGH THIS!
+
+ CHILL_MEMBER_EXP_TYPE exptype;
+
+
+ // constructors
+ chillAST_MemberExpr();
+ chillAST_MemberExpr( chillAST_node *bas, const char *mem, chillAST_node *p, void *unique, CHILL_MEMBER_EXP_TYPE t=CHILL_MEMBER_EXP_DOT);
+
+ // other methods particular to this type of node
+ bool operator!=( const chillAST_MemberExpr& ) ;
+ bool operator==( const chillAST_MemberExpr& ) ;
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void printonly( int indent=0, FILE *fp = stderr );
+ void print( int indent=0, FILE *fp = stderr ) const; // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ char *stringRep( int indent = 0);
+
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+ chillAST_VarDecl* getUnderlyingVarDecl();
+
+ void replaceChild( chillAST_node *old, chillAST_node *newchild );
+
+ void setType( CHILL_MEMBER_EXP_TYPE t ) { exptype = t; };
+ CHILL_MEMBER_EXP_TYPE getType( CHILL_MEMBER_EXP_TYPE t ) { return exptype; };
+
+ chillAST_VarDecl* multibase(); // this one will return the member decl
+ chillAST_node* multibase2(); // this one will return the member expression
+};
+
+
+
+
+class chillAST_IntegerLiteral: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ int value;
+
+ // constructors
+ chillAST_IntegerLiteral(int val, chillAST_node *p = NULL);
+
+ // other methods particular to this type of node
+ int evalAsInt() { return value; }
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool w ){}; // does nothing
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ){}; // does nothing
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls ){}; // does nothing
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls ){}; // does nothing
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls ){}; // does nothing
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls ) {}; // does nothing
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs ) {}; // does nothing
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl) {};
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+ chillAST_node *findref(){return this;}// find the SINGLE constant or data reference at this node or below
+};
+
+
+class chillAST_FloatingLiteral: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ float value;
+ double doublevalue;
+ int float0double1;
+
+ char *allthedigits; // if not NULL, use this as printable representation
+ int precision; // float == 1, double == 2, ???
+
+ // constructors
+ chillAST_FloatingLiteral( float val, chillAST_node *p);
+ chillAST_FloatingLiteral( double val, chillAST_node *p);
+ chillAST_FloatingLiteral( float val, int pre, chillAST_node *p);
+ chillAST_FloatingLiteral( double val, int pre, chillAST_node *p);
+ chillAST_FloatingLiteral( float val, const char *printable, chillAST_node *p);
+ chillAST_FloatingLiteral( float val, int pre, const char *printable, chillAST_node *p);
+ chillAST_FloatingLiteral( chillAST_FloatingLiteral *old );
+
+ // other methods particular to this type of node
+ void setPrecision( int precis ) { precision = precis; };
+ int getPrecision() { return precision; }
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool w ){}; // does nothing
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ){}; // does nothing
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls ){}; // does nothing
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls ){}; // does nothing ;
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls ){}; // does nothing ;
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls ) {}; // does nothing
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs ){}; // does nothing
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl){};
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+ chillAST_node *findref(){return this;};// find the SINGLE constant or data reference at this node or below
+
+ bool isSameAs( chillAST_node *other );
+};
+
+
+
+
+class chillAST_UnaryOperator: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ char *op; // TODO enum
+ bool prefix; // or post
+ chillAST_node *subexpr;
+
+ // constructors
+ chillAST_UnaryOperator( const char *oper, bool pre, chillAST_node *sub, chillAST_node *p );
+
+ // other methods particular to this type of node
+ bool isAssignmentOp() {
+ return( (!strcmp(op, "++")) ||
+ (!strcmp(op, "--")) ); // are there more ??? TODO
+ }
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento ); // chillAST_UnaryOperator
+
+ void gatherVarLHSUsage( std::vector<chillAST_VarDecl*> &decls );
+
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+ int evalAsInt();
+ bool isSameAs( chillAST_node *other );
+
+};
+
+
+
+
+
+class chillAST_ImplicitCastExpr: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *subexpr;
+
+ // constructors
+ chillAST_ImplicitCastExpr(chillAST_node *sub, chillAST_node *p);
+
+ // other methods particular to this type of node
+ bool isNotLeaf() { return true; };
+ bool isLeaf() { return false; };
+
+ // required methods that I can't seem to get to inherit
+ void replaceChild( chillAST_node *old, chillAST_node *newchild );
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void printonly( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ) { print( indent, fp); }; // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+ chillAST_VarDecl *multibase(); // just recurse on subexpr
+
+};
+
+
+
+class chillAST_CStyleCastExpr: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ char * towhat;
+ chillAST_node *subexpr;
+ // constructors
+ chillAST_CStyleCastExpr(const char *to, chillAST_node *sub, chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void replaceChild( chillAST_node *old, chillAST_node *newchild );
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+ chillAST_node *findref(){return subexpr;};// find the SINGLE constant or data reference at this node or below
+
+};
+
+
+class chillAST_CStyleAddressOf: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *subexpr;
+ // constructors
+ chillAST_CStyleAddressOf(chillAST_node *sub, chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+
+};
+
+
+class chillAST_CudaMalloc:public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *devPtr; // Pointer to allocated device memory
+ chillAST_node *sizeinbytes;
+
+ // constructors
+ chillAST_CudaMalloc(chillAST_node *devmemptr, chillAST_node *size, chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ //void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+};
+
+
+class chillAST_CudaFree:public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_VarDecl *variable;
+
+ // constructors
+ chillAST_CudaFree(chillAST_VarDecl *var, chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ //void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+};
+
+
+
+
+
+class chillAST_Malloc:public chillAST_node { // malloc( sizeof(int) * 2048 );
+public:
+ // variables that are special for this type of node
+ char *thing; // to void if this is null , sizeof(thing) if it is not
+ chillAST_node *sizeexpr; // bytes
+
+ // constructors
+ chillAST_Malloc(chillAST_node *size, chillAST_node *p=NULL);
+ chillAST_Malloc(char *thething, chillAST_node *numthings, chillAST_node *p=NULL); // malloc (sizeof(int) *1024)
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ //void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+};
+
+
+
+
+class chillAST_Free:public chillAST_node {
+public:
+
+
+
+
+};
+
+
+
+
+class chillAST_CudaMemcpy:public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_VarDecl *dest; // Pointer to allocated device memory
+ chillAST_VarDecl *src;
+ chillAST_node *size;
+ char *cudaMemcpyKind; // could use the actual enum
+
+ // constructors
+ chillAST_CudaMemcpy(chillAST_VarDecl *d, chillAST_VarDecl *s, chillAST_node *size, char *kind, chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ //void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+};
+
+
+class chillAST_CudaSyncthreads:public chillAST_node {
+public:
+ // variables that are special for this type of node
+
+ // constructors
+ chillAST_CudaSyncthreads(chillAST_node *p=NULL);
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ //chillAST_node* constantFold() {};
+ //chillAST_node* clone();
+ //void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento ){};
+ //void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls ) {}; // does nothing
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls ) {}; // does nothing
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls ) {}; // does nothing
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls ) {}; // does nothing
+ //void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ //bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; };
+
+};
+
+
+
+class chillAST_ReturnStmt: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *returnvalue;
+
+ // constructors
+ chillAST_ReturnStmt( chillAST_node *retval, chillAST_node *p );
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+};
+
+
+
+class chillAST_CallExpr: public chillAST_node { // a function call
+public:
+ // variables that are special for this type of node
+ chillAST_node *callee; // the function declaration (what about builtins?)
+ int numargs;
+ std::vector<class chillAST_node*> args;
+ chillAST_VarDecl *grid;
+ chillAST_VarDecl *block;
+
+ // constructors
+ chillAST_CallExpr(chillAST_node *function, chillAST_node *p );
+ void addArg( chillAST_node *newarg );
+
+ // other methods particular to this type of node
+ // TODO get/set grid, block
+
+ // required methods that I can't seem to get to inherit
+ chillAST_node* constantFold();
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+ chillAST_node* clone();
+};
+
+
+
+class chillAST_ParenExpr: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *subexpr;
+
+ // constructors
+ chillAST_ParenExpr( chillAST_node *sub, chillAST_node *p=NULL );
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+};
+
+
+class chillAST_Sizeof: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ char *thing;
+
+ // constructors
+ chillAST_Sizeof( char *t, chillAST_node *p = NULL );
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl){};
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; }; // no loops under here
+
+};
+
+
+
+class chillAST_NoOp: public chillAST_node {
+public:
+ chillAST_NoOp( chillAST_node *p = NULL ); // { parent = p; };
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ) {}; // print CODE in chill_ast.cc
+ void dump( int indent=0, FILE *fp = stderr ) {}; // print ast in chill_ast.cc
+ chillAST_node* constantFold() {};
+ chillAST_node* clone() { return new chillAST_NoOp( parent ); }; // ??
+
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento ) {};
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) {};
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls ){};
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls ){};
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls ){};
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls ){};
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs ){};
+ void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl){};
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false ){ return false; };//no loops under here
+};
+
+
+
+class chillAST_IfStmt: public chillAST_node {
+public:
+ // variables that are special for this type of node
+ chillAST_node *cond;
+ chillAST_node *thenpart;
+ chillAST_node *elsepart;
+ IR_CONDITION_TYPE conditionoperator; // from ir_code.hh
+
+ // constructors
+ chillAST_IfStmt();
+ chillAST_IfStmt( chillAST_node *c, chillAST_node *t, chillAST_node *e, chillAST_node *p);
+
+ // other methods particular to this type of node
+ chillAST_node *getCond() { return cond; };
+ chillAST_node *getThen() { return thenpart; };
+ chillAST_node *getElse() { return elsepart; };
+
+ void setCond( chillAST_node *b ) { cond = b; if (cond) cond->parent = this; };
+ void setThen( chillAST_node *b ) { thenpart = b; if (thenpart) thenpart->parent = this; };
+ void setElse( chillAST_node *b ) { elsepart = b; if (elsepart) elsepart->parent = this; };
+
+ // required methods that I can't seem to get to inherit
+ void dump( int indent=0, FILE *fp = stderr );
+ void print( int indent=0, FILE *fp = stderr );
+
+ chillAST_node* constantFold();
+ chillAST_node* clone();
+
+ void gatherVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherVarDeclsMore ( std::vector<chillAST_VarDecl*> &decls ) { gatherVarDecls(decls); } ;
+
+ void gatherScalarVarDecls( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayVarDecls ( std::vector<chillAST_VarDecl*> &decls );
+ void gatherArrayRefs( std::vector<chillAST_ArraySubscriptExpr*> &refs, bool writtento );
+ void gatherScalarRefs( std::vector<chillAST_DeclRefExpr*> &refs, bool writtento ) ;
+
+ void gatherVarUsage( std::vector<chillAST_VarDecl*> &decls );
+ void gatherDeclRefExprs( std::vector<chillAST_DeclRefExpr *>&refs );
+ //void replaceVarDecls( chillAST_VarDecl *olddecl, chillAST_VarDecl *newdecl);
+ bool findLoopIndexesToReplace( chillAST_SymbolTable *symtab, bool forcesync=false );
+
+ void gatherStatements( std::vector<chillAST_node*> &statements );
+
+};
+
+
+
+
+
+
+
+
+class chillAST_something: public chillAST_node {
+public:
+ // variables that are special for this type of node
+
+ // constructors
+ chillAST_something();
+
+ // other methods particular to this type of node
+
+
+ // required methods that I can't seem to get to inherit
+ void print( int indent=0, FILE *fp = stderr ); // print CODE in chill_ast.cc
+ //void dump( int indent=0, FILE *fp = stderr ); // print ast in chill_ast.cc
+};
+
+
+
+
+chillAST_FunctionDecl *findFunctionDecl( chillAST_node *node, const char *procname);
+
+
+
+#endif
+
diff --git a/include/chill_error.hh b/include/chill_error.hh
index 88e49fc..c5601ae 100644
--- a/include/chill_error.hh
+++ b/include/chill_error.hh
@@ -21,4 +21,8 @@ struct ir_exp_error: public ir_error {
ir_exp_error(const std::string &msg): ir_error(msg){}
};
+struct omega_error: public std::runtime_error {
+ omega_error(const std::string &msg): std::runtime_error(msg){}
+};
+
#endif
diff --git a/include/dep.hh b/include/dep.hh
index 6c535ce..ed5a83d 100644
--- a/include/dep.hh
+++ b/include/dep.hh
@@ -1,19 +1,6 @@
#ifndef DEP_HH
#define DEP_HH
-/*!
- * \file
- * \brief Data dependence vector and graph.
- *
- * All dependence vectors are normalized, i.e., the first non-zero distance
- * must be positve. Thus the correct dependence meaning can be given based on
- * source/destination pair's read/write type. Suppose for a dependence vector
- * 1, 0~5, -3), we want to permute the first and the second dimension,
- * the result would be two dependence vectors (0, 1, -3) and (1~5, 1, -3).
- * All operations on dependence vectors are non-destructive, i.e., new
- * dependence vectors are returned.
- */
-
#include <omega.h>
#include "graph.hh"
#include "ir_code.hh"
@@ -27,9 +14,12 @@ typedef std::vector<DependenceVector> DependenceList;
struct DependenceVector {
DependenceType type;
IR_Symbol *sym;
-
- bool is_reduction; //!< used to identify a class of flow dependence
- //!< that can be broken
+
+ bool from_same_stmt; // Manu
+ bool is_reduction_cand; // Manu
+
+ bool is_reduction; // used to identify a class of flow dependence
+ // that can be broken
std::vector<omega::coef_t> lbounds;
std::vector<omega::coef_t> ubounds;
@@ -39,11 +29,14 @@ struct DependenceVector {
type = DEP_UNKNOWN;
sym = NULL;
is_reduction = false;
+ from_same_stmt = false; // Manu
+ is_reduction_cand = false; // Manu
quasi = false;
is_scalar_dependence = false;
}
+ // DependenceVector(int size);
DependenceVector(const DependenceVector &that);
- ~DependenceVector() {delete sym;}
+ ~DependenceVector() { delete sym; } // is this legal? TODO
DependenceVector &operator=(const DependenceVector &that);
bool is_data_dependence() const;
@@ -52,7 +45,7 @@ struct DependenceVector {
bool has_been_carried_at(int dim) const;
bool has_been_carried_before(int dim) const;
- // TODO the following functions will be cleaned up or removed later
+ // the following functions will be cleaned up or removed later
bool isZero() const;
bool isPositive() const;
bool isNegative() const;
@@ -67,6 +60,7 @@ struct DependenceVector {
std::vector<DependenceVector> normalize() const;
std::vector<DependenceVector> permute(const std::vector<int> &pi) const;
DependenceVector reverse() const;
+ // std::vector<DependenceVector> matrix(const std::vector<std::vector<int> > &M) const;
DependenceType getType() const;
friend std::ostream& operator<<(std::ostream &os, const DependenceVector &d);
};
@@ -83,8 +77,10 @@ public:
DependenceGraph() { num_dim_ = 0; }
~DependenceGraph() {}
int num_dim() const { return num_dim_; }
+// DependenceGraph permute(const std::vector<int> &pi) const;
DependenceGraph permute(const std::vector<int> &pi,
const std::set<int> &active = std::set<int>()) const;
+ // DependenceGraph matrix(const std::vector<std::vector<int> > &M) const;
DependenceGraph subspace(int dim) const;
bool isPositive() const;
bool hasPositive(int dim) const;
diff --git a/include/ir_clang.hh b/include/ir_clang.hh
new file mode 100755
index 0000000..7c3d451
--- /dev/null
+++ b/include/ir_clang.hh
@@ -0,0 +1,479 @@
+#ifndef IR_CLANG_HH
+#define IR_CLANG_HH
+
+#include <omega.h>
+#include "ir_code.hh"
+//#include <AstInterface_CLANG.h>
+#include "chill_error.hh"
+
+#define __STDC_CONSTANT_MACROS
+#include "clang/AST/Decl.h"
+#include "clang/AST/Expr.h"
+#include "clang/AST/ParentMap.h"
+
+#include "clang/AST/ASTConsumer.h"
+#include "clang/AST/ASTContext.h"
+#include "clang/AST/GlobalDecl.h"
+#include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/Basic/SourceManager.h"
+#include "clang/Frontend/ASTConsumers.h"
+#include "clang/Frontend/FrontendActions.h"
+
+#include "clang/CodeGen/CodeGenAction.h"
+#include "clang/Driver/Compilation.h"
+#include "clang/Driver/Driver.h"
+#include "clang/Driver/Tool.h"
+#include "clang/Frontend/CompilerInvocation.h"
+#include "clang/Frontend/CompilerInstance.h"
+#include "clang/Frontend/FrontendDiagnostic.h"
+#include "clang/Frontend/TextDiagnosticPrinter.h"
+#include "clang/Basic/DiagnosticOptions.h"
+
+#include "chill_ast.hh"
+
+// using namespace clang; // NEVER EVER do this in a header file
+// using namespace clang::driver; // NEVER EVER do this in a header file
+
+extern std::vector<chillAST_VarDecl *> VariableDeclarations; // a global. TODO
+
+typedef llvm::SmallVector<clang::Stmt *, 16> StmtList; // TODO delete
+
+struct IR_chillScalarSymbol: public IR_ScalarSymbol {
+ chillAST_VarDecl *chillvd;
+
+ IR_chillScalarSymbol(const IR_Code *ir, chillAST_VarDecl *vd) {
+ fprintf(stderr, "making scalar symbol %s\n", vd->varname);
+ ir_ = ir;
+ chillvd = vd;
+ }
+
+ std::string name() const;
+ int size() const;
+ bool operator==(const IR_Symbol &that) const;
+ IR_Symbol *clone() const;
+};
+
+
+
+struct IR_chillArraySymbol: public IR_ArraySymbol {
+ //int indirect_; // what was this?
+ int offset_; // what is this?
+ chillAST_VarDecl *chillvd;
+
+ IR_chillArraySymbol(const IR_Code *ir, chillAST_VarDecl *vd, int offset = 0) {
+ //if ( vd == 0 )
+ //fprintf(stderr, "IR_chillArraySymbol::IR_chillArraySymbol (%s) vd 0x%x\n", vd->varname, vd);
+ ir_ = ir;
+ chillvd = vd;
+ //indirect_ = indirect;
+ offset_ = offset;
+ }
+
+
+ // No Fortran support!
+ IR_ARRAY_LAYOUT_TYPE layout_type() const {
+ return IR_ARRAY_LAYOUT_ROW_MAJOR;
+ }
+
+ std::string name() const;
+ int elem_size() const;
+ int n_dim() const;
+ omega::CG_outputRepr *size(int dim) const;
+ bool operator!=(const IR_Symbol &that) const;
+ bool operator==(const IR_Symbol &that) const;
+ IR_Symbol *clone() const;
+ // TODO Hack to pass build
+ IR_CONSTANT_TYPE elem_type() const {return IR_CONSTANT_UNKNOWN;};
+
+};
+
+
+
+struct IR_chillConstantRef: public IR_ConstantRef {
+ union {
+ omega::coef_t i_;
+ double f_;
+ };
+
+ IR_chillConstantRef(const IR_Code *ir, omega::coef_t i) {
+ ir_ = ir;
+ type_ = IR_CONSTANT_INT;
+ i_ = i;
+ }
+ IR_chillConstantRef(const IR_Code *ir, double f) {
+ ir_ = ir;
+ type_ = IR_CONSTANT_FLOAT;
+ f_ = f;
+ }
+ omega::coef_t integer() const {
+ assert(is_integer());
+ return i_;
+ }
+ bool operator==(const IR_Ref &that) const;
+ omega::CG_outputRepr *convert();
+ IR_Ref *clone() const;
+
+};
+
+enum OP_POSITION { OP_DEST =-1, OP_UNKNOWN, OP_SRC };
+#define OP_LEFT OP_DEST
+#define OP_RIGHT OP_SRC
+
+struct IR_chillScalarRef: public IR_ScalarRef {
+ OP_POSITION op_pos_; // -1 means destination operand, 0== unknown, 1 == source operand
+ //chillAST_BinaryOperator *bop; // binary op that contains this scalar?
+ chillAST_DeclRefExpr *dre; // declrefexpr that uses this scalar ref, if that exists
+ chillAST_VarDecl *chillvd; // the vardecl for this scalar
+
+ IR_chillScalarRef(const IR_Code *ir, chillAST_BinaryOperator *ins, OP_POSITION pos) {
+ fprintf(stderr, "\n***** new IR_xxxxScalarRef( ir, ins, pos ) *****\n\n");
+ exit(-1);
+ // this constructor takes a binary operation and an indicator of which side of the op to use,
+ // and finds the scalar in the lhs or rhs of the binary op.
+ ir_ = ir;
+ dre = NULL;
+ //bop = ins; // do we need this?
+ if (pos == OP_LEFT) {
+ chillAST_node *lhs = ins->lhs;
+ if (lhs->isDeclRefExpr()) {
+ chillAST_DeclRefExpr *DRE = (chillAST_DeclRefExpr *) lhs;
+ dre = DRE;
+ chillvd = DRE->getVarDecl();
+ }
+ else if (lhs->isVarDecl()) {
+ chillvd = (chillAST_VarDecl *)lhs;
+ }
+ else {
+ fprintf(stderr, "IR_chillScalarRef constructor, I'm confused\n"); exit(-1);
+ }
+ }
+ else {
+ chillAST_node *rhs = ins->rhs;
+ if (rhs->isDeclRefExpr()) {
+ chillAST_DeclRefExpr *DRE = (chillAST_DeclRefExpr *) rhs;
+ dre = DRE;
+ chillvd = DRE->getVarDecl();
+ }
+ else if (rhs->isVarDecl()) {
+ chillvd = (chillAST_VarDecl *)rhs;
+ }
+ else {
+ fprintf(stderr, "IR_chillScalarRef constructor, I'm confused\n"); exit(-1);
+ }
+ }
+ op_pos_ = pos;
+ }
+
+ IR_chillScalarRef(const IR_Code *ir, chillAST_DeclRefExpr *d) {
+ // fprintf(stderr, "\n***** new IR_xxxxScalarRef( ir, REF EXPR sym %s ) *****\n\n", d->getVarDecl()->varname);
+ //fprintf(stderr, "new IR_chillScalarRef with a DECLREFEXPR (has dre) \n");
+ ir_ = ir;
+ dre = d;
+ //bop = NULL;
+ chillvd = d->getVarDecl();
+ op_pos_ = OP_UNKNOWN;
+
+ //fprintf(stderr, "\nScalarRef has:\n");
+ //fprintf(stderr, "assignment op DOESNT EXIST\n");
+ //fprintf(stderr, "ins_pos %d\n", ins_pos_);
+ //fprintf(stderr, "op_pos %d\n", op_pos_);
+ //fprintf(stderr, "ref expr dre = 0x%x\n", dre);
+ }
+
+ IR_chillScalarRef(const IR_Code *ir, chillAST_VarDecl *vardecl) {
+ fprintf(stderr, "\n***** new IR_xxxxScalarRef( ir, sym 0x1234567 ) ***** THIS SHOULD NEVER HAPPEN\n\n");
+ fprintf(stderr, "vardecl %s\n", vardecl->varname);
+ ir_ = ir;
+ dre = NULL; fprintf(stderr, "new IR_chillScalarRef with a vardecl but no dre\n");
+ //bop = NULL;
+ chillvd = vardecl;
+ op_pos_ = OP_UNKNOWN;
+
+ //fprintf(stderr, "\nScalarRef has:\n");
+ //fprintf(stderr, "assignment op DOESNT EXIST\n");
+ //fprintf(stderr, "ins_pos %d\n", ins_pos_);
+ //fprintf(stderr, "op_pos %d\n", op_pos_);
+ //fprintf(stderr, "ref expr dre = 0x%x\n", dre);
+ }
+
+
+ bool is_write() const;
+ IR_ScalarSymbol *symbol() const;
+ bool operator==(const IR_Ref &that) const;
+ omega::CG_outputRepr *convert();
+ IR_Ref *clone() const;
+};
+
+
+
+struct IR_chillArrayRef: public IR_ArrayRef {
+ //DeclRefExpr *as_;
+ //chillAST_DeclRefExpr *chillDRE;
+ chillAST_ArraySubscriptExpr* chillASE;
+ int iswrite;
+
+
+ // IR_chillArrayRef(const IR_Code *ir, DeclRefExpr *as, ParentMap *pMap = NULL) {
+ // ir_ = ir;
+ // fprintf(stderr, "new IR_chillArrayRef() CLANG ERROR\n"); exit(-1);
+ //}
+
+ IR_chillArrayRef(const IR_Code *ir, chillAST_ArraySubscriptExpr *ase, int write ) {
+ //fprintf(stderr, "IR_chillArrayRef::IR_chillArrayRef() write %d\n", write);
+ ir_ = ir;
+ chillASE = ase;
+ iswrite = write; // ase->imwrittento;
+ }
+
+ bool is_write() const;
+ omega::CG_outputRepr *index(int dim) const;
+ IR_ArraySymbol *symbol() const;
+ bool operator!=(const IR_Ref &that) const;
+ bool operator==(const IR_Ref &that) const;
+ omega::CG_outputRepr *convert();
+ IR_Ref *clone() const;
+ virtual void Dump() const;
+};
+
+
+
+struct IR_chillLoop: public IR_Loop {
+ int step_size_;
+
+ chillAST_DeclRefExpr *chillindex; // the loop index variable (I) // was DeclRefExpr
+ chillAST_ForStmt *chillforstmt;
+ chillAST_node *chilllowerbound;
+ chillAST_node *chillupperbound;
+ chillAST_node *chillbody; // presumably a compound statement, but not guaranteeed
+ IR_CONDITION_TYPE conditionoperator;
+
+ IR_chillLoop(const IR_Code *ir, clang::ForStmt *tf);
+ IR_chillLoop(const IR_Code *ir, chillAST_ForStmt *forstmt);
+
+ ~IR_chillLoop() {}
+ IR_ScalarSymbol *index() const { return new IR_chillScalarSymbol(ir_, chillindex->getVarDecl()); }
+ omega::CG_outputRepr *lower_bound() const;
+ omega::CG_outputRepr *upper_bound() const;
+ IR_CONDITION_TYPE stop_cond() const;
+ IR_Block *body() const;
+
+ // Handle following types of increment expressions:
+
+ // Unary increment/decrement
+ // i += K OR i -= K
+ // i = i + K OR i = i - K
+ // where K is positive
+ int step_size() const { return step_size_; } // K is always an integer ???
+ IR_Control *clone() const;
+ IR_Block *convert() ;
+ virtual void dump() const;
+};
+
+
+
+
+struct IR_chillBlock: public IR_Block { // ONLY ONE OF bDecl or cs_ will be nonNULL ??
+private:
+ //StmtList bDecl_; // declarations in the block??
+ //clang::CompoundStmt *cs_; // will a block always have a compound statement? (no)
+ //StmtList *stmts; // ??
+public:
+
+ // Block is a basic block?? (no, just a chunk of code )
+ std::vector<chillAST_node *>statements;
+ chillAST_node *chillAST; // how about for now we say if there are statements, which is presumably the top level of statements from ... somewhere, otherwise the code is in chillAST
+
+ //IR_chillBlock(const IR_Code *ir, const StmtList& bDecl) : bDecl_(bDecl), cs_(NULL) {
+ // fprintf(stderr, "MISTAKE IR_chillBlock bdecl\n"); exit(-1);
+ // ir_ = ir;
+ //}
+ //IR_chillBlock(const IR_Code *ir, clang::CompoundStmt *cs) : cs_(cs) {
+ // fprintf(stderr, "MISTAKE IR_chillBlock cs\n"); exit(-1);
+ // ir_ = ir;
+ //}
+ IR_chillBlock( const IR_chillBlock *CB ) { // clone existing IR_chillBlock
+ ir_ = CB->ir_;
+ for (int i=0; i<CB->statements.size(); i++) statements.push_back( CB->statements[i] );
+ chillAST = CB->chillAST;
+ }
+
+ IR_chillBlock(const IR_Code *ir, chillAST_node *ast) {
+ ir_ = ir;
+ chillAST = ast;
+ }
+
+ IR_chillBlock(const IR_Code *ir) { // : cs_(NULL), bDecl_(NULL) {
+ chillAST = NULL;
+ ir_ = ir;
+ }
+
+
+ ~IR_chillBlock() {}
+ omega::CG_outputRepr *extract() const;
+ omega::CG_outputRepr *original() const;
+ IR_Control *clone() const;
+ //StmtList getStmtList() const;
+ std::vector<chillAST_node*> getStmtList() const;
+ int numstatements() { return statements.size(); } ;
+ void addStatement( chillAST_node* s );
+
+ void dump() const;
+};
+
+
+#ifdef NOTYET
+struct IR_clangIf: public IR_If {
+ SgNode *ti_;
+
+ IR_clangIf(const IR_Code *ir, SgNode *ti) {
+ ir_ = ir;
+ ti_ = ti;
+ }
+ ~IR_clangIf() {
+ }
+ omega::CG_outputRepr *condition() const;
+ IR_Block *then_body() const;
+ IR_Block *else_body() const;
+ IR_Block *convert();
+ IR_Control *clone() const;
+};
+#endif
+
+
+
+
+
+class aClangCompiler {
+private:
+ //Chill_ASTConsumer *astConsumer_;
+ clang::ASTContext *astContext_;
+
+ clang::DiagnosticOptions *diagnosticOptions;
+ clang::TextDiagnosticPrinter *pTextDiagnosticPrinter;
+ clang::DiagnosticIDs *diagID ;
+ clang::DiagnosticsEngine *diagnosticsEngine;
+ clang::CompilerInstance *Clang;
+ clang::Preprocessor *preprocessor;
+ //FileManager *FileMgr;
+ //clang::CompilerInvocation *CI;
+
+ clang::FileManager *fileManager;
+ clang::SourceManager *sourceManager;
+
+ // UNUSED?
+ clang::Diagnostic *diagnostic;
+ clang::LangOptions *languageOptions;
+ clang::HeaderSearchOptions *headerSearchOptions;
+ //clang::HeaderSearch *headerSearch;
+ clang::TargetOptions *targetOptions;
+ clang::TargetInfo *pTargetInfo;
+ clang::PreprocessorOptions *preprocessorOptions;
+ clang::FrontendOptions *frontendOptions;
+ clang::IdentifierTable *idTable;
+ clang::SelectorTable *selTable;
+ clang::Builtin::Context *builtinContext;
+
+
+public:
+ char *SourceFileName;
+ chillAST_SourceFile * entire_file_AST; // TODO move out of public
+
+ aClangCompiler( char *filename ); // constructor
+ chillAST_FunctionDecl *findprocedurebyname( char *name ); // someday, return the chill AST
+ clang::FunctionDecl *FD;
+ //Chill_ASTConsumer *getASTConsumer() { return astConsumer_; }
+ clang::ASTContext *getASTContext() { return astContext_; }
+ clang::SourceManager *getASTSourceManager() { return sourceManager; };
+};
+
+
+// singleton class for global clang initialization
+// TODO: Add support for multiple files in same script
+class IR_clangCode_Global_Init {
+private:
+ static IR_clangCode_Global_Init *pinstance; // the one and only
+ // protecting the constructor is the SINGLETON PATTERN. a global by any other name
+ // IR_clangCode_Global_Init();
+ ~IR_clangCode_Global_Init(); // is this hidden, too?
+ chillAST_FunctionDecl * chillFD; // the original C code
+
+ clang::ASTContext *astContext_;
+ clang::SourceManager *sourceManager;
+public:
+ clang::ASTContext *getASTContext() { return astContext_; }
+ clang::SourceManager *getSourceManager() { return sourceManager; };
+ static IR_clangCode_Global_Init *Instance(char **argv);
+ static IR_clangCode_Global_Init *Instance() { return pinstance; } ;
+ aClangCompiler *ClangCompiler; // this is the thing we really just want one of
+
+
+ void setCurrentFunction( chillAST_node *F ) { chillFD = (chillAST_FunctionDecl *)F; } ;
+ chillAST_FunctionDecl *getCurrentFunction( ) { return chillFD; } ;
+
+
+ void setCurrentASTContext( clang::ASTContext *A ) { astContext_ = A;};
+ clang::ASTContext *getCurrentASTContext() { return astContext_; } ;
+
+ void setCurrentASTSourceManager( clang::SourceManager *S ) { sourceManager = S; } ;
+ clang::SourceManager *getCurrentASTSourceManager() { return sourceManager; } ;
+};
+
+
+
+class IR_clangCode: public IR_Code{ // for an entire file? A single function?
+protected:
+
+ //
+ char *filename;
+ char *procedurename;
+
+
+ chillAST_node *entire_file_AST;
+ chillAST_FunctionDecl * chillfunc; // the function we're currenly modifying
+
+ std::vector<chillAST_VarDecl> entire_file_symbol_table;
+ // loop symbol table?? for (int i=0; ... ) ??
+
+
+
+ clang::FunctionDecl *func_; // a clang construct the function we're currenly modifying
+ clang::ASTContext *astContext_;
+ clang::SourceManager *sourceManager;
+
+ // firstScope;
+ // symboltable1,2,3 ??
+ // topleveldecls
+ //
+
+public:
+ clang::ASTContext *getASTContext() { return astContext_; } ;
+ clang::SourceManager *getASTSourceManager() { return sourceManager; } ;
+
+ IR_clangCode(const char *filename, char *proc_name);
+ ~IR_clangCode();
+
+ IR_ScalarSymbol *CreateScalarSymbol(const IR_Symbol *sym, int i);
+ IR_ArraySymbol *CreateArraySymbol(const IR_Symbol *sym, std::vector<omega::CG_outputRepr *> &size, int i);
+ IR_ScalarRef *CreateScalarRef(const IR_ScalarSymbol *sym);
+ IR_ArrayRef *CreateArrayRef(const IR_ArraySymbol *sym, std::vector<omega::CG_outputRepr *> &index);
+ int ArrayIndexStartAt() { return 0;} // TODO FORTRAN
+
+ std::vector<IR_ScalarRef *> FindScalarRef(const omega::CG_outputRepr *repr) const;
+ std::vector<IR_ArrayRef *> FindArrayRef(const omega::CG_outputRepr *repr) const;
+ std::vector<IR_Control *> FindOneLevelControlStructure(const IR_Block *block) const;
+ IR_Block *MergeNeighboringControlStructures(const std::vector<IR_Control *> &controls) const;
+ IR_Block *GetCode() const;
+ void ReplaceCode(IR_Control *old, omega::CG_outputRepr *repr);
+ void ReplaceExpression(IR_Ref *old, omega::CG_outputRepr *repr);
+
+ IR_CONDITION_TYPE QueryBooleanExpOperation(const omega::CG_outputRepr*) const;
+ IR_OPERATION_TYPE QueryExpOperation(const omega::CG_outputRepr *repr) const;
+ std::vector<omega::CG_outputRepr *> QueryExpOperand(const omega::CG_outputRepr *repr) const;
+ IR_Ref *Repr2Ref(const omega::CG_outputRepr *) const;
+
+ friend class IR_chillArraySymbol;
+ friend class IR_chillArrayRef;
+};
+
+
+#endif
diff --git a/include/ir_code.hh b/include/ir_code.hh
index d695474..0946610 100644
--- a/include/ir_code.hh
+++ b/include/ir_code.hh
@@ -8,7 +8,7 @@
extra code generation.
.
Notes:
- Unlike CG_outputRepr, IR_Symbol,IR_Ref and IR_Control are place holders
+ Unlike CG_outputRepr, IR_Symbol, IR_Ref and IR_Control are place holders
to the underlying code, thus deleting or duplicating them does not affect
the actual code. Similar to Omega builder's memory allocation strategy,
all non-const pointer parameters of CG_outputRepr/IR_Symbol/IR_Ref/IR_Control
@@ -22,60 +22,49 @@
#ifndef IR_CODE_HH
#define IR_CODE_HH
-/*!
- * \file
- * \brief CHiLL's compiler intermediate representation interface that extends Omega's builder interface to accomodate compiler analyses and extra code generation.
- *
- * Unlike CG_outputRepr, IR_Symbol,IR_Ref and IR_Control are place holders
- * to the underlying code, thus deleting or duplicating them does not affect
- * the actual code. Similar to Omega builder's memory allocation strategy,
- * all non-const pointer parameters of CG_outputRepr/IR_Symbol/IR_Ref/IR_Control
- * are destroyed after the call.
- */
+#include <ir_enums.hh>
+#include <chill_ast.hh>
+
+#include <vector>
+
+// needed for omega::coef_t below
+#include <basic/util.h> // in omega ... why is this not in CG_output*.h?
#include <code_gen/CG_outputRepr.h>
#include <code_gen/CG_outputBuilder.h>
-#include <vector>
-enum IR_OPERATION_TYPE {IR_OP_CONSTANT, IR_OP_VARIABLE,
- IR_OP_PLUS, IR_OP_MINUS, IR_OP_MULTIPLY, IR_OP_DIVIDE,
- IR_OP_POSITIVE, IR_OP_NEGATIVE,
- IR_OP_MIN, IR_OP_MAX,
- IR_OP_ASSIGNMENT,
- IR_OP_NULL, IR_OP_UNKNOWN};
-enum IR_CONTROL_TYPE {IR_CONTROL_LOOP, IR_CONTROL_IF, IR_CONTROL_WHILE, IR_CONTROL_BLOCK};
-enum IR_CONSTANT_TYPE {IR_CONSTANT_INT, IR_CONSTANT_FLOAT,
- IR_CONSTANT_UNKNOWN};
-enum IR_CONDITION_TYPE {IR_COND_LT, IR_COND_LE,
- IR_COND_GT, IR_COND_GE,
- IR_COND_EQ, IR_COND_NE,
- IR_COND_UNKNOWN};
-enum IR_ARRAY_LAYOUT_TYPE {IR_ARRAY_LAYOUT_ROW_MAJOR,
- IR_ARRAY_LAYOUT_COLUMN_MAJOR,
- IR_ARRAY_LAYOUT_SPACE_FILLING};
-
-class IR_Code;
-
-
-//! Base abstract class for scalar and array symbols.
-/*! This is a place holder for related declaration in IR code.*/
+
+
+class IR_Code; // forward declaration
+
+
+// Base abstract class for scalar and array symbols. This is a place
+// holder for related declaration in IR code.
struct IR_Symbol {
const IR_Code *ir_;
-
+
virtual ~IR_Symbol() {/* ir_ is not the responsibility of this object */}
- virtual int n_dim() const = 0;
+ virtual int n_dim() const = 0; // IR_Symbol
virtual std::string name() const = 0;
virtual bool operator==(const IR_Symbol &that) const = 0;
virtual bool operator!=(const IR_Symbol &that) const {return !(*this == that);}
virtual IR_Symbol *clone() const = 0; /* shallow copy */
+
+ virtual bool isScalar() const { return false; } // default
+ virtual bool isArray() const { return false; } // default
+ virtual bool isPointer() const { return false; } // default
+
+ //IR_SYMBOL_TYPE symtype; // base type: int, float, double, struct, .... typedef'd something
+ //IR_SYMBOL_TYPE getDatatype() ;
};
struct IR_ScalarSymbol: public IR_Symbol {
virtual ~IR_ScalarSymbol() {}
- int n_dim() const {return 0;}
+ int n_dim() const {return 0;} // IR_ScalarSymbol
virtual int size() const = 0;
+ bool isScalar() const { return true; }
};
@@ -84,23 +73,33 @@ struct IR_ArraySymbol: public IR_Symbol {
virtual int elem_size() const = 0;
virtual omega::CG_outputRepr *size(int dim) const = 0;
virtual IR_ARRAY_LAYOUT_TYPE layout_type() const = 0;
+ virtual IR_CONSTANT_TYPE elem_type() const = 0;
+ bool isArray() const { return true; }
};
-//! Base abstract class for scalar and array references.
-/*! This is a place holder for related code in IR code. */
+struct IR_PointerSymbol: public IR_Symbol {
+ virtual ~IR_PointerSymbol() {}
+ virtual omega::CG_outputRepr *size(int dim) const = 0;
+ virtual void set_size(int dim, omega::CG_outputRepr*) = 0;
+ virtual IR_CONSTANT_TYPE elem_type() const = 0;
+ bool isPointer() const { return true; }
+};
+
+// Base abstract class for scalar and array references. This is a
+// place holder for related code in IR code.
struct IR_Ref {
const IR_Code *ir_;
virtual ~IR_Ref() {/* ir_ is not the responsibility of this object */}
- virtual int n_dim() const = 0;
+ virtual int n_dim() const = 0; // IR_Ref
virtual bool is_write() const = 0;
virtual std::string name() const = 0;
virtual bool operator==(const IR_Ref &that) const = 0;
virtual bool operator!=(const IR_Ref &that) const {return !(*this == that);}
virtual omega::CG_outputRepr *convert() = 0;
- //! shallow copy
- virtual IR_Ref *clone() const = 0;
+ virtual IR_Ref *clone() const = 0; /* shallow copy */
+ virtual void Dump() const { fprintf(stderr, "some IR_*Ref needs to implement Dump()\n"); int *i=0; int j=i[0]; };
};
@@ -108,7 +107,7 @@ struct IR_ConstantRef: public IR_Ref {
IR_CONSTANT_TYPE type_;
virtual ~IR_ConstantRef() {}
- int n_dim() const {return 0;}
+ int n_dim() const {return 0;} // IR_ConstantRef
bool is_write() const {return false;}
std::string name() const {return std::string();}
virtual bool is_integer() const {return type_ == IR_CONSTANT_INT;}
@@ -118,10 +117,10 @@ struct IR_ConstantRef: public IR_Ref {
struct IR_ScalarRef: public IR_Ref {
virtual ~IR_ScalarRef() {}
- int n_dim() const {return 0;}
+ int n_dim() const {return 0;} // IR_ScalarRef
virtual IR_ScalarSymbol *symbol() const = 0;
std::string name() const {
- IR_ScalarSymbol *sym = symbol();
+ IR_ScalarSymbol *sym = symbol(); // really inefficient. MAKE a symbol, just to get a name
std::string s = sym->name();
delete sym;
return s;
@@ -137,54 +136,78 @@ struct IR_ScalarRef: public IR_Ref {
struct IR_ArrayRef: public IR_Ref {
virtual ~IR_ArrayRef() {}
- int n_dim() const {
+ int n_dim() const { // IR_ArrayRef returns the size of the array
IR_ArraySymbol *sym = symbol();
int n = sym->n_dim();
- delete sym;
+ // ?? delete sym;
return n;
}
virtual omega::CG_outputRepr *index(int dim) const = 0;
virtual IR_ArraySymbol *symbol() const = 0;
- std::string name() const {
- IR_ArraySymbol *sym = symbol();
+ virtual std::string name() const {
+ // makes (constructs!) a symbol, just to copy a string!
+ IR_ArraySymbol *sym = symbol(); // TODO exceedingly wasteful
std::string s = sym->name();
- delete sym;
- return s;
+ // ?? delete sym; (goes out of scope, so deletes itself)
+ return s; // s ALSO goes out of scope but perhaps the info is copied at the other end
}
virtual int elem_size() const {
IR_ArraySymbol *sym = symbol();
int s = sym->elem_size();
- delete sym;
+ // ?? delete sym;
return s;
}
virtual IR_ARRAY_LAYOUT_TYPE layout_type() const {
IR_ArraySymbol *sym = symbol();
IR_ARRAY_LAYOUT_TYPE t = sym->layout_type();
- delete sym;
+ // ?? delete sym;
return t;
}
+ virtual void Dump() const { fprintf(stderr, "IR_ArrayRef needs to implement Dump()\n"); };
+};
+
+struct IR_PointerArrayRef: public IR_Ref {
+
+ const IR_Code *ir_;
+
+ virtual ~IR_PointerArrayRef() {}
+ int n_dim() const { // IR_PointerArrayRef returns size of the ... symbol?
+ IR_PointerSymbol *sym = symbol();
+ int n = sym->n_dim();
+ //Anand: Hack, fix later
+ //delete sym;
+ return n;
+ }
+ virtual omega::CG_outputRepr *index(int dim) const = 0;
+ virtual IR_PointerSymbol *symbol() const = 0;
+ std::string name() const {
+ IR_PointerSymbol *sym = symbol();
+ std::string s = sym->name();
+//Anand: Hack, fix later
+ //delete sym;
+ return s;
+ }
+
+
};
struct IR_Block;
-//! Base abstract class for code structures.
-/*!
- * This is a place holder for the actual structure in the IR code.
- * However, in cases that original source code may be transformed during
- * loop initialization such as converting a while loop to a for loop or
- * reconstructing the loop from low level IR code, the helper loop class (NOT
- * IMPLEMENTED) must contain the transformed code that needs to be
- * freed when out of service.
- */
+// Base abstract class for code structures. This is a place holder
+// for the actual structure in the IR code. However, in cases that
+// original source code may be transformed during loop initialization
+// such as converting a while loop to a for loop or reconstructing the
+// loop from low level IR code, the helper loop class (NOT
+// IMPLEMENTED) must contain the transformed code that needs to be
+// freed when out of service.
struct IR_Control {
- const IR_Code *ir_;
+ const IR_Code *ir_; // hate this
virtual ~IR_Control() {/* ir_ is not the responsibility of this object */}
virtual IR_CONTROL_TYPE type() const = 0;
virtual IR_Block *convert() = 0;
- //! shallow copy
- virtual IR_Control *clone() const = 0;
+ virtual IR_Control *clone() const = 0; /* shallow copy */
};
@@ -224,55 +247,98 @@ struct IR_While: public IR_Control {
};
-//! Abstract class for compiler IR.
+// Abstract class for compiler IR.
class IR_Code {
protected:
- omega::CG_outputBuilder *ocg_;
+ // the only data members in IR_Code are Omega classes
+ omega::CG_outputBuilder *ocg_; // Omega Code Gen
omega::CG_outputRepr *init_code_;
omega::CG_outputRepr *cleanup_code_;
+ // OK, I lied
+ static int ir_pointer_counter;
+ static int ir_array_counter;
+
public:
+
+ int getPointerCounter() { return ir_pointer_counter; }
+ int getArrayCounter() { return ir_array_counter; }
+
+ // TODO can't get the initialize of counters to work !!
+ int getAndIncrementPointerCounter() { if (ir_pointer_counter == 0) ir_pointer_counter= 1; ir_pointer_counter++; return ir_pointer_counter-1; }
+ int getAndIncrementArrayCounter() { if (ir_array_counter == 0) ir_array_counter= 1; ir_array_counter += 1; return ir_array_counter-1; }
+
+ // if all flavors of ir_code use chillAST internally ...
+ chillAST_FunctionDecl * func_defn; // the function we're modifying
+ chillAST_FunctionDecl *GetChillFuncDefinition() { return func_defn; };
+
IR_Code() {ocg_ = NULL; init_code_ = cleanup_code_ = NULL;}
- virtual ~IR_Code() { delete ocg_; delete init_code_; delete cleanup_code_; }
- /* the content of init and cleanup code have already been released in derived classes */
+ virtual ~IR_Code() { delete ocg_; delete init_code_; delete cleanup_code_; } /* the content of init and cleanup code have already been released in derived classes */
- /*!
- * \param memory_type is for differentiating the location of
- * where the new memory is allocated. this is useful for
- * processors with heterogeneous memory hierarchy.
- */
+ omega::CG_outputRepr* init_code(){ return init_code_; }
+
+ // memory_type is for differentiating the location of where the new memory is allocated.
+ // this is useful for processors with heterogeneous memory hierarchy.
virtual IR_ScalarSymbol *CreateScalarSymbol(const IR_Symbol *sym, int memory_type) = 0;
- virtual IR_ArraySymbol *CreateArraySymbol(const IR_Symbol *sym, std::vector<omega::CG_outputRepr *> &size, int memory_type) = 0;
-
+ virtual IR_ScalarSymbol *CreateScalarSymbol(IR_CONSTANT_TYPE type, int memory_type, std::string name="" ) =0;
+
+ virtual IR_ArraySymbol *CreateArraySymbol(const IR_Symbol *sym,
+ std::vector<omega::CG_outputRepr *> &size,
+ int memory_type) = 0;
+ virtual IR_ArraySymbol *CreateArraySymbol(omega::CG_outputRepr *type,
+ std::vector<omega::CG_outputRepr *> &size_repr) =0;
+
+ virtual IR_PointerSymbol *CreatePointerSymbol(const IR_Symbol *sym,
+ std::vector<omega::CG_outputRepr *> &size_repr) =0;
+ virtual IR_PointerSymbol *CreatePointerSymbol(const IR_CONSTANT_TYPE type,
+ std::vector<omega::CG_outputRepr *> &size_repr,
+ std::string name="") =0;
+
+ virtual IR_PointerSymbol *CreatePointerSymbol(omega::CG_outputRepr *type,
+ std::vector<omega::CG_outputRepr *> &size_repr) =0;
+
virtual IR_ScalarRef *CreateScalarRef(const IR_ScalarSymbol *sym) = 0;
virtual IR_ArrayRef *CreateArrayRef(const IR_ArraySymbol *sym, std::vector<omega::CG_outputRepr *> &index) = 0;
+
+ virtual omega::CG_outputRepr* CreateArrayRefRepr(const IR_ArraySymbol *sym,
+ std::vector<omega::CG_outputRepr *> &index) {
+ //IR_ArrayRef *AR = CreateArrayRef(sym, index);
+ //return new omega::CG_outputRepr(AR);
+ fprintf(stderr, "ir_code.hh SOME SUBCLASS OF ir_code did not implement CreateArrayRefRepr()\n");
+ return NULL;
+ }
+
+ virtual IR_PointerArrayRef *CreatePointerArrayRef( IR_PointerSymbol *sym,
+ std::vector<omega::CG_outputRepr *> &index) =0;
virtual int ArrayIndexStartAt() {return 0;}
- /*!
- * Array references should be returned in their accessing order.
- *
- * ~~~
- * e.g. s1: A[i] = A[i-1]
- * s2: B[C[i]] = D[i] + E[i]
- * return A[i-1], A[i], D[i], E[i], C[i], B[C[i]] in this order.
- * ~~~
- */
+ virtual void CreateDefineMacro(std::string s,std::string args, omega::CG_outputRepr *repr) = 0;
+ virtual void CreateDefineMacro(std::string s,std::string args, std::string repr) = 0;
+ virtual void CreateDefineMacro(std::string s, std::vector<std::string> args,omega::CG_outputRepr *repr) {}; // TODO make pure virtual
+
+ virtual omega::CG_outputRepr *CreateArrayType(IR_CONSTANT_TYPE type, omega::CG_outputRepr* size)=0;
+ virtual omega::CG_outputRepr *CreatePointerType(IR_CONSTANT_TYPE type)=0;
+ virtual omega::CG_outputRepr *CreatePointerType(omega::CG_outputRepr *type)=0;
+ virtual omega::CG_outputRepr *CreateScalarType(IR_CONSTANT_TYPE type)=0;
+ // Array references should be returned in their accessing order.
+ // e.g. s1: A[i] = A[i-1]
+ // s2: B[C[i]] = D[i] + E[i]
+ // return A[i-1], A[i], D[i], E[i], C[i], B[C[i]] in this order.
virtual std::vector<IR_ArrayRef *> FindArrayRef(const omega::CG_outputRepr *repr) const = 0;
+ virtual std::vector<IR_PointerArrayRef *> FindPointerArrayRef(const omega::CG_outputRepr *repr) const = 0 ;
virtual std::vector<IR_ScalarRef *> FindScalarRef(const omega::CG_outputRepr *repr) const = 0;
+ virtual bool parent_is_array(IR_ArrayRef *a)=0;
- /*!
- * If there is no sub structure interesting inside the block, return empty,
- * so we know when to stop looking inside.
- */
+ // If there is no sub structure interesting inside the block, return empty,
+ // so we know when to stop looking inside.
virtual std::vector<IR_Control *> FindOneLevelControlStructure(const IR_Block *block) const = 0;
- /*!
- * All controls must be in the same block, at the same level and in
- * contiguous lexical order as appeared in parameter vector.
- */
+ // All controls must be in the same block, at the same level and in
+ // contiguous lexical order as appeared in parameter vector.
virtual IR_Block *MergeNeighboringControlStructures(const std::vector<IR_Control *> &controls) const = 0;
virtual IR_Block *GetCode() const = 0;
+ virtual IR_Control *GetCode(omega::CG_outputRepr *code) const = 0;
virtual void ReplaceCode(IR_Control *old, omega::CG_outputRepr *repr) = 0;
virtual void ReplaceExpression(IR_Ref *old, omega::CG_outputRepr *repr) = 0;
@@ -280,10 +346,34 @@ public:
virtual IR_CONDITION_TYPE QueryBooleanExpOperation(const omega::CG_outputRepr *repr) const = 0;
virtual std::vector<omega::CG_outputRepr *> QueryExpOperand(const omega::CG_outputRepr *repr) const = 0;
virtual IR_Ref *Repr2Ref(const omega::CG_outputRepr *repr) const = 0;
-
- //! Codegen Omega code builder interface
- omega::CG_outputBuilder *builder() const {return ocg_;}
+
+ // Manu:: Added functions required for reduction operation
+ // virtual omega::CG_outputRepr * FromSameStmt(IR_ArrayRef *A, IR_ArrayRef *B) = 0;
+ virtual bool FromSameStmt(IR_ArrayRef *A, IR_ArrayRef *B) = 0;
+ virtual void printStmt(const omega::CG_outputRepr *repr) = 0;
+ virtual int getStmtType(const omega::CG_outputRepr *repr) = 0;
+ virtual IR_OPERATION_TYPE getReductionOp(const omega::CG_outputRepr *repr) = 0;
+ virtual IR_Control * FromForStmt(const omega::CG_outputRepr *repr) = 0;
+
+ // Manu:: Added functions for scalar expansion
+ virtual IR_ArraySymbol *CreateArraySymbol(omega::CG_outputRepr *size, const IR_Symbol *sym) = 0;
+ virtual bool ReplaceRHSExpression(omega::CG_outputRepr *code, IR_Ref *ref) = 0;
+ virtual omega::CG_outputRepr * GetRHSExpression(omega::CG_outputRepr *code) = 0;
+ virtual omega::CG_outputRepr * GetLHSExpression(omega::CG_outputRepr *code) = 0;
+ virtual omega::CG_outputRepr *CreateMalloc(const IR_CONSTANT_TYPE type, std::string lhs,
+ omega::CG_outputRepr * size_repr)=0;
+ virtual omega::CG_outputRepr *CreateMalloc(omega::CG_outputRepr *type, std::string variable,
+ omega::CG_outputRepr * size_repr)=0;
+ virtual omega::CG_outputRepr *CreateFree(omega::CG_outputRepr * exp)=0;
+
+ //void Dump() { ocg_->Dump(); };
+ //---------------------------------------------------------------------------
+ // CC Omega code builder interface here
+
+ //---------------------------------------------------------------------------
+ omega::CG_outputBuilder *builder() const { return ocg_;}
};
#endif
+
diff --git a/include/ir_enums.hh b/include/ir_enums.hh
new file mode 100644
index 0000000..742dbe8
--- /dev/null
+++ b/include/ir_enums.hh
@@ -0,0 +1,24 @@
+
+#ifndef IR_ENUMS_HH
+#define IR_ENUMS_HH
+enum IR_OPERATION_TYPE {IR_OP_CONSTANT, IR_OP_VARIABLE, IR_OP_ARRAY_VARIABLE,
+ IR_OP_PLUS, IR_OP_MINUS, IR_OP_MULTIPLY, IR_OP_DIVIDE,
+ IR_OP_POSITIVE, IR_OP_NEGATIVE, IR_OP_NEQ,
+ IR_OP_MIN, IR_OP_MAX,IR_OP_EQ, IR_OP_LE, IR_OP_GE,
+ IR_OP_ASSIGNMENT, IR_OP_PLUS_ASSIGNMENT, IR_OP_CCAST,
+ IR_OP_NULL, IR_OP_MOD, IR_OP_UNKNOWN};
+enum IR_CONTROL_TYPE {IR_CONTROL_LOOP, IR_CONTROL_IF, IR_CONTROL_WHILE, IR_CONTROL_BLOCK};
+
+enum IR_CONSTANT_TYPE {IR_CONSTANT_INT, IR_CONSTANT_FLOAT, IR_CONSTANT_DOUBLE,
+ IR_CONSTANT_UNKNOWN};
+
+enum IR_SYMBOL_TYPE {IR_SYMBOL_UNKNOWN=0, IR_SYMBOL_INT, IR_SYMBOL_FLOAT, IR_SYMBOL_DOUBLE}; // structs, etc ?? are UNKNOWN
+
+enum IR_CONDITION_TYPE {IR_COND_LT, IR_COND_LE,
+ IR_COND_GT, IR_COND_GE,
+ IR_COND_EQ, IR_COND_NE,
+ IR_COND_UNKNOWN};
+enum IR_ARRAY_LAYOUT_TYPE {IR_ARRAY_LAYOUT_ROW_MAJOR,
+ IR_ARRAY_LAYOUT_COLUMN_MAJOR,
+ IR_ARRAY_LAYOUT_SPACE_FILLING};
+#endif
diff --git a/include/ir_rose.hh b/include/ir_rose.hh
deleted file mode 100644
index 03ea50d..0000000
--- a/include/ir_rose.hh
+++ /dev/null
@@ -1,267 +0,0 @@
-#ifndef IR_ROSE_HH
-#define IR_ROSE_HH
-
-/*!
- * \file
- * \brief CHiLL's rose interface.
- */
-
-#include <omega.h>
-#include "ir_code.hh"
-#include "ir_rose_utils.hh"
-#include <AstInterface_ROSE.h>
-#include "chill_error.hh"
-#include "staticSingleAssignment.h"
-#include "VariableRenaming.h"
-#include "ssaUnfilteredCfg.h"
-#include "virtualCFG.h"
-#include <omega.h>
-
-struct IR_roseScalarSymbol: public IR_ScalarSymbol {
- SgVariableSymbol* vs_;
-
- IR_roseScalarSymbol(const IR_Code *ir, SgVariableSymbol *vs) {
- ir_ = ir;
- vs_ = vs;
- }
-
- std::string name() const;
- int size() const;
- bool operator==(const IR_Symbol &that) const;
- IR_Symbol *clone() const;
-};
-
-struct IR_roseArraySymbol: public IR_ArraySymbol {
-
- SgVariableSymbol* vs_;
-
- IR_roseArraySymbol(const IR_Code *ir, SgVariableSymbol* vs) {
- ir_ = ir;
- vs_ = vs;
- }
- std::string name() const;
- int elem_size() const;
- int n_dim() const;
- omega::CG_outputRepr *size(int dim) const;
- bool operator==(const IR_Symbol &that) const;
- IR_ARRAY_LAYOUT_TYPE layout_type() const;
- IR_Symbol *clone() const;
-
-};
-
-struct IR_roseConstantRef: public IR_ConstantRef {
- union {
- omega::coef_t i_;
- double f_;
- };
-
- IR_roseConstantRef(const IR_Code *ir, omega::coef_t i) {
- ir_ = ir;
- type_ = IR_CONSTANT_INT;
- i_ = i;
- }
- IR_roseConstantRef(const IR_Code *ir, double f) {
- ir_ = ir;
- type_ = IR_CONSTANT_FLOAT;
- f_ = f;
- }
- omega::coef_t integer() const {
- assert(is_integer());
- return i_;
- }
- bool operator==(const IR_Ref &that) const;
- omega::CG_outputRepr *convert();
- IR_Ref *clone() const;
-
-};
-
-struct IR_roseScalarRef: public IR_ScalarRef {
- SgAssignOp *ins_pos_;
- int op_pos_; // -1 means destination operand, otherwise source operand
- SgVarRefExp *vs_;
- int is_write_;
- IR_roseScalarRef(const IR_Code *ir, SgVarRefExp *sym) {
- ir_ = ir;
- ins_pos_ = isSgAssignOp(sym->get_parent());
- op_pos_ = 0;
- if (ins_pos_ != NULL)
- if (sym == isSgVarRefExp(ins_pos_->get_lhs_operand()))
- op_pos_ = -1;
-
- vs_ = sym;
- }
- IR_roseScalarRef(const IR_Code *ir, SgVarRefExp *ins, int pos) {
- ir_ = ir;
- /* ins_pos_ = ins;
- op_pos_ = pos;
- SgExpression* op;
- if (pos == -1)
- op = ins->get_lhs_operand();
- else
- op = ins->get_rhs_operand();
-
- */
-
- is_write_ = pos;
-
- /* if (vs_ == NULL || pos > 0)
- throw ir_error(
- "Src operand not a variable or more than one src operand!!");
- */
-
- vs_ = ins;
-
- }
- bool is_write() const;
- IR_ScalarSymbol *symbol() const;
- bool operator==(const IR_Ref &that) const;
- omega::CG_outputRepr *convert();
- IR_Ref *clone() const;
-};
-
-struct IR_roseArrayRef: public IR_ArrayRef {
-
- SgPntrArrRefExp *ia_;
-
- int is_write_;
- IR_roseArrayRef(const IR_Code *ir, SgPntrArrRefExp *ia, int write) {
- ir_ = ir;
- ia_ = ia;
- is_write_ = write;
- }
- bool is_write() const;
- omega::CG_outputRepr *index(int dim) const;
- IR_ArraySymbol *symbol() const;
- bool operator==(const IR_Ref &that) const;
- omega::CG_outputRepr *convert();
- IR_Ref *clone() const;
-};
-
-struct IR_roseLoop: public IR_Loop {
- SgNode *tf_;
-
- IR_roseLoop(const IR_Code *ir, SgNode *tf) {
- ir_ = ir;
- tf_ = tf;
- }
-
- IR_ScalarSymbol *index() const;
- omega::CG_outputRepr *lower_bound() const;
- omega::CG_outputRepr *upper_bound() const;
- IR_CONDITION_TYPE stop_cond() const;
- IR_Block *body() const;
- IR_Block *convert();
- int step_size() const;
- IR_Control *clone() const;
-};
-
-struct IR_roseBlock: public IR_Block {
- SgNode* tnl_;
- SgNode *start_, *end_;
-
- IR_roseBlock(const IR_Code *ir, SgNode *tnl, SgNode *start, SgNode *end) {
- ir_ = ir;
- tnl_ = tnl;
- start_ = start;
- end_ = end;
- }
-
- IR_roseBlock(const IR_Code *ir, SgNode *tnl) {
- ir_ = ir;
- tnl_ = tnl;
- start_ = tnl_->get_traversalSuccessorByIndex(0);
- end_ = tnl_->get_traversalSuccessorByIndex(
- (tnl_->get_numberOfTraversalSuccessors()) - 1);
- }
- omega::CG_outputRepr *extract() const;
- omega::CG_outputRepr *original() const;
- IR_Control *clone() const;
-};
-
-struct IR_roseIf: public IR_If {
- SgNode *ti_;
-
- IR_roseIf(const IR_Code *ir, SgNode *ti) {
- ir_ = ir;
- ti_ = ti;
- }
- ~IR_roseIf() {
- }
- omega::CG_outputRepr *condition() const;
- IR_Block *then_body() const;
- IR_Block *else_body() const;
- IR_Block *convert();
- IR_Control *clone() const;
-};
-
-class IR_roseCode_Global_Init {
-private:
- static IR_roseCode_Global_Init *pinstance;
-public:
- SgProject* project;
- static IR_roseCode_Global_Init *Instance(char** argv);
-};
-
-class IR_roseCode: public IR_Code {
-protected:
- SgSourceFile* file;
- SgGlobal *root;
- SgGlobal *firstScope;
- SgSymbolTable* symtab_;
- SgSymbolTable* symtab2_;
- SgSymbolTable* symtab3_;
- SgDeclarationStatementPtrList::iterator p;
- SgFunctionDeclaration *func;
- bool is_fortran_;
- int i_;
- StaticSingleAssignment *ssa_for_scalar;
- ssa_unfiltered_cfg::SSA_UnfilteredCfg *main_ssa;
- VariableRenaming *varRenaming_for_scalar;
-public:
- IR_roseCode(const char *filename, const char* proc_name);
- ~IR_roseCode();
-
- IR_ScalarSymbol *CreateScalarSymbol(const IR_Symbol *sym, int memory_type =
- 0);
- IR_ArraySymbol *CreateArraySymbol(const IR_Symbol *sym,
- std::vector<omega::CG_outputRepr *> &size, int memory_type = 0);
- IR_ScalarRef *CreateScalarRef(const IR_ScalarSymbol *sym);
- IR_ArrayRef *CreateArrayRef(const IR_ArraySymbol *sym,
- std::vector<omega::CG_outputRepr *> &index);
- int ArrayIndexStartAt() {
- if (is_fortran_)
- return 1;
- else
- return 0;
- }
-
- void populateLists(SgNode* tnl_1, SgStatementPtrList* list_1,
- SgStatementPtrList& output_list_1);
- void populateScalars(const omega::CG_outputRepr *repr1,
- std::map<SgVarRefExp*, IR_ScalarRef*> &read_scalars_1,
- std::map<SgVarRefExp*, IR_ScalarRef*> &write_scalars_1,
- std::set<std::string> &indices, std::vector<std::string> &index);
- std::vector<IR_ArrayRef *> FindArrayRef(
- const omega::CG_outputRepr *repr) const;
- std::vector<IR_ScalarRef *> FindScalarRef(
- const omega::CG_outputRepr *repr) const;
- std::vector<IR_Control *> FindOneLevelControlStructure(
- const IR_Block *block) const;
- IR_Block *MergeNeighboringControlStructures(
- const std::vector<IR_Control *> &controls) const;
- IR_Block *GetCode() const;
- void ReplaceCode(IR_Control *old, omega::CG_outputRepr *repr);
- void ReplaceExpression(IR_Ref *old, omega::CG_outputRepr *repr);
-
- IR_OPERATION_TYPE QueryExpOperation(const omega::CG_outputRepr *repr) const;
- IR_CONDITION_TYPE QueryBooleanExpOperation(
- const omega::CG_outputRepr *repr) const;
- std::vector<omega::CG_outputRepr *> QueryExpOperand(
- const omega::CG_outputRepr *repr) const;
- IR_Ref *Repr2Ref(const omega::CG_outputRepr *) const;
- void finalizeRose();
- friend class IR_roseArraySymbol;
- friend class IR_roseArrayRef;
-};
-
-#endif
diff --git a/include/ir_rose_utils.hh b/include/ir_rose_utils.hh
deleted file mode 100644
index 350aa24..0000000
--- a/include/ir_rose_utils.hh
+++ /dev/null
@@ -1,19 +0,0 @@
-#ifndef IR_ROSE_UTILS_HH
-#define IR_ROSE_UTILS_HH
-
-/*!
- * \file
- * \brief ROSE interface utilities
- */
-#include <vector>
-#include <rose.h>
-#include <sageBuilder.h>
-
-std::vector<SgForStatement *> find_deepest_loops(SgNode *tnl);
-std::vector<SgForStatement *> find_loops(SgNode *tnl);
-
-SgNode* loop_body_at_level(SgNode* tnl, int level);
-SgNode* loop_body_at_level(SgForStatement* loop, int level);
-void swap_node_for_node_list(SgNode* tn, SgNode* new_tnl);
-
-#endif
diff --git a/include/irtools.hh b/include/irtools.hh
index 6648847..d85b76b 100644
--- a/include/irtools.hh
+++ b/include/irtools.hh
@@ -1,29 +1,22 @@
#ifndef IRTOOLS_HH
#define IRTOOLS_HH
-/*!
- * \file
- * \brief Useful tools to analyze code in compiler IR format.
- */
-
#include <vector>
#include <omega.h>
#include <code_gen/CG_outputRepr.h>
#include "ir_code.hh"
#include "dep.hh"
+#define DEP_DEBUG 0
-//! It is used to initialize a loop.
+// IR tree is used to initialize a loop. For a loop node, payload is
+// its mapped iteration space dimension. For a simple block node,
+// payload is its mapped statement number. Normal if-else is split
+// into two nodes where the one with odd payload represents then-part and
+// the one with even payload represents else-part.
struct ir_tree_node {
IR_Control *content;
ir_tree_node *parent;
std::vector<ir_tree_node *> children;
-/*!
- * * For a loop node, payload is its mapped iteration space dimension.
- * * For a simple block node, payload is its mapped statement number.
- * * Normal if-else is splitted into two nodes
- * * the one with odd payload represents then-part and
- * * the one with even payload represents else-part.
- */
int payload;
~ir_tree_node() {
@@ -33,45 +26,32 @@ struct ir_tree_node {
}
};
-//! Build IR tree from the source code
-/*!
- * Block type node can only be leaf, i.e., there is no further stuctures inside a block allowed
- * @param control
- * @param parent
- * @return
- */
std::vector<ir_tree_node *> build_ir_tree(IR_Control *control,
ir_tree_node *parent = NULL);
-//! Extract statements from IR tree
-/*!
- * Statements returned are ordered in lexical order in the source code
- * @param ir_tree
- * @return
- */
std::vector<ir_tree_node *> extract_ir_stmts(
const std::vector<ir_tree_node *> &ir_tree);
bool is_dependence_valid(ir_tree_node *src_node, ir_tree_node *dst_node,
const DependenceVector &dv, bool before);
-//! test data dependeces between two statements
-/*!
- * The first statement in parameter must be lexically before the second statement in parameter.
- * Returned dependences are all lexicographically positive
- * @param ir
- * @param repr1
- * @param IS1
- * @param repr2
- * @param IS2
- * @param freevar
- * @param index
- * @param i
- * @param j
- * @return Dependecies between the two statements. First vector is dependencies from first to second,
- * second vector is the reverse
- */
std::pair<std::vector<DependenceVector>, std::vector<DependenceVector> > test_data_dependences(
IR_Code *ir, const omega::CG_outputRepr *repr1,
const omega::Relation &IS1, const omega::CG_outputRepr *repr2,
const omega::Relation &IS2, std::vector<omega::Free_Var_Decl*> &freevar,
- std::vector<std::string> index, int i, int j);
+ std::vector<std::string> index, int i, int j,
+ std::map<std::string, std::vector<omega::CG_outputRepr * > > &uninterpreted_symbols,
+ std::map<std::string, std::vector<omega::CG_outputRepr * > > &uninterpreted_symbols_stringrepr);
+
+std::vector<omega::CG_outputRepr *> collect_loop_inductive_and_conditionals(ir_tree_node * stmt_node);
+
+// Manu
+typedef std::map<int, std::pair<std::vector<DependenceVector>, std::vector<DependenceVector> > > tempResultMap;
+typedef std::pair<std::vector<DependenceVector>, std::vector<DependenceVector> > DVPair;
+
+// Manu:: this function is required for reduction operation
+//omega::CG_outputRepr * from_same_statement(IR_Code *ir, IR_ArrayRef *a, IR_ArrayRef *b);
+bool from_same_statement(IR_Code *ir, IR_ArrayRef *a, IR_ArrayRef *b);
+int stmtType(IR_Code *ir, const omega::CG_outputRepr *repr);
+IR_OPERATION_TYPE getReductionOperator(IR_Code *ir, const omega::CG_outputRepr *repr);
+void mapRefstoStatements(IR_Code *ir,std::vector<IR_ArrayRef *> access, int ref2Stmt[], std::map<int,std::set<int> >& rMap, std::set<int>& tnrStmts, std::set<int>& nrStmts);
+void checkReductionDependence(int i, int j, int nestLeveli, omega::coef_t lbound[], omega::coef_t ubound[], int ref2Stmt[], std::map<int,std::set<int> >& rMap, DVPair& dv, tempResultMap& trMap, std::set<int> nrStmts );
#endif
diff --git a/include/loop.hh b/include/loop.hh
index e02b7e9..3909e6c 100644
--- a/include/loop.hh
+++ b/include/loop.hh
@@ -1,24 +1,6 @@
#ifndef LOOP_HH
#define LOOP_HH
-/*!
- * \file
- * \brief Core loop transformation functionality.
- *
- * "level" (starting from 1) means loop level and it corresponds to "dim"
- * (starting from 0) in transformed iteration space [c_1,l_1,c_2,l_2,....,
- * c_n,l_n,c_(n+1)], e.g., l_2 is loop level 2 in generated code, dim 3
- * in transformed iteration space, and variable 4 in Omega relation.
- * All c's are constant numbers only and they will not show up as actual loops.
- *
- * Formula:
- *
- * ~~~
- * dim = 2*level - 1
- * var = dim + 1
- * ~~~
- */
-
#include <omega.h>
#include <code_gen/codegen.h>
@@ -29,6 +11,10 @@
#include "dep.hh"
#include "ir_code.hh"
#include "irtools.hh"
+#include <code_gen/CG_stringBuilder.h>
+
+#include "stencil.hh"
+
class IR_Code;
@@ -36,30 +22,35 @@ enum TilingMethodType { StridedTile, CountedTile };
enum LoopLevelType { LoopLevelOriginal, LoopLevelTile, LoopLevelUnknown };
-//! Describes properties of each loop level of a statement.
+// Describes properties of each loop level of a statement. "payload"
+// for LoopLevelOriginal means iteration space dimension, for
+// LoopLevelTile means tiled loop level. Special value -1 for
+// LoopLevelTile means purely derived loop. For dependence dimension
+// payloads, the values must be in an increasing order.
+// "parallel_level" will be used by code generation to support
+// multi-level parallelization (default 0 means sequential loop under
+// the current parallelization level).
struct LoopLevel {
LoopLevelType type;
-/*!
- * For LoopLevelOriginal means iteration space dimension
- * For LoopLevelTile means tiled loop level. Special value -1 for
- * LoopLevelTile means purely derived loop. For dependence dimension
- * payloads, the values must be in an increasing order.
- */
int payload;
-/*!
- * Used by code generation to support
- * multi-level parallelization (default 0 means sequential loop under
- * the current parallelization level).
- */
int parallel_level;
+ bool segreducible;
+ std::string segment_descriptor;
};
+
struct Statement {
omega::CG_outputRepr *code;
omega::Relation IS;
omega::Relation xform;
std::vector<LoopLevel> loop_level;
ir_tree_node *ir_stmt_node;
+ bool has_inspector;
+ int reduction; // Manu:: 0 == reduction not possible, 1 == reduction possible, 2 == reduction with some processing
+ IR_OPERATION_TYPE reductionOp; // Manu
+
+ class stencilInfo *statementStencil;
+
//protonu--temporarily putting this back here
//omega::Tuple<int> nonSplitLevels;
//end--protonu.
@@ -72,22 +63,30 @@ protected:
static const std::string tmp_loop_var_name_prefix;
int overflow_var_name_counter;
static const std::string overflow_var_name_prefix;
- std::vector<int> stmt_nesting_level_;
+ std::vector<int> stmt_nesting_level_; // UNDERLINE
std::vector<std::string> index;
std::map<int, omega::CG_outputRepr *> replace;
+ std::map<int, std::pair<int, std::string> > reduced_statements;
public:
+ void debugRelations() const;
IR_Code *ir;
+ std::vector<IR_PointerSymbol *> ptr_variables;
std::vector<omega::Free_Var_Decl*> freevar;
std::vector<Statement> stmt;
+ std::vector<omega::CG_outputRepr*> actual_code; // ?????
std::vector<ir_tree_node *> ir_stmt;
std::vector<ir_tree_node *> ir_tree;
+ std::set<std::string> reduced_write_refs;
+ std::map<std::string, int> array_dims;
DependenceGraph dep;
int num_dep_dim;
omega::Relation known;
omega::CG_outputRepr *init_code;
omega::CG_outputRepr *cleanup_code;
std::map<int, std::vector<omega::Free_Var_Decl *> > overflow;
+ std::vector<std::map<std::string, std::vector<omega::CG_outputRepr * > > > uninterpreted_symbols;
+ std::vector<std::map<std::string, std::vector<omega::CG_outputRepr * > > >uninterpreted_symbols_stringrepr;
protected:
@@ -110,15 +109,20 @@ protected:
void apply_xform(std::set<int> &active);
void apply_xform();
std::set<int> getSubLoopNest(int stmt_num, int level) const;
-
+ int getMinLexValue(std::set<int> stmts, int level);
+
public:
Loop() { ir = NULL; tmp_loop_var_name_counter = 1; init_code = NULL; }
Loop(const IR_Control *control);
~Loop();
- omega::CG_outputRepr *getCode(int effort = 1) const;
- void printCode(int effort = 1) const;
+ omega::CG_outputRepr *getCode(int effort = 3) const; // TODO was 1
+ //chillAST_node* LoopCuda::getCode(int effort, std::set<int> stmts) const;
+
+ void stencilASEPadded(int stmt_num);
+
+ void printCode(int effort = 3) const;
void addKnown(const omega::Relation &cond);
void print_internal_loop_structure() const;
bool isInitialized() const;
@@ -130,16 +134,14 @@ public:
std::vector<std::set <int > > sort_by_same_loops(std::set<int > active, int level);
//
- //! legacy unimodular transformations for perfectly nested loops
- /*!
- * e.g. \f$M*(i,j)^T = (i',j')^T or M*(i,j,1)^T = (i',j')^T\f$
- */
+ // legacy unimodular transformations for perfectly nested loops
+ // e.g. M*(i,j)^T = (i',j')^T or M*(i,j,1)^T = (i',j')^T
+ //
bool nonsingular(const std::vector<std::vector<int> > &M);
- /*!
- * \defgroup hltrans High-level loop transformations
- * @{
- */
+ //
+ // high-level loop transformations
+ //
void permute(const std::set<int> &active, const std::vector<int> &pi);
void permute(int stmt_num, int level, const std::vector<int> &pi);
void permute(const std::vector<int> &pi);
@@ -148,46 +150,19 @@ public:
void tile(int stmt_num, int level, int tile_size, int outer_level = 1, TilingMethodType method = StridedTile, int alignment_offset = 0, int alignment_multiple = 1);
std::set<int> split(int stmt_num, int level, const omega::Relation &cond);
std::set<int> unroll(int stmt_num, int level, int unroll_amount, std::vector< std::vector<std::string> >idxNames= std::vector< std::vector<std::string> >(), int cleanup_split_level = 0);
-
- //! Datacopy function by reffering arrays by numbers
- /*!
- * for example
- * ~~~
- * A[i] = A[i-1] + B[i];
- * ~~~
- * parameter array_ref_num=[0,2] means to copy data touched by A[i-1] and A[i]
- *
- * @param array_ref_nums
- * @param level
- * @param allow_extra_read
- * @param fastest_changing_dimension
- * @param padding_stride
- * @param padding_alignment
- * @param memory_type
- * @return
- */
+
bool datacopy(const std::vector<std::pair<int, std::vector<int> > > &array_ref_nums, int level, bool allow_extra_read = false, int fastest_changing_dimension = -1, int padding_stride = 1, int padding_alignment = 4, int memory_type = 0);
- //! Datacopy function by reffering arrays by name
- /*!
- * parameter array_name=A means to copy data touched by A[i-1] and A[i]
- * @param stmt_num
- * @param level
- * @param array_name
- * @param allow_extra_read
- * @param fastest_changing_dimension
- * @param padding_stride
- * @param padding_alignment
- * @param memory_type
- * @return
- */
bool datacopy(int stmt_num, int level, const std::string &array_name, bool allow_extra_read = false, int fastest_changing_dimension = -1, int padding_stride = 1, int padding_alignment = 4, int memory_type = 0);
bool datacopy_privatized(int stmt_num, int level, const std::string &array_name, const std::vector<int> &privatized_levels, bool allow_extra_read = false, int fastest_changing_dimension = -1, int padding_stride = 1, int padding_alignment = 1, int memory_type = 0);
bool datacopy_privatized(const std::vector<std::pair<int, std::vector<int> > > &array_ref_nums, int level, const std::vector<int> &privatized_levels, bool allow_extra_read = false, int fastest_changing_dimension = -1, int padding_stride = 1, int padding_alignment = 1, int memory_type = 0);
bool datacopy_privatized(const std::vector<std::pair<int, std::vector<IR_ArrayRef *> > > &stmt_refs, int level, const std::vector<int> &privatized_levels, bool allow_extra_read, int fastest_changing_dimension, int padding_stride, int padding_alignment, int memory_type = 0);
+ //std::set<int> scalar_replacement_inner(int stmt_num);
+ bool find_stencil_shape( int stmt_num );
Graph<std::set<int>, bool> construct_induced_graph_at_level(std::vector<std::set<int> > s, DependenceGraph dep, int dep_dim);
- std::vector<std::set<int> > typed_fusion(Graph<std::set<int>, bool> g);
+ std::vector<std::set<int> > typed_fusion(Graph<std::set<int>, bool> g, std::vector<bool> &types);
+
void fuse(const std::set<int> &stmt_nums, int level);
void distribute(const std::set<int> &stmt_nums, int level);
void skew(const std::set<int> &stmt_nums, int level, const std::vector<int> &skew_amount);
@@ -195,36 +170,44 @@ public:
void scale(const std::set<int> &stmt_nums, int level, int scale_amount);
void reverse(const std::set<int> &stmt_nums, int level);
void peel(int stmt_num, int level, int peel_amount = 1);
- /*!
- * \defgroup hlfancy fancy loop transformations
- * @{
- */
+ //
+ // more fancy loop transformations
+ //
void modular_shift(int stmt_num, int level, int shift_amount) {}
void diagonal_map(int stmt_num, const std::pair<int, int> &levels, int offset) {}
void modular_partition(int stmt_num, int level, int stride) {}
- /*! @} */
+ void flatten(int stmt_num, std::string index_name, std::vector<int> &loop_levels, std::string inspector_name);
+ void normalize(int stmt_num, int loop_level);
+
+ void generate_ghostcells_v2(std::vector<int> stmt, int loop_num, int ghost_value, int hold_inner_loop_constant=0 );
- /*!
- * \defgroup hlderived derived loop transformations
- * @{
- */
+
+ //
+ // derived loop transformations
+ //
void shift_to(int stmt_num, int level, int absolute_position);
std::set<int> unroll_extra(int stmt_num, int level, int unroll_amount, int cleanup_split_level = 0);
bool is_dependence_valid_based_on_lex_order(int i, int j,
- const DependenceVector &dv, bool before);
- /*! @} */
+ const DependenceVector &dv, bool before);
+ void split_with_alignment(int stmt_num, int level, int alignment,
+ int direction=0);
+
+ // Manu:: reduction operation
+ void reduce(int stmt_num, std::vector<int> &level, int param, std::string func_name, std::vector<int> &seq_levels, std::vector<int> cudaized_levels = std::vector<int>(), int bound_level = -1);
+ void scalar_expand(int stmt_num, const std::vector<int> &levels, std::string arrName, int memory_type =0, int padding_alignment=0, int assign_then_accumulate = 1, int padding_stride = 0);
+ void ELLify(int stmt_num, std::vector<std::string> arrays_to_pad, int pad_to, bool dense_pad = false, std::string dense_pad_pos_array = "");
+ void compact(int stmt_num, int level, std::string new_array, int zero,
+ std::string data_array);
+ void make_dense(int stmt_num, int loop_level, std::string new_loop_index);
+ void set_array_size(std::string name, int size );
+ omega::CG_outputRepr * iegen_parser(std::string &str, std::vector<std::string> &index_names);
- /*!
- * \defgroup hlother other public operations
- * @{
- */
+ //
+ // other public operations
+ //
void pragma(int stmt_num, int level, const std::string &pragmaText);
void prefetch(int stmt_num, int level, const std::string &arrName, int hint);
- /*! @} */
-
- /*! @} */
+ //void prefetch(int stmt_num, int level, const std::string &arrName, const std::string &indexName, int offset, int hint);
};
-
-
#endif
diff --git a/include/omegatools.hh b/include/omegatools.hh
index b51b2bd..2f56971 100644
--- a/include/omegatools.hh
+++ b/include/omegatools.hh
@@ -1,11 +1,6 @@
#ifndef OMEGATOOLS_HH
#define OMEGATOOLS_HH
-/*!
- * \file
- * \brief Useful tools involving Omega manipulation.
- */
-
#include <string>
#include <omega.h>
#include "dep.hh"
@@ -13,81 +8,92 @@
std::string tmp_e();
-//! Convert expression tree to omega relation.
-/*!
- * \param destroy shallow deallocation of "repr", not freeing the actual code inside.
- */
-void exp2formula(IR_Code *ir, omega::Relation &r, omega::F_And *f_root,
+void exp2formula(IR_Code *ir,
+ omega::Relation &r,
+ omega::F_And *f_root,
std::vector<omega::Free_Var_Decl *> &freevars,
- omega::CG_outputRepr *repr, omega::Variable_ID lhs, char side,
- IR_CONDITION_TYPE rel, bool destroy);
+ omega::CG_outputRepr *repr,
+ omega::Variable_ID lhs,
+ char side,
+ IR_CONDITION_TYPE rel,
+ bool destroy,
+ std::map<std::string, std::vector<omega::CG_outputRepr * > > &uninterpreted_symbols,
+ std::map<std::string, std::vector<omega::CG_outputRepr * > > &uninterpreted_symbols_stringrepr);
-//! Build dependence relation for two array references.
omega::Relation arrays2relation(IR_Code *ir, std::vector<omega::Free_Var_Decl*> &freevars,
const IR_ArrayRef *ref_src, const omega::Relation &IS_w,
- const IR_ArrayRef *ref_dst, const omega::Relation &IS_r);
-//! Convert array dependence relation into set of dependence vectors
-/*!
- * assuming ref_w is lexicographically before ref_r in the source code.
- */
+ const IR_ArrayRef *ref_dst, const omega::Relation &IS_r,
+ std::map<std::string, std::vector<omega::CG_outputRepr * > > &uninterpreted_symbols,
+ std::map<std::string, std::vector<omega::CG_outputRepr * > > &uninterpreted_symbols_stringrepr);
+
std::pair<std::vector<DependenceVector>, std::vector<DependenceVector> > relation2dependences(
- const IR_ArrayRef *ref_src, const IR_ArrayRef *ref_dst, const omega::Relation &r);
+ const IR_ArrayRef *ref_src, const IR_ArrayRef *ref_dst, const omega::Relation &r);
-//! Convert a boolean expression to omega relation.
-/*!
- * \param destroy shallow deallocation of "repr", not freeing the actual code inside.
- */
void exp2constraint(IR_Code *ir, omega::Relation &r, omega::F_And *f_root,
std::vector<omega::Free_Var_Decl *> &freevars,
- omega::CG_outputRepr *repr, bool destroy);
+ omega::CG_outputRepr *repr, bool destroy,
+ std::map<std::string, std::vector<omega::CG_outputRepr * > > &uninterpreted_symbols,
+ std::map<std::string, std::vector<omega::CG_outputRepr * > > &uninterpreted_symbols_stringrepr);
+
bool is_single_iteration(const omega::Relation &r, int dim);
-//! Set/get the value of a variable which is know to be constant.
void assign_const(omega::Relation &r, int dim, int val);
-
int get_const(const omega::Relation &r, int dim, omega::Var_Kind type);
-
-//! Find the position index variable in a Relation by name.
omega::Variable_ID find_index(omega::Relation &r, const std::string &s, char side);
-
-//! Generate mapping relation for permuation.
omega::Relation permute_relation(const std::vector<int> &pi);
-
omega::Relation get_loop_bound(const omega::Relation &r, int dim);
-
-//! Determine whether the loop (starting from 0) in the iteration space has only one iteration.
bool is_single_loop_iteration(const omega::Relation &r, int level, const omega::Relation &known);
-//! Get the bound for a specific loop.
omega::Relation get_loop_bound(const omega::Relation &r, int level, const omega::Relation &known);
omega::Relation get_max_loop_bound(const std::vector<omega::Relation> &r, int dim);
omega::Relation get_min_loop_bound(const std::vector<omega::Relation> &r, int dim);
-
-//! Add strident to a loop.
-/*!
- * Issues:
- *
- * * Don't work with relations with multiple disjuncts.
- * * Omega's dealing with max lower bound is awkward.
- */
void add_loop_stride(omega::Relation &r, const omega::Relation &bound, int dim, int stride);
bool is_inner_loop_depend_on_level(const omega::Relation &r, int level, const omega::Relation &known);
-/*!
- * Suppose loop dim is i. Replace i with i+adjustment in loop bounds.
- *
- * ~~~
- * do i = 1, n
- * do j = i, n
- * ~~~
- *
- * after call with dim = 0 and adjustment = 1:
- *
- * ~~~
- * do i = 1, n
- * do j = i+1, n
- * ~~~
- */
+// void adjust_loop_bound(omega::Relation &r, int dim, int adjustment, std::vector<omega::Free_Var_Decl *> globals = std::vector<omega::Free_Var_Decl *>());
omega::Relation adjust_loop_bound(const omega::Relation &r, int level, int adjustment);
+bool lowerBoundIsZero(const omega::Relation &bound, int dim);
+omega::Relation and_with_relation_and_replace_var(const omega::Relation &R, omega::Variable_ID v1,
+ omega::Relation &g);
+omega::Relation replicate_IS_and_add_at_pos(const omega::Relation &R, int level, omega::Relation &bound);
+omega::Relation replicate_IS_and_add_bound(const omega::Relation &R, int level, omega::Relation &bound) ;
+
+omega::CG_outputRepr * construct_int_floor(omega::CG_outputBuilder * ocg, const omega::Relation &R, const omega::GEQ_Handle &h, omega::Variable_ID v,const std::vector<std::pair<omega::CG_outputRepr *, int> > &assigned_on_the_fly,
+ std::map<std::string, std::vector<omega::CG_outputRepr *> > unin);
+//omega::CG_outputRepr * modified_output_subs_repr(omega::CG_outputBuilder * ocg, const omega::Relation &R, const omega::EQ_Handle &h, omega::Variable_ID v,const std::vector<std::pair<omega::CG_outputRepr *, int> > &assigned_on_the_fly,
+// std::map<std::string, std::vector<omega::CG_outputRepr *> > unin);
+std::pair<omega::Relation, bool> replace_set_var_as_existential(const omega::Relation &R,int pos, std::vector<omega::Relation> &bound);
+omega::Relation replace_set_var_as_Global(const omega::Relation &R,int pos,std::vector<omega::Relation> &bound);
+omega::Relation replace_set_var_as_another_set_var(const omega::Relation &old_relation, const omega::Relation &new_relation, int old_pos, int new_pos);
+omega::Relation extract_upper_bound(const omega::Relation &R, omega::Variable_ID v1);
+
+// void adjust_loop_bound(Relation &r, int dim, int adjustment);
+// void adjust_loop_bound(Relation &r, int dim, Free_Var_Decl *global_var, int adjustment);
+// boolean is_private_statement(const omega::Relation &r, int dim);
+
+// coef_t mod(const Relation &r, Variable_ID v, int dividend);
+
enum LexicalOrderType {LEX_MATCH, LEX_BEFORE, LEX_AFTER, LEX_UNKNOWN};
+// template <typename T>
+// LexicalOrderType lexical_order(const std::vector<T> &a, const std::vector<T> &b) {
+// int size = min(a.size(), b.size());
+// for (int i = 0; i < size; i++) {
+// if (a[i] < b[i])
+// return LEX_BEFORE;
+// else if (b[i] < a[i])
+// return LEX_AFTER;
+// }
+// if (a.size() < b.size())
+// return LEX_BEFORE;
+// else if (b.size() < a.size())
+// return LEX_AFTER;
+// else
+// return LEX_MATCH;
+// }
+
+// struct LoopException {
+// std::string descr;
+// LoopException(const std::string &s): descr(s) {};
+// };
+
#endif
diff --git a/include/stencil.hh b/include/stencil.hh
new file mode 100644
index 0000000..69eec42
--- /dev/null
+++ b/include/stencil.hh
@@ -0,0 +1,55 @@
+
+#pragma once
+
+#include "ir_clang.hh"
+#include "chill_ast.hh"
+
+//#include <omega.h>
+#include <omegatools.hh>
+#include "code_gen/CG_chillRepr.h"
+
+// #include "loop.hh"
+#include "ir_code.hh"
+
+
+
+// *extern Loop *myloop; // ???
+
+
+//void stencil( chillAST_node *topstatement, class stencilInfo &SI ); // stencilInfo was originally a struct, and this was the way it got filled in. TODO remove
+
+
+class stencilInfo {
+public:
+ uint dimensions; // number of dimensions in the stencil
+ uint elements; // total number of elements in the stencil
+ chillAST_VarDecl* srcArrayVariable; // the variable of the source array
+ chillAST_VarDecl* dstArrayVariable; // the variable of the destination array
+
+ int minOffset[3]; // the minimum offset for each dimension NOTE hardcoded to 3 dimensions TODO
+ int maxOffset[3]; // the maximum offset for each dimension
+
+ chillAST_VarDecl* indexVariables[3]; // the variable used for indexing each dimension
+
+ std::vector< std::vector< int > > offsets; // k+1, j-47, i+0 etc
+
+
+ std::vector< std::vector<chillAST_node*> > coefficients; // the coefficients of for each element, NOT IN ANY SET ORDER
+ chillAST_node* find_coefficient( int i, int j, int k ) ; // hardcoded 3 dimensions TODO
+
+ // constructors
+ stencilInfo();
+ stencilInfo( chillAST_node *topstatement ) ;
+
+ void walktree( chillAST_node *node, std::vector< chillAST_node * > &coeffstohere );
+
+ void print( FILE *fp=stdout );
+
+ int radius();
+
+ bool isSymmetric();
+
+}; // class stencilInfo
+
+
+