SgFunctionCallExp *
OpenCL::getSetKernelArgumentCallExpression (SgScopeStatement * scope,
SgVarRefExp * openCLKernel, int argumentIndex, SgType * sizeOfArgument,
SgExpression * argument)
{
using namespace SageBuilder;
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (openCLKernel);
actualParameters->append_expression (buildIntVal (argumentIndex));
actualParameters->append_expression (buildSizeOfOp (sizeOfArgument));
if (argument == NULL)
{
actualParameters->append_expression (buildIntVal (0));
}
else
{
actualParameters->append_expression (buildAddressOfOp (argument));
}
return buildFunctionCallExp ("clSetKernelArg", buildIntType (),
actualParameters, scope);
}
void insertChecksum(SgPntrArrRefExp *arrayReference,
SgExprStatement *curExprStatement, bool isWrite) {
if (arrayReference == NULL)
return;
// Create parameter list for function call
SgExprListExp* parameters = new SgExprListExp(NEW_FILE_INFO);
SgAddOp *addOp = new SgAddOp(NEW_FILE_INFO,
arrayReference->get_lhs_operand(),
arrayReference->get_rhs_operand(), arrayReference->get_type());
parameters->append_expression(addOp); // Address
//parameters->append_expression(arrayReference); // Value
parameters->append_expression(
buildStringVal(arrayReference->get_type()->unparseToString())); // type
// Insert Function Call
SgName name1;
SgExprStatement* callStmt;
if (!isWrite) {
name1 = SgName(VERIFY_CALL);
callStmt = buildFunctionCallStmt(name1, buildVoidType(), parameters,
curExprStatement->get_scope());
insertStatementBefore(curExprStatement, callStmt);
} else {
name1 = SgName(UPDATE_CALL);
callStmt = buildFunctionCallStmt(name1, buildVoidType(), parameters,
curExprStatement->get_scope());
insertStatementAfter(curExprStatement, callStmt);
}
}
void
CPPCUDAHostSubroutineDirectLoop::createKernelFunctionCallStatement (
SgScopeStatement * scope)
{
using namespace SageInterface;
using namespace SageBuilder;
using namespace OP2VariableNames;
using namespace OP2::RunTimeVariableNames;
Debug::getInstance ()->debugMessage (
"Creating statement to call CUDA kernel", Debug::FUNCTION_LEVEL,
__FILE__, __LINE__);
SgExprListExp * actualParameters = buildExprListExp ();
for (unsigned int i = 1; i <= parallelLoop->getNumberOfOpDatArgumentGroups (); ++i)
{
if (parallelLoop->isDuplicateOpDat (i) == false)
{
if (parallelLoop->isDirect (i) || parallelLoop->isReductionRequired (i))
{
SgDotExp * dotExpression = buildDotExp (
variableDeclarations->getReference (getOpDatName (i)),
buildOpaqueVarRefExp (data_d, subroutineScope));
SgCastExp * castExpression = buildCastExp (dotExpression,
buildPointerType (parallelLoop->getOpDatBaseType (i)));
actualParameters->append_expression (castExpression);
}
}
}
actualParameters->append_expression (variableDeclarations->getReference (
sharedMemoryOffset));
SgArrowExp * arrowExpression = buildArrowExp (
variableDeclarations->getReference (getOpSetName ()),
buildOpaqueVarRefExp (size, subroutineScope));
actualParameters->append_expression (arrowExpression);
SgCudaKernelExecConfig * kernelConfiguration = new SgCudaKernelExecConfig (
RoseHelper::getFileInfo (), variableDeclarations->getReference (
CUDA::blocksPerGrid), variableDeclarations->getReference (
CUDA::threadsPerBlock), variableDeclarations->getReference (
CUDA::sharedMemorySize));
kernelConfiguration->set_endOfConstruct (RoseHelper::getFileInfo ());
SgCudaKernelCallExp * kernelCallExpression = new SgCudaKernelCallExp (
RoseHelper::getFileInfo (), buildFunctionRefExp (
calleeSubroutine->getSubroutineName (), subroutineScope),
actualParameters, kernelConfiguration);
kernelCallExpression->set_endOfConstruct (RoseHelper::getFileInfo ());
appendStatement (buildExprStatement (kernelCallExpression), scope);
}
SgFunctionCallExp *
OpenCL::getEnqueueKernelCallExpression (SgScopeStatement * scope,
SgVarRefExp * commandQueue, SgVarRefExp * openCLKernel,
SgVarRefExp * globalWorkSize, SgVarRefExp * localWorkSize,
SgVarRefExp * event)
{
using namespace SageBuilder;
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (commandQueue);
actualParameters->append_expression (openCLKernel);
actualParameters->append_expression (buildIntVal (1));
actualParameters->append_expression (buildOpaqueVarRefExp("NULL", scope));
actualParameters->append_expression (buildAddressOfOp (globalWorkSize));
actualParameters->append_expression (buildAddressOfOp (localWorkSize));
actualParameters->append_expression (buildIntVal (0));
actualParameters->append_expression (buildOpaqueVarRefExp("NULL", scope));
actualParameters->append_expression (buildAddressOfOp (event));
return buildFunctionCallExp ("clEnqueueNDRangeKernel", buildIntType (),
actualParameters, scope);
}
SgFunctionCallExp *
OpenCL::OP2RuntimeSupport::getAssertMessage (SgScopeStatement * scope,
SgExpression * assertExpression, SgStringVal * message)
{
using namespace SageBuilder;
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (assertExpression);
actualParameters->append_expression (message);
return buildFunctionCallExp ("assert_m", buildVoidType (), actualParameters,
scope);
}
/*
* Fix OP function calls and inject debug names
*/
void OPSource::fixOpStructs(SgNode *n)
{
SgInitializedName* initname = isSgInitializedName(n);
if(initname)
{
string var_name = initname->get_name().getString();
SgConstructorInitializer *initer = isSgConstructorInitializer(initname->get_initializer());
if(initer)
{
string class_name = initer->get_class_decl()->get_name().getString();
if(class_name.find("op_dat") != string::npos
|| class_name.find("op_dat_gbl") != string::npos
|| class_name.compare("_op_ptr") == 0
|| class_name.compare("_op_set") == 0
|| class_name.compare("_op_dat_const") == 0)
{
cout << "---Injecting Debug Name: " << var_name << "---" << endl;
SgExprListExp* list = initer->get_args();
SgExpressionPtrList &exprs = list->get_expressions();
if( isSgStringVal(exprs.back()) == NULL )
{
list->append_expression(buildStringVal(var_name));
}
}
}
}
}
/*
* Fix op structure calls and inject debug names
*/
void OPSource::fixOpFunctions(SgNode *n)
{
SgName var_name;
SgFunctionCallExp *fn = isSgFunctionCallExp(n);
if(fn != NULL)
{
string fn_name = fn->getAssociatedFunctionDeclaration()->get_name().getString();
if(fn_name.compare("op_decl_const")==0)
{
SgExprListExp* exprList = fn->get_args();
SgExpressionPtrList &exprs = exprList->get_expressions();
if( isSgStringVal(exprs.back()) == NULL )
{
SgVarRefExp* varExp = isSgVarRefExp(exprs[1]);
if(!varExp)
{
varExp = isSgVarRefExp(isSgAddressOfOp(exprs[1])->get_operand_i());
}
if(varExp)
{
var_name = varExp->get_symbol()->get_name();
}
cout << "---Injecting Debug Name for const: " << var_name.getString() << "---" << endl;
exprList->append_expression(buildStringVal(var_name));
}
}
}
}
SgFunctionCallExp *
OpenCL::getSetKernelArgumentCallBufferExpression (SgScopeStatement * scope,
SgVarRefExp * openCLKernel, int argumentIndex, SgExpression * bufferRef)
{
using namespace SageBuilder;
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (openCLKernel);
actualParameters->append_expression (buildIntVal (argumentIndex));
actualParameters->append_expression (bufferRef);
actualParameters->append_expression (buildOpaqueVarRefExp("NULL", scope));
return buildFunctionCallExp ("clSetKernelArg", buildIntType (),
actualParameters, scope);
}
SgFunctionCallExp *
OpenCL::getWorkGroupIDCallStatement (SgScopeStatement * scope,
SgExpression * expression)
{
using namespace SageBuilder;
SgExprListExp * actualParameters = buildExprListExp ();
if (expression == NULL)
{
actualParameters->append_expression (buildIntVal (0));
}
else
{
actualParameters->append_expression (expression);
}
return buildFunctionCallExp ("get_global_id", buildIntType (),
actualParameters, scope);
}
SgFunctionCallExp *
OpenCL::OP2RuntimeSupport::getKernel (SgScopeStatement * scope,
std::string const & kernelName)
{
using namespace SageBuilder;
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (buildStringVal (kernelName));
return buildFunctionCallExp ("getKernel", buildVoidType (), actualParameters,
scope);
}
SgFunctionCallExp *
OpenCL::getFinishCommandQueueCallExpression (SgScopeStatement * scope,
SgVarRefExp * commandQueue)
{
using namespace SageBuilder;
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (commandQueue);
return buildFunctionCallExp ("clFinish", buildIntType (), actualParameters,
scope);
}
SgFunctionCallExp *
OpenCL::createWorkItemsSynchronisationCallStatement (SgScopeStatement * scope)
{
using namespace SageBuilder;
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (buildOpaqueVarRefExp (
CLK_LOCAL_MEM_FENCE, scope));
return buildFunctionCallExp ("barrier", buildVoidType (), actualParameters,
scope);
}
SgStatement *
CPPOpenCLHostSubroutineIndirectLoop::createPlanFunctionCallStatement ()
{
using namespace SageBuilder;
using namespace OP2VariableNames;
using namespace PlanFunctionVariableNames;
Debug::getInstance ()->debugMessage (
"Creating statement to call plan function", Debug::FUNCTION_LEVEL,
__FILE__, __LINE__);
SgExprListExp * actualParamaters = buildExprListExp ();
actualParamaters->append_expression (variableDeclarations->getReference (
getUserSubroutineName ()));
actualParamaters->append_expression (variableDeclarations->getReference (
getOpSetName ()));
actualParamaters->append_expression (variableDeclarations->getReference (
getPartitionSizeVariableName (parallelLoop->getUserSubroutineName ())));
actualParamaters->append_expression (buildIntVal (
parallelLoop->getNumberOfOpDatArgumentGroups ()));
actualParamaters->append_expression (variableDeclarations->getReference (
opDatArray));
actualParamaters->append_expression (buildIntVal (
parallelLoop->getNumberOfDistinctIndirectOpDats ()));
actualParamaters->append_expression (variableDeclarations->getReference (
indirectionDescriptorArray));
SgFunctionCallExp * functionCallExpression = buildFunctionCallExp (
OP2::OP_PLAN_GET, buildVoidType (), actualParamaters, subroutineScope);
SgExprStatement * assignmentStatement = buildAssignStatement (
variableDeclarations->getReference (planRet), functionCallExpression);
return assignmentStatement;
}
void
CPPOpenMPHostSubroutineIndirectLoop::createKernelFunctionCallStatement (
SgScopeStatement * scope)
{
using namespace SageInterface;
using namespace SageBuilder;
using namespace OP2VariableNames;
using namespace PlanFunctionVariableNames;
using namespace OP2::RunTimeVariableNames;
using namespace ReductionVariableNames;
using namespace LoopVariableNames;
Debug::getInstance ()->debugMessage (
"Creating statement to call OpenMP kernel", Debug::FUNCTION_LEVEL,
__FILE__, __LINE__);
SgExprListExp * actualParameters = buildExprListExp ();
for (unsigned int i = 1; i <= parallelLoop->getNumberOfOpDatArgumentGroups (); ++i)
{
if (parallelLoop->isDuplicateOpDat (i) == false)
{
if (parallelLoop->isReductionRequired (i))
{
SgMultiplyOp * multiplyExpression = buildMultiplyOp (
variableDeclarations->getReference (
getIterationCounterVariableName (1)), buildIntVal (64));
SgAddOp * addExpression = buildAddOp (
variableDeclarations->getReference (getReductionArrayHostName (i)),
multiplyExpression);
actualParameters->append_expression (addExpression);
}
else
{
actualParameters->append_expression (
variableDeclarations->getReference (getOpDatLocalName (i)));
}
}
}
unsigned int arrayIndex = 0;
for (unsigned int i = 1; i <= parallelLoop->getNumberOfOpDatArgumentGroups (); ++i)
{
if (parallelLoop->isDuplicateOpDat (i) == false)
{
if (parallelLoop->isIndirect (i))
{
SgVariableDeclaration * variableDeclaration = buildVariableDeclaration (
ind_maps, buildArrayType (buildIntType ()), NULL, subroutineScope);
SgPntrArrRefExp * arrayExpression = buildPntrArrRefExp (buildVarRefExp (
variableDeclaration), buildIntVal (arrayIndex));
SgArrowExp * arrowExpression = buildArrowExp (
variableDeclarations->getReference (planRet), arrayExpression);
actualParameters->append_expression (arrowExpression);
arrayIndex++;
}
}
}
for (unsigned int i = 1; i <= parallelLoop->getNumberOfOpDatArgumentGroups (); ++i)
{
if (parallelLoop->isIndirect (i))
{
SgVariableDeclaration * variableDeclaration = buildVariableDeclaration (
loc_maps, buildArrayType (buildIntType ()), NULL, subroutineScope);
SgPntrArrRefExp * arrayExpression = buildPntrArrRefExp (buildVarRefExp (
variableDeclaration), buildIntVal (i - 1));
SgArrowExp * arrowExpression = buildArrowExp (
variableDeclarations->getReference (planRet), arrayExpression);
actualParameters->append_expression (arrowExpression);
}
}
actualParameters->append_expression (buildArrowExp (
variableDeclarations->getReference (planRet), buildOpaqueVarRefExp (
ind_sizes, subroutineScope)));
actualParameters->append_expression (buildArrowExp (
variableDeclarations->getReference (planRet), buildOpaqueVarRefExp (
ind_offs, subroutineScope)));
actualParameters->append_expression (buildArrowExp (
variableDeclarations->getReference (planRet), buildOpaqueVarRefExp (
blkmap, subroutineScope)));
actualParameters->append_expression (buildArrowExp (
variableDeclarations->getReference (planRet), buildOpaqueVarRefExp (
offset, subroutineScope)));
actualParameters->append_expression (buildArrowExp (
variableDeclarations->getReference (planRet), buildOpaqueVarRefExp (
nelems, subroutineScope)));
actualParameters->append_expression (buildArrowExp (
variableDeclarations->getReference (planRet), buildOpaqueVarRefExp (
thrcol, subroutineScope)));
actualParameters->append_expression (buildArrowExp (
variableDeclarations->getReference (planRet), buildOpaqueVarRefExp (
nthrcol, subroutineScope)));
actualParameters->append_expression (variableDeclarations->getReference (
blockOffset));
actualParameters->append_expression (variableDeclarations->getReference (
blockID));
SgFunctionCallExp * functionCallExpression = buildFunctionCallExp (
calleeSubroutine->getSubroutineName (), buildVoidType (),
actualParameters, subroutineScope);
appendStatement (buildExprStatement (functionCallExpression), scope);
}
SgFunctionCallExp *
OpenCL::getAllocateConstantExpression(SgScopeStatement * scope,
SgVarRefExp * constant, SgType * constantType)
{
using namespace SageBuilder;
using std::string;
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (buildAddressOfOp (constant));
SgType * type = constantType;
string typeName = constantType->class_name ();
long position = typeName.find ("[");
int factor = 0;
if (position != string::npos)
{
string new_type = typeName.substr (position + 1);
string::iterator it;
for (it = new_type.begin (); it < new_type.end (); ++it)
{
if (!isdigit (*it))
{
new_type.erase (it);
}
}
//factor = strtol (new_type, NULL, 10);
factor = atoi (new_type.c_str ());
if (type->isIntegerType ())
{
new_type = "int";
} else if (type->isFloatType ())
{
new_type = "float";
}
type = buildOpaqueType (new_type, scope);
}
SgMultiplyOp * factorMult =
buildMultiplyOp (buildIntVal (factor), buildSizeOfOp (type));
actualParameters->append_expression (factorMult);
return buildFunctionCallExp ("op_allocate_constant",
OpenCL::getMemoryType(scope), actualParameters, scope);
}
void
FortranCUDAHostSubroutineDirectLoop::createKernelFunctionCallStatement (
SgScopeStatement * scope)
{
using namespace SageInterface;
using namespace SageBuilder;
using namespace OP2VariableNames;
using namespace ReductionVariableNames;
using namespace OP2::RunTimeVariableNames;
Debug::getInstance ()->debugMessage (
"Creating statement to call CUDA kernel", Debug::FUNCTION_LEVEL,
__FILE__, __LINE__);
SgExprListExp * actualParameters = buildExprListExp ();
actualParameters->append_expression (variableDeclarations->getReference (
opDatDimensions));
actualParameters->append_expression (variableDeclarations->getReference (
opDatCardinalities));
for (unsigned int i = 1; i <= parallelLoop->getNumberOfOpDatArgumentGroups (); ++i)
{
if (parallelLoop->isDuplicateOpDat (i) == false)
{
if (parallelLoop->isReductionRequired (i))
{
actualParameters->append_expression (
variableDeclarations->getReference (getReductionArrayDeviceName (i)));
}
else if (parallelLoop->isDirect (i))
{
/* Carlo: no more opDatNDevice as arguments */
/* actualParameters->append_expression (
moduleDeclarations->getDeclarations()->getReference (getOpDatDeviceName (i) +
parallelLoop->getUserSubroutineName () + postfixName));*/
}
else if (parallelLoop->isRead (i))
{
if (parallelLoop->isArray (i))
{
actualParameters->append_expression (
variableDeclarations->getReference (getOpDatDeviceName (i)));
}
else
{
actualParameters->append_expression (
variableDeclarations->getReference (getOpDatHostName (i)));
}
}
}
}
SgExpression * dotExpression = buildDotExp (
variableDeclarations->getReference (getOpSetName ()),
buildDotExp ( buildOpaqueVarRefExp (Fortran::setPtr, scope),
buildOpaqueVarRefExp (size, scope)));
actualParameters->append_expression (dotExpression);
actualParameters->append_expression (variableDeclarations->getReference (
warpSize));
actualParameters->append_expression (variableDeclarations->getReference (
sharedMemoryOffset));
SgCudaKernelExecConfig * kernelConfiguration = new SgCudaKernelExecConfig (
RoseHelper::getFileInfo (), variableDeclarations->getReference (
CUDA::blocksPerGrid), variableDeclarations->getReference (
CUDA::threadsPerBlock), variableDeclarations->getReference (
CUDA::sharedMemorySize), buildNullExpression ());
kernelConfiguration->set_endOfConstruct (RoseHelper::getFileInfo ());
SgCudaKernelCallExp * kernelCallExpression = new SgCudaKernelCallExp (
RoseHelper::getFileInfo (), buildFunctionRefExp (
calleeSubroutine->getSubroutineName (), subroutineScope),
actualParameters, kernelConfiguration);
kernelCallExpression->set_endOfConstruct (RoseHelper::getFileInfo ());
appendStatement (buildExprStatement (kernelCallExpression), scope);
}
void mlmTransform::transformCallExp(SgCallExpression* callExp)
{
ROSE_ASSERT(callExp);
SgFunctionRefExp* funcName = isSgFunctionRefExp(callExp->get_function());
if(!funcName)
return;
SgExprListExp* funcArgList = callExp->get_args();
SgExpressionPtrList argList = funcArgList->get_expressions();
SgScopeStatement* scope = getScope(callExp);
//cout << funcName->get_symbol()->get_name() << endl;
/** if it is malloc, search for the mlm attribute and append the memory level **/
if(strncmp("malloc",funcName->get_symbol()->get_name().str(),6) == 0)
{
if(argList.size() != 1) return;
SgExprListExp* funcArgList = callExp->get_args();
// check if LHS of malloc has an attribute assigned
SgNode* parentNode = callExp->get_parent();
// parent node can be a casting expression
if(isSgCastExp(parentNode))
{
parentNode = parentNode->get_parent();
}
// the mlm attribute
AstAttribute* attr = NULL;
// So far we spot two candidates for parentNode that we need to transform
if(isSgAssignOp(parentNode))
{
SgAssignOp* assignOp = isSgAssignOp(parentNode);
SgExpression* lhs = isSgExpression(assignOp->get_lhs_operand());
if(!isSgVarRefExp(lhs))
{
//cout << "lhs:" << assignOp->get_lhs_operand()->class_name() << endl;
// if pointer is packaged inside a struct, then we need to look down in lhs.
if(isSgDotExp(lhs))
{
lhs = isSgDotExp(lhs)->get_rhs_operand();
}
}
SgVarRefExp* lhsVarRef = isSgVarRefExp(lhs);
ROSE_ASSERT(lhsVarRef);
SgSymbol* symbol = lhsVarRef->get_symbol();
ROSE_ASSERT(symbol);
//retrieve the attribute from symbol
attr = symbol->getAttribute("mlmAttribute");
//cout << "LHS symbol name: " << symbol->get_name() << endl;
}
else if(isSgAssignInitializer(parentNode))
{
SgInitializedName* initName = isSgInitializedName(parentNode->get_parent());
ROSE_ASSERT(initName);
SgSymbol* symbol = initName->get_symbol_from_symbol_table();
if(!symbol) return;
ROSE_ASSERT(symbol);
//retrieve the attribute from symbol
attr = symbol->getAttribute("mlmAttribute");
//cout << "Initialized symbol name: " << symbol->get_name() << endl;
}
else
{
// do nothing because no attribute assigned or we always set to default
}
// if there is a mlm attribute attached to the symbol, then create new malloc
if(attr)
{
mlmAttribute* mlmAttr = dynamic_cast<mlmAttribute*> (attr);
SgExprListExp* funcArgList = callExp->get_args();
funcArgList->append_expression(buildIntVal(mlmAttr->getMemType()));
replaceExpression(callExp, buildFunctionCallExp("mlm_malloc",deepCopy(callExp->get_type()),deepCopy(funcArgList),getScope(callExp)));
}
}
else if(strncmp("memcpy",funcName->get_symbol()->get_name().str(),6) == 0)
{
// cout << "replacing memcpy" << endl;
if(argList.size() != 3) return;
Rose_STL_Container<SgNode*> varList = NodeQuery::querySubTree(funcArgList, V_SgVarRefExp);
SgVarRefExp* dst = isSgVarRefExp(varList[0]);
SgVarRefExp* src = isSgVarRefExp(varList[1]);
AstAttribute* attrDst = dst->get_symbol()->getAttribute("mlmAttribute");
AstAttribute* attrSrc = src->get_symbol()->getAttribute("mlmAttribute");
mlmAttribute* mlmAttrDst = dynamic_cast<mlmAttribute*>(attrDst);
mlmAttribute* mlmAttrSrc = dynamic_cast<mlmAttribute*>(attrSrc);
// if((mlmAttrDst && !mlmAttrSrc) || (mlmAttrDst && mlmAttrSrc && (mlmAttrDst->getMemType() < mlmAttrDst->getMemType())))
// {
// replaceExpression(callExp, buildFunctionCallExp("mlm_memcpy",deepCopy(callExp->get_type()),deepCopy(funcArgList),scope),true);
// DeletepragmasList2.push_back(callExp);
// }
//
// else if((!mlmAttrDst && mlmAttrSrc) || (mlmAttrDst && mlmAttrSrc && (mlmAttrDst->getMemType() > mlmAttrDst->getMemType())))
// 09/30/14 Following Simon's suggestion, we always insert wait for the mlm_memcpy
// {
string tagName = generateUniqueVariableName(scope,"copy_tag");
SgVariableDeclaration* newDecl = buildVariableDeclaration(tagName,
buildOpaqueType("mlm_Tag",getGlobalScope(callExp)),
buildAssignInitializer(buildFunctionCallExp("mlm_memcpy",deepCopy(callExp->get_type()),deepCopy(funcArgList),scope)));
SgExprStatement* waitStmt = buildFunctionCallStmt("mlm_waitComplete",
buildVoidType(),
buildExprListExp(buildVarRefExp(tagName,scope)),
scope);
insertStatement(getEnclosingStatement(callExp),newDecl,true);
insertStatement(getEnclosingStatement(callExp),waitStmt,true);
removeStatement(getEnclosingStatement(callExp));
// }
}
else if(strncmp("free",funcName->get_symbol()->get_name().str(),4) == 0)
{
// cout << "replacing free" << endl;
if(argList.size() != 1) return;
SgExpression* varExp = isSgExpression(argList[0]);
//cout << "exp:" << varExp->class_name() << endl;
if(!isSgVarRefExp(varExp))
{
if(isSgCastExp(varExp))
{
varExp = isSgCastExp(varExp)->get_operand_i();
}
// if pointer is packaged inside a struct, then we need to look down in lhs.
if(isSgDotExp(varExp))
{
varExp = isSgDotExp(varExp)->get_rhs_operand();
}
}
SgVarRefExp* varRef = isSgVarRefExp(varExp);
ROSE_ASSERT(varRef);
AstAttribute* attr = varRef->get_symbol()->getAttribute("mlmAttribute");
if(attr)
{
replaceExpression(callExp, buildFunctionCallExp("mlm_free",deepCopy(callExp->get_type()),deepCopy(funcArgList),scope),false);
}
}
}
