diff --git a/hpvm/include/SupportHPVM/DFG2LLVM.h b/hpvm/include/SupportHPVM/DFG2LLVM.h
index e517a7d542ccde98d4e212c59803226045fce632..c2270de3cffc4767b1220e0ea7cbaf11b8c9dc61 100644
--- a/hpvm/include/SupportHPVM/DFG2LLVM.h
+++ b/hpvm/include/SupportHPVM/DFG2LLVM.h
@@ -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) {
diff --git a/hpvm/lib/Transforms/DFG2LLVM_CPU/DFG2LLVM_CPU.cpp b/hpvm/lib/Transforms/DFG2LLVM_CPU/DFG2LLVM_CPU.cpp
index 3655279a99ceecf462ca4aab46c25f82cf238ef0..de9c025c0e7e996b6abfaa8748adf6688d04d10d 100644
--- a/hpvm/lib/Transforms/DFG2LLVM_CPU/DFG2LLVM_CPU.cpp
+++ b/hpvm/lib/Transforms/DFG2LLVM_CPU/DFG2LLVM_CPU.cpp
@@ -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
diff --git a/hpvm/lib/Transforms/DFG2LLVM_OpenCL/DFG2LLVM_OpenCL.cpp b/hpvm/lib/Transforms/DFG2LLVM_OpenCL/DFG2LLVM_OpenCL.cpp
index 2d9a07500f355f7fd805f74c668814d905842fed..5a58f272b3042a4ebfc2e3c7bb3606b5c19e8d84 100644
--- a/hpvm/lib/Transforms/DFG2LLVM_OpenCL/DFG2LLVM_OpenCL.cpp
+++ b/hpvm/lib/Transforms/DFG2LLVM_OpenCL/DFG2LLVM_OpenCL.cpp
@@ -1,4 +1,4 @@
-//=== 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;
diff --git a/hpvm/lib/Transforms/GenHPVM/GenHPVM.cpp b/hpvm/lib/Transforms/GenHPVM/GenHPVM.cpp
index 864dad58cb507fb49ddf3bdc9fcafefa57f866d7..6c3dcd75f6bc5f81b2e834d89ae080c731df210e 100644
--- a/hpvm/lib/Transforms/GenHPVM/GenHPVM.cpp
+++ b/hpvm/lib/Transforms/GenHPVM/GenHPVM.cpp
@@ -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;
diff --git a/hpvm/projects/hpvm-rt/hpvm-rt.cpp b/hpvm/projects/hpvm-rt/hpvm-rt.cpp
index b6273ec2cca712469269f68f538ce437e9b062ec..b42668f2b46e007e0892908fbcbda09f67774f0d 100644
--- a/hpvm/projects/hpvm-rt/hpvm-rt.cpp
+++ b/hpvm/projects/hpvm-rt/hpvm-rt.cpp
@@ -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);