From 833e38a87672636a97824e13847c4503a3dda124 Mon Sep 17 00:00:00 2001
From: Prakalp Srivastava <prakalps@gmail.com>
Date: Sun, 2 Aug 2015 09:22:23 -0500
Subject: [PATCH] Improved Debug messages in Passes
---
llvm/lib/Transforms/BuildDFG/BuildDFG.cpp | 31 ++++++-------
llvm/lib/Transforms/ClearDFG/ClearDFG.cpp | 5 +--
.../DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp | 45 +++++++++----------
.../Transforms/DFG2LLVM_X86/DFG2LLVM_X86.cpp | 14 +++---
llvm/lib/Transforms/GenVISC/GenVISC.cpp | 45 ++++++++++---------
5 files changed, 66 insertions(+), 74 deletions(-)
diff --git a/llvm/lib/Transforms/BuildDFG/BuildDFG.cpp b/llvm/lib/Transforms/BuildDFG/BuildDFG.cpp
index 4c938b5f01..15d5ee4605 100644
--- a/llvm/lib/Transforms/BuildDFG/BuildDFG.cpp
+++ b/llvm/lib/Transforms/BuildDFG/BuildDFG.cpp
@@ -18,12 +18,10 @@
using namespace llvm;
-STATISTIC(IntrinsicCounter, "Counts number of visc intrinsics greeted");
-
namespace builddfg {
static visc::Target getPreferredTarget(Function* F) {
- errs() << "Finding preferred target for " << F->getName() << "\n";
+ DEBUG(errs() << "Finding preferred target for " << F->getName() << "\n");
Module* M = F->getParent();
NamedMDNode* HintNode = M->getOrInsertNamedMetadata("visc_hint_gpu");
for(unsigned i = 0; i < HintNode->getNumOperands(); i++) {
@@ -35,20 +33,20 @@ static visc::Target getPreferredTarget(Function* F) {
}
bool BuildDFG::runOnModule(Module &M) {
-
- errs() << "-------- Searching for launch sites ----------\n";
+ errs() << "\nBUILDDFG PASS\n";
+ DEBUG(errs() << "-------- Searching for launch sites ----------\n");
IntrinsicInst* II;
// Iterate over all functions in the module
for (Module::iterator mi = M.begin(), me = M.end(); mi != me; ++mi) {
Function* f = &*mi;
- errs() << "Function: " << f->getName() << "\n";
+ DEBUG(errs() << "Function: " << f->getName() << "\n");
for (inst_iterator i = inst_begin(f), e = inst_end(f); i != e ; ++i) {
Instruction* I = &*i; // Grab pointer to Instruction
if (isViscLaunchIntrinsic(I)) {
- errs() << "------------ Found launch site --------------\n";
+ DEBUG(errs() << "------------ Found launch site --------------\n");
II = cast<IntrinsicInst>(I);
assert(II && "Launch intrinsic not recognized.");
@@ -62,13 +60,13 @@ bool BuildDFG::runOnModule(Module &M) {
for(DFGraph::children_iterator i = Root->getChildGraph()->begin(),
e = Root->getChildGraph()->end(); i!=e; i++) {
DFNode* N = *i;
- errs() << "\t" << N->getFuncPointer()->getName() << "\n";
+ DEBUG(errs() << "\t" << N->getFuncPointer()->getName() << "\n");
}
Root->getChildGraph()->sortChildren();
for(DFGraph::children_iterator i = Root->getChildGraph()->begin(),
e = Root->getChildGraph()->end(); i!=e; i++) {
DFNode* N = *i;
- errs() << "\t" << N->getFuncPointer()->getName() << "\n";
+ DEBUG(errs() << "\t" << N->getFuncPointer()->getName() << "\n");
}
viewDFGraph(Root->getChildGraph());
@@ -327,9 +325,9 @@ void BuildDFG::BuildGraph (DFInternalNode* N, Function *F) {
// intrinsics.
for (inst_iterator i = inst_begin(F), e = inst_end(F); i != e; ++i) {
Instruction* I = &*i; // Grab pointer to instruction reference
- errs() << *I << "\n";
+ DEBUG(errs() << *I << "\n");
if(IntrinsicInst* II = dyn_cast<IntrinsicInst>(I)) {
- errs() << "IntrinsicID = " << II->getIntrinsicID() << ": " << II->getCalledFunction()->getName()<<"\n";
+ DEBUG(errs() << "IntrinsicID = " << II->getIntrinsicID() << ": " << II->getCalledFunction()->getName()<<"\n");
switch(II->getIntrinsicID()) {
case Intrinsic::visc_createNode:
@@ -352,17 +350,20 @@ void BuildDFG::BuildGraph (DFInternalNode* N, Function *F) {
//TODO: Reconsider launch within a dataflow graph (recursion?)
case Intrinsic::visc_wait:
case Intrinsic::visc_launch:
- errs() << "Error: Launch/wait intrinsic used within a dataflow graph\n";
+ errs() << "Error: Launch/wait intrinsic used within a dataflow graph\n\t" << *II << "\n";
break;
default:
- errs() << "Error: Invalid VISC Intrinsic inside Internal node!\n";
+ errs() << "Error: Invalid VISC Intrinsic inside Internal node!\n\t" << *II << "\n";
break;
}
}
- else {
- errs() << "Non-intrinsic instruction\n";
+ else if(!isa<ReturnInst>(I)) {
+ errs() << "Non-intrinsic instruction: " << *I << "\n";
+ llvm_unreachable("Found non-intrinsic instruction inside an internal node. Only return instruction is allowed!");
+
}
+
}
}
diff --git a/llvm/lib/Transforms/ClearDFG/ClearDFG.cpp b/llvm/lib/Transforms/ClearDFG/ClearDFG.cpp
index b3ea7d17b2..6d04074b96 100644
--- a/llvm/lib/Transforms/ClearDFG/ClearDFG.cpp
+++ b/llvm/lib/Transforms/ClearDFG/ClearDFG.cpp
@@ -80,13 +80,13 @@ public:
virtual void visit(DFLeafNode* N) {
DEBUG(errs() << "Erasing Node (L) - " << N->getFuncPointer()->getName() << "\n");
deleteNode(N);
- errs() << "DONE" << "\n";
+ DEBUG(errs() << "DONE" << "\n");
}
};
bool ClearDFG::runOnModule(Module &M) {
-
+ errs() << "\nCLEARDFG PASS\n";
// Get the BuildDFG Analysis Results:
// - Dataflow graph
// - Maps from i8* hansles to DFNode and DFEdge
@@ -130,7 +130,6 @@ void TreeTraversal::deleteNode(DFNode* N) {
IntrinsicInst* LI = N->getInstruction();
LI->replaceAllUsesWith(UndefValue::get(LI->getType()));
LI->eraseFromParent();
-
}
} // End of namespace
diff --git a/llvm/lib/Transforms/DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp b/llvm/lib/Transforms/DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp
index ecaf0a4139..dd4cc894e4 100644
--- a/llvm/lib/Transforms/DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp
+++ b/llvm/lib/Transforms/DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp
@@ -35,12 +35,6 @@ using namespace llvm;
using namespace builddfg;
using namespace dfg2llvm;
-//STATISTIC(IntrinsicCounter, "Counts number of visc intrinsics greeted");
-
-#define DECLARE(X) X = M.getOrInsertFunction(#X, \
- runtimeModule->getFunction(#X)->getFunctionType()); \
- DEBUG(errs() << *X)
-
// VISC Command line option to use timer or not
static cl::opt<bool>
VISCTimer_NVPTX("visc-timers-ptx", cl::desc("Enable visc timers"));
@@ -254,9 +248,9 @@ void CGT_NVPTX::initRuntimeAPI() {
DECLARE(llvm_visc_ptx_free);
DECLARE(llvm_visc_ptx_getOutput);
DECLARE(llvm_visc_ptx_executeNode);
- DECLARE(llvm_visc_initializeTimerSet);
- DECLARE(llvm_visc_switchToTimer);
- DECLARE(llvm_visc_printTimerSet);
+
+ // Get or insert timerAPI functions as well if you plan to use timers
+ initTimerAPI();
// Insert init context in main
DEBUG(errs() << "Gen Code to initialize NVPTX Timer\n");
@@ -272,7 +266,7 @@ void CGT_NVPTX::initRuntimeAPI() {
// Insert print instruction at visc exit
DEBUG(errs() << "Gen Code to print NVPTX Timer\n");
Function* VC = M.getFunction("llvm.visc.cleanup");
- errs() << *VC << "\n";
+ DEBUG(errs() << *VC << "\n");
assert(VC->getNumUses() == 1 && "__visc__clear should only be used once");
CleanupCall = cast<Instruction>(*VC->use_begin());
@@ -392,7 +386,7 @@ void CGT_NVPTX::insertRuntimeCalls(DFInternalNode* N, Kernel* K, const Twine& Fi
GlobalValue::CommonLinkage,
Constant::getNullValue(NVPTX_Ctx->getType()),
"graph"+KF->getName()+".addr");
- errs() << "Store at: " << *GraphIDAddr << "\n";
+ DEBUG(errs() << "Store at: " << *GraphIDAddr << "\n");
StoreInst* SI = new StoreInst(NVPTX_Ctx, GraphIDAddr, InitCall);
DEBUG(errs() << *SI << "\n");
switchToTimer(visc_TimerID_NONE, InitCall);
@@ -551,13 +545,13 @@ void CGT_NVPTX::insertRuntimeCalls(DFInternalNode* N, Kernel* K, const Twine& Fi
// Read all the pointer arguments which had side effects i.e., had out
// attribute
- errs() << "Output Pointers : " << OutputPointers.size() << "\n";
+ DEBUG(errs() << "Output Pointers : " << OutputPointers.size() << "\n");
// FIXME: Not reading output pointers anymore as we read them when data is
// actually requested
/*for(auto output: OutputPointers) {
- errs() << "Read: " << *output.d_ptr << "\n";
- errs() << "\tTo: " << *output.h_ptr << "\n";
- errs() << "\t#bytes: " << *output.bytes << "\n";
+ DEBUG(errs() << "Read: " << *output.d_ptr << "\n");
+ DEBUG(errs() << "\tTo: " << *output.h_ptr << "\n");
+ DEBUG(errs() << "\t#bytes: " << *output.bytes << "\n");
Value* GetOutputArgs[] = {GraphID, output.h_ptr, output.d_ptr, output.bytes};
CallInst* CI = CallInst::Create(llvm_visc_ptx_getOutput,
@@ -658,7 +652,7 @@ void CGT_NVPTX::codeGen(DFInternalNode* N) {
// Now the remaining nodes to be visited should be ignored
KernelLaunchNode = NULL;
- errs() << "Insert Runtime calls\n";
+ DEBUG(errs() << "Insert Runtime calls\n");
insertRuntimeCalls(N, kernel, getPTXFilename(M));
} else {
@@ -747,7 +741,7 @@ void CGT_NVPTX::codeGen(DFLeafNode* N) {
}
else {
// Converting a 2-level DFG to opencl kernel
- errs() << "*************** Kernel Gen: 2-Level Hierarchy **************\n";
+ DEBUG(errs() << "*************** Kernel Gen: 2-Level Hierarchy **************\n");
KernelLaunchNode = PNode->getParent();
assert((PNode->getNumOfDim() == N->getNumOfDim()) && "Dimension number must match");
// Contains the instructions generating the kernel configuration parameters
@@ -882,8 +876,8 @@ void CGT_NVPTX::codeGen(DFLeafNode* N) {
Function * OpenCLFunction;
int parentLevel = N->getParent()->getLevel();
int parentReplFactor = N->getParent()->getNumOfDim();
- errs() << "Parent Level = " << parentLevel << "\n";
- errs() << "Parent Repl factor = " << parentReplFactor << "\n";
+ DEBUG(errs() << "Parent Level = " << parentLevel << "\n");
+ DEBUG(errs() << "Parent Repl factor = " << parentReplFactor << "\n");
if ((!parentLevel || !parentReplFactor) && ArgDFNode == N) {
// We only have one level in the hierarchy or the parent node is not
@@ -891,8 +885,8 @@ void CGT_NVPTX::codeGen(DFLeafNode* N) {
// launch, so we need to specify a global id.
// We can translate this only if the argument is the current node
// itself
- errs() << "Substitute with get_global_id()\n";
- errs() << *II << "\n";
+ DEBUG(errs() << "Substitute with get_global_id()\n");
+ DEBUG(errs() << *II << "\n");
FunctionType* FT =
FunctionType::get(Type::getInt32Ty(getGlobalContext() /*KernelM.getContext()*/),
std::vector<Type*>(1, Type::getInt32Ty(getGlobalContext() /*KernelM.getContext()*/)),
@@ -1047,6 +1041,7 @@ void CGT_NVPTX::codeGen(DFLeafNode* N) {
}
bool DFG2LLVM_NVPTX::runOnModule(Module &M) {
+ errs() << "\nDFG2LLVM_NVPTX PASS\n";
// Get the BuildDFG Analysis Results:
// - Dataflow graph
@@ -1315,12 +1310,12 @@ static Value* genWorkGroupPtr(std::vector<Value*> WGSize, ValueToValueMapTy& VMa
// generate code to extend it to i64. We need to use the mapped value in
// the new generated function, hence the use of VMap
// FIXME: Why are we changing the kernel WGSize vector here?
- errs() << "Not i64. Zero extend required.\n";
- errs() << *WGSize[i] << "\n";
+ DEBUG(errs() << "Not i64. Zero extend required.\n");
+ DEBUG(errs() << *WGSize[i] << "\n");
CastInst* CI = BitCastInst::CreateIntegerCast(WGSize[i], Int64Ty, true, "", IB);
- errs() << "Bitcast done.\n";
+ DEBUG(errs() << "Bitcast done.\n");
StoreInst* SI = new StoreInst(CI, nextDim, IB);
- errs() << "Zero extend done.\n";
+ DEBUG(errs() << "Zero extend done.\n");
DEBUG(errs() << "\tZero extended work group size: " << *SI << "\n");
} else {
// Store the value representing work group size in ith dimension on
diff --git a/llvm/lib/Transforms/DFG2LLVM_X86/DFG2LLVM_X86.cpp b/llvm/lib/Transforms/DFG2LLVM_X86/DFG2LLVM_X86.cpp
index 1419b4446e..9141723f08 100644
--- a/llvm/lib/Transforms/DFG2LLVM_X86/DFG2LLVM_X86.cpp
+++ b/llvm/lib/Transforms/DFG2LLVM_X86/DFG2LLVM_X86.cpp
@@ -25,11 +25,6 @@ using namespace llvm;
using namespace builddfg;
using namespace dfg2llvm;
-//STATISTIC(IntrinsicCounter, "Counts number of visc intrinsics greeted");
-#define DECLARE(X) X = M.getOrInsertFunction(#X, \
- runtimeModule->getFunction(#X)->getFunctionType()); \
- DEBUG(errs() << *X)
-
// VISC Command line option to use timer or not
static cl::opt<bool>
VISCTimer_X86("visc-timers-x86", cl::desc("Enable visc timers"));
@@ -95,6 +90,7 @@ public:
};
bool DFG2LLVM_X86::runOnModule(Module &M) {
+ errs() << "\nDFG2LLVM_X86 PASS\n";
// Get the BuildDFG Analysis Results:
// - Dataflow graph
@@ -150,9 +146,9 @@ void CGT_X86::initRuntimeAPI() {
DECLARE(llvm_visc_x86_pop);
DECLARE(llvm_visc_x86_getDimLimit);
DECLARE(llvm_visc_x86_getDimInstance);
- DECLARE(llvm_visc_initializeTimerSet);
- DECLARE(llvm_visc_switchToTimer);
- DECLARE(llvm_visc_printTimerSet);
+
+ // Get or insert timerAPI functions as well if you plan to use timers
+ initTimerAPI();
// Insert init context in main
Function* VI = M.getFunction("llvm.visc.init");
@@ -489,7 +485,7 @@ void CGT_X86::invokeChild_X86(DFNode* C, Function* F_X86,
CallInst* Pop = CallInst::Create(llvm_visc_x86_pop, None, "", NextI);
DEBUG(errs() << "Pop from stack: " << *Pop << "\n");
- errs() << *CI->getParent()->getParent();
+ DEBUG(errs() << *CI->getParent()->getParent());
}
void CGT_X86::codeGen(DFInternalNode* N) {
diff --git a/llvm/lib/Transforms/GenVISC/GenVISC.cpp b/llvm/lib/Transforms/GenVISC/GenVISC.cpp
index 97e0a56600..68f5ca56a3 100644
--- a/llvm/lib/Transforms/GenVISC/GenVISC.cpp
+++ b/llvm/lib/Transforms/GenVISC/GenVISC.cpp
@@ -96,7 +96,7 @@ static void addArgs(Function* F, unsigned numArgs, std::string names[]) {
static void addHint(Function* F, visc::Target T) {
// Get Module
Module* M = F->getParent();
- errs() << "Set preferred target for " << F->getName() << ": " << T << "\n";
+ DEBUG(errs() << "Set preferred target for " << F->getName() << ": " << T << "\n");
//assert(isa<ConstantInt>(CI->getArgOperand(0))
//&& "Argument to hint must be constant integer!");
@@ -117,7 +117,7 @@ static void addHint(Function* F, visc::Target T) {
}
static visc::Target getPreferredTarget(Function* F) {
- errs() << "Finding preferred target for " << F->getName() << "\n";
+ DEBUG(errs() << "Finding preferred target for " << F->getName() << "\n");
Module* M = F->getParent();
NamedMDNode* HintNode = M->getOrInsertNamedMetadata("visc_hint_gpu");
for(unsigned i = 0; i < HintNode->getNumOperands(); i++) {
@@ -245,7 +245,7 @@ static void handleVISCAttributes(Function* F, CallInst* CI) {
assert(CI->getNumArgOperands() > offset
&& "Too few arguments for __visc__attributes call!");
unsigned numInPtrs = getNumericValue(CI->getArgOperand(offset));
- errs() << "\tNum of in pointers = " << numInPtrs << "\n";
+ DEBUG(errs() << "\tNum of in pointers = " << numInPtrs << "\n");
for(unsigned i = offset+1; i< offset+1+numInPtrs; i++) {
Value* V = CI->getArgOperand(i);
@@ -262,7 +262,7 @@ static void handleVISCAttributes(Function* F, CallInst* CI) {
assert(CI->getNumArgOperands() > offset
&& "Too few arguments for __visc__attributes call!");
unsigned numOutPtrs = getNumericValue(CI->getOperand(offset));
- errs() << "\tNum of out Pointers = " << numOutPtrs << "\n";
+ DEBUG(errs() << "\tNum of out Pointers = " << numOutPtrs << "\n");
for(unsigned i = offset+1; i< offset+1+numOutPtrs; i++) {
Value* V = CI->getArgOperand(i);
if(Argument* arg = dyn_cast<Argument>(V)) {
@@ -414,6 +414,7 @@ static Function* genInternalNode(Function* KernelF, unsigned level,
// Public Functions of GenVISC pass
bool GenVISC::runOnModule(Module &M) {
+ errs() << "\nGENVISC PASS\n";
this->M = &M;
// Load Runtime API Module
@@ -487,12 +488,12 @@ bool GenVISC::runOnModule(Module &M) {
if(isVISCinitCall(I)) {
CallInst* CI = cast<CallInst>(I);
Function* InitF = Intrinsic::getDeclaration(&M, Intrinsic::visc_init);
- errs() << *InitF << "\n";
+ DEBUG(errs() << *InitF << "\n");
CallInst* InitInst = CallInst::Create(InitF,
None, "", CI);
toBeErased.push_back(CI);
- errs() << "Found visc init call: " << *CI << "\n";
- errs() << "\tSubstitute with: " << *InitInst << "\n";
+ DEBUG(errs() << "Found visc init call: " << *CI << "\n");
+ DEBUG(errs() << "\tSubstitute with: " << *InitInst << "\n");
}
if(isVISCcleanupCall(I)) {
CallInst* CI = cast<CallInst>(I);
@@ -500,20 +501,20 @@ bool GenVISC::runOnModule(Module &M) {
CallInst* CleanupInst = CallInst::Create(CleanupF,
None,
"", CI);
- errs() << "Found visc cleanup call: " << *CI << "\n";
- errs() << "\tSubstitute with: " << *CleanupInst << "\n";
+ DEBUG(errs() << "Found visc cleanup call: " << *CI << "\n");
+ DEBUG(errs() << "\tSubstitute with: " << *CleanupInst << "\n");
toBeErased.push_back(CI);
}
if(isVISCwaitCall(I)) {
CallInst* CI = cast<CallInst>(I);
Function* WaitF = Intrinsic::getDeclaration(&M, Intrinsic::visc_wait);
- errs() << *WaitF << "\n";
- errs() << *CI->getArgOperand(1) << "\n";
+ DEBUG(errs() << *WaitF << "\n");
+ DEBUG(errs() << *CI->getArgOperand(1) << "\n");
CallInst* WaitInst = CallInst::Create(WaitF,
ArrayRef<Value*>(CI->getArgOperand(0)),
"", CI);
- errs() << "Found visc wait call: " << *CI << "\n";
- errs() << "\tSubstitute with: " << *WaitInst << "\n";
+ DEBUG(errs() << "Found visc wait call: " << *CI << "\n");
+ DEBUG(errs() << "\tSubstitute with: " << *WaitInst << "\n");
toBeErased.push_back(CI);
}
if(isVISCtrackMemoryCall(I)) {
@@ -523,8 +524,8 @@ bool GenVISC::runOnModule(Module &M) {
CallInst* TrackMemInst = CallInst::Create(TrackMemoryF,
ArrayRef<Value*>(TrackMemArgs,2),
"", CI);
- errs() << "Found visc track memory call: " << *CI << "\n";
- errs() << "\tSubstitute with: " << *TrackMemInst << "\n";
+ DEBUG(errs() << "Found visc track memory call: " << *CI << "\n");
+ DEBUG(errs() << "\tSubstitute with: " << *TrackMemInst << "\n");
}
if(isVISCuntrackMemoryCall(I)) {
CallInst* CI = cast<CallInst>(I);
@@ -532,8 +533,8 @@ bool GenVISC::runOnModule(Module &M) {
CallInst* UntrackMemInst = CallInst::Create(UntrackMemoryF,
ArrayRef<Value*>(CI->getArgOperand(0)),
"", CI);
- errs() << "Found visc *un*track memory call: " << *CI << "\n";
- errs() << "\tSubstitute with: " << *UntrackMemInst << "\n";
+ DEBUG(errs() << "Found visc *un*track memory call: " << *CI << "\n");
+ DEBUG(errs() << "\tSubstitute with: " << *UntrackMemInst << "\n");
}
if(isVISCrequestMemoryCall(I)) {
CallInst* CI = cast<CallInst>(I);
@@ -542,8 +543,8 @@ bool GenVISC::runOnModule(Module &M) {
CallInst* RequestMemInst = CallInst::Create(RequestMemoryF,
ArrayRef<Value*>(RequestMemArgs,2),
"", CI);
- errs() << "Found visc request memory call: " << *CI << "\n";
- errs() << "\tSubstitute with: " << *RequestMemInst << "\n";
+ DEBUG(errs() << "Found visc request memory call: " << *CI << "\n");
+ DEBUG(errs() << "\tSubstitute with: " << *RequestMemInst << "\n");
}
if(isVISChintCall(I)) {
CallInst* CI = cast<CallInst>(I);
@@ -554,16 +555,16 @@ bool GenVISC::runOnModule(Module &M) {
visc::Target t = hint->equalsInt(visc::GPU_TARGET)?
visc::GPU_TARGET : visc::CPU_TARGET;
addHint(CI->getParent()->getParent(), t);
- errs() << "Found visc hint call: " << *CI << "\n";
+ DEBUG(errs() << "Found visc hint call: " << *CI << "\n");
toBeErased.push_back(CI);
}
}
}
// Erase the __visc__node calls
- errs() << "Erase Statements:\n";
+ DEBUG(errs() << "Erase Statements:\n");
for(auto I: toBeErased) {
- errs() << *I << "\n";
+ DEBUG(errs() << *I << "\n");
I->eraseFromParent();
}
--
GitLab