Commit 957c5a34 authored by Akash Kothari's avatar Akash Kothari
Browse files

Fix warnings in most of HPVM

parent 6e8dfead
......@@ -38,12 +38,10 @@ using namespace builddfg;
#define DECLARE(X) \
X = M.getOrInsertFunction( \
#X, runtimeModule->getFunction(#X)->getFunctionType()); \
// DEBUG(errs() << *X)
namespace dfg2llvm {
// Helper Functions
static inline ConstantInt *getTimerID(Module &, enum hpvm_TimerID);
static inline ConstantInt *getTimerID(Module &, enum hpvm::Target);
bool hasAttribute(Function *, unsigned, Attribute::AttrKind);
......@@ -261,29 +259,6 @@ Value *CodeGenTraversal::getStringPointer(const Twine &S, Instruction *IB,
return SPtr;
}
// Add an argument of type Ty to the given function F
// void CodeGenTraversal::addArgument(Function* F, Type* Ty, const Twine& name)
// {
// // Add the argument to argument list
// new Argument(Ty, name, F);
//
// // Create the argument type list with added argument types
// std::vector<Type*> ArgTypes;
// for(Function::const_arg_iterator ai = F->arg_begin(), ae = F->arg_end();
// ai != ae; ++ai) {
// ArgTypes.push_back(ai->getType());
// }
// // Adding new arguments to the function argument list, would not change the
// // function type. We need to change the type of this function to reflect the
// // added arguments
// FunctionType* FTy = FunctionType::get(F->getReturnType(), ArgTypes,
// F->isVarArg()); PointerType* PTy = PointerType::get(FTy,
// cast<PointerType>(F->getType())->getAddressSpace());
//
// // Change the function type
// F->mutateType(PTy);
//}
void renameNewArgument(Function *newF, const Twine &argName) {
// Get Last argument in Function Arg List and rename it to given name
Argument *lastArg = &*(newF->arg_end() - 1);
......@@ -323,15 +298,6 @@ Function *CodeGenTraversal::addArgument(Function *F, Type *Ty,
return newF;
}
// Change the argument list of function F to add index and limit arguments
// void CodeGenTraversal::addIdxDimArgs(Function* F) {
// // Add Index and Dim arguments
// std::string names[] = {"idx_x", "idx_y", "idx_z", "dim_x", "dim_y",
// "dim_z"}; for (int i = 0; i < 6; ++i) {
// addArgument(F, Type::getInt32Ty(F->getContext()), names[i]);
// }
//}
// Return new function with additional index and limit arguments.
// The original function is removed from the module and erased.
Function *CodeGenTraversal::addIdxDimArgs(Function *F) {
......
......@@ -1152,11 +1152,8 @@ Function *CGT_CPU::createFunctionFilter(DFNode *C) {
// Add loop around the basic block, which exits the loop if isLastInput is
// false Pointers to keep the created loop structure
BasicBlock *EntryBB, *CondBB, *BodyBB;
Instruction *CondStartI = cast<Instruction>(isLastInputPop);
Instruction *BodyStartI = cast<Instruction>(Cond)->getNextNode();
EntryBB = CondStartI->getParent();
addWhileLoop(CondStartI, BodyStartI, RI, Cond);
// Return the Function pointer
......
//=== DFG2LLVM_OpenCL.cpp ===//
//===----------------------- DFG2LLVM_OpenCL.cpp ---------------------------===//
//
// The LLVM Compiler Infrastructure
//
......@@ -148,14 +148,11 @@ static Value *genWorkGroupPtr(Module &M, std::vector<Value *>,
ValueToValueMapTy &, Instruction *,
const Twine &WGName = "WGSize");
static std::string getPTXFilename(const Module &);
static std::string getFilenameFromModule(const Module &M);
static void changeDataLayout(Module &);
static void changeTargetTriple(Module &);
static void findReturnInst(Function *, std::vector<ReturnInst *> &);
static void findIntrinsicInst(Function *, Intrinsic::ID,
std::vector<IntrinsicInst *> &);
static AtomicRMWInst::BinOp getAtomicOp(Intrinsic::ID);
static std::string getAtomicOpName(Intrinsic::ID);
// DFG2LLVM_OpenCL - The first implementation.
struct DFG2LLVM_OpenCL : public DFG2LLVM {
......@@ -538,7 +535,7 @@ void CGT_OpenCL::insertRuntimeCalls(DFInternalNode *N, Kernel *K,
// location
AllocaInst *inputValPtr = new AllocaInst(
inputVal->getType(), 0, inputVal->getName() + ".ptr", RI);
StoreInst *SI = new StoreInst(inputVal, inputValPtr, RI);
new StoreInst(inputVal, inputValPtr, RI);
Value *inputValI8Ptr = CastInst::CreatePointerCast(
inputValPtr, Type::getInt8PtrTy(M.getContext()),
......@@ -592,7 +589,7 @@ void CGT_OpenCL::insertRuntimeCalls(DFInternalNode *N, Kernel *K,
AllocaInst *allocSizePtr =
new AllocaInst(allocSize->getType(), 0,
allocSize->getName() + ".sharedMem.ptr", RI);
StoreInst *SI = new StoreInst(allocSize, allocSizePtr, RI);
new StoreInst(allocSize, allocSizePtr, RI);
Value *allocSizeI8Ptr = CastInst::CreatePointerCast(
allocSizePtr, Type::getInt8PtrTy(M.getContext()),
......@@ -648,7 +645,7 @@ void CGT_OpenCL::insertRuntimeCalls(DFInternalNode *N, Kernel *K,
AllocaInst *allocSizePtr =
new AllocaInst(allocSize->getType(), 0,
allocSize->getName() + ".sharedMem.ptr", RI);
StoreInst *SI = new StoreInst(allocSize, allocSizePtr, RI);
new StoreInst(allocSize, allocSizePtr, RI);
Value *allocSizeI8Ptr = CastInst::CreatePointerCast(
allocSizePtr, Type::getInt8PtrTy(M.getContext()),
......@@ -1492,7 +1489,7 @@ void CGT_OpenCL::codeGen(DFLeafNode *N) {
dyn_cast<GetElementPtrInst>(Destination)) {
Value *SourcePtrOperand = sourceGEPI->getPointerOperand();
Value *DestPtrOperand = destGEPI->getPointerOperand();
for (int i = 0; i < memcpy_count; ++i) {
for (unsigned i = 0; i < memcpy_count; ++i) {
Constant *increment;
LoadInst *newLoadI;
StoreInst *newStoreI;
......@@ -1727,7 +1724,6 @@ void CGT_OpenCL::codeGen(DFLeafNode *N) {
continue; // not in pattern
}
DEBUG("HERE!\n");
// check that we load from pointer we got from bitcast - assert - the unique
// argument must be the use we found it from
assert(LoadI->getPointerOperand() == BitCastI &&
......@@ -1895,9 +1891,6 @@ bool DFG2LLVM_OpenCL::runOnModule(Module &M) {
// DFInternalNode *Root = DFG.getRoot();
std::vector<DFInternalNode *> Roots = DFG.getRoots();
// BuildDFG::HandleToDFNode &HandleToDFNodeMap =
// DFG.getHandleToDFNodeMap(); BuildDFG::HandleToDFEdge &HandleToDFEdgeMap
// = DFG.getHandleToDFEdgeMap();
// Visitor for Code Generation Graph Traversal
CGT_OpenCL *CGTVisitor = new CGT_OpenCL(M, DFG);
......@@ -1917,11 +1910,6 @@ bool DFG2LLVM_OpenCL::runOnModule(Module &M) {
}
std::string CGT_OpenCL::getKernelsModuleName(Module &M) {
/*SmallString<128> currentDir;
llvm::sys::fs::current_path(currentDir);
std::string fileName = getFilenameFromModule(M);
Twine output = Twine(currentDir) + "/Output/" + fileName + "";
return output.str().append(".kernels.ll");*/
std::string mid = M.getModuleIdentifier();
return mid.append(".kernels.ll");
}
......@@ -2364,12 +2352,6 @@ static std::string getPTXFilename(const Module &M) {
return moduleID;
}
// Get the name of the input file from module ID
static std::string getFilenameFromModule(const Module &M) {
std::string moduleID = M.getModuleIdentifier();
return moduleID.substr(moduleID.find_last_of("/") + 1);
}
// Changes the data layout of the Module to be compiled with OpenCL backend
// TODO: Figure out when to call it, probably after duplicating the modules
static void changeDataLayout(Module &M) {
......@@ -2422,55 +2404,6 @@ static void findIntrinsicInst(Function *F, Intrinsic::ID IntrinsicID,
}
}
// Helper funtion, returns the atomicrmw op, corresponding to intrinsic atomic
// op
static AtomicRMWInst::BinOp getAtomicOp(Intrinsic::ID ID) {
switch (ID) {
case Intrinsic::hpvm_atomic_add:
return AtomicRMWInst::Add;
case Intrinsic::hpvm_atomic_sub:
return AtomicRMWInst::Sub;
case Intrinsic::hpvm_atomic_min:
return AtomicRMWInst::Min;
case Intrinsic::hpvm_atomic_max:
return AtomicRMWInst::Max;
case Intrinsic::hpvm_atomic_xchg:
return AtomicRMWInst::Xchg;
case Intrinsic::hpvm_atomic_and:
return AtomicRMWInst::And;
case Intrinsic::hpvm_atomic_or:
return AtomicRMWInst::Or;
case Intrinsic::hpvm_atomic_xor:
return AtomicRMWInst::Xor;
default:
llvm_unreachable("Unsupported atomic intrinsic!");
};
}
// Helper funtion, returns the OpenCL function name, corresponding to atomic op
static std::string getAtomicOpName(Intrinsic::ID ID) {
switch (ID) {
case Intrinsic::hpvm_atomic_add:
return "atom_add";
case Intrinsic::hpvm_atomic_sub:
return "atom_sub";
case Intrinsic::hpvm_atomic_min:
return "atom_min";
case Intrinsic::hpvm_atomic_max:
return "atom_max";
case Intrinsic::hpvm_atomic_xchg:
return "atom_xchg";
case Intrinsic::hpvm_atomic_and:
return "atom_and";
case Intrinsic::hpvm_atomic_or:
return "atom_or";
case Intrinsic::hpvm_atomic_xor:
return "atom_xor";
default:
llvm_unreachable("Unsupported atomic intrinsic!");
};
}
} // End of namespace
char DFG2LLVM_OpenCL::ID = 0;
......
......@@ -132,10 +132,6 @@ static void ReplaceCallWithIntrinsic(Instruction *I, Intrinsic::ID IntrinsicID,
IS_HPVM_CALL(launch) /* Exists but not required */
IS_HPVM_CALL(edge) /* Exists but not required */
IS_HPVM_CALL(createNodeND)
// IS_HPVM_CALL(createNode)
// IS_HPVM_CALL(createNode1D)
// IS_HPVM_CALL(createNode2D)
// IS_HPVM_CALL(createNode3D)
IS_HPVM_CALL(bindIn)
IS_HPVM_CALL(bindOut)
IS_HPVM_CALL(push)
......@@ -152,23 +148,15 @@ IS_HPVM_CALL(getNumNodeInstances_x)
IS_HPVM_CALL(getNumNodeInstances_y)
IS_HPVM_CALL(getNumNodeInstances_z)
// Atomics
IS_HPVM_CALL(atomic_cmpxchg)
IS_HPVM_CALL(atomic_add)
IS_HPVM_CALL(atomic_sub)
IS_HPVM_CALL(atomic_xchg)
IS_HPVM_CALL(atomic_inc)
IS_HPVM_CALL(atomic_dec)
IS_HPVM_CALL(atomic_min)
IS_HPVM_CALL(atomic_max)
IS_HPVM_CALL(atomic_umin)
IS_HPVM_CALL(atomic_umax)
IS_HPVM_CALL(atomic_and)
IS_HPVM_CALL(atomic_or)
IS_HPVM_CALL(atomic_xor)
// Misc Fn
IS_HPVM_CALL(floor)
IS_HPVM_CALL(rsqrt)
IS_HPVM_CALL(sqrt)
IS_HPVM_CALL(sin)
IS_HPVM_CALL(cos)
......@@ -183,8 +171,7 @@ IS_HPVM_CALL(hint)
// Return the constant integer represented by value V
static unsigned getNumericValue(Value *V) {
assert(
isa<ConstantInt>(V) &&
assert(isa<ConstantInt>(V) &&
"Value indicating the number of arguments should be a constant integer");
return cast<ConstantInt>(V)->getZExtValue();
}
......@@ -892,6 +879,7 @@ static Type *getReturnTypeFromReturnInst(Function *F) {
return RI->getReturnValue()->getType();
}
}
return NULL;
}
char genhpvm::GenHPVM::ID = 0;
......
......@@ -1671,7 +1671,6 @@ void *llvm_hpvm_ocl_executeNode(void *graphID, unsigned workDim,
// TODO: Would like to use event to ensure better scheduling of kernels.
// Currently passing the event paratemeter results in seg fault with
// clEnqueueNDRangeKernel.
cl_event *event;
DEBUG(cout << "Enqueuing kernel:\n");
DEBUG(cout << "\tCommand Queue: " << Context->clCommandQue << flush << "\n");
DEBUG(cout << "\tKernel: " << Context->clKernel << flush << "\n");
......@@ -1707,7 +1706,7 @@ void *llvm_hpvm_ocl_executeNode(void *graphID, unsigned workDim,
hpvm_SwitchToTimer(&kernel_timer, hpvm_TimerID_NONE);
pthread_mutex_unlock(&ocl_mtx);
return event;
return NULL;
}
//////////////////////////////////////////////////////////////////////////////
......@@ -1781,7 +1780,6 @@ void *llvm_hpvm_ocl_launch(const char *FileName, const char *KernelName) {
checkErr(programSource != NULL, 1 /*bool true*/,
"Failure to load Program Binary");
cl_int binaryStatus;
// pthread_mutex_lock(&ocl_mtx);
Context->clProgram = clCreateProgramWithSource(
Context->clOCLContext, 1, (const char **)&programSource, NULL, &errcode);
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment