Alexander_KS/SAPFOR
4
2
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fixed dead flag for functions

Browse Source
This commit is contained in:
ALEXks
2025-02-09 20:48:06 +03:00
committed by Dudarenko
parent c6b09ad285
commit 68c779790d
19 changed files with 897 additions and 241 deletions
Download Patch File Download Diff File Expand all files Collapse all files

254
dvm/fdvm/trunk/fdvm/acc.cpp
View File

@@ -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<char, const char*>('\n', "\\n\"\n\""));
SpecialSymbols.insert(std::pair<char, const char*>('"', "\\\""));
@@ -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)
{
@@ -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; i<Rank(s); i++)
elist = AddListToList(elist, new SgExprListExp(*Calculate(LowerBound(s,i-1))));
}
*/
for (int i=0; i<Rank(s); i++)
elist = AddListToList(elist, new SgExprListExp(*new SgValueExp(0)));
elist->addAttribute(NULL_SUBSCRIPTS, (void*)1, 0);
return elist;
}
SgSymbol *DummyDvmHeaderSymbol(SgSymbol *ar, SgStatement *st_hedr)
@@ -8503,7 +8590,7 @@ SgStatement *CreateLoopKernel(SgSymbol *skernel, SgType *indexTypeInKernel)
ReplaceExitCycleGoto(block, stk);
last = cur_st;
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;
SgExpression *arg_list = NULL;
if (!options.isOn(C_CUDA) || !PrivateArrayClassUse(sizeOfPrivateArraysInBytes())) // !sizeOfPrivateArraysInBytes())
return NULL;
for (el = private_list; el; el = el->rhs())
@@ -9984,6 +10071,67 @@ void DeclareInternalPrivateVars()
st->addComment("! Internal private variables\n");
}
}
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()
{
@@ -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 (HEADER(var->lhs()->symbol())) continue; // dvm-array declared as dummy argument
if (!options.isOn(C_CUDA) || !IS_ARRAY(s) || !PrivateArrayClassUse(e_all_private_size))
{
st = Declaration_Statement(SymbolInKernel(s));
@@ -10009,17 +10162,19 @@ void DeclarePrivateVars(SgType *idxTypeInKernel)
st_first = st;
}
else
{
SgSymbol *s_dims=NULL;
st = new SgStatement(PRIVATE_AR_DECL);
{
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_first = st;
SgSymbol *s_dims=NULL;
if (Rank(s)>1)
{
@@ -10039,17 +10194,14 @@ void DeclarePrivateVars(SgType *idxTypeInKernel)
}
}
else
{
for (int i=Rank(s)-1; i; 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;
}
if (s_dims)
{
@@ -10058,11 +10210,11 @@ void DeclarePrivateVars(SgType *idxTypeInKernel)
SgSymbol **satr = (SgSymbol **) var->lhs()->attributeValue(0, PRIVATE_POINTER);
if (satr)
{
SgSymbol *sp = *satr;
efc->addArg(*new SgVarRefExp(sp)); //e->setLhs(new SgExprListExp(*new SgVarRefExp(sp)));
}
SgSymbol *sp = *satr;
efc->addArg(*new SgVarRefExp(sp));
}
}
}
if (!st_first)
return;
@@ -14000,7 +14152,7 @@ SgStatement *Create_C_Adapter_Function(SgSymbol *sadapter)
/* --------- add argument list to kernel call ----*/
for (sg = hgpu_first, sb = base_first, sl = acc_array_list, ln = 0; ln<num; sg = sg->next(), 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));
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);
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)
sdev = sdev->next();
}
e_all_private_size = sizeOfPrivateArraysInBytes();
e_all_private_size = sizeOfPrivateArraysInBytes();
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)
stmt = new SgCExpStmt(SgAssignOp(*new SgVarRefExp(*s_max_blocks), *getProp));
st_end->insertStmtBefore(*stmt, *st_hedr);
// insert code for big private arrays
// insert code for big private arrays
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);
@@ -14215,7 +14367,7 @@ SgStatement *Create_C_Adapter_Function(SgSymbol *sadapter)
SgSymbol *s_total_threads = new SgSymbol(VARIABLE_NAME, TestAndCorrectName("totalThreads"), *C_DvmType(), *st_hedr);
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_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 ? 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)
if (red_list)
InsertFinishReductionCalls(st_end, s_loop_ref, s_red_num);
// to dispose private arrays
// to dispose private arrays
if (options.isOn(C_CUDA) && PrivateArrayClassUse(e_all_private_size))
for (s = private_first, ln = 0; ln < lnp; s = s->next(), ln++) // private arrays
{
@@ -14583,6 +14735,19 @@ SgExpression *sizeOfElementInBytes(SgSymbol *symb)
return (new SgValueExp(isz));
}
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;
@@ -14614,19 +14779,18 @@ SgExpression *sizeOfPrivateArraysInBytes()
// e_size = &( *ArrayLengthInElems(symb, dvm_parallel_dir, 1) * *sizeOfElementInBytes(symb));
}
}
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;
if (e_size && e_size->isInteger()) // calculating length if it is possible
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)
{

13
dvm/fdvm/trunk/fdvm/acc_across.cpp
View File

@@ -1638,7 +1638,7 @@ vector<ArgsForKernel> 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<ArgsForKernel> 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<ArgsForKernel> 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<ArgsForKernel> 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<ArgsForKernel> 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));

141
dvm/fdvm/trunk/fdvm/acc_f2c.cpp
View File

@@ -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<SgExpression*> 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 (j<dims)
{
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)));
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<SgType*> *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<SgStatement*, SgStatement*> &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)
{

452
dvm/fdvm/trunk/fdvm/calls.cpp
View File

@@ -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 <string, vector<vector<SgType*> > > 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))

56
dvm/fdvm/trunk/fdvm/dvm.cpp
View File

@@ -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;

3
dvm/fdvm/trunk/fdvm/stmt.cpp
View File

@@ -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;
}

6
dvm/fdvm/trunk/include/calls.h
View File

@@ -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
};
};

35
dvm/fdvm/trunk/include/dvm.h
View File

@@ -205,6 +205,11 @@ struct local_part_list {
SgSymbol *local_part;
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;
@@ -265,6 +270,9 @@ const int L_BOUNDS = 1054; /*ACC*/
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
@@ -599,6 +607,7 @@ EXTERN int in_checksection,undefined_Tcuda, cuda_functions; /*ACC*/
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 */
@@ -635,7 +644,8 @@ int Rank(SgSymbol *s);
int Rank(SgSymbol *s);
SgExpression *dvm_array_ref();
SgExpression *dvm_ref(int n);
int DeleteDArFromList(SgStatement *stmt);
int DeleteDArFromList(SgStatement *stmt);
//void ChangeArg_DistArrayRef(SgExpression *e);
void ChangeArg_DistArrayRef(SgExpression *ele, SgSymbol *fsym, int i);
void ChangeDistArrayRef(SgExpression *e);
void ChangeDistArrayRef_Left(SgExpression *e);
@@ -1436,6 +1446,7 @@ void ExtractCopy(SgExpression *elist);
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);
@@ -1453,8 +1464,12 @@ SgExpression *CreatePrivateDummyList();
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 */
@@ -1563,6 +1578,7 @@ SgStatement *else_dir();
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);
@@ -2118,6 +2134,7 @@ void Translate_Fortran_To_C(SgStatement *stat, SgStatement *last, int countOfCop
void Translate_Fortran_To_C(SgStatement *stat, SgStatement *last, int countOfCopy, SgStatement *st_header);
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);
@@ -2181,12 +2198,13 @@ void Call_Site (SgSymbol *s, int inlined, SgStatement *stat, SgExpression *e);
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();
void InsertCalledProcedureCopies();
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);
@@ -2197,7 +2215,9 @@ int IsRecursiveProcedure(SgSymbol *s);
int IsRecursiveProcedure(SgSymbol *s);
int IsNoBodyProcedure(SgSymbol *s);
int isUserFunction(SgSymbol *s);
int IsInternalProcedure(SgSymbol *s);
int IsInternalProcedure(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);
@@ -2212,14 +2232,15 @@ SgStatement *FunctionPrototype(SgSymbol *sf);
SgStatement *FunctionPrototype(SgSymbol *sf);
bool CreateIntefacePrototype(SgStatement *header);
SgStatement *hasInterface(SgSymbol *s);
void SaveInterface(SgSymbol *s, SgStatement *interface);
void SaveInterface(SgSymbol *s, SgStatement *interface);
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);
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(...);
@@ -2267,7 +2288,7 @@ void ConvertLoopWithLabelToEnddoLoop (SgStatement *stat); /*OMP*/
// options on FDVM converter
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,
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_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

2
dvm/fdvm/trunk/parser/cftn.c
View File

@@ -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"))

69
dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.DVMH_cuda.cu_opt
View File

@@ -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<<<blocks, threads, 0, stream>>>((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<<<blocks, threads, 0, stream>>>((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<<<blocks, threads, 0, stream>>>((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<<<blocks, threads, 0, stream>>>((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<<<blocks, threads, 0, stream>>>((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<<<blocks, threads, 0, stream>>>((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<<<blocks, threads, 0, stream>>>((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<<<blocks, threads, 0, stream>>>((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;
}
}
}

12
dvm/tools/tester/trunk/test-suite/Performance/NPB/FDVMH.fdv/CG/cg.fdv
View File

@@ -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