diff --git a/dvm/fdvm/trunk/fdvm/acc.cpp b/dvm/fdvm/trunk/fdvm/acc.cpp index 303e8d8..5762e0a 100644 --- a/dvm/fdvm/trunk/fdvm/acc.cpp +++ b/dvm/fdvm/trunk/fdvm/acc.cpp @@ -16,7 +16,7 @@ local_part_list *lpart_list; static int dvmh_targets, has_io_stmt; static int targets[Ndev]; -static int has_region, in_arg_list, analyzing, has_max_minloc, for_shadow_compute; +static int has_region, in_arg_list, analyzing, has_max_minloc, for_shadow_compute, private_array_arg; //static char *fname_gpu; static SgStatement *cur_in_block, *cur_in_source, *mod_gpu_end; @@ -37,6 +37,7 @@ static SgSymbol *s_end[MAX_LOOP_LEVEL], *s_blocksS_k[MAX_LOOP_LEVEL], *s_loopSte static SgType *type_DvmType, *type_CudaIndexType, *type_with_len_DvmType, *type_FortranDvmType, *CudaIndexType_k; static int loopIndexCount; + //------ C ---------- static const char *red_kernel_func_names[] = { NULL, @@ -83,6 +84,7 @@ void InitializeACC() declaration_cmnt = NULL; indexType_int = indexType_long = indexType_llong = NULL; dvmh_targets = options.isOn(NO_CUDA) ? HOST_DEVICE : HOST_DEVICE | CUDA_DEVICE; + private_array_class = new SgSymbol(TYPE_NAME, "PrivateArray", *(current_file->firstStatement())); SpecialSymbols.insert(std::pair('\n', "\\n\"\n\"")); SpecialSymbols.insert(std::pair('"', "\\\"")); @@ -287,6 +289,7 @@ void InitializeInFuncACC() acc_registered_list = NULL; /*ACC*/ registered_uses_list = NULL; /*ACC*/ acc_declared_list = NULL; /*ACC*/ + } int GeneratedForCuda() @@ -1015,7 +1018,7 @@ void EnterDataRegion(SgExpression *ale,SgStatement *stmt) void ExitDataRegion(SgExpression *ale,SgStatement *stmt) { SgExpression *e,*size; SgSymbol *ar,*ar2; - + e = &(ale->copy()); if(isSgRecordRefExp(e)) { @@ -1218,7 +1221,6 @@ void ExitDataRegionForLocalVariables(SgStatement *st, int is) } } - void testScopeOfDeclaredVariables(SgStatement *stmt) { SgExpression *el; @@ -1235,7 +1237,7 @@ void testDeclareDirectives(SgStatement *first_dvm_exec) { SgStatement *stmt; for (stmt = cur_func->lexNext(); stmt && (stmt != first_dvm_exec); stmt = stmt->lastNodeOfStmt()->lexNext()) - { + { if (stmt->variant()==ACC_DECLARE_DIR) { if (IN_MODULE) @@ -1484,7 +1486,7 @@ SgStatement *ACC_Directive(SgStatement *stmt) void ACC_DECLARE_Directive(SgStatement *stmt) { - if (ACC_program) + if (ACC_program) acc_declared_list = ExpressionListsUnion(acc_declared_list, &(stmt->expr(0)->copy())); } @@ -1504,15 +1506,43 @@ void ACC_ROUTINE_Directive(SgStatement *stmt) return; } if (!mod_gpu_symb) - CreateGPUModule(); - int targets = stmt->expr(0) ? TargetsList(stmt->expr(0)->lhs()) : dvmh_targets; + CreateGPUModule(); + + SgExpression *targets_spec= NULL, *private_spec = NULL, *el; + + for (el=stmt->expr(0); el; el=el->rhs()) + { + switch (el->lhs()->variant()) + { + case ACC_TARGETS_OP: + if (!targets_spec) + { + targets_spec = el->lhs(); + } else + err("Double TARGETS clause",669,stmt); + break; + case ACC_PRIVATE_OP: + if (!private_spec) + { + private_spec = el->lhs(); + } else + err("Double PRIVATE clause",607,stmt); + break; + } + } + int targets = targets_spec ? TargetsList(targets_spec->lhs()) : dvmh_targets; //stmt->expr(0) ? TargetsList(stmt->expr(0)->lhs()) : dvmh_targets; targets = targets & dvmh_targets; SgSymbol *s = stmt->controlParent()->symbol(); if(!s) return; - if(targets & CUDA_DEVICE) + if (targets & CUDA_DEVICE) + { MarkAsCalled(s); + if (private_spec) + MarkPrivateArgumentsOfRoutine(s, private_spec->lhs()); + } MarkAsRoutine(s); + return; } @@ -2846,6 +2876,7 @@ void ACC_CreateParallelLoop(int ipl, SgStatement *first_do, int nloop, SgStateme // creating private_list private_list = clause[PRIVATE_] ? clause[PRIVATE_]->lhs() : NULL; + private_array_arg = 0; dost = InnerMostLoop(first_do, nloop); @@ -4616,6 +4647,17 @@ void Argument(SgExpression *e, int i, SgSymbol *s) } else if (isSgArrayRefExp(e)) { + if (analyzing && e->lhs() && isSgArrayType(e->type())) // case of array section + { + Warning("Array section of %s in a region", e->symbol()->identifier(), 667, cur_st); + doNotForCuda(); + } + if (!analyzing && isPrivate(e->symbol()) && isArrayParameter(ProcedureSymbol(s),i)) + { // scheme with PrivateArray Class + private_array_arg++; // += isArrayParameter(ProcedureSymbol(s),i); + if (!FromOtherFile(s)) + addArgumentNumber(i, s); + } RefInExpr(e, _READ_WRITE_); return; } @@ -4632,7 +4674,6 @@ void Argument(SgExpression *e, int i, SgSymbol *s) } } - void Call(SgSymbol *s, SgExpression *e) { SgExpression *el; @@ -4650,7 +4691,7 @@ void Call(SgSymbol *s, SgExpression *e) } if (IsInternalProcedure(s) && analyzing) Error(" Call of the procedure %s in a region, which is internal/module procedure", s->identifier(), 580, cur_st); - + if (!isUserFunction(s) && (s->attributes() & INTRINSIC_BIT || isIntrinsicFunctionName(s->identifier()))) //IsNoBodyProcedure(s) { RefInExpr(e, _READ_); @@ -7174,6 +7215,24 @@ int ExplicitShape(SgExpression *eShape) return 1; } +int AssumedShape(SgExpression *eShape) +{ + SgExpression *el; + SgSubscriptExp *sbe; + for(el=eShape; el; el=el->rhs()) + { + //SgExpression *uBound = (sbe=isSgSubscriptExp(el->lhs())) ? sbe->ubound() : el->lhs(); + sbe=isSgSubscriptExp(el->lhs()); + if(sbe && !sbe->ubound()) + //if(!uBound) + continue; + else + return 0; + } + return 1; +} + + int TestArrayShape(SgSymbol *ar) { int i; @@ -7356,6 +7415,34 @@ SgExpression *FirstArrayElementSubscriptsForHandler(SgSymbol *ar) return(el); } +SgExpression *FirstArrayElementSubscriptsOfPrivateArray(SgSymbol *s) +{//generating reference AR(L1,...,Ln), where Li - lower bound of i-th dimension for kernel in C_Cuda + // Li - is constant or dummy argument reference + SgExpression *elist = NULL, *var; +/* + if (!TestArrayShape(s)) + { + var = ElementOfPrivateList(s); + SgExpression **eatr = (SgExpression **) var->lhs()->attributeValue(0, L_BOUNDS); + SgExpression *ela; + for (ela = *eatr; ela->rhs(); ela = ela->rhs()) + { + SgExpression *ed = new SgVarRefExp(ela->lhs()->lhs()->symbol()); + elist = AddListToList(new SgExprListExp(*ed), elist); + } + } + else + { + for (int i=0; iaddAttribute(NULL_SUBSCRIPTS, (void*)1, 0); + return elist; +} SgSymbol *DummyDvmHeaderSymbol(SgSymbol *ar, SgStatement *st_hedr) { @@ -8503,7 +8590,7 @@ SgStatement *CreateLoopKernel(SgSymbol *skernel, SgType *indexTypeInKernel) last = cur_st; - TranslateBlock(cur_in_kernel); + TranslateBlock(cur_in_kernel); if (options.isOn(C_CUDA)) { @@ -9534,7 +9621,7 @@ SgExpression *CreatePrivateDummyList() SgSymbol *s_dummy, *s; SgExpression *el, *ae; SgExpression *arg_list = NULL; - if (!options.isOn(C_CUDA) || !sizeOfPrivateArraysInBytes()) + if (!options.isOn(C_CUDA) || !PrivateArrayClassUse(sizeOfPrivateArraysInBytes())) // !sizeOfPrivateArraysInBytes()) return NULL; for (el = private_list; el; el = el->rhs()) { @@ -9984,6 +10071,67 @@ void DeclareInternalPrivateVars() } } +SgStatement *makeClassObjectDeclaration(SgSymbol *s, SgSymbol *sp, SgStatement *header_st, SgType *idxType, SgExpression *dim_list, int flag_true) +{ + SgStatement *st = new SgStatement(VAR_DECL); + SgSymbol *s_new = & s->copy(); + SYMB_SCOPE(s_new->thesymb) = header_st->thebif; + SgExpression *e = new SgExprListExp(*new SgTypeRefExp(*C_Type(s_new->type()))); + SgDerivedTemplateType *tp = new SgDerivedTemplateType(e, private_array_class); + tp->addArg(new SgValueExp(Rank(s))); + s_new->setType(tp); + SgFunctionCallExp *efc = new SgFunctionCallExp(*s_new); + efc->setType(tp); + st->setExpression(0, *new SgExprListExp(*efc)); + header_st->insertStmtAfter(*st); + + SgSymbol *s_dims=NULL; + SgStatement *st_dims = NULL; + if (Rank(s)>1) + { + char *name = new char[strlen(s->identifier())+7]; + sprintf(name, "_%s_dims", s->identifier()); + s_dims = ArraySymbol(name, idxType, new SgValueExp(Rank(s)-1), header_st); + SgExpression *einit = new SgExpression(INIT_LIST); +/* SgExpression *elist = NULL; + + if (for_kernel && !TestArrayShape(s)) + { + SgExpression **eatr = (SgExpression **) var->lhs()->attributeValue(0, DIM_SIZES); + SgExpression *ela; + for (ela = *eatr; ela->rhs(); ela = ela->rhs()) + { + SgExpression *ed = new SgVarRefExp(ela->lhs()->lhs()->symbol()); + elist = AddListToList(new SgExprListExp(*ed), elist); + } + } + else + { + for (int i=Rank(s)-1; i; i--) + elist = AddListToList(elist, new SgExprListExp(*Calculate(ArrayDimSize(s,i)))); + } + + einit->setLhs(elist); +*/ + einit->setLhs(dim_list); + SgStatement *st_dims = makeSymbolDeclarationWithInit(s_dims, einit); + header_st->insertStmtAfter(*st_dims); + //st_first = st_dims; + } + if (s_dims) + efc->addArg(*new SgVarRefExp(s_dims)); + + //SgSymbol **satr = (SgSymbol **) var->lhs()->attributeValue(0, PRIVATE_POINTER); + if (sp) + // { + // SgSymbol *sp = *satr; + efc->addArg(*new SgVarRefExp(sp)); + // } + if (flag_true) + efc->addArg(*new SgKeywordValExp("true")); + return (st_dims ? st_dims : st); +} + void DeclarePrivateVars() { DeclarePrivateVars(C_UnsignedLongLongType()); @@ -9994,14 +10142,19 @@ void DeclarePrivateVars(SgType *idxTypeInKernel) SgStatement *st = NULL, *st_first=NULL; SgExpression *var = NULL, *e; SgSymbol *s; + + if(!private_list) return; + SgExpression *e_all_private_size = sizeOfPrivateArraysInBytes(); + //SgSymbol *class_name = new SgSymbol(TYPE_NAME, "PrivateArray"); + // declare private variables for (var = private_list; var; var = var->rhs()) { s = var->lhs()->symbol(); if (isParDoIndexVar(s)) continue; // declared as index variable of parallel loop //if (HEADER(var->lhs()->symbol())) continue; // dvm-array declared as dummy argument - if (!options.isOn(C_CUDA) || !IS_ARRAY(s) || !e_all_private_size ) + if (!options.isOn(C_CUDA) || !IS_ARRAY(s) || !PrivateArrayClassUse(e_all_private_size)) { st = Declaration_Statement(SymbolInKernel(s)); kernel_st->insertStmtAfter(*st); @@ -10009,17 +10162,19 @@ void DeclarePrivateVars(SgType *idxTypeInKernel) } else { - SgSymbol *s_dims=NULL; - st = new SgStatement(PRIVATE_AR_DECL); - kernel_st->insertStmtAfter(*st); + SgStatement *st = new SgStatement(VAR_DECL); + SgSymbol *s_new = & s->copy(); + SYMB_SCOPE(s_new->thesymb) = kernel_st->thebif; + e = new SgExprListExp(*new SgTypeRefExp(*C_Type(s_new->type()))); + SgDerivedTemplateType *tp = new SgDerivedTemplateType(e, private_array_class); + tp->addArg(new SgValueExp(Rank(s))); + s_new->setType(tp); + SgFunctionCallExp *efc = new SgFunctionCallExp(*s_new); + efc->setType(tp); + st->setExpression(0, *new SgExprListExp(*efc)); + kernel_st->insertStmtAfter(*st); st_first = st; - - e = new SgExpression(TYPE_OP); - e->setType(C_Type(s->type()->baseType())); - st->setExpression(0, e); - - e = new SgValueExp(Rank(s)); - st->setExpression(1, e); + SgSymbol *s_dims=NULL; if (Rank(s)>1) { char *name = new char[strlen(s->identifier())+7]; @@ -10039,17 +10194,14 @@ void DeclarePrivateVars(SgType *idxTypeInKernel) } else { - for (int i=Rank(s)-1; i; i--) - elist = AddListToList(elist, new SgExprListExp(*Calculate(ArrayDimSize(s,i)))); + for (int i=Rank(s)-1; i; i--) + elist = AddListToList(elist, new SgExprListExp(*Calculate(ArrayDimSize(s,i)))); } einit->setLhs(elist); SgStatement *st_dims = makeSymbolDeclarationWithInit(s_dims, einit);//Declaration_Statement(s_dims); kernel_st->insertStmtAfter(*st_dims); st_first = st_dims; } - SgSymbol *s_new = & s->copy(); - SYMB_SCOPE(s_new->thesymb) = kernel_st->thebif; - SgFunctionCallExp *efc = new SgFunctionCallExp(*s_new); if (s_dims) { efc->addArg(*new SgVarRefExp(s_dims)); @@ -10058,11 +10210,11 @@ void DeclarePrivateVars(SgType *idxTypeInKernel) if (satr) { SgSymbol *sp = *satr; - efc->addArg(*new SgVarRefExp(sp)); //e->setLhs(new SgExprListExp(*new SgVarRefExp(sp))); - } - st->setExpression(2, efc); + efc->addArg(*new SgVarRefExp(sp)); + } } } + if (!st_first) return; @@ -14000,7 +14152,7 @@ SgStatement *Create_C_Adapter_Function(SgSymbol *sadapter) for (sg = hgpu_first, sb = base_first, sl = acc_array_list, ln = 0; lnnext(), sb = sb->next(), sl = sl->next, ln++) { e = new SgCastExp(*C_PointerType(options.isOn(C_CUDA) ? C_Type(sl->symb->type()) : new SgDescriptType(*SgTypeChar(), BIT_SIGNED)), *new SgVarRefExp(sb)); - fcall->addArg(*e); + fcall->addArg(*e); for (i = NumberOfCoeffs(sg); i>0; i--) fcall->addArg(*new SgArrayRefExp(*sg, *new SgValueExp(i))); } @@ -14105,7 +14257,7 @@ SgStatement *Create_C_Adapter_Function(SgSymbol *sadapter) } e_all_private_size = sizeOfPrivateArraysInBytes(); - if (options.isOn(C_CUDA) && e_all_private_size) + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) { for (el=private_list, lnp=0; el; el=el->rhs()) { @@ -14206,7 +14358,7 @@ SgStatement *Create_C_Adapter_Function(SgSymbol *sadapter) st_end->insertStmtBefore(*stmt, *st_hedr); // insert code for big private arrays - if (options.isOn(C_CUDA) && e_all_private_size) //(e_size = sizeOfPrivateArraysInBytes())) + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) //(e_size = sizeOfPrivateArraysInBytes())) { SgSymbol *s_private_size = new SgSymbol(VARIABLE_NAME, TestAndCorrectName("privateSizeForBlock"), *C_DvmType(), *st_hedr); stmt = makeSymbolDeclaration(s_private_size); @@ -14215,7 +14367,7 @@ SgStatement *Create_C_Adapter_Function(SgSymbol *sadapter) addDeclExpList(s_total_threads, stmt->expr(0)); SgExpression *e_threads = &(*new SgRecordRefExp(*s_threads, "x") * *new SgRecordRefExp(*s_threads, "y") * *new SgRecordRefExp(*s_threads, "z")); - SgExpression *e_private_size_for_block = &(*e_threads * *e_all_private_size); + SgExpression *e_private_size_for_block = &(*e_threads * *(e_all_private_size ? e_all_private_size : CalculateSizeOfPrivateArraysInBytes())); stmt = new SgCExpStmt(SgAssignOp(*new SgVarRefExp(*s_private_size), *e_private_size_for_block)); st_end->insertStmtBefore(*stmt, *st_hedr); @@ -14254,7 +14406,7 @@ SgStatement *Create_C_Adapter_Function(SgSymbol *sadapter) InsertFinishReductionCalls(st_end, s_loop_ref, s_red_num); // to dispose private arrays - if (options.isOn(C_CUDA) && e_all_private_size) + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) for (s = private_first, ln = 0; ln < lnp; s = s->next(), ln++) // private arrays { stmt = new SgCExpStmt(*DisposePrivateArray(s_loop_ref, s)); @@ -14583,6 +14735,19 @@ SgExpression *sizeOfElementInBytes(SgSymbol *symb) } SgExpression *sizeOfPrivateArraysInBytes() +{ + SgExpression *e_size = CalculateSizeOfPrivateArraysInBytes(); + if (e_size && e_size->isInteger()) // calculating length if it is possible + { + if (options.isOn(BIG_PRIVATES)) + return e_size; + else + return NULL; + } + return e_size; +} + +SgExpression *CalculateSizeOfPrivateArraysInBytes() { SgExpression *el, *e_size = NULL; int isize = 0; @@ -14614,19 +14779,18 @@ SgExpression *sizeOfPrivateArraysInBytes() } } if (e_size && e_size->isInteger()) // calculating length if it is possible - { - int i_size = e_size->valueInteger(); - e_size = new SgValueExp(i_size); - //TODO: need to add option - /*if (i_size > 2048) - return e_size; - else */ - return NULL; - } + e_size = new SgValueExp(e_size->valueInteger()); return e_size; } +int PrivateArrayClassUse(SgExpression *e_all_private_size) +{ + if (private_array_arg || e_all_private_size) + return 1; + return 0; +} + SgExpression *ProductOfDimSizeArgs(SgExpression *esizes) { SgExpression *el, *eprod = NULL; diff --git a/dvm/fdvm/trunk/fdvm/acc_across.cpp b/dvm/fdvm/trunk/fdvm/acc_across.cpp index d4bc926..43142dd 100644 --- a/dvm/fdvm/trunk/fdvm/acc_across.cpp +++ b/dvm/fdvm/trunk/fdvm/acc_across.cpp @@ -1638,7 +1638,7 @@ vector Create_C_Adapter_Function_Across_OneThread(SgSymbol *sadap } e_all_private_size = sizeOfPrivateArraysInBytes(); - if (options.isOn(C_CUDA) && e_all_private_size) + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) { for (el=private_list, lnp=0; el; el=el->rhs()) { @@ -1710,7 +1710,7 @@ vector Create_C_Adapter_Function_Across_OneThread(SgSymbol *sadap } } // insert code for big private arrays - if (options.isOn(C_CUDA) && e_all_private_size) + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) { GetMemoryForPrivateArrays(private_first, s_loop_ref, lnp, st_where, st_hedr, new SgValueExp(1)); @@ -1796,7 +1796,8 @@ static inline void insertReductionArgs(SgSymbol **reduction_ptr, SgSymbol **redu static void createPrivatePointers(SgSymbol* &private_first, int &lnp, SgStatement* st_hedr, SgExpression* &e_all_private_size) { private_first = NULL; - if (options.isOn(C_CUDA) && (e_all_private_size=sizeOfPrivateArraysInBytes())) + e_all_private_size = sizeOfPrivateArraysInBytes(); + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) { SgExpression *el, *ae; SgSymbol *sarg; @@ -3020,7 +3021,7 @@ vector Create_C_Adapter_Function_Across_variants(SgSymbol *sadapt } e_all_private_size = sizeOfPrivateArraysInBytes(); - if (options.isOn(C_CUDA) && e_all_private_size) + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) { for (el=private_list, lnp=0; el; el=el->rhs()) { @@ -3085,7 +3086,7 @@ vector Create_C_Adapter_Function_Across_variants(SgSymbol *sadapt stmt = simple; } stmt->addComment("// GPU execution"); - if (options.isOn(C_CUDA) && e_all_private_size) + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) { e_totalThreads = &(*new SgRecordRefExp(*s_blocks, "x") * *new SgRecordRefExp(*s_blocks, "y") * *new SgRecordRefExp(*s_blocks, "z") * *new SgRecordRefExp(*s_threads, "x") * *new SgRecordRefExp(*s_threads, "y") * *new SgRecordRefExp(*s_threads, "z")); GetMemoryForPrivateArrays(private_first, s_loop_ref, lnp, stmt, st_hedr, e_totalThreads); @@ -4067,7 +4068,7 @@ vector Create_C_Adapter_Function_Across_variants(SgSymbol *sadapt } } // to dispose private arrays - if (options.isOn(C_CUDA) && e_all_private_size) + if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size)) for (s = private_first, ln = 0; ln < lnp; s = s->next(), ln++) // private arrays { stmt = new SgCExpStmt(*DisposePrivateArray(s_loop_ref, s)); diff --git a/dvm/fdvm/trunk/fdvm/acc_f2c.cpp b/dvm/fdvm/trunk/fdvm/acc_f2c.cpp index 4ed299e..e64fd5f 100644 --- a/dvm/fdvm/trunk/fdvm/acc_f2c.cpp +++ b/dvm/fdvm/trunk/fdvm/acc_f2c.cpp @@ -103,7 +103,7 @@ static int lvl_convert_st = 0; // functions void convertExpr(SgExpression*, SgExpression*&); void createNewFCall(SgExpression *expr, SgExpression *&retExp, const char *name, int nArgs); - +static bool isPrivate(const string& array); #if TRACE void printfSpaces(int num) @@ -191,6 +191,14 @@ static bool inNewVars(const char *name) return ret; } +static bool isNullSubscripts(SgExpression *subs) +{ + if (subs && subs->attributeValue(0, NULL_SUBSCRIPTS)) + return true; + else + return false; +} + static void addInListIfNeed(SgSymbol *tmp, int type, reduction_operation_list *tmpR) { stack allArraySub; @@ -285,6 +293,14 @@ static void addRandStateIfNeeded(const string& name) private_list = e; } +void swapDimentionsInprivateList(SgExpression *pList) +{ + private_list = pList; + red_struct_list = NULL; + swapDimentionsInprivateList(); + private_list = NULL; +} + void swapDimentionsInprivateList() { SgExpression *tmp = private_list; @@ -639,46 +655,53 @@ SgExpression* switchArgumentsByKeyword(const string& name, SgExpression* funcCal if (argDims != dims) { char buf[256]; - sprintf(buf, "Dimention of the %d formal and actual parameters of '%s' call is not equal", i, name.c_str()); + sprintf(buf, "Rank of the %d dummy and actual arguments of '%s' call is not equal", i, name.c_str()); Error(buf, "", 651, first_do_par); } SgExpression* argList = NULL; - for (int j = 6; j >= 0; --j) + for (int j = MAX_DIMS; j >= 0; --j) { if (argInfo->elem(j) == NULL) continue; - //TODO: not checked!! - SgExpression* val = Calculate(&(*UpperBound(resultExprCall[i]->symbol(), j) - *LowerBound(resultExprCall[i]->symbol(), j) + *LowerBound(s->parameter(i), j))); - if (val != NULL) - tmp = new SgExprListExp(*val); - else - tmp = new SgExprListExp(*new SgValueExp(int(0))); + if (jsymbol(), j) - *LowerBound(resultExprCall[i]->symbol(), j) + *LowerBound(s->parameter(i), j))); + if (val != NULL) + tmp = new SgExprListExp(*val); + else + tmp = new SgExprListExp(*new SgValueExp(int(0))); - tmp->setRhs(argList); - argList = tmp; - val = LowerBound(s->parameter(i), j); - if (val != NULL) - tmp = new SgExprListExp(*val); - else - tmp = new SgExprListExp(*new SgValueExp(int(0))); - tmp->setRhs(argList); - argList = tmp; + tmp->setRhs(argList); + argList = tmp; + val = LowerBound(s->parameter(i), j); + if (val != NULL) + tmp = new SgExprListExp(*val); + else + tmp = new SgExprListExp(*new SgValueExp(int(0))); + tmp->setRhs(argList); + argList = tmp; + } } + if (isPrivate(resultExprCall[i]->symbol()->identifier())) //isPrivateArrayDummy==1 + { + resultExprCall[i] = new SgArrayRefExp(*resultExprCall[i]->symbol()); + } + else + { + SgArrayRefExp* arrRef = new SgArrayRefExp(*resultExprCall[i]->symbol()); + for (int j = 0; j < dims; ++j) + arrRef->addSubscript(*new SgValueExp(0)); - SgArrayRefExp* arrRef = new SgArrayRefExp(*resultExprCall[i]->symbol()); - for (int j = 0; j < dims; ++j) - arrRef->addSubscript(*new SgValueExp(0)); - - tmp = new SgExprListExp(SgAddrOp(*arrRef)); - tmp->setRhs(argList); - argList = tmp; - SgSymbol* aa = s->parameter(i); - - SgTypeRefExp* typeExpr = new SgTypeRefExp(*C_Type(s->parameter(i)->type())); - resultExprCall[i] = new SgFunctionCallExp(*((new SgDerivedTemplateType(typeExpr, new SgSymbol(TYPE_NAME, "s_array")))->typeName()), *argList); - resultExprCall[i]->setRhs(typeExpr); + tmp = new SgExprListExp(SgAddrOp(*arrRef)); + tmp->setRhs(argList); + argList = tmp; + SgSymbol* aa = s->parameter(i); + SgTypeRefExp* typeExpr = new SgTypeRefExp(*C_Type(s->parameter(i)->type())); + resultExprCall[i] = new SgFunctionCallExp(*((new SgDerivedTemplateType(typeExpr, new SgSymbol(TYPE_NAME, "s_array")))->typeName()), *argList); + resultExprCall[i]->setRhs(typeExpr); + } } } } @@ -925,7 +948,7 @@ static bool isPrivate(const string& array) static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isFunction) { bool ret = true; - + bool casePrivateArray = false; const string name(funcSymb->identifier()); vector *prototype = NULL; @@ -1072,15 +1095,29 @@ static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isF } else if (inCall->dimension() != inProt->dimension()) { - typeInCall = NULL; + if (isPrivate(argInCall->lhs()->symbol()->identifier()) && isPrivateArrayDummy(argInCall->lhs()->symbol()) != 1) + typeInCall = typeInProt; + else + typeInCall = NULL; + #ifdef DEB printf("typeInCall NULL 2\n"); #endif } else + { typeInCall = typeInProt; + if (for_kernel && isPrivate(argInCall->lhs()->symbol()->identifier()) || isPrivateArrayDummy(argInCall->lhs()->symbol())==1) + { + typeInCall = NULL; + casePrivateArray = true; +#ifdef DEB + printf("typeInCall NULL 2_p\n"); +#endif + } + } } - else + else // countOfSubscrInCall != 0 { //TODO: not supported yet if (inCall && inProt) @@ -1092,12 +1129,12 @@ static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isF } else { - if (options.isOn(O_PL2) && dvm_parallel_dir->expr(0) == NULL) + if (options.isOn(O_PL2) && dvm_parallel_dir && dvm_parallel_dir->expr(0) == NULL) dimSizeInProt = inCall->dimension(); const int arrayDim = isPrivate(argInCall->lhs()->symbol()->identifier()) ? inCall->dimension() : 1; - if (isSgArrayType(typeInProt) && (!options.isOn(O_PL2) || dvm_parallel_dir->expr(0) != NULL)) // inconsistency + if (isSgArrayType(typeInProt) && (!options.isOn(O_PL2) || !for_kernel || dvm_parallel_dir && dvm_parallel_dir->expr(0) != NULL)) // inconsistency { if (inCall->dimension() == inProt->dimension()) { @@ -1208,7 +1245,7 @@ static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isF } } else - ; //TODO + printf("typeInCall NULL 11\n"); //TODO } } @@ -1228,14 +1265,15 @@ static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isF if (CompareKind(typeInProt, typeInCall)) typeInCall = typeInProt; } - } - + } // end of type analysis + //---------------------------------------------------------------------------------------------------- if (typeInProt != typeInCall) { char buf[256]; sprintf(buf, "Can not match the %d argument of '%s' procedure", i + 1, name.c_str()); - Error(buf, "", 656, first_do_par); - ret = false; + if (!casePrivateArray) + Error(buf, "", 656, first_do_par); + //ret = false; } else if (argInCall->lhs()->variant() == ARRAY_REF) { @@ -1253,7 +1291,7 @@ static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isF { if (dimSizeInProt == 0) { - if (isFunction) + //if (isFunction) //04.02.25 podd { SgExpression* arrayRef = argInCall->lhs(); convertExpr(arrayRef, arrayRef); @@ -1307,7 +1345,9 @@ static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isF } else { - SgExpression* arr = argInCall->lhs(); + SgExpression* arr = argInCall->lhs(); + if (!isNullSubscripts(arr->lhs())) + convertExpr(arr, arr); if (options.isOn(O_PL2)) { @@ -1316,14 +1356,23 @@ static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isF cast = C_PointerType(C_Type(typeInProtSave->baseType())); else cast = C_PointerType(C_Type(typeInProtSave)); + if (for_kernel && isPrivate(arr->symbol()->identifier()) || isPrivateArrayDummy(arr->symbol())==2) + { + cast = C_PointerType(C_VoidType()); + } argInCall->setLhs(*new SgCastExp(*cast, SgAddrOp(*arr))); } else - argInCall->setLhs(SgAddrOp(*arr)); + { + if (for_kernel && isPrivate(arr->symbol()->identifier()) || isPrivateArrayDummy(arr->symbol())==2) + argInCall->setLhs(*new SgCastExp(*C_PointerType(C_VoidType()), SgAddrOp(*arr))); + else + argInCall->setLhs(SgAddrOp(*arr)); + } } } } - } + } //end of ARRAY_REF else { SgExpression* arg = argInCall->lhs(); @@ -1337,7 +1386,7 @@ static bool matchPrototype(SgSymbol *funcSymb, SgExpression *&listArgs, bool isF arg->setType(typeCopy); } - if (isFunction) + //if (isFunction) // 04.02.25 podd convertExpr(arg, arg); if (selector) @@ -2752,7 +2801,7 @@ static bool convertStmt(SgStatement* &st, pair &retS lvl_convert_st += 2; #endif SgExpression *lhs = st->expr(0); - convertExpr(lhs, lhs); + //convertExpr(lhs, lhs); // !!!! 04.02.25 podd if (lhs == NULL || SAPFOR_CONV) { diff --git a/dvm/fdvm/trunk/fdvm/calls.cpp b/dvm/fdvm/trunk/fdvm/calls.cpp index 7a08ba3..bd37e4c 100644 --- a/dvm/fdvm/trunk/fdvm/calls.cpp +++ b/dvm/fdvm/trunk/fdvm/calls.cpp @@ -41,7 +41,9 @@ int isStatementFunction(SgSymbol *s); int isHeaderNode(graph_node *gnode); int isDeadNode(graph_node *gnode); int isNoBodyNode(graph_node *gnode); -void PrototypeOfFunctionFromOtherFile(graph_node *node, SgStatement *after); +void InsertPrototypesOfFunctionFromOtherFile(graph_node *node, SgStatement *after); +void PrototypeOfFunctionFromOtherFile(graph_node *node, SgStatement *after, argument_numbers *arg_numbs); +void InsertCopiesOfProcedure(graph_node *ndl, SgStatement *after); graph_node_list *addToNodeList(graph_node_list *pnode, graph_node *gnode); graph_node_list *delFromNodeList(graph_node_list *pnode, graph_node *gnode); graph_node_list *isInNodeList(graph_node_list *pnode, graph_node *gnode); @@ -58,6 +60,7 @@ void ScanSymbolTable(SgFile *f); void ScanTypeTable(SgFile *f); void printSymb(SgSymbol *s); void printType(SgType *t); +void replaceVectorRef(SgExpression *e); //------------------------------------------------------------------------------------- extern SgExpression *private_list; extern map > > interfaceProcedures; @@ -103,7 +106,7 @@ int IsInternalProcedure(SgSymbol *s) SgStatement *hasInterface(SgSymbol *s) { - return (ATTR_NODE(s) ? GRAPHNODE(s)->st_interface : NULL); + return (ATTR_NODE(s) ? GRAPHNODE(s)->st_interface : NULL); } void SaveInterface(SgSymbol *s, SgStatement *interface) @@ -140,6 +143,25 @@ int isInParameter(SgSymbol *s, int i) { return (s && ((SgFunctionSymb *) s)->parameter(i) && (((SgFunctionSymb *) s)->parameter(i)->attributes() & IN_BIT) ? 1 : 0); } + +int isArrayParameter(SgSymbol *s, int i) +{ + return (s && ((SgFunctionSymb *) s)->parameter(i) && (((SgFunctionSymb *) s)->parameter(i)->attributes() & DIMENSION_BIT) ? 1 : 0); +} + +int isArrayParameterWithAssumedShape(SgSymbol *s, int i) +{ + return (isArrayParameter(s,i) && AssumedShape(isSgArrayType(((SgFunctionSymb *) s)->parameter(i)->type())->getDimList())); +} + +int isPrivateArrayDummy(SgSymbol *s) +{ + int *private_attr = (int *) s->attributeValue(0, DUMMY_PRIVATE_AR); + if (!private_attr) + return 0; + else + return *private_attr; +} SgSymbol *ProcedureSymbol(SgSymbol *s) { @@ -204,6 +226,7 @@ void MarkAsCalled(SgSymbol *s) { graph_node *gnode; edge *gedge; + if (!ATTR_NODE(s)) return; gnode = GRAPHNODE(s); @@ -215,6 +238,21 @@ void MarkAsCalled(SgSymbol *s) } +void MarkPrivateArgumentsOfRoutine(SgSymbol *s, SgExpression *private_args) +{ + SgExpression *el; + for (el=private_args; el; el=el->rhs()) + { + SgSymbol *arg = el->lhs()->symbol(); + if (IS_ARRAY(arg) && !IS_DVM_ARRAY(arg)) + { + int i = findParameterNumber(s,arg->identifier()); + if (i>=0) + addArgumentNumber(findParameterNumber(s,arg->identifier()), s); + } + } +} + void MakeFunctionCopy(SgSymbol *s) { SgSymbol *s_header; @@ -262,11 +300,13 @@ void InsertCalledProcedureCopies() { if (ndl->st_header && current_file_id == ndl->file_id) //procedure from current_file { - ndl->st_copy = InsertProcedureCopy(ndl->st_header, ndl->st_header->symbol(), ndl->is_routine, after); //C_Cuda ? mod_gpu : mod_gpu->lexNext()); + InsertCopiesOfProcedure(ndl, after); n++; } else //procedure from other file - PrototypeOfFunctionFromOtherFile(ndl,after); + { + InsertPrototypesOfFunctionFromOtherFile(ndl,after); + } ndl->count = 0; ndl->st_interface = NULL; @@ -345,11 +385,181 @@ int HasDerivedTypeVariables(SgStatement *header) return 0; } -SgStatement *InsertProcedureCopy(SgStatement *st_header, SgSymbol *sproc, int is_routine, SgStatement *after) +argument_numbers *GetFirstLenPlus(argument_numbers *source, int source_len, int list_len) { + // copy first (list_len+1) elements of source + if (list_len == source_len) + return NULL; + argument_numbers *cur_list, *source_list, *new_list, *new_elem; + new_list = new argument_numbers; + new_list->number = source->number; + new_list->next = NULL; + int i; + for (i=2, cur_list=new_list, source_list = source->next; i<=list_len+1; i++, source_list = source_list->next) + { + new_elem = new argument_numbers; + new_elem->number = source_list->number; + new_elem->next=NULL; + cur_list->next = new_elem; + cur_list = new_elem; + } + return new_list; +} + +argument_numbers *elementByValue(int numb, argument_numbers *nlist) +{ + for (; nlist; nlist=nlist->next) + if (nlist->number == numb) + return nlist; + return NULL; +} + +argument_numbers *element(int n, argument_numbers *nlist) +{ + for (int i=1; nlist; nlist=nlist->next, i++) + if (i == n) + return nlist; + return NULL; +} + +int numberOfElements(argument_numbers *nlist) +{ + int i; + for (i=0; nlist; nlist=nlist->next, i++) + ; + return i; +} + +void printValueList(argument_numbers *nlist) +{ + printf(" ("); + for (; nlist; nlist=nlist->next) + printf("%d ",nlist->number); + printf(") "); +} + +argument_numbers *GetNextWithChange(argument_numbers *source, int source_len, argument_numbers *nlist, int list_len) +{ + int i; + argument_numbers *elem, *source_elem; + for (i=1, elem=nlist; elem; i++, elem=elem->next) + if ( elem->number == (source_elem=element(source_len+i-list_len, source))->number ) + break; + + if (i == 1) return NULL; + elem = element(i-1, nlist); //element with serial number i-1 + int numb = elem->number; + source_elem = elementByValue(numb, source)->next; + for (int j=i-1; j<=list_len; j++, elem=elem->next, source_elem=source_elem->next) + elem->number = source_elem->number; + + return nlist; +} + +argument_numbers *GetNextNumberList(argument_numbers *source, argument_numbers *nlist) +{ + if (!source) return NULL; + if (!nlist) + { + nlist = new argument_numbers; + nlist->next = NULL; + nlist->number = source->number; + return nlist; + } + int source_len = numberOfElements(source); + int list_len = numberOfElements(nlist); + argument_numbers * last_elem = element(list_len, nlist); + argument_numbers *last_in_source = element(source_len, source); + + if (list_len == source_len) return NULL; + + argument_numbers *elem_in_source = elementByValue(last_elem->number, source); + if (elem_in_source != last_in_source) + { //get next in row + last_elem->number = elem_in_source->next->number; + return nlist; + } + else if ((nlist = GetNextWithChange(source, source_len, nlist, list_len))) + return nlist; + else + return GetFirstLenPlus(source, source_len, list_len); +} + +argument_numbers *correctArgList(argument_numbers *arg_numbs, SgStatement *st_header) +{ + SgSymbol *s = st_header->symbol(); + int i; + + argument_numbers *numb_list=NULL, *elem; + for (i=0; arg_numbs; arg_numbs=arg_numbs->next, i++) + { + if ( !isArrayParameterWithAssumedShape(s, arg_numbs->number) ) + { + elem = new argument_numbers; + elem->number = arg_numbs->number; + + if (numb_list) + { + elem->next = numb_list; + numb_list =elem; + } + else + elem->next = NULL; + numb_list = elem; + } + } + return numb_list; +} + +void InsertCopiesOfProcedure(graph_node *ndl, SgStatement *after) +{ + //insert copies of procedure after statement 'after' + argument_numbers *numb_list = NULL; + ndl->st_copy = InsertProcedureCopy(ndl->st_header, ndl->st_header->symbol(), ndl->is_routine, numb_list, after); + ndl->st_copy_first = ndl->st_copy; + + if (ndl->arg_numbs) + { + argument_numbers *arg_numbs = correctArgList(ndl->arg_numbs, ndl->st_header); + while ((numb_list = GetNextNumberList(arg_numbs, numb_list))) + ndl->st_copy = InsertProcedureCopy(ndl->st_header, ndl->st_header->symbol(), ndl->is_routine, numb_list, after); + } +} + +SgExpression *PrivateArrayDummyList(SgStatement *new_header, argument_numbers *arg_numbs) +{ + SgSymbol *s = new_header->symbol(); + SgExpression *pList = NULL; + SgExpression *ae; + int *id; + int n = ((SgFunctionSymb *)s)->numberOfParameters(); + for (int i = 0; i < n; i++) + { + SgSymbol *sarg = ((SgFunctionSymb *)s)->parameter(i); + if (isArrayParameterWithAssumedShape(s, i)) + { + id = new int; + *id = 1; + } + else if (arg_numbs && elementByValue(i, arg_numbs)) + { + id = new int; + *id = 2; + } + else + continue; + sarg->addAttribute(DUMMY_PRIVATE_AR, (void *)id, sizeof(int)); + ae = new SgArrayRefExp(*sarg); + ae ->setType(sarg->type()); + pList = AddListToList(pList, new SgExprListExp(*ae)); + } + return pList; +} + +SgStatement *InsertProcedureCopy(SgStatement *st_header, SgSymbol *sproc, int is_routine, argument_numbers *arg_numbs, SgStatement *after) +{ //InsertProcedureCopy(ndl->st_header, ndl->st_header->symbol(), ndl->is_routine, after) //insert copy of procedure after statement 'after' SgStatement *new_header, *end_st; - SgSymbol *new_sproc = &sproc->copySubprogram(*after); new_header = after->lexNext(); // new procedure header //new_sproc->body() SYMB_SCOPE(new_sproc->thesymb) = mod_gpu->thebif; @@ -362,13 +572,15 @@ SgStatement *InsertProcedureCopy(SgStatement *st_header, SgSymbol *sproc, int is int flagHasDerivedTypeVariables = HasDerivedTypeVariables(new_header); end_st = new_header->lastNodeOfStmt(); - ConvertArrayReferences(new_header->lexNext(), end_st); //!!!! + + private_list = PrivateArrayDummyList(new_header,arg_numbs); + ConvertArrayReferences(new_header->lexNext(), end_st); + + TranslateProcedureHeader_To_C(new_header,arg_numbs); - TranslateProcedureHeader_To_C(new_header); + // extract specification statements and add local arrays to private_list + ExtractDeclarationStatements(new_header); - private_list = NULL; - - ExtractDeclarationStatements(new_header); SgSymbol *s_last = LastSymbolOfFunction(new_header); if (sproc->variant() == FUNCTION_NAME) { @@ -378,20 +590,19 @@ SgStatement *InsertProcedureCopy(SgStatement *st_header, SgSymbol *sproc, int is InsertReturnBeforeEnd(new_header, end_st); } - swapDimentionsInprivateList(); - //std::vector < std::stack < SgStatement*> > zero = std::vector < std::stack < SgStatement*> >(0); - //cur_func = after; - Translate_Fortran_To_C(new_header, end_st, 0, st_header); //TranslateProcedure_Fortran_To_C(after->lexNext()); - cur_func = after; + + Translate_Fortran_To_C(new_header, end_st, 0, st_header); + cur_func = after; if (sproc->variant() == FUNCTION_NAME) { new_header->insertStmtAfter(*Declaration_Statement(new_sproc), *new_header); ChangeReturnStmts(new_header, end_st, returnSymbol); } + if(!flagHasDerivedTypeVariables) //!!! derived data type is not supported MakeFunctionDeclarations(new_header, s_last); - + newVars.clear(); private_list = NULL; // generate prototype of function and insert it before 'after' @@ -405,7 +616,7 @@ SgStatement *InsertProcedureCopy(SgStatement *st_header, SgSymbol *sproc, int is new_header->addComment("\n"); // add comment (empty line) to new procedure header ReplaceInterfaceBlocks(new_header); } - + return(new_header); } @@ -434,13 +645,15 @@ void doPrototype(SgStatement *func_hedr, SgStatement *block_header, int static_f block_header->insertStmtAfter(*st, *block_header); //before->insertStmtAfter(*st,*before->controlParent()); } -SgStatement *TranslateProcedureHeader_To_C(SgStatement *new_header) +SgStatement *TranslateProcedureHeader_To_C(SgStatement *new_header, argument_numbers *arg_numbs) { SgSymbol *new_sproc = new_header->symbol(); SgFunctionRefExp *fe = new SgFunctionRefExp(*new_sproc); fe->setSymbol(*new_sproc); new_header->setExpression(0, *fe); + SgSymbol *returnSymbol = getReturnSymbol(new_header, new_sproc); + if (new_sproc->variant() == PROCEDURE_NAME) new_sproc->setType(C_VoidType()); else // FUNCTION_NAME @@ -449,14 +662,30 @@ SgStatement *TranslateProcedureHeader_To_C(SgStatement *new_header) new_sproc->setType(C_Type(returnSymbol->type())); } fe->setType(new_sproc->type()); - fe->setLhs(FunctionDummyList(new_sproc)); + fe->setLhs(FunctionDummyList(new_sproc, new_header, arg_numbs)); BIF_LL3(new_header->thebif) = NULL; new_header->addDeclSpec(BIT_CUDA_DEVICE); new_header->setVariant(FUNC_HEDR); + return new_header; } -void PrototypeOfFunctionFromOtherFile(graph_node *node, SgStatement *after) +void InsertPrototypesOfFunctionFromOtherFile(graph_node *node, SgStatement *after) +{ + if (options.isOn(RTC)) return; + //insert prototypes of procedure after statement 'after' + argument_numbers *numb_list = NULL; + PrototypeOfFunctionFromOtherFile(node, after, numb_list); + + if (node->arg_numbs) + { + argument_numbers *arg_numbs = correctArgList(node->arg_numbs, node->st_header); + while ((numb_list = GetNextNumberList(arg_numbs, numb_list))) + PrototypeOfFunctionFromOtherFile(node, after, numb_list); + } +} + +void PrototypeOfFunctionFromOtherFile(graph_node *node, SgStatement *after, argument_numbers *arg_numbs) { if (options.isOn(RTC)) return; if(!node->st_interface) return; @@ -469,9 +698,9 @@ void PrototypeOfFunctionFromOtherFile(graph_node *node, SgStatement *after) SYMB_SCOPE(sh->thesymb) = current_file->firstStatement()->thebif; SgStatement *new_hedr = &(interface->copy()); new_hedr->setSymbol(*sh); - TranslateProcedureHeader_To_C(new_hedr); + TranslateProcedureHeader_To_C(new_hedr, arg_numbs); doPrototype(new_hedr, mod_gpu, !STATIC); - + //current_file->firstStatement()->insertStmtAfter(*new_hedr, *current_file->firstStatement()); //SYMB_FUNC_HEDR(sh->thesymb) = new_hedr->thebif; @@ -483,12 +712,11 @@ void PrototypeOfFunctionFromOtherFile(graph_node *node, SgStatement *after) return; } -SgExpression *FunctionDummyList(SgSymbol *s) +SgExpression *FunctionDummyList(SgSymbol *s, SgStatement *st_header, argument_numbers *arg_numbs) { SgExpression *arg_list = NULL, *ae = NULL; int n = ((SgFunctionSymb *)s)->numberOfParameters(); - //insert at 0-th position inf-argument //check for optional arguments, if some argunemt exist with optional then add argument-mask @@ -548,21 +776,73 @@ SgExpression *FunctionDummyList(SgSymbol *s) } dimList = dimList->rhs(); } + int rank = Rank(sarg); + SgType *ar_type = sarg->type(); + SgType *tbase = C_Type(sarg->type()->baseType()); + SgType *t = C_PointerType(tbase); + SgSymbol *new_arg = &sarg->copy();//new SgVariableSymb(sarg->identifier(), *t, *st_header); + //new_arg->thesymb->entry.var_decl.local = IO; // is new dummy argument + new_arg->setType(t); + ae = new SgVarRefExp(new_arg); + //ae->setType(t); - SgType *t = C_PointerType(C_Type(sarg->type()->baseType())); - sarg->setType(t); - ae = new SgVarRefExp(sarg); - ae->setType(t); if (needChanged) - { - sarg->setType(new SgDerivedTemplateType(new SgTypeRefExp(*t), new SgSymbol(TYPE_NAME, "s_array"))); - ae = new SgVarRefExp(sarg); + { + SgExpression *ce = new SgExprListExp(*new SgTypeRefExp(*tbase)); + SgDerivedTemplateType *tpc = new SgDerivedTemplateType(ce, private_array_class); + tpc->addArg(new SgValueExp(rank)); + new_arg->setType(tpc); + //int *id = new int; + //*id = 1; + //sarg->addAttribute(DUMMY_PRIVATE_AR, (void *)id, sizeof(int)); + ae = &SgAddrOp(*new SgVarRefExp(new_arg)); + //} + //else + //{ else + + // sarg->setType(new SgDerivedTemplateType(new SgTypeRefExp(*t), new SgSymbol(TYPE_NAME, "s_array"))); + // ae = new SgVarRefExp(sarg); + //} ae = new SgExprListExp(*ae); arg_list = AddListToList(arg_list, ae); continue; } - //ae->setType(C_ReferenceType(sarg->type())); + if (arg_numbs && elementByValue(i, arg_numbs)) + { + //SgType *tp = C_Type(sarg->type()->baseType()); + t = C_PointerType(C_VoidType()); + SgSymbol *sp = new SgSymbol(VARIABLE_NAME, PointerNameForPrivateArray(sarg), *t, *st_header); + //sp->thesymb->entry.var_decl.local = IO; + SgSymbol **symb = new (SgSymbol *); + *symb= sarg; + sp->addAttribute(FUNCTION_AR_DUMMY, (void*) symb, sizeof(SgSymbol *)); + ae->setSymbol(*sp); + //ae->setType(t); + //sarg->setType(ar_type); // restoration of argument type + //int *id = new int; + //*id = 2; + //sarg->addAttribute(DUMMY_PRIVATE_AR, (void *)id, sizeof(int)); + + //ae->setType(C_ReferenceType(t)); + //ae = new SgPointerDerefExp(*ae); + //ae = new SgCastExp(*C_PointerType(C_VoidType()), *ae); + //arg_list = AddListToList( int *id = new int; + //continue; + + //ae = new SgCastExp(*C_PointerType( t), *new SgVarRefExp(sp)); + // ae = new SgVarRefExp(sp); + //ae = new SgExprListExp(*ae); + //arg_list = AddListToList(arg_list, ae); + //continue; + + } + + //sarg->setType(t); + //ae = new SgVarRefExp(sarg); + //ae->setType(t); + + ae->setType(C_ReferenceType(t));//(sarg->type())); //t ae = new SgExprListExp(*new SgPointerDerefExp(*ae)); arg_list = AddListToList(arg_list, ae); //SgSymbol *arr_info = new SgSymbol(VAR_REF, ("inf_" + std::string(sarg->identifier())).c_str()); @@ -679,11 +959,11 @@ void ExtractDeclarationStatements(SgStatement *header) if(stmt->variant()==CONTROL_END) return; - while (stmt && !isSgExecutableStatement(stmt)) //is Fortran specification statement - { + while (stmt && (!isSgExecutableStatement(stmt) || stmt->variant()==ACC_ROUTINE_DIR)) //is Fortran specification statement or ROUTINE directive + { cur_st = stmt; stmt = stmt->lexNext(); - if(cur_st->variant() == INTERFACE_STMT || cur_st->variant() == INTERFACE_ASSIGNMENT || cur_st->variant() == INTERFACE_OPERATOR) + if (cur_st->variant() == INTERFACE_STMT || cur_st->variant() == INTERFACE_ASSIGNMENT || cur_st->variant() == INTERFACE_OPERATOR) { SgStatement *last = cur_st->lastNodeOfStmt(); SgStatement *start = cur_st; @@ -722,7 +1002,7 @@ void ExtractDeclarationStatements(SgStatement *header) } //if(cur_st->variant()==IMPL_DECL || cur_st->variant()==DATA_DECL || cur_st->variant()==USE_STMT || cur_st->variant()==FORMAT_STAT || cur_st->variant()==ENTRY_STAT || cur_st->variant()==COMM_STAT || cur_st->variant()==STMTFN_STAT ) if(!isSgVarDeclStmt(cur_st) && !isSgVarListDeclStmt(cur_st)) - { + { cur_st->extractStmt(); continue; } @@ -858,6 +1138,27 @@ void CorrectSubscript(SgExpression *e) e->setLhs((new SgExprListExp(*line))); } +void replaceArgument(SgSymbol *fsym, SgExpression *arg, int i) +{ + if (isSgArrayRefExp(arg) && !arg->lhs()) //argument is whole array (array name) + { + if (fsym && isPrivateArrayDummy(arg->symbol()) && !isArrayParameterWithAssumedShape(ProcedureSymbol(fsym),i)) + arg->setLhs(FirstArrayElementSubscriptsOfPrivateArray(arg->symbol())); + return; + } + replaceVectorRef(arg); + return; +} + +void replaceArgumentList(SgSymbol *fsym, SgExpression *arg_list) +{ + if (!arg_list) return; + int i; + SgExpression *el; + for (el=arg_list, i=0; el; el=el->rhs(),i++) + replaceArgument(fsym, el->lhs(), i); +} + void replaceVectorRef(SgExpression *e) { SgType *type; @@ -868,11 +1169,16 @@ void replaceVectorRef(SgExpression *e) type = isSgArrayType(e->symbol()->type()); if (IS_DUMMY(e->symbol()) && type) { - CorrectSubscript(e); + if (!isPrivateArrayDummy(e->symbol())) //isPrivate(e->symbol()->identifier()) + CorrectSubscript(e); } return; } - + if (isSgFunctionCallExp(e)) + { + replaceArgumentList(e->symbol(),e->lhs()); + return; + } replaceVectorRef(e->lhs()); replaceVectorRef(e->rhs()); } @@ -882,6 +1188,16 @@ void ConvertArrayReferences(SgStatement *first, SgStatement *last) SgStatement *st; for (st = first; st != last; st = st->lexNext()) { + if (isInterfaceStatement(st)) + { + st = st->lastNodeOfStmt(); + continue; + } + if (st->variant() == PROC_STAT) // call statement + { + replaceArgumentList(st->symbol(), st->expr(0)); + continue; + } if (st->expr(0)) replaceVectorRef(st->expr(0)); if (st->expr(1)) @@ -939,7 +1255,7 @@ void MakeFunctionDeclarations(SgStatement *header, SgSymbol *s_last) int flags = s->attributes(); if (IS_DUMMY(s)) - { + { if (flags & (IN_BIT | OUT_BIT | INOUT_BIT)) ; else if(!options.isOn(NO_PURE_FUNC)) @@ -957,13 +1273,23 @@ void MakeFunctionDeclarations(SgStatement *header, SgSymbol *s_last) //printf("%s: %d \n",s->identifier(),s->scope()->variant()); //printf("%s: %d %s \n",s->identifier(),s->scope()->variant(),s->scope()->symbol()->identifier()); continue; } + + SgSymbol **sarg = (SgSymbol **) s->attributeValue(0, FUNCTION_AR_DUMMY); + if (sarg) // pointer for PrivateArray class object + { + SgExpression *elist = NULL; + int rank = Rank(*sarg); + if (rank > 1) + for (int i=rank-1; i; i--) + elist = AddListToList(elist, new SgExprListExp(*Calculate(ArrayDimSize(*sarg,i)))); + makeClassObjectDeclaration(*sarg, s, header, C_UnsignedLongLongType(), elist, 1); //makeSymbolDeclaration(s); //MakePrivateArrayDeclaration(*sarg, s); + continue; + } if (!isSgArrayType(s->type())) //scalar variable s->setType(C_Type(s->type())); else - { continue; - } - + if (isSgConstantSymb(s)) { SgExpression *ce = ((SgConstantSymb *)s)->constantValue(); @@ -978,9 +1304,10 @@ void MakeFunctionDeclarations(SgStatement *header, SgSymbol *s_last) cur_stat->insertStmtAfter(*st); cur_stat = st; } - //printf("\n"); if(private_list) private_list->unparsestdout(); printf("\n"); + for (el = private_list; el; el = el->rhs()) { + if (IS_DUMMY(el->lhs()->symbol())) return; convertArrayDecl(el->lhs()->symbol()); st = makeSymbolDeclaration(el->lhs()->symbol()); cur_stat->insertStmtAfter(*st); @@ -1005,6 +1332,7 @@ void ProjectStructure(SgProject &project) int n = project.numberOfFiles(); SgFile *file; int i; + // building program structure // looking through the file list of project (first time) for (i = n - 1; i >= 0; i--) @@ -1246,6 +1574,7 @@ END_: PrintGraphNode(cur_node); in_routine = 0; return(last); + } void FunctionCallSearch(SgExpression *e) @@ -1609,13 +1938,13 @@ void Call_Site(SgSymbol *s, int inlined, SgStatement *stat, SgExpression *e) graph_node * gnode, *node_by_attr = NULL; SgSymbol *s_new = s; SgStatement *interface_st = NULL; - //printf("\n%s id= %d \n", s->identifier(), s->id()); + //printf("\n%s id= %d type= %d\n", s->identifier(), s->id(), s->type() ? s->type()->variant() : 0); if (!do_dummy && isDummyArgument(s)) return; if (!do_stmtfn && isStatementFunction(s)) return; // if(isIntrinsicFunction(s)) return; //printf("\nLINE %d", cur_st->lineNumber()); - - if(s->variant() == INTERFACE_NAME && in_region) + + if(s->variant() == INTERFACE_NAME && (in_region || in_routine)) { //printf("INTERFACE_NAME %s\n",s->identifier()); interface_st = getGenericInterface(s, stat ? stat->expr(0) : e->lhs()); @@ -1816,7 +2145,8 @@ graph_node *NewGraphNode(SgSymbol *s, SgStatement *header_st) gnode->clone = 0; gnode->count = 0; gnode->is_routine = 0; - gnode->st_interface = NULL; + gnode->st_interface = NULL; + gnode->arg_numbs = NULL; //printf("%s --- %d %d\n",gnode->name,gnode->id,gnode->type); return(gnode); } @@ -1981,6 +2311,24 @@ graph_node_list *isInNodeList(graph_node_list *pnode, graph_node *gnode) return (NULL); } +void addArgumentNumber(int i, SgSymbol *s) +{ + if (!ATTR_NODE(s)) + return; + graph_node *gnode = GRAPHNODE(s); + argument_numbers *nl; + for (nl=gnode->arg_numbs; nl; nl=nl->next) + if(i == nl->number) return; + nl = new argument_numbers; + nl->number = i; + if (gnode->arg_numbs) + { + nl->next = gnode->arg_numbs; + gnode->arg_numbs = nl; + } + else + gnode->arg_numbs = nl; +} void PrintGraphNode(graph_node *gnode) { @@ -2167,6 +2515,14 @@ void printSymb(SgSymbol *s) head = isHeaderStmtSymbol(s) ? "HEADER " : " "; printf("SYMB[%3d] scope=STMT[%3d] : %s %s", s->id(), (s->scope()) ? (s->scope())->id() : -1, s->identifier(), head); printType(s->type()); + if(s->variant() == CONST_NAME) + { + printf(" CONST_NAME "); + if(IS_BY_USE(s)) + printf(" BY_USE"); + printf("\n"); + return; + } if(IS_BY_USE(s)) printf(" BY_USE %s", ORIGINAL_SYMBOL(s)->scope()->symbol()->identifier()); if(ATTR_NODE(s)) diff --git a/dvm/fdvm/trunk/fdvm/dvm.cpp b/dvm/fdvm/trunk/fdvm/dvm.cpp index 7447afc..61c47f0 100644 --- a/dvm/fdvm/trunk/fdvm/dvm.cpp +++ b/dvm/fdvm/trunk/fdvm/dvm.cpp @@ -303,6 +303,8 @@ int main(int argc, char *argv[]) if ((*argv)[12] != '\0' && (isz = is_integer_value(*argv + 12))) UnparserBufSize = isz * 1024 * 1024; } + else if (!strcmp(argv[0], "-bigPrivates")) /*ACC*/ + options.setOn(BIG_PRIVATES); else if (!strcmp(argv[0], "-ioRTS")) options.setOn(IO_RTS); else if (!strcmp(argv[0], "-read_all")) @@ -1971,11 +1973,11 @@ void TranslateFileDVM(SgFile *f) } if(ACC_program) - { InsertCalledProcedureCopies(); + { InsertCalledProcedureCopies(); AddExternStmtToBlock_C(); GenerateEndIfDir(); GenerateDeclarationDir(); - GenerateStmtsForInfoFile(); + GenerateStmtsForInfoFile(); } } @@ -2093,8 +2095,8 @@ void TransFunc(SgStatement *func,SgStatement* &end_of_unit) { // all directives of F-DVM { //!!!debug - // printVariantName(stmt->variant()); //for debug - // printf("\n"); + //printVariantName(stmt->variant()); //for debug + //printf("\n"); //discovering distributed arrays in COMMON-blocks if(stmt->variant()==COMM_STAT) { @@ -2681,7 +2683,7 @@ void TransFunc(SgStatement *func,SgStatement* &end_of_unit) { // current statement is executable (F77/DVM) break; - } + } // checking semantics of DECLARE directives testDeclareDirectives(stmt); @@ -2884,6 +2886,7 @@ void TransFunc(SgStatement *func,SgStatement* &end_of_unit) { // testing procedure // if(dvm_debug && dbg_if_regim>1 && ((func->variant() == PROC_HEDR) || (func->variant() == FUNC_HEDR)) && !pstmt)// && !hasParallelDir(first_exec,func)) // copy_proc=1; + for(;pstmt; pstmt= pstmt->next) Extract_Stmt(pstmt->st);// extracting DVM Specification Directives @@ -3237,8 +3240,9 @@ EXEC_PART_: } else { // looking through the arguments list SgExpression * el; - for(el=stmt->expr(0); el; el=el->rhs()) - ChangeArg_DistArrayRef(el); // argument + int i; + for(el=stmt->expr(0),i=0; el; el=el->rhs(),i++) + ChangeArg_DistArrayRef(el,stmt->symbol(),i); // argument } break; case ALLOCATE_STMT: @@ -4243,8 +4247,8 @@ END_: // end of program unit cur_st = first_dvm_exec; if(last_dvm_entry) lentry = last_dvm_entry->lexNext(); // lentry - statement following first_dvm_exec or last generated dvm-initialization statement(before first_exec) - // before first_exec may be new statements generated for first_exec - if(!IN_MODULE) { + // before first_exec may be new statements generated for first_exec + if(!IN_MODULE) { if(has_contains) MarkCoeffsAsUsed(); InitBaseCoeffs(); @@ -7821,9 +7825,10 @@ void ChangeDistArrayRef(SgExpression *e) return; } if(isSgFunctionCallExp(e)) { + int i; ReplaceFuncCall(e); - for(el=e->lhs(); el; el=el->rhs()) - ChangeArg_DistArrayRef(el); + for(el=e->lhs(), i=0; el; el=el->rhs(),i++) + ChangeArg_DistArrayRef(el,e->symbol(),i); return; } @@ -7887,7 +7892,7 @@ void ChangeDistArrayRef_Left(SgExpression *e) return; } -void ChangeArg_DistArrayRef(SgExpression *ele) +void ChangeArg_DistArrayRef(SgExpression *ele, SgSymbol *fsym, int i) {//ele is SgExprListExp SgExpression *el, *e; e = ele->lhs(); @@ -7904,15 +7909,23 @@ void ChangeArg_DistArrayRef(SgExpression *ele) if(IS_POINTER(e->symbol())) Error("Illegal POINTER reference: '%s'",e->symbol()->identifier(),138,cur_st); if((inparloop && parloop_by_handler || IN_COMPUTE_REGION) ) + { if(DUMMY_FOR_ARRAY(e->symbol()) && isIn_acc_array_list(*DUMMY_FOR_ARRAY(e ->symbol())) ) { e->setLhs(FirstArrayElementSubscriptsForHandler(e->symbol())); //changed by first array element reference if(!for_host) - DistArrayRef(e,0,cur_st); + DistArrayRef(e,0,cur_st); } - if(HEADER(e->symbol()) && for_host) - e->setSymbol(*HeaderSymbolForHandler(e->symbol())); - return; + else if(options.isOn(C_CUDA) && for_kernel && isPrivate(e->symbol())) // && PrivateArrayClassUse(sizeOfPrivateArraysInBytes()))) + { + if(fsym && !isArrayParameterWithAssumedShape(ProcedureSymbol(fsym),i)) + e->setLhs(FirstArrayElementSubscriptsOfPrivateArray(e->symbol())); + } + } + if(HEADER(e->symbol()) && for_host) + e->setSymbol(*HeaderSymbolForHandler(e->symbol())); + + return; } el=e->lhs()->lhs(); //first subscript of argument //testing: is first subscript of ArrayRef a POINTER @@ -10604,8 +10617,8 @@ void InsertDebugStat(SgStatement *func, SgStatement* &end_of_unit) //including the DVM specification directive to list pstmt = addToStmtList(pstmt, stmt); continue; - case(ACC_ROUTINE_DIR): - case(ACC_DECLARE_DIR): + case(ACC_ROUTINE_DIR): + case(ACC_DECLARE_DIR): case(HPF_PROCESSORS_STAT): case(HPF_TEMPLATE_STAT): case(DVM_DYNAMIC_DIR): @@ -14253,10 +14266,11 @@ void TranslateFromTo(SgStatement *first, SgStatement *last, int error_msg) break; case PROC_STAT: // CALL - {SgExpression *el; + {SgExpression *el; + int i; // looking through the arguments list - for(el=stmt->expr(0); el; el=el->rhs()) - ChangeArg_DistArrayRef(el); // argument + for(el=stmt->expr(0), i=0; el; el=el->rhs(), i++) + ChangeArg_DistArrayRef(el, stmt->symbol(), i); // argument } break; diff --git a/dvm/fdvm/trunk/fdvm/stmt.cpp b/dvm/fdvm/trunk/fdvm/stmt.cpp index 4b4e46c..59cb720 100644 --- a/dvm/fdvm/trunk/fdvm/stmt.cpp +++ b/dvm/fdvm/trunk/fdvm/stmt.cpp @@ -1034,7 +1034,7 @@ RE: st = st->lastNodeOfStmt(); SgStatement * lastStmtOf(SgStatement *st) { SgStatement *last; - if(st->variant() == LOGIF_NODE) + if(st->variant() == LOGIF_NODE || st->variant() == FORALL_STAT) last = st->lexNext(); else if((st->variant() == FOR_NODE) || (st->variant() == WHILE_NODE)) last = lastStmtOfDo(st); @@ -1543,6 +1543,7 @@ int isDvmSpecification (SgStatement * st) { case DVM_CONSISTENT_GROUP_DIR: case DVM_CONSISTENT_DIR: case ACC_ROUTINE_DIR: + case ACC_DECLARE_DIR: return 1; break; } diff --git a/dvm/fdvm/trunk/include/calls.h b/dvm/fdvm/trunk/include/calls.h index 3449327..1a23a13 100644 --- a/dvm/fdvm/trunk/include/calls.h +++ b/dvm/fdvm/trunk/include/calls.h @@ -17,6 +17,7 @@ struct graph_node { SgStatement *st_header; SgStatement *st_last; SgStatement *st_copy; + SgStatement *st_copy_first; SgStatement *st_interface; SgSymbol *symb; //??? st_header->symbol() char *name; @@ -30,6 +31,7 @@ struct graph_node { int count; //counter of inline expansions or calls int is_routine;// has ROUTINE attribute - 1, else - 0 int samenamed; // flag - there is samenamed symbol + struct argument_numbers *arg_numbs; #if __SPF graph_node() { addToCollection(__LINE__, __FILE__, this, 1); } @@ -67,4 +69,6 @@ struct edge_list { edge_list() { addToCollection(__LINE__, __FILE__, this, 1); } ~edge_list() { removeFromCollection(this); } #endif -}; \ No newline at end of file +}; + + diff --git a/dvm/fdvm/trunk/include/dvm.h b/dvm/fdvm/trunk/include/dvm.h index 3d3be5c..45ed055 100644 --- a/dvm/fdvm/trunk/include/dvm.h +++ b/dvm/fdvm/trunk/include/dvm.h @@ -205,6 +205,11 @@ struct local_part_list { local_part_list *next; }; /*ACC*/ +struct argument_numbers { // numbers of dummy arguments that correspond to a private array when calling a procedure + argument_numbers *next; + int number; +}; /*ACC*/ + const int ROOT = 1; const int NODE = 2; const int ALIGN_TREE = 1000; @@ -265,6 +270,9 @@ const int L_BOUNDS = 1054; /*ACC*/ const int DIM_SIZES = 1055; /*ACC*/ const int PRIVATE_ARRAY = 1056; /*ACC*/ const int PRIVATE_POINTER = 1057; /*ACC*/ +const int FUNCTION_AR_DUMMY = 1058; /*ACC*/ +const int DUMMY_PRIVATE_AR = 1059; /*ACC*/ +const int NULL_SUBSCRIPTS = 1060; /*ACC*/ const int MAX_LOOP_LEVEL = 20; // 7 - maximal number of loops in parallel loop nest const int MAX_LOOP_NEST = 25; // maximal number of nested loops @@ -599,6 +607,7 @@ EXTERN int in_checksection,undefined_Tcuda, cuda_functions; /*ACC*/ EXTERN symb_list *RGname_list; /*ACC*/ EXTERN int parloop_by_handler; //set to 1 by option -Opl and /*ACC*/ // to 2 by option -Opl2 +EXTERN SgSymbol *private_array_class; //--------------------------------------------------------------------- /* dvm.cpp */ void TranslateFileDVM(SgFile *f); @@ -635,7 +644,8 @@ int Rank(SgSymbol *s); SgExpression *dvm_array_ref(); SgExpression *dvm_ref(int n); int DeleteDArFromList(SgStatement *stmt); -void ChangeArg_DistArrayRef(SgExpression *e); +//void ChangeArg_DistArrayRef(SgExpression *e); +void ChangeArg_DistArrayRef(SgExpression *ele, SgSymbol *fsym, int i); void ChangeDistArrayRef(SgExpression *e); void ChangeDistArrayRef_Left(SgExpression *e); SgExpression *SearchDistArrayField(SgExpression *e); @@ -1436,6 +1446,7 @@ void ExtractCopy(SgExpression *elist); void CleanAllocatedList(); SgStatement *CreateIndirectDistributionProcedure(SgSymbol *sProc,symb_list *paramList,symb_list *dummy_index_list,SgExpression *derived_elem_list,int flag); SgExpression *FirstArrayElementSubscriptsForHandler(SgSymbol *ar); +SgExpression *FirstArrayElementSubscriptsOfPrivateArray(SgSymbol *s); SgSymbol *HeaderSymbolForHandler(SgSymbol *ar); void TestRoutineAttribute(SgSymbol *s, SgStatement *routine_interface); int LookForRoutineDir(SgStatement *interfaceFunc); @@ -1453,8 +1464,12 @@ SgExpression *CreatePrivateDummyList(); char *PointerNameForPrivateArray(SgSymbol *symb); void GetMemoryForPrivateArrays(SgSymbol *private_first, SgSymbol *s_loop_ref, int nump, SgStatement *st_end, SgStatement *st_hedr, SgExpression *e_totalThreads); SgSymbol *LocRedVariableSymbolInKernel(reduction_operation_list *rsl); +int PrivateArrayClassUse(SgExpression *e_all_private_size); +SgExpression *CalculateSizeOfPrivateArraysInBytes(); +SgExpression *ElementOfPrivateList(SgSymbol *ar); void testDeclareDirectives(SgStatement *first_dvm_exec); void ACC_DECLARE_Directive(SgStatement *stmt); +SgStatement *makeClassObjectDeclaration(SgSymbol *s, SgSymbol *sp, SgStatement *header_st, SgType *idxType, SgExpression *dim_list, int flag_true); /* acc_analyzer.cpp */ //void Private_Vars_Analyzer(SgStatement *firstSt, SgStatement *lastSt); @@ -1563,6 +1578,7 @@ SgStatement *else_dir(); SgExpression *CalculateArrayBound(SgExpression *edim,SgSymbol *ar, int flag_private); void ReplaceArrayBoundsInDeclaration(SgExpression *e); int ExplicitShape(SgExpression *eShape); +int AssumedShape(SgExpression *eShape); SgSymbol *ArraySymbolInHostHandler(SgSymbol *ar,SgStatement *scope); SgSymbol *DeclareSymbolInHostHandler(SgSymbol *var, SgStatement *st_hedr, SgSymbol *loc_var); char *RegisterConstName(); @@ -2118,6 +2134,7 @@ void Translate_Fortran_To_C(SgStatement *stat, SgStatement *last, int countOfCop SgStatement* Translate_Fortran_To_C(SgStatement* Stmt, bool isSapforConv = false); SgSymbol* createNewFunctionSymbol(const char *name); void swapDimentionsInprivateList(void); +void swapDimentionsInprivateList(SgExpression *pList); void createNewFCall(SgExpression*, SgExpression*&, const char*, int); SgFunctionCallExp* createNewFCall(const char *name); void convertExpr(SgExpression*, SgExpression*&); @@ -2181,12 +2198,13 @@ void Call_Site (SgSymbol *s, int inlined, SgStatement *stat, SgExpression *e); SgSymbol * GetProcedureHeaderSymbol(SgSymbol *s); void MarkAsRoutine(SgSymbol *s); void MarkAsCalled(SgSymbol *s); +void MarkPrivateArgumentsOfRoutine(SgSymbol *s, SgExpression *private_args); void MarkAsUserProcedure(SgSymbol *s); void MarkAsExternalProcedure(SgSymbol *s); void MakeFunctionCopy(SgSymbol *s); SgStatement *HeaderStatement(SgSymbol *s); void InsertCalledProcedureCopies(); -SgStatement *InsertProcedureCopy(SgStatement *st_header, SgSymbol *sproc, int is_routine, SgStatement *after); +SgStatement *InsertProcedureCopy(SgStatement *st_header, SgSymbol *sproc, int is_routine, argument_numbers *arg_numbs, SgStatement *after); int FromOtherFile(SgSymbol *s); int findParameterNumber(SgSymbol *s, char *name); int isInParameter(SgSymbol *s, int i); @@ -2197,7 +2215,9 @@ int IsRecursiveProcedure(SgSymbol *s); int IsNoBodyProcedure(SgSymbol *s); int isUserFunction(SgSymbol *s); int IsInternalProcedure(SgSymbol *s); -SgExpression *FunctionDummyList(SgSymbol *s); +int isArrayParameter(SgSymbol *s, int i); +int isArrayParameterWithAssumedShape(SgSymbol *s, int i); +SgExpression *FunctionDummyList(SgSymbol *s, SgStatement *st_header, argument_numbers *arg_numbs); char *FunctionResultIdentifier(SgSymbol *sfun); SgSymbol *isSameNameInProcedure(char *name, SgSymbol *sfun); char *NameCheck(char *name, SgSymbol *sfun); @@ -2212,14 +2232,15 @@ SgStatement *FunctionPrototype(SgSymbol *sf); bool CreateIntefacePrototype(SgStatement *header); SgStatement *hasInterface(SgSymbol *s); void SaveInterface(SgSymbol *s, SgStatement *interface); -SgStatement *TranslateProcedureHeader_To_C(SgStatement *new_header); +SgStatement *TranslateProcedureHeader_To_C(SgStatement *new_header, argument_numbers *arg_numbs); SgStatement *getInterface(SgSymbol *s); SgStatement *getGenericInterface(SgSymbol *s, SgExpression *arg_list); int CompareKind(SgType* type_arg, SgType* type_dummy); SgExpression* TypeKindExpr(SgType* t); SgFunctionSymb *SymbolForIntrinsicFunction(const char *name, int i, SgType *tp, SgStatement *func); - - +void addArgumentNumber(int i, SgSymbol *s); +argument_numbers *GetNextNumberList(argument_numbers *source, argument_numbers *list); +int isPrivateArrayDummy(SgSymbol *s); //----------------------------------------------------------------------- extern "C" char* funparse_bfnd(...); extern "C" char* Tool_Unparse2_LLnode(...); @@ -2267,7 +2288,7 @@ void ConvertLoopWithLabelToEnddoLoop (SgStatement *stat); /*OMP*/ enum OPTIONS { AUTO_TFM = 0, ONE_THREAD, SPEED_TEST_L0, SPEED_TEST_L1, GPU_O0, GPU_O1, RTC, C_CUDA, OPT_EXP_COMP, O_HOST, NO_CUDA, NO_BL_INFO, LOOP_ANALYSIS, PRIVATE_ANALYSIS, IO_RTS, READ_ALL, NO_REMOTE, NO_PURE_FUNC, - GPU_IRR_ACC, O_PL, O_PL2, BIG_P, NUM_OPT}; + GPU_IRR_ACC, O_PL, O_PL2, BIG_PRIVATES, NUM_OPT}; // ONE_THREAD - compile one thread CUDA-kernels only for across (TODO for all CUDA-kernels) // SPEED_TEST_L0, SPEED_TEST_L1 - debug options for speed testof CUDA-kernels for across // RTC - enable CUDA run-time compilation of all CUDA-kernels diff --git a/dvm/fdvm/trunk/parser/cftn.c b/dvm/fdvm/trunk/parser/cftn.c index 871c9d4..3dae8f8 100644 --- a/dvm/fdvm/trunk/parser/cftn.c +++ b/dvm/fdvm/trunk/parser/cftn.c @@ -465,6 +465,8 @@ int main(int argc, char *argv[]) if ((*argv)[12] == '\0' || (!is_integer_value(*argv + 12))) goto ERR; } + else if (!strcmp(argv[0], "-bigPrivates")) + ; else if (!strcmp(argv[0], "-ioRTS")) ; else if (!strcmp(argv[0], "-read_all")) diff --git a/dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.DVMH_cuda.cu_opt b/dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.DVMH_cuda.cu_opt index 90ed1a4..1e80d19 100644 --- a/dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.DVMH_cuda.cu_opt +++ b/dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.DVMH_cuda.cu_opt @@ -487,7 +487,7 @@ typedef long long __indexTypeLLong; //--------------------- Kernel for loop on line 558 --------------------- - __global__ void loop_cg_558_cuda_kernel_int(double _p_rma[], double _q[], int _colidx[], double _a[], int _rowstr[], __indexTypeInt begin_1, __indexTypeInt end_1, __indexTypeInt add_blocks) + __global__ void loop_cg_558_cuda_kernel_int(double _q[], double _p[], int _colidx[], double _a[], int _rowstr[], __indexTypeInt begin_1, __indexTypeInt end_1, __indexTypeInt add_blocks) { // Private variables @@ -496,8 +496,8 @@ typedef long long __indexTypeLLong; int cond_0; int __k; int gid = blockIdx.x * blockDim.x + threadIdx.x; - int tid = gid / warpSize; - int lid = gid % warpSize; + int tid = gid / 32; + int lid = gid % 32; // Local needs __indexTypeInt _j; @@ -520,7 +520,7 @@ typedef long long __indexTypeLLong; __k < cond_0 ; _k = _k + warpSize, __k = __k + warpSize) { - _sum = _p_rma[_colidx[_k]] * _a[_k] + _sum; + _sum = _p[_colidx[_k]] * _a[_k] + _sum; } _sum = __dvmh_warpReduceSum(_sum); if (lid == 0) { @@ -532,7 +532,7 @@ typedef long long __indexTypeLLong; //--------------------- Kernel for loop on line 558 --------------------- - __global__ void loop_cg_558_cuda_kernel_llong(double _p_rma[], double _q[], int _colidx[], double _a[], int _rowstr[], __indexTypeLLong begin_1, __indexTypeLLong end_1, __indexTypeLLong add_blocks) + __global__ void loop_cg_558_cuda_kernel_llong(double _q[], double _p[], int _colidx[], double _a[], int _rowstr[], __indexTypeLLong begin_1, __indexTypeLLong end_1, __indexTypeLLong add_blocks) { // Private variables @@ -541,8 +541,8 @@ typedef long long __indexTypeLLong; int cond_0; int __k; int gid = blockIdx.x * blockDim.x + threadIdx.x; - int tid = gid / warpSize; - int lid = gid % warpSize; + int tid = gid / 32; + int lid = gid % 32; // Local needs __indexTypeLLong _j; __indexTypeLLong rest_blocks, cur_blocks; @@ -564,7 +564,7 @@ typedef long long __indexTypeLLong; __k < cond_0 ; _k = _k + warpSize, __k = __k + warpSize) { - _sum = _p_rma[_colidx[_k]] * _a[_k] + _sum; + _sum = _p[_colidx[_k]] * _a[_k] + _sum; } _sum = __dvmh_warpReduceSum(_sum); if (lid == 0) { @@ -752,7 +752,7 @@ typedef long long __indexTypeLLong; //--------------------- Kernel for loop on line 605 --------------------- - __global__ void loop_cg_605_cuda_kernel_int(double _z_rma[], double _r[], int _colidx[], double _a[], int _rowstr[], __indexTypeInt begin_1, __indexTypeInt end_1, __indexTypeInt add_blocks) + __global__ void loop_cg_605_cuda_kernel_int(double _r[], double _z[], int _colidx[], double _a[], int _rowstr[], __indexTypeInt begin_1, __indexTypeInt end_1, __indexTypeInt add_blocks) { // Private variables @@ -761,8 +761,8 @@ typedef long long __indexTypeLLong; int cond_0; int __k; int gid = blockIdx.x * blockDim.x + threadIdx.x; - int tid = gid / warpSize; - int lid = gid % warpSize; + int tid = gid / 32; + int lid = gid % 32; // Local needs __indexTypeInt _j; __indexTypeInt rest_blocks, cur_blocks; @@ -784,7 +784,7 @@ typedef long long __indexTypeLLong; __k < cond_0 ; _k = _k + warpSize, __k = __k + warpSize) { - _d = _z_rma[_colidx[_k]] * _a[_k] + _d; + _d = _z[_colidx[_k]] * _a[_k] + _d; } _d = __dvmh_warpReduceSum(_d); if (lid == 0) { @@ -796,7 +796,7 @@ typedef long long __indexTypeLLong; //--------------------- Kernel for loop on line 605 --------------------- - __global__ void loop_cg_605_cuda_kernel_llong(double _z_rma[], double _r[], int _colidx[], double _a[], int _rowstr[], __indexTypeLLong begin_1, __indexTypeLLong end_1, __indexTypeLLong add_blocks) + __global__ void loop_cg_605_cuda_kernel_llong(double _r[], double _z[], int _colidx[], double _a[], int _rowstr[], __indexTypeLLong begin_1, __indexTypeLLong end_1, __indexTypeLLong add_blocks) { // Private variables @@ -805,8 +805,8 @@ typedef long long __indexTypeLLong; int cond_0; int __k; int gid = blockIdx.x * blockDim.x + threadIdx.x; - int tid = gid / warpSize; - int lid = gid % warpSize; + int tid = gid / 32; + int lid = gid % 32; // Local needs __indexTypeLLong _j; __indexTypeLLong rest_blocks, cur_blocks; @@ -828,7 +828,7 @@ typedef long long __indexTypeLLong; __k < cond_0 ; _k = _k + warpSize, __k = __k + warpSize) { - _d = _z_rma[_colidx[_k]] * _a[_k] + _d; + _d = _z[_colidx[_k]] * _a[_k] + _d; } _d = __dvmh_warpReduceSum(_d); if (lid == 0) { @@ -1727,10 +1727,11 @@ extern "C" { // CUDA handler for loop on line 558 - void loop_cg_558_cuda_(DvmType *loop_ref, DvmType _p_rma[], DvmType _q[], DvmType _colidx[], DvmType _a[], DvmType _rowstr[]) + void loop_cg_558_cuda_(DvmType *loop_ref, DvmType _q[], DvmType _p[], DvmType _colidx[], DvmType _a[], DvmType _rowstr[]) { - void *p_rma_base, *q_base, *colidx_base, *a_base, *rowstr_base; - DvmType d_p_rma[4], d_q[4], d_colidx[4], d_a[4], d_rowstr[4]; + + void *q_base, *p_base, *colidx_base, *a_base, *rowstr_base; + DvmType d_q[4], d_p[4], d_colidx[4], d_a[4], d_rowstr[4]; DvmType idxTypeInKernel; dim3 blocks, threads; cudaStream_t stream; @@ -1742,15 +1743,15 @@ extern "C" { device_num = loop_get_device_num_(loop_ref); // Get 'natural' bases - p_rma_base = dvmh_get_natural_base(&device_num, _p_rma); q_base = dvmh_get_natural_base(&device_num, _q); + p_base = dvmh_get_natural_base(&device_num, _p); colidx_base = dvmh_get_natural_base(&device_num, _colidx); a_base = dvmh_get_natural_base(&device_num, _a); rowstr_base = dvmh_get_natural_base(&device_num, _rowstr); // Fill 'device' headers - dvmh_fill_header_(&device_num, p_rma_base, _p_rma, d_p_rma); dvmh_fill_header_(&device_num, q_base, _q, d_q); + dvmh_fill_header_(&device_num, p_base, _p, d_p); dvmh_fill_header_(&device_num, colidx_base, _colidx, d_colidx); dvmh_fill_header_(&device_num, a_base, _a, d_a); dvmh_fill_header_(&device_num, rowstr_base, _rowstr, d_rowstr); @@ -1778,8 +1779,8 @@ extern "C" { } loop_fill_bounds_(loop_ref, idxL, idxH, loopSteps); blocksS[0] = (idxH[0] - idxL[0] + threads.x) / threads.x; - overallBlocks = blocksS[0]; - restBlocks = overallBlocks * dvmh_get_warp_size(loop_ref); + overallBlocks = blocksS[0]* dvmh_get_warp_size(loop_ref);; + restBlocks = overallBlocks; addBlocks = 0; blocks = dim3(1, 1, 1); maxBlocks = loop_cuda_get_device_prop(loop_ref, CUDA_MAX_GRID_X); @@ -1798,11 +1799,11 @@ extern "C" { } if (idxTypeInKernel == rt_INT) { - loop_cg_558_cuda_kernel_int<<>>((double *)p_rma_base, (double *)q_base, (int *)colidx_base, (double *)a_base, (int *)rowstr_base, idxL[0], idxH[0], addBlocks); + loop_cg_558_cuda_kernel_int<<>>((double *)q_base, (double *)p_base, (int *)colidx_base, (double *)a_base, (int *)rowstr_base, idxL[0], idxH[0], addBlocks); } else { - loop_cg_558_cuda_kernel_llong<<>>((double *)p_rma_base, (double *)q_base, (int *)colidx_base, (double *)a_base, (int *)rowstr_base, idxL[0], idxH[0], addBlocks); + loop_cg_558_cuda_kernel_llong<<>>((double *)q_base, (double *)p_base, (int *)colidx_base, (double *)a_base, (int *)rowstr_base, idxL[0], idxH[0], addBlocks); } addBlocks += blocks.x; restBlocks -= blocks.x; @@ -2098,10 +2099,10 @@ extern "C" { // CUDA handler for loop on line 605 - void loop_cg_605_cuda_(DvmType *loop_ref, DvmType _z_rma[], DvmType _r[], DvmType _colidx[], DvmType _a[], DvmType _rowstr[]) + void loop_cg_605_cuda_(DvmType *loop_ref, DvmType _r[], DvmType _z[], DvmType _colidx[], DvmType _a[], DvmType _rowstr[]) { - void *z_rma_base, *r_base, *colidx_base, *a_base, *rowstr_base; - DvmType d_z_rma[4], d_r[4], d_colidx[4], d_a[4], d_rowstr[4]; + void *r_base, *z_base, *colidx_base, *a_base, *rowstr_base; + DvmType d_r[4], d_z[4], d_colidx[4], d_a[4], d_rowstr[4]; DvmType idxTypeInKernel; dim3 blocks, threads; cudaStream_t stream; @@ -2113,15 +2114,15 @@ extern "C" { device_num = loop_get_device_num_(loop_ref); // Get 'natural' bases - z_rma_base = dvmh_get_natural_base(&device_num, _z_rma); r_base = dvmh_get_natural_base(&device_num, _r); + z_base = dvmh_get_natural_base(&device_num, _z); colidx_base = dvmh_get_natural_base(&device_num, _colidx); a_base = dvmh_get_natural_base(&device_num, _a); rowstr_base = dvmh_get_natural_base(&device_num, _rowstr); // Fill 'device' headers - dvmh_fill_header_(&device_num, z_rma_base, _z_rma, d_z_rma); dvmh_fill_header_(&device_num, r_base, _r, d_r); + dvmh_fill_header_(&device_num, z_base, _z, d_z); dvmh_fill_header_(&device_num, colidx_base, _colidx, d_colidx); dvmh_fill_header_(&device_num, a_base, _a, d_a); dvmh_fill_header_(&device_num, rowstr_base, _rowstr, d_rowstr); @@ -2155,7 +2156,7 @@ extern "C" { blocks = dim3(1, 1, 1); maxBlocks = loop_cuda_get_device_prop(loop_ref, CUDA_MAX_GRID_X); maxBlocks = maxBlocks / dvmh_get_warp_size(loop_ref) * dvmh_get_warp_size(loop_ref); - + // GPU execution while (restBlocks > 0) { @@ -2169,15 +2170,15 @@ extern "C" { } if (idxTypeInKernel == rt_INT) { - loop_cg_605_cuda_kernel_int<<>>((double *)z_rma_base, (double *)r_base, (int *)colidx_base, (double *)a_base, (int *)rowstr_base, idxL[0], idxH[0], addBlocks); + loop_cg_605_cuda_kernel_int<<>>((double *)r_base, (double *)z_base, (int *)colidx_base, (double *)a_base, (int *)rowstr_base, idxL[0], idxH[0], addBlocks); } else { - loop_cg_605_cuda_kernel_llong<<>>((double *)z_rma_base, (double *)r_base, (int *)colidx_base, (double *)a_base, (int *)rowstr_base, idxL[0], idxH[0], addBlocks); + loop_cg_605_cuda_kernel_llong<<>>((double *)r_base, (double *)z_base, (int *)colidx_base, (double *)a_base, (int *)rowstr_base, idxL[0], idxH[0], addBlocks); } addBlocks += blocks.x; restBlocks -= blocks.x; - } + } } diff --git a/dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.fdv b/dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.fdv index f077345..1f6e535 100644 --- a/dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.fdv +++ b/dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.fdv @@ -550,11 +550,11 @@ c The conj grad iteration loop c----> c--------------------------------------------------------------------- do cgit = 1, cgitmax - d = 0.0d0 +! DVM$ interval 11 CDVM$ region - -CDVM$ parallel (j) on p(j), private(sum,k), remote_access(p(:)) +!WANR for many process, remote_access(p(:)) is needed +CDVM$ parallel (j) on p(j), private(sum,k) do j=1,lastrow-firstrow+1 sum = 0.d0 do k=rowstr(j),rowstr(j+1)-1 @@ -570,7 +570,7 @@ CDVM$ parallel (j) on q(j), reduction(SUM(d)) CDVM$ end region alpha = rho / d rho0 = rho - +! DVM$ end interval rho = 0.0d0 CDVM$ region CDVM$ parallel (j) on r(j), private(d), reduction(SUM(rho)) @@ -598,10 +598,10 @@ c Compute residual norm explicitly: ||r|| = ||x - A.z|| c First, form A.z c The partition submatrix-vector multiply c--------------------------------------------------------------------- - +!WANR for many process, remote_access(z(:)) is needed sum = 0.0d0 CDVM$ region -CDVM$ parallel (j) on r(j), private(d,k),remote_access(z(:)) +CDVM$ parallel (j) on r(j), private(d,k) do j=1,lastrow-firstrow+1 d = 0.d0 do k=rowstr(j),rowstr(j+1)-1 diff --git a/sapfor/experts/Sapfor_2017/_src/DvmhRegions/DvmhRegionInserter.cpp b/sapfor/experts/Sapfor_2017/_src/DvmhRegions/DvmhRegionInserter.cpp index 6096dbf..70549c4 100644 --- a/sapfor/experts/Sapfor_2017/_src/DvmhRegions/DvmhRegionInserter.cpp +++ b/sapfor/experts/Sapfor_2017/_src/DvmhRegions/DvmhRegionInserter.cpp @@ -1561,12 +1561,25 @@ static set for (auto& decl : decls) { declStat = SgStatement::getStatementByFileAndLine(decl.first, decl.second); - checkNull(declStat, convertFileName(__FILE__).c_str(), __LINE__); + if (declStat == NULL) // check in inlcudes + { + for (auto st = main; st != main->lastNodeOfStmt() && !declStat; st = st->lexNext()) + { + if (st->fileName() == decl.first && st->lineNumber() == decl.second) + declStat = st; + } - if (declStat != main) + if (declStat) + break; + } + else { - declStat = NULL; - continue; + declStat = getFuncStat(declStat); + if (declStat != main) + { + declStat = NULL; + continue; + } } } } diff --git a/sapfor/experts/Sapfor_2017/_src/GraphCall/graph_calls.cpp b/sapfor/experts/Sapfor_2017/_src/GraphCall/graph_calls.cpp index a44543c..945ae77 100644 --- a/sapfor/experts/Sapfor_2017/_src/GraphCall/graph_calls.cpp +++ b/sapfor/experts/Sapfor_2017/_src/GraphCall/graph_calls.cpp @@ -2235,10 +2235,10 @@ static bool hasRecursionChain(vector currentChainCalls, const FuncInf currentChainCalls.push_back(itF->second); const string &chain = printChainRec(currentChainCalls); - __spf_print(1, "For function on line %d found recursive chain calls: %s\n", currentChainCalls[0]->linesNum.first, chain.c_str()); + __spf_print(1, " For function on line %d found recursive chain calls: %s\n", currentChainCalls[0]->linesNum.first, chain.c_str()); wstring bufE, bufR; - __spf_printToLongBuf(bufE, L"Found recursive chain calls: %s, this function will be ignored", to_wstring(chain).c_str()); + __spf_printToLongBuf(bufE, L" Found recursive chain calls: %s, this function will be ignored", to_wstring(chain).c_str()); __spf_printToLongBuf(bufR, R46, to_wstring(chain).c_str()); messagesForFile.push_back(Messages(NOTE, currentChainCalls[0]->linesNum.first, bufR, bufE, 1014)); @@ -2273,7 +2273,7 @@ void checkForRecursion(SgFile *file, map> &allFuncInfo for (int i = 0; i < itCurrFuncs->second.size(); ++i) { - __spf_print(1, "run for func %s\n", itCurrFuncs->second[i]->funcName.c_str()); + __spf_print(1, " run for func %s\n", itCurrFuncs->second[i]->funcName.c_str()); if (hasRecursionChain( { itCurrFuncs->second[i] }, itCurrFuncs->second[i], mapFuncInfo, messagesForFile)) itCurrFuncs->second[i]->doNotAnalyze = true; } diff --git a/sapfor/experts/Sapfor_2017/_src/GraphCall/graph_calls_base.cpp b/sapfor/experts/Sapfor_2017/_src/GraphCall/graph_calls_base.cpp index 824e8f5..47f8fef 100644 --- a/sapfor/experts/Sapfor_2017/_src/GraphCall/graph_calls_base.cpp +++ b/sapfor/experts/Sapfor_2017/_src/GraphCall/graph_calls_base.cpp @@ -404,6 +404,7 @@ void findDeadFunctionsAndFillCalls(map> &allFuncInfo, set allExternalCalls; set allChildCalls; + for (auto &it : mapFuncInfo) { FuncInfo *currInfo = it.second; @@ -448,9 +449,37 @@ void findDeadFunctionsAndFillCalls(map> &allFuncInfo, } } } + + FuncInfo* main = NULL; + for (auto& it : mapFuncInfo) + if (it.second->isMain) + main = it.second; + + checkNull(main, convertFileName(__FILE__).c_str(), __LINE__); + + set liveFunctions; + liveFunctions.insert(main); + for (auto& callFrom : main->callsFromV) + liveFunctions.insert(callFrom); + + //find live functions + bool changes = true; + while (changes) + { + changes = false; + for (auto& currInfo : liveFunctions) + { + for (auto& callFrom : currInfo->callsFromV) { + if (liveFunctions.find(callFrom) == liveFunctions.end()) { + changes = true; + liveFunctions.insert(callFrom); + } + } + } + } // propagate 'deadFunction' status for all 'CallsFrom' from dead functions - bool changes = true; + changes = true; while (changes) { changes = false; @@ -466,10 +495,11 @@ void findDeadFunctionsAndFillCalls(map> &allFuncInfo, auto itFrom = mapFuncInfo.find(callFrom); if (itFrom != mapFuncInfo.end()) { - if (!itFrom->second->deadFunction) + auto func = itFrom->second; + if (!func->deadFunction && liveFunctions.find(func) == liveFunctions.end()) { changes = true; - itFrom->second->deadFunction = itFrom->second->doNotAnalyze = true; + func->deadFunction = func->doNotAnalyze = true; } } } diff --git a/sapfor/experts/Sapfor_2017/_src/GraphLoop/graph_loops_base.cpp b/sapfor/experts/Sapfor_2017/_src/GraphLoop/graph_loops_base.cpp index b1f9e35..98e13ca 100644 --- a/sapfor/experts/Sapfor_2017/_src/GraphLoop/graph_loops_base.cpp +++ b/sapfor/experts/Sapfor_2017/_src/GraphLoop/graph_loops_base.cpp @@ -778,7 +778,7 @@ static void isAllOk(const vector &loops, vector &currMessa if (loops[i]->countOfIters == 0 && loops[i]->region && loops[i]->isFor) { wstring bufE, bufR; - __spf_printToLongBuf(bufE, L"Can not calculate count of iterations for this loop, information about iterations in all loops in parallel regions '%s' will be ignored", + __spf_printToLongBuf(bufE, L" Can not calculate count of iterations for this loop, information about iterations in all loops in parallel regions '%s' will be ignored", to_wstring(loops[i]->region->GetName()).c_str()); auto itM = uniqMessages.find(bufE); @@ -789,7 +789,7 @@ static void isAllOk(const vector &loops, vector &currMessa __spf_printToLongBuf(bufR, R48, to_wstring(loops[i]->region->GetName()).c_str()); currMessages.push_back(Messages(NOTE, loops[i]->lineNum, bufR, bufE, 1016)); - __spf_print(1, "Can not calculate count of iterations for loop on line %d, information about iterations in all loops in parallel regions '%s' will be ignored\n", loops[i]->lineNum, loops[i]->region->GetName().c_str()); + __spf_print(1, " Can not calculate count of iterations for loop on line %d, information about iterations in all loops in parallel regions '%s' will be ignored\n", loops[i]->lineNum, loops[i]->region->GetName().c_str()); } isNotOkey.insert(loops[i]->region); } @@ -1085,7 +1085,7 @@ static void checkArraysMapping(vector &loopList, mapIsDimDepracated(z)) { std::wstring bufw, bufR; - __spf_printToLongBuf(bufw, L"Array '%s' can not be distributed due to different writes to %d dimension, this dimension will deprecated", + __spf_printToLongBuf(bufw, L" Array '%s' can not be distributed due to different writes to %d dimension, this dimension will deprecated", to_wstring(elem.first->GetShortName()).c_str(), z + 1); __spf_printToLongBuf(bufR, R85, z + 1,to_wstring(elem.first->GetShortName()).c_str()); @@ -1122,7 +1122,7 @@ void checkArraysMapping(const map> &loopGraph, mapIsAllDeprecated()) { wstring bufw, bufR; - __spf_printToLongBuf(bufw, L"Array '%s' can not be distributed due to all dimensions will deprecated", to_wstring(elem->GetShortName()).c_str()); + __spf_printToLongBuf(bufw, L" Array '%s' can not be distributed due to all dimensions will deprecated", to_wstring(elem->GetShortName()).c_str()); __spf_printToLongBuf(bufR, R86, to_wstring(elem->GetShortName()).c_str()); for (auto &decl : elem->GetDeclInfo()) @@ -1210,7 +1210,7 @@ static void filterArrayInCSRGraph(vector &loops, const mapsecond < 0) { wstring bufw, bufR; - __spf_printToLongBuf(bufw, L"Array '%s' can not be distributed", to_wstring(array->GetShortName()).c_str()); + __spf_printToLongBuf(bufw, L" Array '%s' can not be distributed", to_wstring(array->GetShortName()).c_str()); __spf_printToLongBuf(bufR, R87, to_wstring(array->GetShortName()).c_str()); getObjectForFileFromMap(loop->fileName.c_str(), messages).push_back(Messages(NOTE, loop->lineNum, bufR, bufw, 1047)); @@ -1251,7 +1251,7 @@ static void filterArrayInCSRGraph(vector &loops, const mapsecond != treeNum) { wstring bufw, bufR; - __spf_printToLongBuf(bufw, L"Array '%s' can not be distributed", to_wstring(array->GetShortName()).c_str()); + __spf_printToLongBuf(bufw, L" Array '%s' can not be distributed", to_wstring(array->GetShortName()).c_str()); __spf_printToLongBuf(bufR, R88, to_wstring(array->GetShortName()).c_str()); getObjectForFileFromMap(loop->fileName.c_str(), messages).push_back(Messages(NOTE, loop->lineNum, bufR, bufw, 1047)); @@ -1283,7 +1283,7 @@ static void filterArrayInCSRGraph(vector &loops, const mapGetShortName()).c_str()); + __spf_printToLongBuf(bufw, L" Array '%s' can not be distributed", to_wstring(inCall->GetShortName()).c_str()); __spf_printToLongBuf(bufR, R89, to_wstring(inCall->GetShortName()).c_str()); getObjectForFileFromMap(loop->fileName.c_str(), messages).push_back(Messages(NOTE, loop->lineNum, bufR, bufw, 1047)); diff --git a/sapfor/experts/Sapfor_2017/_src/ParallelizationRegions/ParRegions.cpp b/sapfor/experts/Sapfor_2017/_src/ParallelizationRegions/ParRegions.cpp index d22d115..85d1953 100644 --- a/sapfor/experts/Sapfor_2017/_src/ParallelizationRegions/ParRegions.cpp +++ b/sapfor/experts/Sapfor_2017/_src/ParallelizationRegions/ParRegions.cpp @@ -836,13 +836,13 @@ bool buildGraphFromUserDirectives(const vector &userDvmAlignDirs, DI string tmp; for (auto& elem : realAlignArrayRefsSet) tmp += elem->GetName() + " "; - __spf_print(1, "align array%s '%s' from user dir in line %d\n", (realAlignArrayRefsSet.size() == 1 ? "" : "s"), tmp.c_str(), dir->lineNumber()); - __spf_print(1, "template align:\n"); + __spf_print(1, " align array%s '%s' from user dir in line %d\n", (realAlignArrayRefsSet.size() == 1 ? "" : "s"), tmp.c_str(), dir->lineNumber()); + __spf_print(1, " template align:\n"); for (int i = 0; i < alignTemplate.size(); ++i) - __spf_print(1, "-- %d: %s -- [%d, %d]\n", i, alignTemplate[i].first.c_str(), alignTemplate[i].second.first, alignTemplate[i].second.second); - __spf_print(1, "template align with:\n"); + __spf_print(1, " -- %d: %s -- [%d, %d]\n", i, alignTemplate[i].first.c_str(), alignTemplate[i].second.first, alignTemplate[i].second.second); + __spf_print(1, " template align with:\n"); for (int i = 0; i < alignWithTemplate.size(); ++i) - __spf_print(1, "-- %d: %s -- [%d, %d]\n", i, alignWithTemplate[i].first.c_str(), alignWithTemplate[i].second.first, alignWithTemplate[i].second.second); + __spf_print(1, " -- %d: %s -- [%d, %d]\n", i, alignWithTemplate[i].first.c_str(), alignWithTemplate[i].second.first, alignWithTemplate[i].second.second); for (int i = 0; i < alignTemplate.size(); ++i) { @@ -891,6 +891,6 @@ void calculateLinesOfCode(vector &allRegions) for (auto &lineV : line.second) lineCounter += (lineV.lines.second - lineV.lines.first); - __spf_print(1, "Count of lines in region '%s' = %d\n", elem->GetName().c_str(), lineCounter); + __spf_print(1, " Count of lines in region '%s' = %d\n", elem->GetName().c_str(), lineCounter); } } \ No newline at end of file diff --git a/sapfor/experts/Sapfor_2017/_src/Sapfor.cpp b/sapfor/experts/Sapfor_2017/_src/Sapfor.cpp index c4d9961..c221a1c 100644 --- a/sapfor/experts/Sapfor_2017/_src/Sapfor.cpp +++ b/sapfor/experts/Sapfor_2017/_src/Sapfor.cpp @@ -1283,7 +1283,7 @@ static bool runAnalysis(SgProject &project, const int curr_regime, const bool ne if (loop.second->hasLimitsToParallel()) { loop.second->addConflictMessages(&SPF_messages[loop.second->fileName]); - __spf_print(1, "added conflict messages to loop on line %d\n", loop.second->lineNum); + __spf_print(1, " added conflict messages to loop on line %d\n", loop.second->lineNum); } } } diff --git a/sapfor/experts/Sapfor_2017/_src/Utils/utils.cpp b/sapfor/experts/Sapfor_2017/_src/Utils/utils.cpp index 13d4081..a5f9cdc 100644 --- a/sapfor/experts/Sapfor_2017/_src/Utils/utils.cpp +++ b/sapfor/experts/Sapfor_2017/_src/Utils/utils.cpp @@ -434,7 +434,7 @@ static map> removeCopies(map> __spf_print(1, "%s\n", tmp.c_str());*/ uniq[key] = &message; } - __spf_print(1, "messages filtering for file '%s': count before %d, count after %d\n", byFile.first.c_str(), byFile.second.size(), uniq.size()); + __spf_print(1, " messages filtering for file '%s': count before %d, count after %d\n", byFile.first.c_str(), byFile.second.size(), uniq.size()); vector uniqV; for (auto& elem : uniq) { diff --git a/sapfor/experts/Sapfor_2017/_src/Utils/version.h b/sapfor/experts/Sapfor_2017/_src/Utils/version.h index 9bc5377..8182a93 100644 --- a/sapfor/experts/Sapfor_2017/_src/Utils/version.h +++ b/sapfor/experts/Sapfor_2017/_src/Utils/version.h @@ -1,3 +1,3 @@ #pragma once -#define VERSION_SPF "2384" +#define VERSION_SPF "2386"