diff --git a/Compiler/BackEnd/BackendDAECreate.mo b/Compiler/BackEnd/BackendDAECreate.mo index 88aa394cbcd..fbc31827d9c 100644 --- a/Compiler/BackEnd/BackendDAECreate.mo +++ b/Compiler/BackEnd/BackendDAECreate.mo @@ -1827,6 +1827,8 @@ algorithm DAE.Ident iteratorName; DAE.Exp e,iteratorExp; list iteratorexps, expExpLst; + list> loopPrlVars "list of parallel variables used/referenced in a parfor loop"; + case (v,_,{},_) then v; case (v,knv,(DAE.STMT_ASSIGN(exp1 =DAE.CREF(componentRef = cr)) :: xs),true) equation @@ -1868,6 +1870,19 @@ algorithm v_2 = detectImplicitDiscreteAlgsStatemens(v_1,knv, xs,true); then v_2; + +/* + case (v,knv,(DAE.STMT_PARFOR(type_= tp, iter = iteratorName, range = e,statementLst = statementLst,loopPrlVars=loopPrlVars) :: xs),true) + equation + cr = ComponentReference.makeCrefIdent(iteratorName, tp, {}); + iteratorExp = Expression.crefExp(cr); + iteratorexps = extendRange(e,knv); + v_1 = detectImplicitDiscreteAlgsStatemensFor(iteratorExp,iteratorexps,v,knv,statementLst,true); + v_2 = detectImplicitDiscreteAlgsStatemens(v_1,knv, xs,true); + then + v_2; +*/ + case (v,knv,(DAE.STMT_WHEN(statementLst = statementLst,elseWhen = NONE()) :: xs),_) equation v_1 = detectImplicitDiscreteAlgsStatemens(v,knv,statementLst,true); @@ -2740,6 +2755,7 @@ algorithm DAE.ElementSource source; Algorithm.Else algElse; tuple, DAE.Exp, tuple,Integer>, tuple> extraArg; + list> loopPrlVars "list of parallel variables used/referenced in the parfor loop"; case ({},_,extraArg,knvars) then (({},extraArg)); @@ -2791,6 +2807,15 @@ algorithm ((xs_1, extraArg)) = traverseStmtsExps(xs, func, extraArg, knvars); then ((DAE.STMT_FOR(tp,b1,id1,ix,e,stmts2,source) :: xs_1,extraArg)); + case (((x as DAE.STMT_PARFOR(type_=tp,iterIsArray=b1,iter=id1,index=ix,range=e,statementLst=stmts, loopPrlVars= loopPrlVars, source = source)) :: xs),func, extraArg, knvars) + equation + cr = ComponentReference.makeCrefIdent(id1, tp, {}); + iteratorExp = Expression.crefExp(cr); + iteratorexps = extendRange(e,knvars); + ((stmts2, extraArg)) = traverseStmtsForExps(iteratorExp, iteratorexps,e, stmts, knvars, func, extraArg); + ((xs_1, extraArg)) = traverseStmtsExps(xs, func, extraArg, knvars); + then ((DAE.STMT_PARFOR(tp,b1,id1,ix,e,stmts2,loopPrlVars,source) :: xs_1,extraArg)); + case (((x as DAE.STMT_WHILE(exp = e,statementLst=stmts, source = source)) :: xs),func, extraArg, knvars) equation ((stmts2, extraArg)) = traverseStmtsExps(stmts,func, extraArg, knvars); diff --git a/Compiler/BackEnd/BackendDAEUtil.mo b/Compiler/BackEnd/BackendDAEUtil.mo index 915f9ce3f29..b7403087cc0 100644 --- a/Compiler/BackEnd/BackendDAEUtil.mo +++ b/Compiler/BackEnd/BackendDAEUtil.mo @@ -5478,6 +5478,15 @@ algorithm extraArg = traverseStmts(stmts,func,extraArg); then traverseStmts(xs, func, extraArg); + + case (((x as DAE.STMT_PARFOR(type_=tp,iterIsArray=b1,iter=id1,range=e,statementLst=stmts)) :: xs),func,extraArg) + equation + ((_, extraArg)) = func((e, extraArg)); + cr = ComponentReference.makeCrefIdent(id1, tp, {}); + (stmts,_) = DAEUtil.traverseDAEEquationsStmts(stmts,Expression.traverseSubexpressionsHelper,(Expression.replaceCref,(cr,e))); + extraArg = traverseStmts(stmts,func,extraArg); + then + traverseStmts(xs, func, extraArg); case (((x as DAE.STMT_WHILE(exp = e,statementLst=stmts)) :: xs),func,extraArg) equation diff --git a/Compiler/BackEnd/BackendVarTransform.mo b/Compiler/BackEnd/BackendVarTransform.mo index 1c2f72568c2..ee3816f779f 100644 --- a/Compiler/BackEnd/BackendVarTransform.mo +++ b/Compiler/BackEnd/BackendVarTransform.mo @@ -1314,6 +1314,7 @@ algorithm list helpVarIndices; Integer index; Boolean b,b1,b2,b3; + list> loopPrlVars "list of parallel variables used/referenced in the parfor loop"; case ({},_,_,_) then (listReverse(inAcc),inBAcc); @@ -1380,6 +1381,18 @@ algorithm (es_1,b) = replaceStatementLst(es, repl,DAE.STMT_FOR(type_,iterIsArray,ident,index,e1_2,statementLst_1,source)::inAcc,true); then ( es_1,b); + + case ((DAE.STMT_PARFOR(type_=type_,iterIsArray=iterIsArray,iter=ident,index=index,range=e1,statementLst=statementLst,loopPrlVars=loopPrlVars,source=source)::es),repl,_,_) + equation + (statementLst_1,b1) = replaceStatementLst(statementLst, repl,{},false); + (e1_1,b2) = replaceExp(e1, repl,NONE()); + true = b1 or b2; + source = DAEUtil.addSymbolicTransformationSubstitution(b2,source,e1,e1_1); + (e1_2,b1) = ExpressionSimplify.condsimplify(b2,e1_1); + source = DAEUtil.addSymbolicTransformationSimplify(b1,source,e1_1,e1_2); + (es_1,b) = replaceStatementLst(es, repl,DAE.STMT_PARFOR(type_,iterIsArray,ident,index,e1_2,statementLst_1,loopPrlVars,source)::inAcc,true); + then + ( es_1,b); case ((DAE.STMT_WHILE(exp=e1,statementLst=statementLst,source=source)::es),repl,_,_) equation diff --git a/Compiler/BackEnd/DAEQuery.mo b/Compiler/BackEnd/DAEQuery.mo index 78b67d7e20e..160892fed58 100644 --- a/Compiler/BackEnd/DAEQuery.mo +++ b/Compiler/BackEnd/DAEQuery.mo @@ -609,6 +609,12 @@ algorithm print("- DAEQuery.incidenceRowStmts on FOR not implemented\n"); then {}; + + case ((DAE.STMT_PARFOR(type_ = _) :: rest),vars) + equation + print("- DAEQuery.incidenceRowStmts on PARFOR not implemented\n"); + then + {}; case ((DAE.STMT_WHILE(exp = _) :: rest),vars) equation diff --git a/Compiler/BackEnd/PartFn.mo b/Compiler/BackEnd/PartFn.mo index 4f3e324efea..947270dc349 100644 --- a/Compiler/BackEnd/PartFn.mo +++ b/Compiler/BackEnd/PartFn.mo @@ -462,6 +462,8 @@ algorithm DAE.Statement stmt,stmt_1; DAE.ElementSource source; Integer ix; + list> loopPrlVars "list of parallel variables used/referenced in the parfor loop"; + case({},dae) then ({},dae); case(DAE.STMT_ASSIGN(ty,e1,e2,source) :: cdr,dae) equation @@ -492,6 +494,13 @@ algorithm (cdr_1,dae) = elabStmts(cdr,dae); then (DAE.STMT_FOR(ty,b,i,ix,e_1,stmts_1,source) :: cdr_1,dae); + case(DAE.STMT_PARFOR(ty,b,i,ix,e,stmts,loopPrlVars,source) :: cdr,dae) + equation + ((e_1,dae)) = Expression.traverseExp(e,elabExp,dae); + (stmts_1,dae) = elabStmts(stmts,dae); + (cdr_1,dae) = elabStmts(cdr,dae); + then + (DAE.STMT_PARFOR(ty,b,i,ix,e_1,stmts_1,loopPrlVars,source) :: cdr_1,dae); case(DAE.STMT_WHILE(e,stmts,source) :: cdr,dae) equation ((e_1,dae)) = Expression.traverseExp(e,elabExp,dae); @@ -1178,6 +1187,8 @@ algorithm list elst,elst_1; DAE.ElementSource source; Integer ix; + list> loopPrlVars "list of parallel variables used/referenced in the parfor loop"; + case({},_,_,_,_) then {}; case(DAE.STMT_ASSIGN(ty,e1,e2,source) :: cdr,dae,p,inputs,current) equation @@ -1214,6 +1225,13 @@ algorithm cdr_1 = fixCallsAlg(cdr,dae,p,inputs,current); then DAE.STMT_FOR(ty,b,i,ix,e_1,stmts_1,source) :: cdr_1; + case(DAE.STMT_PARFOR(ty,b,i,ix,e,stmts,loopPrlVars,source) :: cdr,dae,p,inputs,current) + equation + ((e_1,_)) = Expression.traverseExp(e,fixCall,(p,inputs,dae,current)); + stmts_1 = fixCallsAlg(stmts,dae,p,inputs,current); + cdr_1 = fixCallsAlg(cdr,dae,p,inputs,current); + then + DAE.STMT_PARFOR(ty,b,i,ix,e_1,stmts_1,loopPrlVars,source) :: cdr_1; case(DAE.STMT_WHILE(e,stmts,source) :: cdr,dae,p,inputs,current) equation ((e_1,_)) = Expression.traverseExp(e,fixCall,(p,inputs,dae,current)); diff --git a/Compiler/BackEnd/SimCode.mo b/Compiler/BackEnd/SimCode.mo index ac9a9f6bf17..e09f7860165 100644 --- a/Compiler/BackEnd/SimCode.mo +++ b/Compiler/BackEnd/SimCode.mo @@ -12316,6 +12316,14 @@ algorithm files = getFilesFromStatements(rest, files); then files; + + case (DAE.STMT_PARFOR(source = source, statementLst = stmts)::rest, files) + equation + files = getFilesFromDAEElementSource(source, files); + files = getFilesFromStatements(stmts, files); + files = getFilesFromStatements(rest, files); + then + files; case (DAE.STMT_WHILE(source = source, statementLst = stmts)::rest, files) equation diff --git a/Compiler/FrontEnd/Algorithm.mo b/Compiler/FrontEnd/Algorithm.mo index 280cbf226e7..38967a75b3c 100644 --- a/Compiler/FrontEnd/Algorithm.mo +++ b/Compiler/FrontEnd/Algorithm.mo @@ -605,6 +605,44 @@ algorithm end matchcontinue; end makeFor; +public function makeParFor "function: makeParFor + This function creates a DAE.STMT_PARFOR construct, checking + that the types of the parts are correct." + input Ident inIdent; + input DAE.Exp inExp; + input DAE.Properties inProperties; + input list inStatementLst; + input list> inLoopPrlVars; + input DAE.ElementSource source; + output Statement outStatement; +algorithm + outStatement := matchcontinue (inIdent,inExp,inProperties,inStatementLst,inLoopPrlVars,source) + local + Boolean isArray; + DAE.Type et; + Ident i,e_str,t_str; + DAE.Exp e; + DAE.Type t; + list stmts; + DAE.Dimensions dims; + + case (i,e,DAE.PROP(type_ = DAE.T_ARRAY(ty = t, dims = dims)),stmts,_,source) + equation + isArray = Types.isArray(t, dims); + et = Types.simplifyType(t); + then + DAE.STMT_PARFOR(t,isArray,i,-1,e,stmts,inLoopPrlVars,source); + + case (_,e,DAE.PROP(type_ = t),_,_,source) + equation + e_str = ExpressionDump.printExpStr(e); + t_str = Types.unparseType(t); + Error.addSourceMessage(Error.FOR_EXPRESSION_TYPE_ERROR, {e_str,t_str}, DAEUtil.getElementSourceFileInfo(source)); + then + fail(); + end matchcontinue; +end makeParFor; + public function makeWhile "function: makeWhile This function creates a DAE.STMT_WHILE construct, checking that the types of the parts are correct." @@ -794,6 +832,7 @@ algorithm case DAE.STMT_ASSIGN_ARR(source=source) then source; case DAE.STMT_IF(source=source) then source; case DAE.STMT_FOR(source=source) then source; + case DAE.STMT_PARFOR(source=source) then source; case DAE.STMT_WHILE(source=source) then source; case DAE.STMT_WHEN(source=source) then source; case DAE.STMT_ASSERT(source=source) then source; diff --git a/Compiler/FrontEnd/DAE.mo b/Compiler/FrontEnd/DAE.mo index 1d229de1e6b..6a6ba28da17 100644 --- a/Compiler/FrontEnd/DAE.mo +++ b/Compiler/FrontEnd/DAE.mo @@ -587,6 +587,17 @@ uniontype Statement "There are four kinds of statements. Assignments (`a := b;\ list statementLst; ElementSource source "the origin of the component/equation/algorithm"; end STMT_FOR; + + record STMT_PARFOR + Type type_ "this is the type of the iterator"; + Boolean iterIsArray "True if the iterator has an array type, otherwise false."; + Ident iter "the iterator variable"; + Integer index "the index of the iterator variable, to make it unique; used by the new inst"; + Exp range "range for the loop"; + list statementLst; + list> loopPrlVars "list of parallel variables used/referenced in the parfor loop"; + ElementSource source "the origin of the component/equation/algorithm"; + end STMT_PARFOR; record STMT_WHILE Exp exp; diff --git a/Compiler/FrontEnd/DAEDump.mo b/Compiler/FrontEnd/DAEDump.mo index 2299963858b..47a03400fbd 100644 --- a/Compiler/FrontEnd/DAEDump.mo +++ b/Compiler/FrontEnd/DAEDump.mo @@ -1642,6 +1642,22 @@ algorithm Print.printBuf("end for;\n"); then (); + + case (DAE.STMT_PARFOR(iter = id,index=index,range = e,statementLst = stmts),i) + equation + indent(i); + Print.printBuf("parfor "); + Print.printBuf(id); + Debug.bcall(index <> -1, Print.printBuf, " /* iter index " +& intString(index) +& " */"); + Print.printBuf(" in "); + ExpressionDump.printExp(e); + Print.printBuf(" loop\n"); + i_1 = i + 2; + ppStmtList(stmts, i_1); + indent(i); + Print.printBuf("end parfor;\n"); + then + (); case (DAE.STMT_WHILE(exp = e,statementLst = stmts),i) equation diff --git a/Compiler/FrontEnd/DAEUtil.mo b/Compiler/FrontEnd/DAEUtil.mo index 43ef21ef5ef..ee84e1d0ad9 100644 --- a/Compiler/FrontEnd/DAEUtil.mo +++ b/Compiler/FrontEnd/DAEUtil.mo @@ -4128,6 +4128,7 @@ algorithm DAE.ElementSource source; Algorithm.Else algElse; Type_a extraArg; + list> loopPrlVars "list of parallel variables used/referenced in the parfor loop"; case ({},_,extraArg) then ({},extraArg); @@ -4179,6 +4180,13 @@ algorithm ((e_1, extraArg)) = func((e, extraArg)); (xs_1, extraArg) = traverseDAEEquationsStmts(xs, func, extraArg); then (DAE.STMT_FOR(tp,b1,id1,ix,e_1,stmts2,source) :: xs_1,extraArg); + + case (((x as DAE.STMT_PARFOR(type_=tp,iterIsArray=b1,iter=id1,index=ix,range=e,statementLst=stmts, loopPrlVars=loopPrlVars, source = source)) :: xs),func,extraArg) + equation + (stmts2, extraArg) = traverseDAEEquationsStmts(stmts,func,extraArg); + ((e_1, extraArg)) = func((e, extraArg)); + (xs_1, extraArg) = traverseDAEEquationsStmts(xs, func, extraArg); + then (DAE.STMT_PARFOR(tp,b1,id1,ix,e_1,stmts2,loopPrlVars,source) :: xs_1,extraArg); case (((x as DAE.STMT_WHILE(exp = e,statementLst=stmts, source = source)) :: xs),func,extraArg) equation diff --git a/Compiler/FrontEnd/InstSection.mo b/Compiler/FrontEnd/InstSection.mo index bee0930b86d..a9a7ba9c5c7 100644 --- a/Compiler/FrontEnd/InstSection.mo +++ b/Compiler/FrontEnd/InstSection.mo @@ -4853,6 +4853,7 @@ algorithm InstanceHierarchy ih; DAE.ElementSource source; list> loopPrlVars; + DAE.ComponentRef parforIter; // one iterator case (cache,env,ih,pre,ci_state,i,SOME(e),sl,info,source,initial_,impl,unrollForLoops) @@ -4868,13 +4869,18 @@ algorithm // situations. Start with empty list and collect all variables cref'ed // in the loop body. loopPrlVars = collectParallelVariables({},sl_1); + + // Remove the parfor loop iterator from the list(implicitly declared). + parforIter = DAE.CREF_IDENT(i, t,{}); + (loopPrlVars,_) = List.deleteMemberOnTrue(parforIter,loopPrlVars,crefInfoListCrefsEqual); + // Check the cref's in the list one by one to make // sure that they are parallel variables. // checkParallelVariables(cache,env_1,loopPrlVars); List.map2_0(loopPrlVars, isCrefParGlobalOrForIterator, cache, env_1); source = DAEUtil.addElementSourceFileInfo(source,info); - stmt = Algorithm.makeFor(i, e_2, prop, sl_1, source); + stmt = Algorithm.makeParFor(i, e_2, prop, sl_1, loopPrlVars, source); then (cache,{stmt}); @@ -4936,12 +4942,16 @@ algorithm // is it parglobal var? isParglobal = SCode.parallelismEqual(prl, SCode.PARGLOBAL()); + + // Now the iterator is already removed. No need for this. // is it the iterator of the parfor loop(implicitly declared)? - isForiterator = Util.isSome(cnstForRange); + // isForiterator = Util.isSome(cnstForRange); //is it either a parglobal var or for iterator - true = isParglobal or isForiterator; + //true = isParglobal or isForiterator; + true = isParglobal; + then (); case((cref,info),_,_) diff --git a/Compiler/Template/CodegenC.tpl b/Compiler/Template/CodegenC.tpl index d5f0253fe2f..8e5e98e854f 100644 --- a/Compiler/Template/CodegenC.tpl +++ b/Compiler/Template/CodegenC.tpl @@ -2224,8 +2224,12 @@ template simulationFunctionsFile(String filePrefix, list functions, li extern "C" { #endif - /* the OpenCL program. Made global to avoid repeated builds - cl_program ocl_kernels_program = ocl_build_p_from_src("<%filePrefix%>_kernels.cl", true); */ + <%if acceptParModelicaGrammar() then + << /* the OpenCL program. Made global to avoid repeated builds */ + cl_program ocl_kernels_program = ocl_build_p_from_src("<%filePrefix%>_kernels.cl", true); + >> + %> + <%literals |> literal hasindex i0 fromindex 0 => literalExpConst(literal,i0) ; separator="\n"%> <%functionBodies(functions)%> @@ -2241,6 +2245,10 @@ end simulationFunctionsFile; template parModelicaKernelsFile(String filePrefix, list functions, list literals) "Generates the content of the C file for functions in the simulation case." ::= + + /* Reset the parfor loop id counter to 0*/ + let()= System.parForTickReset(0) + << #include "OCLRuntimeUtil_new.cl" @@ -2249,7 +2257,6 @@ template parModelicaKernelsFile(String filePrefix, list functions, lis // Headers finish here. - <%literals |> literal hasindex i0 fromindex 0 => literalExpConst(literal,i0) ; separator="\n"%> <%functionBodiesParModelica(functions)%> @@ -2774,8 +2781,7 @@ template functionHeader(Function fn, Boolean inFunc) >> case KERNEL_FUNCTION(__) then << - <%functionHeaderNormal(underscorePath(name), functionArguments, outVars, inFunc, false)%> - <%functionHeaderBoxed(underscorePath(name), functionArguments, outVars, isBoxedFunction(fn), false)%> + <%functionHeaderKernelFunctionInterface(underscorePath(name), functionArguments, outVars)%> >> case EXTERNAL_FUNCTION(dynamicLoad=true) then << @@ -2970,6 +2976,33 @@ template functionHeaderImpl(String fname, list fargs, list o >> end functionHeaderImpl; +template functionHeaderKernelFunctionInterface(String fname, list fargs, list outVars) + "Generates function header for a ParModelica Kernel function interface." +::= + let fargsStr = (fargs |> var => funArgDefinitionKernelFunctionInterface(var) ;separator=", ") + + if outVars then << + typedef struct <%fname%>_rettype_s { + <%outVars |> var hasindex i1 fromindex 1 => + match var + case VARIABLE(__) then + let dimStr = match ty case T_ARRAY(__) then + '[<%dims |> dim => dimension(dim) ;separator=", "%>]' + let typeStr = varType(var) + '<%typeStr%> c<%i1%>; /* <%crefStr(name)%><%dimStr%> */' + case FUNCTION_PTR(__) then + 'modelica_fnptr c<%i1%>; /* <%name%> */' + ;separator="\n" + %> + } <%fname%>_rettype; + + <%fname%>_rettype _<%fname%>(<%fargsStr%>); + >> else << + + void _<%fname%>(<%fargsStr%>); + >> +end functionHeaderKernelFunctionInterface; + template funArgName(Variable var) ::= match var @@ -2984,49 +3017,104 @@ template funArgDefinition(Variable var) case FUNCTION_PTR(__) then 'modelica_fnptr <%name%>' end funArgDefinition; -template parFunArgDefinition2(Variable var, Text &parArgList /*BUFPA*/) +template funArgDefinitionKernelFunctionInterface(Variable var) +::= + match var + case VARIABLE(ty=T_ARRAY(__)) then 'device_<%varType(var)%> <%contextCref(name,contextFunction)%>' + case VARIABLE(__) then '<%varType(var)%> <%contextCref(name,contextFunction)%>' + else 'Invalid function argument to Kernel function Interface.' +end funArgDefinitionKernelFunctionInterface; + +template funArgDefinitionKernelFunctionBody(Variable var) "Generates code to initialize variables. Does not return anything: just appends declarations to buffers." ::= match var //function args will have nill instdims even if they are arrays. handled here case var as VARIABLE(ty=T_ARRAY(__)) then - let varName = '<%contextCref(var.name,contextFunction)%>' - let &parArgList += ',<%\n%> __global modelica_<%expTypeShort(var.ty)%>* data_<%varName%>,' - let &parArgList += '<%\n%> __global modelica_integer* info_<%varName%>' - "" + let varName = '<%contextCref(var.name,contextParallelFunction)%>' + '__global modelica_<%expTypeShort(var.ty)%>* data_<%varName%>,<%\n%> __global modelica_integer* info_<%varName%>' + case var as VARIABLE(__) then - let varName = '<%contextCref(var.name,contextFunction)%>' + let varName = '<%contextCref(var.name,contextParallelFunction)%>' if instDims then - let &parArgList += ',<%\n%> __global modelica_<%expTypeShort(var.ty)%>* data_<%varName%>,' - let &parArgList += '<%\n%> __global modelica_integer* info_<%varName%>' - " " + '__global modelica_<%expTypeShort(var.ty)%>* data_<%varName%>,<%\n%> __global modelica_integer* info_<%varName%>' + else - let &parArgList += ',<%\n%> modelica_<%expTypeShort(var.ty)%> <%varName%>' - "" -else let &parArgList += ' #error Unknown variable type in as function argument parFunArgDefinition<%\n%>' "" -end parFunArgDefinition2; + 'modelica_<%expTypeShort(var.ty)%> <%varName%>' + +else '#error Unknown variable type in as function argument funArgDefinitionKernelFunctionBody<%\n%>' +end funArgDefinitionKernelFunctionBody; -template parFunArgDefinition(Variable var) +template funArgDefinitionKernelFunctionBody2(Variable var, Text &parArgList /*BUFPA*/) "Generates code to initialize variables. Does not return anything: just appends declarations to buffers." ::= match var //function args will have nill instdims even if they are arrays. handled here case var as VARIABLE(ty=T_ARRAY(__)) then - let varName = '<%contextCref(var.name,contextFunction)%>' - '__global modelica_<%expTypeShort(var.ty)%>* data_<%varName%>,<%\n%> __global modelica_integer* info_<%varName%>' - + let varName = '<%contextCref(var.name,contextParallelFunction)%>' + let &parArgList += ',<%\n%> __global modelica_<%expTypeShort(var.ty)%>* data_<%varName%>,' + let &parArgList += '<%\n%> __global modelica_integer* info_<%varName%>' + "" case var as VARIABLE(__) then - let varName = '<%contextCref(var.name,contextFunction)%>' + let varName = '<%contextCref(var.name,contextParallelFunction)%>' if instDims then - '__global modelica_<%expTypeShort(var.ty)%>* data_<%varName%>,<%\n%> __global modelica_integer* info_<%varName%>' - + let &parArgList += ',<%\n%> __global modelica_<%expTypeShort(var.ty)%>* data_<%varName%>,' + let &parArgList += '<%\n%> __global modelica_integer* info_<%varName%>' + " " else - 'modelica_<%expTypeShort(var.ty)%> <%varName%>' + let &parArgList += ',<%\n%> modelica_<%expTypeShort(var.ty)%> <%varName%>' + "" +else let &parArgList += ' #error Unknown variable type in as function argument funArgDefinitionKernelFunctionBody2<%\n%>' "" +end funArgDefinitionKernelFunctionBody2; + +template parFunArgDefinitionFromLooptupleVar(tuple tupleVar) +::= +match tupleVar +case tupleVar as ((cref as CREF_IDENT(identType = T_ARRAY(__)),_)) then + let varName = contextArrayCref(cref,contextParallelFunction) + match cref.identType + case identType as T_ARRAY(ty = T_INTEGER(__)) then + '__global modelica_integer* data_<%varName%>,<%\n%>__global modelica_integer* info_<%varName%>' + case identType as T_ARRAY(ty = T_REAL(__)) then + '__global modelica_integer* data_<%varName%>,<%\n%>__global modelica_integer* info_<%varName%>' + + else 'Tempalte error in parFunArgDefinitionFromLooptupleVar' + +case tupleVar as ((cref as CREF_IDENT(__),_)) then + let varName = contextArrayCref(cref,contextParallelFunction) + match cref.identType + case identType as T_INTEGER(__) then + 'modelica_integer <%varName%>' + case identType as T_REAL(__) then + 'modelica_real <%varName%>' + + else 'Tempalte error in parFunArgDefinitionFromLooptupleVar' -else '#error Unknown variable type in as function argument parFunArgDefinition<%\n%>' -end parFunArgDefinition; +end parFunArgDefinitionFromLooptupleVar; + +template reconstructKernelArraysFromLooptupleVars(tuple tupleVar, Text &reconstructedArrs) + "reconstructs modelica arrays in the kernels." +::= +match tupleVar +case tupleVar as ((cref as CREF_IDENT(identType = T_ARRAY(__)),_)) then + let varName = contextArrayCref(cref,contextParallelFunction) + match cref.identType + case identType as T_ARRAY(ty = T_INTEGER(__)) then + let &reconstructedArrs += 'integer_array <%varName%>; <%\n%>' + let &reconstructedArrs += '<%varName%>.data = data_<%varName%>; <%\n%>' + let &reconstructedArrs += '<%varName%>.ndims = info_<%varName%>[0]; <%\n%>' + let &reconstructedArrs += '<%varName%>.dim_size = info_<%varName%> + 1; <%\n%>' + "" + case identType as T_ARRAY(ty = T_REAL(__)) then + let &reconstructedArrs += 'real_array <%varName%>; <%\n%>' + let &reconstructedArrs += '<%varName%>.data = data_<%varName%>; <%\n%>' + let &reconstructedArrs += '<%varName%>.ndims = info_<%varName%>[0]; <%\n%>' + let &reconstructedArrs += '<%varName%>.dim_size = info_<%varName%> + 1; <%\n%>' + "" +else let &reconstructedArrs += '#wiered variable in kerenl reconstruction of arrays<%\n%>' "" +end reconstructKernelArraysFromLooptupleVars; template reconstructKernelArrays(Variable var, Text &reconstructedArrs) "reconstructs modelica arrays in the kernels." @@ -3034,14 +3122,14 @@ template reconstructKernelArrays(Variable var, Text &reconstructedArrs) match var //function args will have nill instdims even if they are arrays. handled here case var as VARIABLE(ty=T_ARRAY(__)) then - let varName = '<%contextCref(var.name,contextFunction)%>' + let varName = '<%contextCref(var.name,contextParallelFunction)%>' let &reconstructedArrs += '<%expTypeShort(var.ty)%>_array <%varName%>; <%\n%>' - let &reconstructedArrs += '<%varName%>.data = data_<%varName%>; <%\n%>' - let &reconstructedArrs += '<%varName%>.ndims = info_<%varName%>[0]; <%\n%>' - let &reconstructedArrs += '<%varName%>.dim_size = info_<%varName%> + 1; <%\n%>' + let &reconstructedArrs += '<%varName%>.data = data_<%varName%>; <%\n%>' + let &reconstructedArrs += '<%varName%>.ndims = info_<%varName%>[0]; <%\n%>' + let &reconstructedArrs += '<%varName%>.dim_size = info_<%varName%> + 1; <%\n%>' "" case var as VARIABLE(__) then - let varName = '<%contextCref(var.name,contextFunction)%>' + let varName = '<%contextCref(var.name,contextParallelFunction)%>' if instDims then let &reconstructedArrs += '<%expTypeShort(var.ty)%>_array <%varName%>; <%\n%>' let &reconstructedArrs += '<%varName%>.data = data_<%varName%>; <%\n%>' @@ -3228,6 +3316,7 @@ template functionBodyParModelica(Function fn, Boolean inFunc) "Generates the body for a function." ::= match fn + case fn as FUNCTION(__) then extractParforBodies(fn, inFunc) case fn as KERNEL_FUNCTION(__) then functionBodyKernelFunction(fn, inFunc) case fn as PARALLEL_FUNCTION(__) then functionBodyParallelFunction(fn, inFunc) end functionBodyParModelica; @@ -3248,6 +3337,24 @@ template addRoots(Variable var) end match end addRoots; +template extractParforBodies(Function fn, Boolean inFunc) + "Generates the body for a Modelica/MetaModelica function." +::= +match fn +case FUNCTION(__) then + let()= System.tmpTickReset(1) + let()= System.tmpTickResetIndex(0,1) /* Boxed array indices */ + + let &varDecls = buffer "" /*BUFD*/ + + let bodyPart = (body |> stmt => extractParFors(stmt, &varDecls /*BUFD*/) ;separator="\n") + + + << + <%bodyPart%> + >> +end extractParforBodies; + template functionBodyRegularFunction(Function fn, Boolean inFunc) "Generates the body for a Modelica/MetaModelica function." ::= @@ -3363,10 +3470,10 @@ case KERNEL_FUNCTION(__) then // This odd arrangment and call is to get the commas in the right places // between the argumetns. // This puts correct comma placment even when the 'outvar' list is empty - let argStr = (functionArguments |> var => '<%parFunArgDefinition(var)%>' ;separator=", \n ") + let argStr = (functionArguments |> var => '<%funArgDefinitionKernelFunctionBody(var)%>' ;separator=", \n ") //let &argStr += (outVars |> var => '<%parFunArgDefinition(var)%>' ;separator=", \n") let _ = (outVars |> var => - parFunArgDefinition2(var, &argStr /*BUFP*/) ;separator=",\n") + funArgDefinitionKernelFunctionBody2(var, &argStr /*BUFP*/) ;separator=",\n") // Reconstruct array arguments to structures in the kernels let &reconstrucedArrays = buffer "" @@ -3377,21 +3484,22 @@ case KERNEL_FUNCTION(__) then reconstructKernelArrays(var, &reconstrucedArrays /*BUFP*/) ) - let bodyPart = (body |> stmt => parallelfunStatement(stmt, &varDecls /*BUFD*/) ;separator="\n") + let bodyPart = (body |> stmt => parModelicafunStatement(stmt, &varDecls /*BUFD*/) ;separator="\n") /* Needs to be done last as it messes with the tmp ticks :) */ let &varDecls += addRootsTempArray() << + __kernel void _<%fname%>( <%\t%><%\t%><%argStr%> ) { /* functionBodyKernelFunction: Reconstruct Arrays */ - <%reconstrucedArrays%> + <%reconstrucedArrays%> - /* functionBodyKernelFunction: locals */ + /* functionBodyKernelFunction: locals */ <%varDecls%> /* functionBodyKernelFunction: state in */ @@ -3405,7 +3513,7 @@ case KERNEL_FUNCTION(__) then restore_memory_state(<%stateVar%>); /* Free GPU/OpenCL CPU memory */ - <%varFrees%> + <%varFrees%> } >> @@ -3430,7 +3538,7 @@ case PARALLEL_FUNCTION(__) then ) let funArgs = (functionArguments |> var => functionArg(var, &varInits) ;separator="\n") - let bodyPart = (body |> stmt => parallelfunStatement(stmt, &varDecls) ;separator="\n") + let bodyPart = (body |> stmt => parModelicafunStatement(stmt, &varDecls) ;separator="\n") let &outVarInits = buffer "" let &outVarCopy = buffer "" let &outVarAssign = buffer "" @@ -3441,7 +3549,7 @@ case PARALLEL_FUNCTION(__) then << - <%retType%> _<%fname%>(<%functionArguments |> var => funArgDefinition(var) ;separator=", "%>, memory* memory_state) + <%retType%> _<%fname%>(<%functionArguments |> var => funArgDefinition(var) ;separator=", "%>) { <%funArgs%> <%varDecls%> @@ -3453,13 +3561,12 @@ case PARALLEL_FUNCTION(__) then <%bodyPart%> - //_return: <%outVarCopy%> restore_memory_state(<%stateVar%>); <%outVarAssign%> - /**mahge: Free unwanted meomory allocated**/ - <%varFrees%> + /*mahge: Free unwanted meomory allocated*/ + <%varFrees%> return<%if outVars then ' <%retVar%>' %>; } @@ -3504,7 +3611,7 @@ case KERNEL_FUNCTION(__) then let cl_kernelVar = tempDecl("cl_kernel", &varDecls) - let kernel_arg_number = "ker_arg_nr" + let kernel_arg_number = '<%fname%>_arg_nr' let &kernelArgSets = buffer "" let _ = (functionArguments |> var => @@ -3526,8 +3633,8 @@ case KERNEL_FUNCTION(__) then << - //Interface function to _<%fname%> defined in parallelFunctions.cl file. - <%retType%> _<%fname%>(<%functionArguments |> var => funArgDefinition(var) ;separator=", "%>) + /* Interface function to <%fname%> defined in parallelFunctions.cl file. */ + <%retType%> _<%fname%>(<%functionArguments |> var => funArgDefinitionKernelFunctionInterface(var) ;separator=", "%>) { <%funArgs%> <%varDecls%> @@ -3537,24 +3644,21 @@ case KERNEL_FUNCTION(__) then <%varInits%> - ///////////////// parallel operations/////////////////////// + /* functionBodyKernelFunctionInterface : <%fname%> Kernel creation and execution */ + int <%kernel_arg_number%> = 0; <%cl_kernelVar%> = ocl_create_kernel(ocl_kernels_program, "_<%fname%>"); - - int <%kernel_arg_number%> = 0; <%kernelArgSets%> - ocl_execute_kernel(<%cl_kernelVar%>); clReleaseKernel(<%cl_kernelVar%>); - ////////////////////////////////////////////////////////// + /*functionBodyKernelFunctionInterface : <%fname%> kernel execution ends here.*/ - //_return: <%outVarCopy%> <%if not acceptMetaModelicaGrammar() then 'restore_memory_state(<%stateVar%>);'%> <%outVarAssign%> - /**mahge: Free unwanted meomory allocated**/ - <%varFrees%> + /*mahge: Free unwanted meomory allocated*/ + <%varFrees%> return<%if outVars then ' <%retVar%>' %>; } @@ -3583,6 +3687,22 @@ case var as VARIABLE(__) then end setKernelArg_ith; +template setKernelArgFormTupleLoopVars_ith(tuple tupleVar, Text &KernelName, Text &argNr, Text &parVarList, Context context /*BUFPA*/) +::= +match tupleVar +//function args will have nill instdims even if they are arrays. handled here +case tupleVar as ((cref as CREF_IDENT(identType = T_ARRAY(__)),_)) then + let varName = contextArrayCref(cref,context) + let &parVarList += 'ocl_set_kernel_arg(<%KernelName%>, <%argNr%>, <%varName%>.data); ++<%argNr%>; <%\n%>' + let &parVarList += 'ocl_set_kernel_arg(<%KernelName%>, <%argNr%>, <%varName%>.info_dev); ++<%argNr%>; <%\n%>' + "" +case tupleVar as ((cref as CREF_IDENT(__),_)) then + let varName = contextArrayCref(cref,context) + let &parVarList += 'ocl_set_kernel_arg(<%KernelName%>, <%argNr%>, <%varName%>); ++<%argNr%>; <%\n%>' + "" +end setKernelArgFormTupleLoopVars_ith; + + template functionBodyExternalFunction(Function fn, Boolean inFunc) "Generates the body for an external function (just a wrapper)." ::= @@ -4585,7 +4705,7 @@ template funStatement(Statement stmt, Text &varDecls /*BUFP*/) "NOT IMPLEMENTED FUN STATEMENT" end funStatement; -template parallelfunStatement(Statement stmt, Text &varDecls) +template parModelicafunStatement(Statement stmt, Text &varDecls) "Generates function statements With PARALLEL context. Similar to Function context. Except in some cases like assignments." ::= @@ -4596,7 +4716,31 @@ template parallelfunStatement(Statement stmt, Text &varDecls) ;separator="\n") else "NOT IMPLEMENTED FUN STATEMENT" -end parallelfunStatement; +end parModelicafunStatement; + +template extractParFors(Statement stmt, Text &varDecls) + "Generates bodies of parfor loops to the kernel file. + The sequential C operations needed to implement the parallel + for loop will be handled by the normal funStatment template." +::= + match stmt + case ALGORITHM(__) then + (statementLst |> stmt => + extractParFors_impl(stmt, contextParallelFunction, &varDecls) + ;separator="\n") + else + "NOT IMPLEMENTED FUN STATEMENT" +end extractParFors; + + +template extractParFors_impl(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/) + "Generates an algorithm statement." +::= + match stmt + case s as STMT_PARFOR(__) then algStmtParForBody(s, contextParallelFunction, &varDecls /*BUFD*/) +end extractParFors_impl; + + template algStatement(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/) "Generates an algorithm statement." @@ -4608,6 +4752,7 @@ template algStatement(DAE.Statement stmt, Context context, Text &varDecls /*BUFP case s as STMT_TUPLE_ASSIGN(__) then algStmtTupleAssign(s, context, &varDecls /*BUFD*/) case s as STMT_IF(__) then algStmtIf(s, context, &varDecls /*BUFD*/) case s as STMT_FOR(__) then algStmtFor(s, context, &varDecls /*BUFD*/) + case s as STMT_PARFOR(__) then algStmtParForInterface(s, context, &varDecls /*BUFD*/) case s as STMT_WHILE(__) then algStmtWhile(s, context, &varDecls /*BUFD*/) case s as STMT_ASSERT(__) then algStmtAssert(s, context, &varDecls /*BUFD*/) case s as STMT_TERMINATE(__) then algStmtTerminate(s, context, &varDecls /*BUFD*/) @@ -4904,6 +5049,140 @@ case STMT_IF(__) then >> end algStmtIf; +template algStmtParForBody(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/) + "Generates a for algorithm statement." +::= + match stmt + case s as STMT_PARFOR(range=rng as RANGE(__)) then + algStmtParForRangeBody(s, context, &varDecls /*BUFD*/) + case s as STMT_PARFOR(__) then + algStmtForGeneric(s, context, &varDecls /*BUFD*/) +end algStmtParForBody; + +template algStmtParForRangeBody(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/) + "Generates a for algorithm statement where range is RANGE." +::= +match stmt +case STMT_PARFOR(range=rng as RANGE(__)) then + let iterName = contextIteratorName(iter, context) + let identType = expType(type_, iterIsArray) + let identTypeShort = expTypeShort(type_) + + let parforKernelName = 'parfor_<%System.parForTick()%>' + + let &loopVarDecls = buffer "" + let body = (statementLst |> stmt => algStatement(stmt, context, &loopVarDecls) + ;separator="\n") + + // Reconstruct array arguments to structures in the kernels + let &reconstrucedArrays = buffer "" + let _ = (loopPrlVars |> var => + reconstructKernelArraysFromLooptupleVars(var, &reconstrucedArrays /*BUFP*/) + ) + + let argStr = (loopPrlVars |> var => '<%parFunArgDefinitionFromLooptupleVar(var)%>' ;separator=", \n") + + << + + __kernel void <%parforKernelName%>( + modelica_integer loop_start, + modelica_integer loop_step, + modelica_integer loop_end, + <%argStr%> + ) { + + + /* algStmtParForRangeBody : Thread managment for parfor loops */ + modelica_integer inner_start = (get_global_id(0) * loop_step) + (loop_start); + modelica_integer stride = get_global_size(0) * loop_step; + + for(modelica_integer <%iterName%> = (modelica_integer) inner_start; in_range_integer(<%iterName%>, loop_start, loop_end); <%iterName%> += stride) + { + + /*algStmtParForRangeBody : Reconstruct Arrays */ + <%reconstrucedArrays%> + + /* algStmtParForRangeBody : locals */ + <%loopVarDecls%> + + <%body%> + + } + + } + >> +end algStmtParForRangeBody; + +template algStmtParForInterface(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/) + "Generates a for algorithm statement." +::= + match stmt + case s as STMT_PARFOR(range=rng as RANGE(__)) then + algStmtParForRangeInterface(s, context, &varDecls /*BUFD*/) + case s as STMT_PARFOR(__) then + algStmtForGeneric(s, context, &varDecls /*BUFD*/) +end algStmtParForInterface; + +template algStmtParForRangeInterface(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/) + "Generates a for algorithm statement where range is RANGE." +::= +match stmt +case STMT_PARFOR(range=rng as RANGE(__)) then + let identType = expType(type_, iterIsArray) + let identTypeShort = expTypeShort(type_) + let stmtStr = (statementLst |> stmt => algStatement(stmt, context, &varDecls) + ;separator="\n") + algStmtParForRangeInterface_impl(rng, iter, identType, identTypeShort, loopPrlVars, stmtStr, context, &varDecls) +end algStmtParForRangeInterface; + +template algStmtParForRangeInterface_impl(Exp range, Ident iterator, String type, String shortType, list> loopPrlVars, Text body, Context context, Text &varDecls) + "The implementation of algStmtParForRangeInterface." +::= +match range +case RANGE(__) then + let iterName = contextIteratorName(iterator, context) + let stateVar = if not acceptMetaModelicaGrammar() then tempDecl("state", &varDecls) + let startVar = tempDecl(type, &varDecls) + let stepVar = tempDecl(type, &varDecls) + let stopVar = tempDecl(type, &varDecls) + let &preExp = buffer "" + let startValue = daeExp(start, context, &preExp, &varDecls) + let stepValue = match step case SOME(eo) then + daeExp(eo, context, &preExp, &varDecls) + else + "(1)" + let stopValue = daeExp(stop, context, &preExp, &varDecls) + + let cl_kernelVar = tempDecl("cl_kernel", &varDecls) + + let parforKernelName = 'parfor_<%System.parForTick()%>' + + let kerArgNr = '<%parforKernelName%>_arg_nr' + + let &kernelArgSets = buffer "" + let _ = (loopPrlVars |> varTuple => + setKernelArgFormTupleLoopVars_ith(varTuple, &cl_kernelVar, &kerArgNr, &kernelArgSets, context /*BUFP*/) + ) + + << + <%preExp%> + <%startVar%> = <%startValue%>; <%stepVar%> = <%stepValue%>; <%stopVar%> = <%stopValue%>; + <%cl_kernelVar%> = ocl_create_kernel(ocl_kernels_program, "<%parforKernelName%>"); + int <%kerArgNr%> = 0; + + ocl_set_kernel_arg(<%cl_kernelVar%>, <%kerArgNr%>, <%startVar%>); ++<%kerArgNr%>; <%\n%> + ocl_set_kernel_arg(<%cl_kernelVar%>, <%kerArgNr%>, <%stepVar%>); ++<%kerArgNr%>; <%\n%> + ocl_set_kernel_arg(<%cl_kernelVar%>, <%kerArgNr%>, <%stopVar%>); ++<%kerArgNr%>; <%\n%> + + <%kernelArgSets%> + + ocl_execute_kernel(<%cl_kernelVar%>); + clReleaseKernel(<%cl_kernelVar%>); + + + >> /* else we're looping over a zero-length range */ +end algStmtParForRangeInterface_impl; + template algStmtFor(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/) "Generates a for algorithm statement." @@ -4915,7 +5194,6 @@ template algStmtFor(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/ algStmtForGeneric(s, context, &varDecls /*BUFD*/) end algStmtFor; - template algStmtForRange(DAE.Statement stmt, Context context, Text &varDecls /*BUFP*/) "Generates a for algorithm statement where range is RANGE." ::= @@ -5454,7 +5732,7 @@ template daeExpCrefRhs2(Exp ecr, Context context, Text &preExp /*BUFP*/, << (*<%arrayType%>_element_addr(&<%arrName%>, <%dimsLenStr%>, <%dimsValuesStr%>)) >> - case PARALLEL_FUNCTION_CONTEXT(__) then + case PARALLEL_FUNCTION_CONTEXT(__) then let dimsValuesStr = (crefSubs(cr) |> INDEX(__) => daeExp(exp, context, &preExp, &varDecls) ;separator=", ") @@ -5639,15 +5917,15 @@ template daeExpCrefLhs2(Exp ecr, Context context, Text &afterExp /*BUFP*/, 'arrayGet(<%arrName%>,<%dimsValuesStr%>) /* DAE.CREF */' else match context - case FUNCTION_CONTEXT(__) then - << - (*<%arrayType%>_element_addr(&<%arrName%>, <%dimsLenStr%>, <%dimsValuesStr%>)) - >> case PARALLEL_FUNCTION_CONTEXT(__) then << (*<%arrayType%>_element_addr_c99_<%dimsLenStr%>(&<%arrName%>, <%dimsLenStr%>, <%dimsValuesStr%>)) >> - else "" + else + << + (*<%arrayType%>_element_addr(&<%arrName%>, <%dimsLenStr%>, <%dimsValuesStr%>)) + >> + else // The array subscript denotes a slice let &afterExp += '/* daeExpCrefLhs2 SLICE(<%ExpressionDump.printExpStr(ecr)%>) afterExp */<%\n%>' @@ -6303,7 +6581,7 @@ template daeExpCall(Exp call, Context context, Text &preExp /*BUFP*/, Text &varD match context case FUNCTION_CONTEXT(__) then "" case PARALLEL_FUNCTION_CONTEXT(__) then "" - else + else << <%\n%>#ifdef _OMC_MEASURE_TIME SIM_PROF_TICK_FN(<%funName%>_index); @@ -6314,7 +6592,7 @@ template daeExpCall(Exp call, Context context, Text &preExp /*BUFP*/, Text &varD match context case FUNCTION_CONTEXT(__) then "" case PARALLEL_FUNCTION_CONTEXT(__) then "" - else + else << <%\n%>#ifdef _OMC_MEASURE_TIME SIM_PROF_ACC_FN(<%funName%>_index); @@ -6909,8 +7187,9 @@ template arrayScalarRhs(Type ty, list subs, String arrName, Context context (*<%arrayType%>_element_addr_c99_<%dimsLenStr%>(&<%arrName%>, <%dimsLenStr%>, <%dimsValuesStr%>)) >> else - '(*<%arrayType%>_element_addr(&<%arrName%>, <%dimsLenStr%>, <%dimsValuesStr%>))' - + << + (*<%arrayType%>_element_addr(&<%arrName%>, <%dimsLenStr%>, <%dimsValuesStr%>)) + >> end arrayScalarRhs; template daeExpList(Exp exp, Context context, Text &preExp /*BUFP*/, diff --git a/Compiler/Template/SimCodeTV.mo b/Compiler/Template/SimCodeTV.mo index fd8ff5f4840..0383c6557c4 100644 --- a/Compiler/Template/SimCodeTV.mo +++ b/Compiler/Template/SimCodeTV.mo @@ -690,6 +690,17 @@ package System input Integer index; output Integer tickNo; end tmpTickIndex; + + function parForTick + "returns a tick that can be reset. used to uniquely identify parallel for loops. + parfor loops are generated in two rounds. We need to match them." + output Integer tickNo; + end parForTick; + + function parForTickReset + "Resets parfor loop id for second round" + input Integer start; + end parForTickReset; function tmpTickIndexReserve input Integer index; @@ -1199,6 +1210,15 @@ package DAE list statementLst; ElementSource source; end STMT_FOR; + record STMT_PARFOR + Type type_; + Boolean iterIsArray; + Ident iter; + Exp range; + list statementLst; + list> loopPrlVars "list of parallel variables used/referenced in the parfor loop"; + ElementSource source; + end STMT_PARFOR; record STMT_WHILE Exp exp; list statementLst; diff --git a/Compiler/Util/System.mo b/Compiler/Util/System.mo index bacd8dfa1b2..715ee7109e3 100644 --- a/Compiler/Util/System.mo +++ b/Compiler/Util/System.mo @@ -597,6 +597,18 @@ public function tmpTickReset external "C" SystemImpl_tmpTickReset(start) annotation(Library = "omcruntime"); end tmpTickReset; +public function parForTick + "returns a tick that can be reset. used to uniquely identify parallel for loops" + output Integer loopNo; + external "C" loopNo = SystemImpl_parForTick() annotation(Library = "omcruntime"); +end parForTick; + +public function parForTickReset + "Resets parfor loop id for second round" + input Integer start; + external "C" SystemImpl_parForTickReset(start) annotation(Library = "omcruntime"); +end parForTickReset; + public function tmpTickIndex "returns a tick that can be reset. TODO: remove me when bootstrapped (default argument index=0)" input Integer index; diff --git a/Compiler/runtime/System_rml.c b/Compiler/runtime/System_rml.c index f74f8605dba..94c34bfbd68 100644 --- a/Compiler/runtime/System_rml.c +++ b/Compiler/runtime/System_rml.c @@ -751,6 +751,20 @@ RML_BEGIN_LABEL(System__tmpTickReset) } RML_END_LABEL +RML_BEGIN_LABEL(System__parForTick) +{ + rmlA0 = (void*) mk_icon(SystemImpl_parForTick()); + RML_TAILCALLK(rmlSC); +} +RML_END_LABEL + +RML_BEGIN_LABEL(System__parForTickReset) +{ + SystemImpl_parForTickReset(RML_UNTAGFIXNUM(rmlA0)); + RML_TAILCALLK(rmlSC); +} +RML_END_LABEL + RML_BEGIN_LABEL(System__tmpTickIndex) { rmlA0 = (void*) mk_icon(SystemImpl_tmpTickIndex(RML_UNTAGFIXNUM(rmlA0))); diff --git a/Compiler/runtime/systemimpl.c b/Compiler/runtime/systemimpl.c index a9e53823fcd..eb0d588af11 100644 --- a/Compiler/runtime/systemimpl.c +++ b/Compiler/runtime/systemimpl.c @@ -1556,6 +1556,7 @@ int SystemImpl__getLoadModelPath(const char *name, void *prios, void *mps, const #define MAX_TMP_TICK 16 static modelica_integer tmp_tick_no[MAX_TMP_TICK] = {0}; +static modelica_integer parfor_tick_no = 0; extern int SystemImpl_tmpTick() { @@ -1567,6 +1568,16 @@ extern void SystemImpl_tmpTickReset(int start) tmp_tick_no[0] = start; } +extern int SystemImpl_parForTick() +{ + return parfor_tick_no++; +} + +extern void SystemImpl_parForTickReset(int start) +{ + parfor_tick_no = start; +} + extern int SystemImpl_tmpTickIndex(int index) { assert(index < MAX_TMP_TICK && index >= 0); diff --git a/SimulationRuntime/ParModelica/OpenCLRuntime/omc_ocl_util.c b/SimulationRuntime/ParModelica/OpenCLRuntime/omc_ocl_util.c index 68e02c39687..25ebcb64e95 100644 --- a/SimulationRuntime/ParModelica/OpenCLRuntime/omc_ocl_util.c +++ b/SimulationRuntime/ParModelica/OpenCLRuntime/omc_ocl_util.c @@ -163,6 +163,8 @@ void ocl_initialize(){ void ocl_create_context_and_comm_queue(){ // Create a context to run OpenCL on the OCL-enabled Device cl_int err; + + printf("--- Creating OpenCL context"); device_context = clCreateContext(0, 1, &ocl_device, NULL, NULL, &err); ocl_error_check(OCL_CREATE_CONTEXT, err); @@ -174,6 +176,7 @@ void ocl_create_context_and_comm_queue(){ clGetContextInfo(device_context, CL_CONTEXT_DEVICES, ParmDataBytes, OCL_Devices, NULL); // Create a command-queue on the first OCL_ device + printf("--- Creating OpenCL command queue"); device_comm_queue = clCreateCommandQueue(device_context, OCL_Devices[0], 0, &err); @@ -196,24 +199,27 @@ cl_program ocl_build_p_from_src(const char* source, int isfile){ else program_source = source; + + printf("--- Creating OpenCL program"); cl_program ocl_program = clCreateProgramWithSource(device_context, 1, (const char**)&program_source, NULL, NULL); - printf("********** program created.\n"); + printf("\t\t\t - OK.\n"); free((void*)program_source); // Build the program (OpenCL JIT compilation) + printf("--- Building OpenCL program \n"); char options[100]; - const char* flags = "-I "; + const char* flags = "-I\""; const char* OMHOME = getenv("OPENMODELICAHOME"); - const char* OMEXT = "/include/omc/"; + const char* OMEXT = "/include/omc/\""; if ( OMHOME != NULL ) { strcpy(options, flags); strcat(options, OMHOME); strcat(options, OMEXT); - printf("Building OpenCL code with flags %s\n",options); + printf("\t :Using flags %s\n",options); cl_int err; err = clBuildProgram(ocl_program, 0, NULL, options, NULL, NULL); ocl_error_check(OCL_BUILD_PROGRAM, err); @@ -223,14 +229,16 @@ cl_program ocl_build_p_from_src(const char* source, int isfile){ 0, NULL, &size); char * log = (char*)malloc(size); clGetProgramBuildInfo(ocl_program,ocl_device,CL_PROGRAM_BUILD_LOG,size,log, NULL); - printf("\t\tCL_PROGRAM_BUILD_LOG: \t%s\n", log); + printf("CL_PROGRAM_BUILD_LOG: \n%s\n", log); free(log); if(err){ - printf("Errors detected in compilation of OpenCL code:\n"); + printf("Build failed: Errors detected in compilation of OpenCL code:\n"); exit(1); } + + return ocl_program; } @@ -443,7 +451,7 @@ void ocl_error_check(int operation, cl_int error_code){ printf("CL_OUT_OF_HOST_MEMORY \n"); break; case CL_SUCCESS: - printf("********** program built(JIT compilation).\n"); + printf("\t\t\t\t\t\t - OK.\n"); break; default: printf("Possible unknown error in : OCL_BUILD_PROGRAM\n"); @@ -543,7 +551,7 @@ void ocl_error_check(int operation, cl_int error_code){ printf("CL_OUT_OF_HOST_MEMORY \n"); break; case CL_SUCCESS: - printf("********** Context created.\n"); + printf("\t\t\t - OK.\n"); break; default: printf("Possible unknown error in : OCL_CREATE_CONTEXT\n"); @@ -573,7 +581,7 @@ void ocl_error_check(int operation, cl_int error_code){ printf("CL_OUT_OF_HOST_MEMORY \n"); break; case CL_SUCCESS: - printf("********** Command queue created.\n"); + printf("\t\t - OK.\n"); break; default: printf("Possible unknown error in : OCL_CREATE_COMMAND_QUEUE\n");