diff --git a/CREDITS.txt b/CREDITS.TXT
similarity index 95%
rename from CREDITS.txt
rename to CREDITS.TXT
index 1c996e47ee6ddd6c7b6e7b2c1ed31ab9fe23cf8a..1e130c355c301fc42da1df20f5034f9b6a9f0681 100644
--- a/CREDITS.txt
+++ b/CREDITS.TXT
@@ -1,5 +1,4 @@
-This file is a partial list of people who have contributed to the HPVM-CAVA
-pilot project.
+This file is a partial list of people who have contributed to the HPVM project.
As with LLVM's credits, the list is sorted by surname and formatted to allow
easy grepping and beautification by scripts. The fields are: name (N) and
diff --git a/LICENSE b/LICENSE.TXT
similarity index 100%
rename from LICENSE
rename to LICENSE.TXT
diff --git a/README.md b/README.md
index 5061f0df3f8742e73d1c7327119a44a825b3f028..2e73dc27600a27e843171099d2506af0114f9958 100644
--- a/README.md
+++ b/README.md
@@ -25,22 +25,35 @@ The following components are required to be installed on your machine to build H
* CMake (>=3.4.3)
* Python (>=2.7)
* GNU Make (>=3.79.1)
-* CUDA (>= 9.1) (only required for GPU support)
+* OpenCL (>=1.0.0) or CUDA (>=9.1, only required for GPU support)
+
+
+## Supported Targets
+Supported/tested CPU architectures:
+* Intel Xeon E5-2640
+* Intel Xeon W-2135
+* ARM Cortex A-57
+
+Supported/tested GPU architectures:
+* Nvidia Quadro P1000
+* Nvidia GeForce GTX 1080
+
+HPVM has not been tested but might work on other CPUs supported by LLVM Backend, and GPUs supported by OpenCL such as Intel, AMD, etc.
+
## Getting source code and building HPVM
Checkout HPVM:
```shell
git clone https://gitlab.engr.illinois.edu/llvm/hpvm-release.git/
-cd hpvm-release
+cd hpvm-release/hpvm
```
-HPVM installer script can be used to download, configure and build HPVM along with LLVM and other subprojects including Clang.
+HPVM installer script can be used to download, configure and build HPVM along with LLVM and Clang.
```shell
-cd hpvm
bash install.sh
```
-Specifically, the HPVM installer downloads LLVM, Clang, libcxx and libcxxabi, copies HPVM source into
+Specifically, the HPVM installer downloads LLVM, and Clang, copies HPVM source into
llvm/tools and builds the entire tree. It also builds a modified LLVM C-Backend, based on the one maintained by [Julia Computing](https://github.com/JuliaComputing/llvm-cbe), as a part of HPVM and is currently used
to generate OpenCL kernels for GPUs.
@@ -48,7 +61,7 @@ In the beginning of the building process, the installer provides users the choic
If HPVM is selected to be built automatically, the installer allows users to type in the number of threads they want to use.
The default number of threads used to build HPVM is two.
-Alternatively, CMake can be run manually.
+Alternatively, CMake can be run manually using the following steps in ./hpvm-release/hpvm directory.
```shell
mkdir build
cd build
@@ -62,15 +75,14 @@ Some common options that can be used with CMake are:
* -DLLVM_ENABLE_ASSERTIONS=On --- Compile with assertion checks enabled (default is Yes for Debug builds, No for all other build types).
-In order to manually build and install HPVM, GNU Make can be run using the following.
+In order to manually build and install HPVM, GNU Make can be run using the following in the build directory.
```shell
make -j<number of threads>
make install
```
In the end of the installation process, the installer automatically runs all the regression tests to ensure that the installation is
-successful. If HPVM is built and installed manually, the tests can be automatically run by executing the following step from the
-current directory.
+successful. If HPVM is built and installed manually, the tests can be automatically run by executing the following step from the ./hpvm-release/hpvm directory.
```shell
bash scripts/automate_tests.sh
```
@@ -87,4 +99,7 @@ Benchmark descriptions and instructions on how to compile and run them are [here
We are also providing [unit tests](/hpvm/test/unitTests) and [regression tests](/hpvm/test/regressionTests).
+## Support
+All questions can be directed to [hpvm-dev@lists.cs.illinois.edu](mailto:hpvm-dev@lists.cs.illinois.edu).
+
diff --git a/hpvm/CMakeLists.txt b/hpvm/CMakeLists.txt
index 9fc8bd884b4562d3575943e76c4bf35ec4a39f95..3d304fa99b61884a8bb9c26d57ae0141ad551bdd 100644
--- a/hpvm/CMakeLists.txt
+++ b/hpvm/CMakeLists.txt
@@ -1,7 +1,7 @@
include_directories(./include/)
add_subdirectory(lib)
add_subdirectory(projects)
-
+add_subdirectory(tools)
add_subdirectory(test)
diff --git a/hpvm/include/BuildDFG/BuildDFG.h b/hpvm/include/BuildDFG/BuildDFG.h
index f568ffc7c3083050b6d5c40d7c2779951aad73d9..fd36a2593e4f0200e70a2244a2c3b95c6a6d0823 100644
--- a/hpvm/include/BuildDFG/BuildDFG.h
+++ b/hpvm/include/BuildDFG/BuildDFG.h
@@ -11,8 +11,8 @@
//===----------------------------------------------------------------------===//
//
// This pass defines the BuildDFG pass which uses LLVM IR with HPVM intrinsics
-// to infer information about dataflow graph hierarchy and structure to construct
-// HPVM IR.
+// to infer information about dataflow graph hierarchy and structure to
+// construct HPVM IR.
//
//===----------------------------------------------------------------------===//
diff --git a/hpvm/include/GenHPVM/GenHPVM.h b/hpvm/include/GenHPVM/GenHPVM.h
index 78afb325e7a37efeae1a57d2aaa59951783e9da8..f61d4a7c90dff2e4bff5f781c6c9cc92e3246232 100644
--- a/hpvm/include/GenHPVM/GenHPVM.h
+++ b/hpvm/include/GenHPVM/GenHPVM.h
@@ -7,12 +7,11 @@
//
//===----------------------------------------------------------------------===//
//
-// This file defines the GenHPVM pass responsible for converting HPVM-C to
+// This file defines the GenHPVM pass responsible for converting HPVM-C to
// HPVM intrinsics. Note that this pass relies on memory-to-register optimiza-
-// tion pass to execute before this executes.
+// tion pass to execute before this executes.
//
//===----------------------------------------------------------------------===//
-
#include "SupportHPVM/HPVMTimer.h"
#include "llvm/IR/DerivedTypes.h"
diff --git a/hpvm/include/SupportHPVM/DFG2LLVM.h b/hpvm/include/SupportHPVM/DFG2LLVM.h
index 47bd242520c7d2ddca0e0bc7bbbab38fa9fdc543..c2270de3cffc4767b1220e0ea7cbaf11b8c9dc61 100644
--- a/hpvm/include/SupportHPVM/DFG2LLVM.h
+++ b/hpvm/include/SupportHPVM/DFG2LLVM.h
@@ -10,7 +10,7 @@
//
//===----------------------------------------------------------------------===//
//
-// This defines different classes for traversing Dataflow Graph for code
+// This defines different classes for traversing Dataflow Graph for code
// generation for different nodes for different targets.
//
//===----------------------------------------------------------------------===//
@@ -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/include/SupportHPVM/HPVMUtils.h b/hpvm/include/SupportHPVM/HPVMUtils.h
index 3060c6c0321dc197023f24cd8fefe106e18562b7..537a92caec4a5c63b0fcc06b3714bddefc0a3fde 100644
--- a/hpvm/include/SupportHPVM/HPVMUtils.h
+++ b/hpvm/include/SupportHPVM/HPVMUtils.h
@@ -8,8 +8,8 @@
//
//===----------------------------------------------------------------------===//
//
-// This filed defines utility functions used for target-specific code generation
-// for different nodes of dataflow graphs.
+// This filed defines utility functions used for target-specific code generation
+// for different nodes of dataflow graphs.
//
//===----------------------------------------------------------------------===//
diff --git a/hpvm/install.sh b/hpvm/install.sh
index 776b8aa6d11c0d92af507a161fb902b183da7672..692c6195d28b23c418f37d8a574d91ef744e2b33 100644
--- a/hpvm/install.sh
+++ b/hpvm/install.sh
@@ -8,4 +8,4 @@ BASH=/bin/bash
$BASH $SCRIPTS_DIR/llvm_installer.sh
# Run the tests
-$BASH $SCRIPTS_DIR/automate_tests.sh
+$BASH $SCRIPTS_DIR/automated_tests.sh
diff --git a/hpvm/lib/Transforms/BuildDFG/BuildDFG.cpp b/hpvm/lib/Transforms/BuildDFG/BuildDFG.cpp
index 25f5cd0870dc343c2cabe9e872f2b341f4fa899e..3177f860057b48dc1b22bc61940b53da848dca09 100644
--- a/hpvm/lib/Transforms/BuildDFG/BuildDFG.cpp
+++ b/hpvm/lib/Transforms/BuildDFG/BuildDFG.cpp
@@ -10,8 +10,8 @@
// BuildDFG pass is responsible for constructing dataflow graph from a textual
// representation of HPVM IR with HPVM intrinsics from GenHPVM pass. This pass
// makes use of three crutial abstractions: graph itself, dataflow nodes repre-
-// -senting functions and data edges representing tranfer of data between
-// the functions (or nodes in the graph). This pass is part of HPVM frontend
+// -senting functions and data edges representing tranfer of data between
+// the functions (or nodes in the graph). This pass is part of HPVM frontend
// and does not make any changes to the textual representation of the IR.
//
//===----------------------------------------------------------------------===//
diff --git a/hpvm/lib/Transforms/ClearDFG/ClearDFG.cpp b/hpvm/lib/Transforms/ClearDFG/ClearDFG.cpp
index b90b641a9d4a224279af064eafcc397985277168..c905f745e658426951ae06fbe2a1d85685dfec74 100644
--- a/hpvm/lib/Transforms/ClearDFG/ClearDFG.cpp
+++ b/hpvm/lib/Transforms/ClearDFG/ClearDFG.cpp
@@ -7,7 +7,7 @@
//
//===----------------------------------------------------------------------===//
//
-// This pass HPVM intrinsics from HPVM IR. This pass is the final pass that
+// This pass HPVM intrinsics from HPVM IR. This pass is the final pass that
// runs as a part of clean up after construction of dataflowgraph and LLVM
// code generation for different targets from the dataflow graph.
//
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 18b851c9461bcebb67d941ae6f8a8fba56fcc9f7..6c3dcd75f6bc5f81b2e834d89ae080c731df210e 100644
--- a/hpvm/lib/Transforms/GenHPVM/GenHPVM.cpp
+++ b/hpvm/lib/Transforms/GenHPVM/GenHPVM.cpp
@@ -8,7 +8,7 @@
//===----------------------------------------------------------------------===//
//
// This pass takes LLVM IR with HPVM-C functions to generate textual representa-
-// -tion for HPVM IR consisting of HPVM intrinsics. Memory-to-register
+// -tion for HPVM IR consisting of HPVM intrinsics. Memory-to-register
// optimization pass is expected to execute prior to execution of this pass.
//
//===----------------------------------------------------------------------===//
@@ -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/lib/Transforms/LocalMem/LocalMem.cpp b/hpvm/lib/Transforms/LocalMem/LocalMem.cpp
index 5cc134874c5c21da36fe290b30b49a26cc9e271f..c0823274266feba5c78df3cead41c86e80b8a542 100644
--- a/hpvm/lib/Transforms/LocalMem/LocalMem.cpp
+++ b/hpvm/lib/Transforms/LocalMem/LocalMem.cpp
@@ -7,8 +7,8 @@
//
//===----------------------------------------------------------------------===//
//
-// This pass traverses the dataflow graph to recognize the allocation nodes which
-// allocate scratch memory. This pass does not make changes to the textual
+// This pass traverses the dataflow graph to recognize the allocation nodes
+// which allocate scratch memory. This pass does not make changes to the textual
// representation of HPVM IR.
//
//===----------------------------------------------------------------------===//
diff --git a/hpvm/projects/hpvm-rt/hpvm-rt.cpp b/hpvm/projects/hpvm-rt/hpvm-rt.cpp
index b6273ec2cca712469269f68f538ce437e9b062ec..56a84a2711c0ac70c46be56f5c50cb2d8a07766f 100644
--- a/hpvm/projects/hpvm-rt/hpvm-rt.cpp
+++ b/hpvm/projects/hpvm-rt/hpvm-rt.cpp
@@ -31,7 +31,6 @@ typedef struct {
std::vector<pthread_t> *threads;
// Map from InputPort to Size
std::map<unsigned, uint64_t> *ArgInPortSizeMap;
- // std::vector<uint64_t>* BindInSizes;
std::vector<unsigned> *BindInSourcePort;
std::vector<uint64_t> *BindOutSizes;
std::vector<uint64_t> *EdgeSizes;
@@ -326,18 +325,14 @@ static void *llvm_hpvm_ocl_request_mem(void *ptr, size_t size,
clFlags = CL_MEM_READ_ONLY;
hpvm_SwitchToTimer(&kernel_timer, hpvm_TimerID_COPY);
- // pthread_mutex_lock(&ocl_mtx);
cl_mem d_input =
clCreateBuffer(Context->clOCLContext, clFlags, size, NULL, &errcode);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to allocate memory on device");
DEBUG(cout << "\nMemory allocated on device: " << d_input << flush << "\n");
if (isInput) {
DEBUG(cout << "\tCopying ...");
- // pthread_mutex_lock(&ocl_mtx);
errcode = clEnqueueWriteBuffer(Context->clCommandQue, d_input, CL_TRUE, 0,
size, MTE->getAddress(), 0, NULL, NULL);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to copy memory to device");
}
@@ -435,30 +430,14 @@ static void insert_marker(struct hpvm_TimerSet *tset, enum hpvm_TimerID timer) {
*new_event = (struct hpvm_async_time_marker_list *)malloc(
sizeof(struct hpvm_async_time_marker_list));
(*new_event)->marker = calloc(1, sizeof(cl_event));
- /*
- // I don't think this is needed at all. I believe clEnqueueMarker 'creates'
-the event #if ( __OPENCL_VERSION__ >= CL_VERSION_1_1 ) fprintf(stderr, "Creating
-Marker [%d]\n", timer);
- *((cl_event *)((*new_event)->marker)) = clCreateUserEvent(*clContextPtr,
-&ciErrNum); if (ciErrNum != CL_SUCCESS) { fprintf(stderr, "Error Creating User
-Event Object!\n");
- }
- ciErrNum = clSetUserEventStatus(*((cl_event *)((*new_event)->marker)),
-CL_QUEUED); if (ciErrNum != CL_SUCCESS) { fprintf(stderr, "Error Setting User
-Event Status!\n");
- }
-#endif
-*/
(*new_event)->next = NULL;
}
/* valid event handle now aquired: insert the event record */
(*new_event)->label = NULL;
(*new_event)->timerID = timer;
- // pthread_mutex_lock(&ocl_mtx);
ciErrNum =
clEnqueueMarker(globalCommandQue, (cl_event *)(*new_event)->marker);
- // pthread_mutex_unlock(&ocl_mtx);
if (ciErrNum != CL_SUCCESS) {
fprintf(stderr, "Error Enqueueing Marker!\n");
}
@@ -477,29 +456,14 @@ static void insert_submarker(struct hpvm_TimerSet *tset, char *label,
*new_event = (struct hpvm_async_time_marker_list *)malloc(
sizeof(struct hpvm_async_time_marker_list));
(*new_event)->marker = calloc(1, sizeof(cl_event));
- /*
-#if ( __OPENCL_VERSION__ >= CL_VERSION_1_1 )
-fprintf(stderr, "Creating SubMarker %s[%d]\n", label, timer);
- *((cl_event *)((*new_event)->marker)) = clCreateUserEvent(*clContextPtr,
-&ciErrNum); if (ciErrNum != CL_SUCCESS) { fprintf(stderr, "Error Creating User
-Event Object!\n");
- }
- ciErrNum = clSetUserEventStatus(*((cl_event *)((*new_event)->marker)),
-CL_QUEUED); if (ciErrNum != CL_SUCCESS) { fprintf(stderr, "Error Setting User
-Event Status!\n");
- }
-#endif
-*/
(*new_event)->next = NULL;
}
/* valid event handle now aquired: insert the event record */
(*new_event)->label = label;
(*new_event)->timerID = timer;
- // pthread_mutex_lock(&ocl_mtx);
ciErrNum =
clEnqueueMarker(globalCommandQue, (cl_event *)(*new_event)->marker);
- // pthread_mutex_unlock(&ocl_mtx);
if (ciErrNum != CL_SUCCESS) {
fprintf(stderr, "Error Enqueueing Marker!\n");
}
@@ -617,7 +581,6 @@ void hpvm_StartTimerAndSubTimer(struct hpvm_Timer *timer,
numNotStopped &= 0x2; // Zero out 2^0
}
if (numNotStopped == 0x0) {
- // fputs("Ignoring attempt to start running timer and subtimer\n", stderr);
return;
}
@@ -681,7 +644,6 @@ void hpvm_StopTimerAndSubTimer(struct hpvm_Timer *timer,
numNotRunning &= 0x2; // Zero out 2^0
}
if (numNotRunning == 0x0) {
- // fputs("Ignoring attempt to stop stopped timer and subtimer\n", stderr);
return;
}
@@ -1127,7 +1089,6 @@ void hpvm_DestroyTimerSet(struct hpvm_TimerSet *timers) {
free(event);
- // (*event) = NULL;
event = next;
}
@@ -1157,7 +1118,6 @@ void *llvm_hpvm_streamLaunch(void (*LaunchFunc)(void *, void *), void *args) {
Context->threads = new std::vector<pthread_t>();
Context->ArgInPortSizeMap = new std::map<unsigned, uint64_t>();
- // Context->BindInSizes = new std::vector<uint64_t>();
Context->BindInSourcePort = new std::vector<unsigned>();
Context->BindOutSizes = new std::vector<uint64_t>();
Context->EdgeSizes = new std::vector<uint64_t>();
@@ -1185,7 +1145,6 @@ void llvm_hpvm_streamPush(void *graphID, void *args) {
for (unsigned j = 0; j < Ctx->BindInputBuffers->size(); j++) {
if (Ctx->BindInSourcePort->at(j) == i) {
// Push to all bind buffers connected to parent node at this port
- // DEBUG(cout << "\tPushing Value " << element << " to buffer\n");
llvm_hpvm_bufferPush(Ctx->BindInputBuffers->at(j), element);
}
}
@@ -1206,7 +1165,6 @@ void *llvm_hpvm_streamPop(void *graphID) {
unsigned offset = 0;
for (unsigned i = 0; i < Ctx->BindOutputBuffers->size(); i++) {
uint64_t element = llvm_hpvm_bufferPop(Ctx->BindOutputBuffers->at(i));
- // DEBUG(cout << "\tPopped Value " << element << " from buffer\n");
memcpy((char *)output + offset, &element, Ctx->BindOutSizes->at(i));
offset += Ctx->BindOutSizes->at(i);
}
@@ -1220,7 +1178,6 @@ void llvm_hpvm_streamWait(void *graphID) {
// Push garbage to all other input buffers
for (unsigned i = 0; i < Ctx->BindInputBuffers->size(); i++) {
uint64_t element = 0;
- // DEBUG(cout << "\tPushing Value " << element << " to buffer\n");
llvm_hpvm_bufferPush(Ctx->BindInputBuffers->at(i), element);
}
// Push 1 in isLastInput buffers of all child nodes
@@ -1250,7 +1207,6 @@ void *llvm_hpvm_createBindOutBuffer(void *graphID, uint64_t size) {
DEBUG(cout << "Create BindOutBuffer -- Graph: " << graphID
<< ", Size: " << size << flush << "\n");
DFNodeContext_CPU *Context = (DFNodeContext_CPU *)graphID;
- // Twine name = Twine("Bind.Out.")+Twine(Context->BindOutputBuffers->size());
CircularBuffer<uint64_t> *bufferID =
new CircularBuffer<uint64_t>(BUFFER_SIZE, "BindOut");
DEBUG(cout << "\tNew Buffer: " << bufferID << flush << "\n");
@@ -1262,7 +1218,6 @@ void *llvm_hpvm_createEdgeBuffer(void *graphID, uint64_t size) {
DEBUG(cout << "Create EdgeBuffer -- Graph: " << graphID << ", Size: " << size
<< flush << "\n");
DFNodeContext_CPU *Context = (DFNodeContext_CPU *)graphID;
- // Twine name = Twine("Edge.")+Twine(Context->EdgeBuffers->size());
CircularBuffer<uint64_t> *bufferID =
new CircularBuffer<uint64_t>(BUFFER_SIZE, "Edge");
DEBUG(cout << "\tNew Buffer: " << bufferID << flush << "\n");
@@ -1275,7 +1230,6 @@ void *llvm_hpvm_createLastInputBuffer(void *graphID, uint64_t size) {
DEBUG(cout << "Create isLastInputBuffer -- Graph: " << graphID
<< ", Size: " << size << flush << "\n");
DFNodeContext_CPU *Context = (DFNodeContext_CPU *)graphID;
- // Twine name = Twine("isLastInput.")+Twine(Context->EdgeBuffers->size());
CircularBuffer<uint64_t> *bufferID =
new CircularBuffer<uint64_t>(BUFFER_SIZE, "LastInput");
DEBUG(cout << "\tNew Buffer: " << bufferID << flush << "\n");
@@ -1346,8 +1300,6 @@ void *llvm_hpvm_cpu_launch(void *(*rootFunc)(void *), void *arguments) {
void llvm_hpvm_cpu_wait(void *graphID) {
DEBUG(cout << "Waiting for pthread to finish ...\n");
- // DFNodeContext_CPU* Context = (DFNodeContext_CPU*) graphID;
- // pthread_join(Context->threadID, NULL);
free(graphID);
DEBUG(cout << "\t... pthread Done!\n");
}
@@ -1500,9 +1452,6 @@ void *llvm_hpvm_ocl_initContext(enum hpvm::Target T) {
cl_uint numDevices;
clCreateSubDevices(clDevices[0], props, num_entries, subdevice_id,
&numDevices);
- // printf("Num of devices = %d\n", numDevices);
- // for(unsigned i =0 ; i< numDevices; i++)
- // printf("Subdevice id %d = %p\n", i, subdevice_id[i]);
clDevices[0] = subdevice_id[0];
globalOCLContext =
clCreateContext(properties, 1, clDevices, NULL, NULL, &errcode);
@@ -1527,10 +1476,6 @@ void llvm_hpvm_ocl_clearContext(void *graphID) {
// FIXME: Have separate function to release command queue and clear context.
// Would be useful when a context has multiple command queues
clReleaseKernel(Context->clKernel);
- // clReleaseProgram(Context->clProgram);
- // clReleaseCommandQueue(Context->clCommandQue);
- // clReleaseContext(globalOCLContext);
- // DEBUG(cout << "Released context at: " << globalOCLContext);
free(Context);
DEBUG(cout << "Done with OCL kernel\n");
cout << "Printing HPVM Timer: KernelTimer\n";
@@ -1546,9 +1491,7 @@ void llvm_hpvm_ocl_argument_shared(void *graphID, int arg_index, size_t size) {
DFNodeContext_OCL *Context = (DFNodeContext_OCL *)graphID;
DEBUG(cout << "Using Context: " << Context << flush << "\n");
DEBUG(cout << "Using clKernel: " << Context->clKernel << flush << "\n");
- // pthread_mutex_lock(&ocl_mtx);
cl_int errcode = clSetKernelArg(Context->clKernel, arg_index, size, NULL);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to set shared memory argument");
pthread_mutex_unlock(&ocl_mtx);
}
@@ -1562,9 +1505,7 @@ void llvm_hpvm_ocl_argument_scalar(void *graphID, void *input, int arg_index,
DFNodeContext_OCL *Context = (DFNodeContext_OCL *)graphID;
DEBUG(cout << "Using Context: " << Context << flush << "\n");
DEBUG(cout << "Using clKernel: " << Context->clKernel << flush << "\n");
- // pthread_mutex_lock(&ocl_mtx);
cl_int errcode = clSetKernelArg(Context->clKernel, arg_index, size, input);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to set constant input argument");
pthread_mutex_unlock(&ocl_mtx);
}
@@ -1588,10 +1529,8 @@ void *llvm_hpvm_ocl_argument_ptr(void *graphID, void *input, int arg_index,
pthread_mutex_lock(&ocl_mtx);
// Set Kernel Argument
- // pthread_mutex_lock(&ocl_mtx);
cl_int errcode = clSetKernelArg(Context->clKernel, arg_index, sizeof(cl_mem),
(void *)&d_input);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to set pointer argument");
DEBUG(cout << "\tDevicePtr = " << d_input << flush << "\n");
pthread_mutex_unlock(&ocl_mtx);
@@ -1605,15 +1544,11 @@ void *llvm_hpvm_ocl_output_ptr(void *graphID, int arg_index, size_t size) {
<< flush << "\n");
DFNodeContext_OCL *Context = (DFNodeContext_OCL *)graphID;
cl_int errcode;
- // pthread_mutex_lock(&ocl_mtx);
cl_mem d_output = clCreateBuffer(Context->clOCLContext, CL_MEM_WRITE_ONLY,
size, NULL, &errcode);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to create output buffer on device");
- // pthread_mutex_lock(&ocl_mtx);
errcode = clSetKernelArg(Context->clKernel, arg_index, sizeof(cl_mem),
(void *)&d_output);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to set pointer argument");
DEBUG(cout << "\tDevicePtr = " << d_output << flush << "\n");
pthread_mutex_unlock(&ocl_mtx);
@@ -1621,9 +1556,6 @@ void *llvm_hpvm_ocl_output_ptr(void *graphID, int arg_index, size_t size) {
}
void llvm_hpvm_ocl_free(void *ptr) {
- // DEBUG(cout << "Release Device Pointer: " << ptr << flush << "\n");
- // cl_mem d_ptr = (cl_mem) ptr;
- // clReleaseMemObject(d_ptr);
}
void *llvm_hpvm_ocl_getOutput(void *graphID, void *h_output, void *d_output,
@@ -1635,11 +1567,9 @@ void *llvm_hpvm_ocl_getOutput(void *graphID, void *h_output, void *d_output,
if (h_output == NULL)
h_output = malloc(size);
DFNodeContext_OCL *Context = (DFNodeContext_OCL *)graphID;
- // pthread_mutex_lock(&ocl_mtx);
cl_int errcode =
clEnqueueReadBuffer(Context->clCommandQue, (cl_mem)d_output, CL_TRUE, 0,
size, h_output, 0, NULL, NULL);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "[getOutput] Failure to read output");
pthread_mutex_unlock(&ocl_mtx);
return h_output;
@@ -1671,7 +1601,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");
@@ -1688,26 +1617,17 @@ void *llvm_hpvm_ocl_executeNode(void *graphID, unsigned workDim,
}
DEBUG(cout << ")\n");
}
- // pthread_mutex_lock(&ocl_mtx);
clFinish(Context->clCommandQue);
- // pthread_mutex_unlock(&ocl_mtx);
hpvm_SwitchToTimer(&kernel_timer, hpvm_TimerID_COMPUTATION);
- // for(int i=0 ;i < NUM_TESTS; i++) {
- // cout << "Iteration = " << i << flush << "\n";
- // pthread_mutex_lock(&ocl_mtx);
cl_int errcode = clEnqueueNDRangeKernel(
Context->clCommandQue, Context->clKernel, workDim, NULL, GlobalWG,
(localWorkSize == NULL) ? NULL : LocalWG, 0, NULL, NULL);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to enqueue kernel");
- //}
- // pthread_mutex_lock(&ocl_mtx);
clFinish(Context->clCommandQue);
- // pthread_mutex_unlock(&ocl_mtx);
hpvm_SwitchToTimer(&kernel_timer, hpvm_TimerID_NONE);
pthread_mutex_unlock(&ocl_mtx);
- return event;
+ return NULL;
}
//////////////////////////////////////////////////////////////////////////////
@@ -1769,11 +1689,9 @@ void *llvm_hpvm_ocl_launch(const char *FileName, const char *KernelName) {
Context->clOCLContext = globalOCLContext;
// Create a command-queue
- // pthread_mutex_lock(&ocl_mtx);
Context->clCommandQue = clCreateCommandQueue(
Context->clOCLContext, clDevices[0], CL_QUEUE_PROFILING_ENABLE, &errcode);
globalCommandQue = Context->clCommandQue;
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to create command queue");
DEBUG(cout << "Loading program binary: " << FileName << flush << "\n");
@@ -1781,11 +1699,8 @@ 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);
- // pthread_mutex_unlock(&ocl_mtx);
checkErr(errcode, CL_SUCCESS, "Failure to create program from binary");
DEBUG(cout << "Building kernel - " << KernelName << " from file " << FileName
@@ -1816,7 +1731,6 @@ void *llvm_hpvm_ocl_launch(const char *FileName, const char *KernelName) {
checkErr(errcode, CL_SUCCESS, "Failure to create kernel");
DEBUG(cout << "Kernel ID = " << Context->clKernel << "\n");
- // free(clDevices);
free(programSource);
pthread_mutex_unlock(&ocl_mtx);
@@ -1827,16 +1741,12 @@ void llvm_hpvm_ocl_wait(void *graphID) {
pthread_mutex_lock(&ocl_mtx);
DEBUG(cout << "Wait\n");
DFNodeContext_OCL *Context = (DFNodeContext_OCL *)graphID;
- // pthread_mutex_lock(&ocl_mtx);
clFinish(Context->clCommandQue);
- // pthread_mutex_unlock(&ocl_mtx);
pthread_mutex_unlock(&ocl_mtx);
}
void llvm_hpvm_switchToTimer(void **timerSet, enum hpvm_TimerID timer) {
- // cout << "Switching to timer " << timer << flush << "\n";
pthread_mutex_lock(&ocl_mtx);
- // hpvm_SwitchToTimer((hpvm_TimerSet*)(*timerSet), timer);
pthread_mutex_unlock(&ocl_mtx);
}
void llvm_hpvm_printTimerSet(void **timerSet, char *timerName) {
diff --git a/hpvm/projects/hpvm-rt/hpvm-rt.h b/hpvm/projects/hpvm-rt/hpvm-rt.h
index 94fe5b5ef0d82aca9f7556f7022aa513b9d2cc28..8f8b3eef69beaee37ec383fabcb03e82bf11d9ee 100644
--- a/hpvm/projects/hpvm-rt/hpvm-rt.h
+++ b/hpvm/projects/hpvm-rt/hpvm-rt.h
@@ -11,7 +11,6 @@
#include <pthread.h>
#include <string>
#include <vector>
-//#include <condition_variable>
#include "../../include/SupportHPVM/HPVMHint.h"
#include "../../include/SupportHPVM/HPVMTimer.h"
@@ -206,43 +205,27 @@ public:
template <class ElementType>
bool CircularBuffer<ElementType>::push(ElementType E) {
- // DEBUG(cout << name << " Buffer[" << ID << "]: Push " << E << flush <<
- // "\n"); unique_lock<mutex> lk(mtx);
pthread_mutex_lock(&mtx);
if ((Head + 1) % bufferSize == Tail) {
- // DEBUG(cout << name << " Buffer[" << ID << "]: Push going to sleep
- // ...\n"); cv.wait(lk);
pthread_cond_wait(&cv, &mtx);
- // DEBUG(cout << name << " Buffer[" << ID << "]: Push woke up\n");
}
buffer[Head] = E;
Head = (Head + 1) % bufferSize;
numElements++;
- // DEBUG(cout << name << " Buffer[" << ID << "]: Total Elements = " <<
- // numElements << flush << "\n"); lk.unlock();
pthread_mutex_unlock(&mtx);
- // cv.notify_one();
pthread_cond_signal(&cv);
return true;
}
template <class ElementType> ElementType CircularBuffer<ElementType>::pop() {
- // unique_lock<mutex> lk(mtx);
- // DEBUG(cout << name << " Buffer[" << ID << "]: Pop\n");
pthread_mutex_lock(&mtx);
if (Tail == Head) {
- // DEBUG(cout << name << " Buffer[" << ID << "]: Pop going to sleep ...\n");
- // cv.wait(lk);
pthread_cond_wait(&cv, &mtx);
- // DEBUG(cout << name << " Buffer[" << ID << "]: Pop woke up\n");
}
ElementType E = buffer[Tail];
Tail = (Tail + 1) % bufferSize;
numElements--;
- // DEBUG(cout << name << " Buffer[" << ID << "]: Total Elements = " <<
- // numElements << flush << "\n"); lk.unlock();
pthread_mutex_unlock(&mtx);
- // cv.notify_one();
pthread_cond_signal(&cv);
return E;
}
diff --git a/hpvm/projects/hpvm-rt/policy.h b/hpvm/projects/hpvm-rt/policy.h
index d50e65868b376bfbcc3d4bd00d4919db677722b8..78aacfc94a87c4855e67997fbdadbf10621cbd30 100644
--- a/hpvm/projects/hpvm-rt/policy.h
+++ b/hpvm/projects/hpvm-rt/policy.h
@@ -24,8 +24,6 @@ private:
class NodePolicy : public Policy {
virtual int getVersion(const char *name, int64_t it) override {
std::string s(name);
- // std::string NodeNames[1] = {
- // "_Z9mysgemmNTPfiS_iS_iiff_clonedInternal_level2_cloned" };
std::string NodeNames[] = {
"WrapperGaussianSmoothing_cloned",
"WrapperlaplacianEstimate_cloned",
@@ -34,10 +32,6 @@ class NodePolicy : public Policy {
"WrapperComputeMaxGradient_cloned",
"WrapperRejectZeroCrossings_cloned",
};
- // if (!s.compare(NodeNames[4])) {
- // std::cout << s << ": CPU" << "\n";
- // return 0;
- //}
return 2;
}
};
@@ -54,17 +48,14 @@ class IterationPolicy : public Policy {
class DeviceStatusPolicy : public Policy {
virtual int getVersion(const char *name, int64_t it) override {
if (deviceStatus) {
- // std::cout << "Returning GPU\n";
return 2;
} else {
- // std::cout << "Returning CPU\n";
return 0;
}
}
};
/* ------------------------------------------------------------------------- */
-// Added for the CFAR interactive policy demo.
class InteractivePolicy : public Policy {
private:
diff --git a/hpvm/scripts/automate_tests.sh b/hpvm/scripts/automated_tests.sh
similarity index 91%
rename from hpvm/scripts/automate_tests.sh
rename to hpvm/scripts/automated_tests.sh
index 517f66ad3e1f55aa18e1cf6500743be3198b5507..aece74eaa0a59e2e91807e9444c5c10054126cdd 100644
--- a/hpvm/scripts/automate_tests.sh
+++ b/hpvm/scripts/automated_tests.sh
@@ -7,6 +7,7 @@ HPVM_RT=hpvm-rt/hpvm-rt.bc
if [ -f $BUILD_DIR/tools/hpvm/projects/$HPVM_RT ]; then
true
else
+ echo $BUILD_DIR/tools/hpvm/projects/$HPVM_RT
echo HPVM not installed! Exiting without running tests!.
exit 0
fi
diff --git a/hpvm/scripts/llvm_installer.sh b/hpvm/scripts/llvm_installer.sh
index 2c6f6ebe43ae57585d6f7d5d64337bef68a65c5e..9c71f865c0799c3f380c21c0c45a31b8ba53dcb3 100755
--- a/hpvm/scripts/llvm_installer.sh
+++ b/hpvm/scripts/llvm_installer.sh
@@ -16,8 +16,6 @@ NUM_THREADS=2
SUFFIX=".tar.xz"
CLANG_SRC="cfe-$VERSION.src"
LLVM_SRC="llvm-$VERSION.src"
-LIBCXX_SRC="libcxx-$VERSION.src"
-LIBCXXABI_SRC="libcxxabi-$VERSION.src"
HPVM_RT=hpvm-rt/hpvm-rt.bc
@@ -42,9 +40,9 @@ fi
if [ -d $LLVM_SRC ]; then
- echo Found $LLVM_SRC!
+ echo Found $LLVM_SRC, not dowloading it again!
elif [ -d llvm ]; then
- echo Found llvm, not downloading $LLVM_SRC!
+ echo Found LLVM, not downloading it again!
else
echo $WGET $URL/$VERSION/$LLVM_SRC$SUFFIX
$WGET $URL/$VERSION/$LLVM_SRC$SUFFIX
@@ -58,7 +56,7 @@ if [ -d $LLVM_SRC ]; then
elif [ -d llvm ]; then
echo Everything looks sane.
else
- echo Install had problems. Quitting.
+ echo Problem with LLVM download. Exiting!
exit
fi
@@ -68,7 +66,7 @@ if [ -d $CURRENT_DIR/$LLVM_SRC/tools ]; then
cd $CURRENT_DIR/$LLVM_SRC/tools
echo In tools.
else
- echo Fail! Something is wrong with your $LLVM_SRC checkout!
+ echo Something is wrong with LLVM checkout. Exiting!
exit 1
fi
@@ -82,52 +80,13 @@ else
if [ -d clang ]; then
echo Everything looks sane.
else
- echo Install had problems. Quitting.
+ echo Problem with clang download. Exiting!
exit
fi
fi
cd $CURRENT_DIR
-if [ -d $CURRENT_DIR/$LLVM_SRC/projects ]; then
- cd $CURRENT_DIR/$LLVM_SRC/projects
-else
- echo Fail! Something is wrong with $LLVM_SRC.
- exit 1
-fi
-
-if [ -d libcxx ]; then
- echo Found libcxx! Not downloading libcxx again.
-else
- $WGET $URL/$VERSION/$LIBCXX_SRC$SUFFIX
- tar xf $LIBCXX_SRC$SUFFIX
- rm $LIBCXX_SRC$SUFFIX
- mv $LIBCXX_SRC libcxx
- if [ -d libcxx ]; then
- echo Everything looks sane.
- else
- echo Install had problems. Quitting.
- exit
- fi
-fi
-
-if [ $LIBCXXABI_SRC != "" ]; then
- if [ -d libcxxabi ]; then
- echo Found libcxxabi! Not downloading libcxx again.
- else
- $WGET $URL/$VERSION/$LIBCXXABI_SRC$SUFFIX
- tar xf $LIBCXXABI_SRC$SUFFIX
- rm $LIBCXXABI_SRC$SUFFIX
- mv $LIBCXXABI_SRC libcxxabi
- if [ -d libcxxabi ]; then
- echo Everything looks sane.
- else
- echo Install had problems. Quitting.
- exit
- fi
- fi
-fi
-
HPVM_DIR=$CURRENT_DIR/$LLVM_SRC/tools/hpvm
if [ ! -d $HPVM_DIR ]; then
@@ -138,6 +97,7 @@ if [ ! -d $HPVM_DIR ]; then
ln -s $CURRENT_DIR/lib $HPVM_DIR/
ln -s $CURRENT_DIR/projects $HPVM_DIR/
ln -s $CURRENT_DIR/test $HPVM_DIR/
+ ln -s $CURRENT_DIR/tools $HPVM_DIR/
else
echo $CURRENT_DIR/$LLVM_SRC/tools/hpvm exists.
fi
@@ -185,6 +145,7 @@ make -j$NUM_THREADS
if [ -f $BUILD_DIR/tools/hpvm/projects/$HPVM_RT ]; then
true
else
+ echo $BUILD_DIR/tools/hpvm/projects/$HPVM_RT
echo HPVM not installed properly.
exit 0
fi
diff --git a/hpvm/test/unitTests/ThreeLevel.ll b/hpvm/test/unitTests/ThreeLevel.ll
index d8bf050234264e55be6af269e40ab5f2ef36a03b..460dd15b6b1f6dd38483a18e899f0d96b68cac08 100644
--- a/hpvm/test/unitTests/ThreeLevel.ll
+++ b/hpvm/test/unitTests/ThreeLevel.ll
@@ -1,70 +1,67 @@
-; RUN: opt - load LLVMGenHPVM.so -S -genhpvm < %s
-; ModuleID = 'TwoLevel.c'
-source_filename = "TwoLevel.c"
+; RUN: opt -load LLVMBuildDFG.so -load LLVMDFG2LLVM_CPU.so -load LLVMClearDFG.so -S -dfg2llvm-cpu -clearDFG < %s | FileCheck %s
+; ModuleID = 'ThreeLevel.ll'
+source_filename = "ThreeLevel.c"
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
%struct.Root = type { i32*, i64, i32*, i64, i32*, i64 }
+%struct.out.Func1 = type <{ i32* }>
+%struct.out.Func3 = type <{ i32* }>
+%struct.out.Func2 = type <{ i32* }>
+%struct.out.PipeRoot = type <{ i32* }>
-; Function Attrs: nounwind uwtable
-define dso_local void @Func1(i32* %In, i64 %Insize, i32* %Out, i64 %Outsize) #0 {
-entry:
- tail call void @__hpvm__hint(i32 1) #3
- tail call void (i32, ...) @__hpvm__attributes(i32 1, i32* %In, i32* %Out, i32 1, i32* %Out) #3
- %0 = load i32, i32* %In, align 4, !tbaa !2
- store i32 %0, i32* %Out, align 4, !tbaa !2
- tail call void (i32, ...) @__hpvm__return(i32 1, i32* %Out) #3
- ret void
-}
-declare dso_local void @__hpvm__hint(i32) local_unnamed_addr #1
+; CHECK-LABEL: i32 @main(
+; CHECK-NOT: call void @llvm.hpvm.init()
+; CHECK: call i8* @llvm_hpvm_cpu_launch(i8* (i8*)* @LaunchDataflowGraph, i8*
+; CHECK-NOT: call i8* @llvm.hpvm.launch(i8*
+; CHECK: call void @llvm_hpvm_cpu_wait(i8*
-declare dso_local void @__hpvm__attributes(i32, ...) local_unnamed_addr #1
+; CHECK-LABEL: @Func1_cloned.1_cloned_cloned_cloned_cloned_cloned_cloned
+; CHECK: call i8* @llvm_hpvm_cpu_argument_ptr(
-declare dso_local void @__hpvm__return(i32, ...) local_unnamed_addr #1
+; CHECK-LABEL: @Func3_cloned.2_cloned_cloned_cloned_cloned_cloned_cloned(
+; CHECK-LABEL: for.body1:
+; CHECK: %index.y = phi i64 [ 0, %for.body ], [ %index.y.inc, %for.body1 ]
+; CHECK-NEXT: call void @llvm_hpvm_cpu_dstack_push(
+; CHECK-NEXT: @Func1_cloned.1_cloned_cloned_cloned_cloned_cloned_cloned(
+; CHECK-NEXT: call void @llvm_hpvm_cpu_dstack_pop()
-; Function Attrs: nounwind uwtable
-define dso_local void @Func2(i32* %In, i64 %Insize, i32* %Out, i64 %Outsize) #0 {
-entry:
- tail call void @__hpvm__hint(i32 1) #3
- tail call void (i32, ...) @__hpvm__attributes(i32 2, i32* %In, i32* %Out, i32 1, i32* %Out) #3
- %call = tail call i8* (i32, ...) @__hpvm__createNodeND(i32 1, void (i32*, i64, i32*, i64)* nonnull @Func1, i64 3) #3
- tail call void @__hpvm__bindIn(i8* %call, i32 0, i32 0, i32 0) #3
- tail call void @__hpvm__bindIn(i8* %call, i32 1, i32 1, i32 0) #3
- tail call void @__hpvm__bindIn(i8* %call, i32 2, i32 2, i32 0) #3
- tail call void @__hpvm__bindIn(i8* %call, i32 3, i32 3, i32 0) #3
- tail call void @__hpvm__bindOut(i8* %call, i32 0, i32 0, i32 0) #3
- ret void
-}
+; CHECK-LABEL: @Func2_cloned.3_cloned_cloned_cloned_cloned_cloned_cloned(
+; CHECK-LABEL: for.body:
+; CHECK-NEXT: %index.x = phi i64 [ 0, %entry ], [ %index.x.inc, %for.body ]
+; CHECK-NEXT: call void @llvm_hpvm_cpu_dstack_push(
+; CHECK-NEXT: @Func3_cloned.2_cloned_cloned_cloned_cloned_cloned_cloned(
+; CHECK-NEXT: call void @llvm_hpvm_cpu_dstack_pop()
-; Function Attrs: argmemonly nounwind
-declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #2
+; CHECK-LABEL: @PipeRoot_cloned.4(
+; CHECK: call void @llvm_hpvm_cpu_dstack_push(
+; CHECK-NEXT: @Func2_cloned.3_cloned_cloned_cloned_cloned_cloned_cloned(
+; CHECK-NEXT: call void @llvm_hpvm_cpu_dstack_pop()
+
+; CHECK-LABEL: @LaunchDataflowGraph(
+; CHECK: call %struct.out.PipeRoot @PipeRoot_cloned.4(
-declare dso_local i8* @__hpvm__createNodeND(i32, ...) local_unnamed_addr #1
+declare dso_local void @__hpvm__hint(i32) local_unnamed_addr #0
-declare dso_local void @__hpvm__bindIn(i8*, i32, i32, i32) local_unnamed_addr #1
+declare dso_local void @__hpvm__attributes(i32, ...) local_unnamed_addr #0
-declare dso_local void @__hpvm__bindOut(i8*, i32, i32, i32) local_unnamed_addr #1
+declare dso_local void @__hpvm__return(i32, ...) local_unnamed_addr #0
; Function Attrs: argmemonly nounwind
-declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #2
+declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
-; Function Attrs: nounwind uwtable
-define dso_local void @PipeRoot(i32* %In1, i64 %Insize1, i32* %In2, i64 %InSize2, i32* %Out, i64 %Outsize) #0 {
-entry:
- tail call void @__hpvm__hint(i32 1) #3
- tail call void (i32, ...) @__hpvm__attributes(i32 3, i32* %In1, i32* %In2, i32* %Out, i32 1, i32* %Out) #3
- %call = tail call i8* (i32, ...) @__hpvm__createNodeND(i32 0, void (i32*, i64, i32*, i64)* nonnull @Func2) #3
- tail call void @__hpvm__bindIn(i8* %call, i32 0, i32 0, i32 0) #3
- tail call void @__hpvm__bindIn(i8* %call, i32 1, i32 1, i32 0) #3
- tail call void @__hpvm__bindIn(i8* %call, i32 2, i32 2, i32 0) #3
- tail call void @__hpvm__bindIn(i8* %call, i32 3, i32 3, i32 0) #3
- tail call void @__hpvm__bindOut(i8* %call, i32 0, i32 0, i32 0) #3
- ret void
-}
+declare dso_local i8* @__hpvm__createNodeND(i32, ...) local_unnamed_addr #0
+
+declare dso_local void @__hpvm__bindIn(i8*, i32, i32, i32) local_unnamed_addr #0
+
+declare dso_local void @__hpvm__bindOut(i8*, i32, i32, i32) local_unnamed_addr #0
+
+; Function Attrs: argmemonly nounwind
+declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
; Function Attrs: nounwind uwtable
-define dso_local i32 @main() local_unnamed_addr #0 {
+define dso_local i32 @main() local_unnamed_addr #2 {
entry:
%In1 = alloca i32, align 4
%In2 = alloca i32, align 4
@@ -72,31 +69,32 @@ entry:
%RootArgs = alloca %struct.Root, align 8
%0 = bitcast i32* %In1 to i8*
call void @llvm.lifetime.start.p0i8(i64 4, i8* nonnull %0) #3
- store i32 1, i32* %In1, align 4, !tbaa !2
+ store i32 1, i32* %In1, align 4, !tbaa !6
%1 = bitcast i32* %In2 to i8*
call void @llvm.lifetime.start.p0i8(i64 4, i8* nonnull %1) #3
- store i32 2, i32* %In2, align 4, !tbaa !2
+ store i32 2, i32* %In2, align 4, !tbaa !6
%2 = bitcast i32* %Out to i8*
call void @llvm.lifetime.start.p0i8(i64 4, i8* nonnull %2) #3
- store i32 0, i32* %Out, align 4, !tbaa !2
+ store i32 0, i32* %Out, align 4, !tbaa !6
%3 = bitcast %struct.Root* %RootArgs to i8*
call void @llvm.lifetime.start.p0i8(i64 48, i8* nonnull %3) #3
%input1 = getelementptr inbounds %struct.Root, %struct.Root* %RootArgs, i64 0, i32 0
- store i32* %In1, i32** %input1, align 8, !tbaa !6
+ store i32* %In1, i32** %input1, align 8, !tbaa !10
%Insize1 = getelementptr inbounds %struct.Root, %struct.Root* %RootArgs, i64 0, i32 1
- store i64 32, i64* %Insize1, align 8, !tbaa !10
+ store i64 32, i64* %Insize1, align 8, !tbaa !14
%input2 = getelementptr inbounds %struct.Root, %struct.Root* %RootArgs, i64 0, i32 2
- store i32* %In2, i32** %input2, align 8, !tbaa !11
+ store i32* %In2, i32** %input2, align 8, !tbaa !15
%Insize2 = getelementptr inbounds %struct.Root, %struct.Root* %RootArgs, i64 0, i32 3
- store i64 32, i64* %Insize2, align 8, !tbaa !12
+ store i64 32, i64* %Insize2, align 8, !tbaa !16
%output = getelementptr inbounds %struct.Root, %struct.Root* %RootArgs, i64 0, i32 4
- store i32* %Out, i32** %output, align 8, !tbaa !13
+ store i32* %Out, i32** %output, align 8, !tbaa !17
%Outsize = getelementptr inbounds %struct.Root, %struct.Root* %RootArgs, i64 0, i32 5
- store i64 32, i64* %Outsize, align 8, !tbaa !14
- call void (...) @__hpvm__init() #3
- %call = call i8* (i32, ...) @__hpvm__launch(i32 0, void (i32*, i64, i32*, i64, i32*, i64)* nonnull @PipeRoot, %struct.Root* nonnull %RootArgs) #3
- call void @__hpvm__wait(i8* %call) #3
- call void (...) @__hpvm__cleanup() #3
+ store i64 32, i64* %Outsize, align 8, !tbaa !18
+ call void @llvm.hpvm.init()
+ %4 = bitcast %struct.Root* %RootArgs to i8*
+ %graphID = call i8* @llvm.hpvm.launch(i8* bitcast (%struct.out.PipeRoot (i32*, i64, i32*, i64, i32*, i64)* @PipeRoot_cloned to i8*), i8* %4, i1 false)
+ call void @llvm.hpvm.wait(i8* %graphID)
+ call void @llvm.hpvm.cleanup()
call void @llvm.lifetime.end.p0i8(i64 48, i8* nonnull %3) #3
call void @llvm.lifetime.end.p0i8(i64 4, i8* nonnull %2) #3
call void @llvm.lifetime.end.p0i8(i64 4, i8* nonnull %1) #3
@@ -104,34 +102,124 @@ entry:
ret i32 0
}
-declare dso_local void @__hpvm__init(...) local_unnamed_addr #1
+declare dso_local void @__hpvm__init(...) local_unnamed_addr #0
+
+declare dso_local i8* @__hpvm__launch(i32, ...) local_unnamed_addr #0
+
+declare dso_local void @__hpvm__wait(i8*) local_unnamed_addr #0
+
+declare dso_local void @__hpvm__cleanup(...) local_unnamed_addr #0
+
+declare i8* @llvm_hpvm_initializeTimerSet()
+
+declare void @llvm_hpvm_switchToTimer(i8**, i32)
+
+declare void @llvm_hpvm_printTimerSet(i8**, i8*)
+
+; Function Attrs: nounwind uwtable
+define dso_local %struct.out.Func1 @Func1_cloned(i32* in %In, i64 %Insize, i32* out %Out, i64 %Outsize) #2 {
+entry:
+ %returnStruct = insertvalue %struct.out.Func1 undef, i32* %Out, 0
+ ret %struct.out.Func1 %returnStruct
+}
+
+; Function Attrs: nounwind
+declare i8* @llvm.hpvm.createNode2D(i8*, i64, i64) #3
+
+; Function Attrs: nounwind
+declare void @llvm.hpvm.bind.input(i8*, i32, i32, i1) #3
+
+; Function Attrs: nounwind
+declare void @llvm.hpvm.bind.output(i8*, i32, i32, i1) #3
+
+; Function Attrs: nounwind uwtable
+define dso_local %struct.out.Func3 @Func3_cloned(i32* in %In, i64 %Insize, i32* out %Out, i64 %Outsize) #2 {
+; CHECK-NOT: @Func3_cloned
+entry:
+ %Func1_cloned.node = call i8* @llvm.hpvm.createNode2D(i8* bitcast (%struct.out.Func1 (i32*, i64, i32*, i64)* @Func1_cloned to i8*), i64 3, i64 5)
+ call void @llvm.hpvm.bind.input(i8* %Func1_cloned.node, i32 0, i32 0, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func1_cloned.node, i32 1, i32 1, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func1_cloned.node, i32 2, i32 2, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func1_cloned.node, i32 3, i32 3, i1 false)
+ call void @llvm.hpvm.bind.output(i8* %Func1_cloned.node, i32 0, i32 0, i1 false)
+ ret %struct.out.Func3 undef
+}
+
+; Function Attrs: nounwind
+declare i8* @llvm.hpvm.createNode1D(i8*, i64) #3
+
+; Function Attrs: nounwind uwtable
+define dso_local %struct.out.Func2 @Func2_cloned(i32* in %In, i64 %Insize, i32* out %Out, i64 %Outsize) #2 {
+; CHECK-NOT: @Func2_cloned
+entry:
+ %Func3_cloned.node = call i8* @llvm.hpvm.createNode1D(i8* bitcast (%struct.out.Func3 (i32*, i64, i32*, i64)* @Func3_cloned to i8*), i64 3)
+ call void @llvm.hpvm.bind.input(i8* %Func3_cloned.node, i32 0, i32 0, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func3_cloned.node, i32 1, i32 1, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func3_cloned.node, i32 2, i32 2, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func3_cloned.node, i32 3, i32 3, i1 false)
+ call void @llvm.hpvm.bind.output(i8* %Func3_cloned.node, i32 0, i32 0, i1 false)
+ ret %struct.out.Func2 undef
+}
+
+; Function Attrs: nounwind
+declare i8* @llvm.hpvm.createNode(i8*) #3
+
+; Function Attrs: nounwind uwtable
+define dso_local %struct.out.PipeRoot @PipeRoot_cloned(i32* in %In1, i64 %Insize1, i32* in %In2, i64 %InSize2, i32* out %Out, i64 %Outsize) #2 {
+; CHECK-NOT: @PipeRoot_cloned
+entry:
+ %Func2_cloned.node = call i8* @llvm.hpvm.createNode(i8* bitcast (%struct.out.Func2 (i32*, i64, i32*, i64)* @Func2_cloned to i8*))
+ call void @llvm.hpvm.bind.input(i8* %Func2_cloned.node, i32 0, i32 0, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func2_cloned.node, i32 1, i32 1, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func2_cloned.node, i32 2, i32 2, i1 false)
+ call void @llvm.hpvm.bind.input(i8* %Func2_cloned.node, i32 3, i32 3, i1 false)
+ call void @llvm.hpvm.bind.output(i8* %Func2_cloned.node, i32 0, i32 0, i1 false)
+ ret %struct.out.PipeRoot undef
+}
+
+; Function Attrs: nounwind
+declare void @llvm.hpvm.init() #3
-declare dso_local i8* @__hpvm__launch(i32, ...) local_unnamed_addr #1
+; Function Attrs: nounwind
+declare i8* @llvm.hpvm.launch(i8*, i8*, i1) #3
-declare dso_local void @__hpvm__wait(i8*) local_unnamed_addr #1
+; Function Attrs: nounwind
+declare void @llvm.hpvm.wait(i8*) #3
-declare dso_local void @__hpvm__cleanup(...) local_unnamed_addr #1
+; Function Attrs: nounwind
+declare void @llvm.hpvm.cleanup() #3
-attributes #0 = { nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="false" "no-infs-fp-math"="true" "no-jump-tables"="false" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="true" "use-soft-float"="false" }
-attributes #1 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="true" "use-soft-float"="false" }
-attributes #2 = { argmemonly nounwind }
+attributes #0 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="cpu-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="true" "use-soft-float"="false" }
+attributes #1 = { argmemonly nounwind }
+attributes #2 = { nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="false" "no-infs-fp-math"="true" "no-jump-tables"="false" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="cpu-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="true" "use-soft-float"="false" }
attributes #3 = { nounwind }
!llvm.module.flags = !{!0}
!llvm.ident = !{!1}
+!hpvm_hint_cpu = !{!2, !3, !4, !5}
+!hpvm_hint_gpu = !{}
+!hpvm_hint_spir = !{}
+!hpvm_hint_cudnn = !{}
+!hpvm_hint_promise = !{}
+!hpvm_hint_cpu_gpu = !{}
+!hpvm_hint_cpu_spir = !{}
!0 = !{i32 1, !"wchar_size", i32 4}
-!1 = !{!"clang version 9.0.0 (https://gitlab.engr.illinois.edu/llvm/hpvm.git 3551132592a00cab6c966df508ab511598269f78)"}
-!2 = !{!3, !3, i64 0}
-!3 = !{!"int", !4, i64 0}
-!4 = !{!"omnipotent char", !5, i64 0}
-!5 = !{!"Simple C/C++ TBAA"}
-!6 = !{!7, !8, i64 0}
-!7 = !{!"Root", !8, i64 0, !9, i64 8, !8, i64 16, !9, i64 24, !8, i64 32, !9, i64 40}
-!8 = !{!"any pointer", !4, i64 0}
-!9 = !{!"long", !4, i64 0}
-!10 = !{!7, !9, i64 8}
-!11 = !{!7, !8, i64 16}
-!12 = !{!7, !9, i64 24}
-!13 = !{!7, !8, i64 32}
-!14 = !{!7, !9, i64 40}
+!1 = !{!"clang version 9.0.0 (https://gitlab.engr.illinois.edu/llvm/hpvm.git 6690f9e7e8b46b96aea222d3e85315cd63545953)"}
+!2 = !{%struct.out.Func1 (i32*, i64, i32*, i64)* @Func1_cloned}
+!3 = !{%struct.out.Func3 (i32*, i64, i32*, i64)* @Func3_cloned}
+!4 = !{%struct.out.Func2 (i32*, i64, i32*, i64)* @Func2_cloned}
+!5 = !{%struct.out.PipeRoot (i32*, i64, i32*, i64, i32*, i64)* @PipeRoot_cloned}
+!6 = !{!7, !7, i64 0}
+!7 = !{!"int", !8, i64 0}
+!8 = !{!"omnipotent char", !9, i64 0}
+!9 = !{!"Simple C/C++ TBAA"}
+!10 = !{!11, !12, i64 0}
+!11 = !{!"Root", !12, i64 0, !13, i64 8, !12, i64 16, !13, i64 24, !12, i64 32, !13, i64 40}
+!12 = !{!"any pointer", !8, i64 0}
+!13 = !{!"long", !8, i64 0}
+!14 = !{!11, !13, i64 8}
+!15 = !{!11, !12, i64 16}
+!16 = !{!11, !13, i64 24}
+!17 = !{!11, !12, i64 32}
+!18 = !{!11, !13, i64 40}
diff --git a/hpvm/tools/CMakeLists.txt b/hpvm/tools/CMakeLists.txt
new file mode 100644
index 0000000000000000000000000000000000000000..c653aebd5047db4540ad6ff018b0377d0a4e82a4
--- /dev/null
+++ b/hpvm/tools/CMakeLists.txt
@@ -0,0 +1 @@
+add_llvm_tool_subdirectory(hpvm-config)
diff --git a/hpvm/tools/README.md b/hpvm/tools/README.md
new file mode 100644
index 0000000000000000000000000000000000000000..4609a86407cdc06b531af1b4cdf89609a4a58050
--- /dev/null
+++ b/hpvm/tools/README.md
@@ -0,0 +1,61 @@
+# hpvm-config -- Print HPVM compilation options
+
+### Synopsis
+hpvm-config option [components…]
+
+### Description
+hpvm-config print the compiler flags, linker flags and object libraries needed to link against HPVM.
+In addtion to printing flags printed by [llvm-config](http://llvm.org/docs/CommandGuide/llvm-config.html), hpvm-config also prints HPVM version and compiler
+flags required to compile HPVM programs.
+
+### Options
+–version: Print the version number of HPVM.
+
+-llvm-version: Print LLVM version.
+
+-hpvm-version: Print HPVM version.
+
+-help: Print a summary of llvm-config arguments.
+
+–prefix: Print the installation prefix for LLVM.
+
+–src-root: Print the source root from which LLVM was built.
+
+–obj-root: Print the object root used to build HPVM.
+
+–bindir: Print the installation directory for LLVM binaries.
+
+–includedir: Print the installation directory for LLVM headers.
+
+–libdir: Print the installation directory for LLVM libraries.
+
+–cxxflags: Print the C++ compiler flags needed to use LLVM headers.
+
+–ldflags: Print the flags needed to link against LLVM libraries.
+
+–libs: Print all the libraries needed to link against the specified LLVM components, including any dependencies.
+
+–libnames: Similar to –libs, but prints the bare filenames of the libraries without -l or pathnames. Useful for linking against a not-yet-installed copy of LLVM.
+
+–libfiles: Similar to –libs, but print the full path to each library file. This is useful when creating makefile dependencies, to ensure that a tool is relinked if any library it uses changes.
+
+–components: Print all valid component names.
+
+–targets-built: Print the component names for all targets supported by this copy of LLVM.
+
+–build-mode: Print the build mode used when LLVM was built (e.g. Debug or Release)
+
+-genHPVM: Generate HPVM textual IR from LLVM IR.
+
+-dfg2llvm-cpu: Generate code for CPU and host.
+
+-dfg2llvm-opencl: Generate kernel code for GPU target in OpenCL.
+
+-clearDFG: Clear dataflow graph for HPVM and extraneous HPVM-specific instructions from IR.
+
+### Exit Status
+If hpvm-config succeeds, it will exit with 0. Otherwise, if an error occurs, it will exit with a non-zero value.
+
+
+
+
diff --git a/hpvm/tools/hpvm-config/BuildVariables.inc.in b/hpvm/tools/hpvm-config/BuildVariables.inc.in
new file mode 100644
index 0000000000000000000000000000000000000000..3a24d3e974e1fb7c97335b8622181305b494aa87
--- /dev/null
+++ b/hpvm/tools/hpvm-config/BuildVariables.inc.in
@@ -0,0 +1,36 @@
+//===-- BuildVariables.inc.in - llvm-config build variables -*- C++ -*-----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file is configured by the build system to define the variables
+// llvm-config wants to report to the user, but which can only be determined at
+// build time.
+//
+// The variant of this file not ending with .in has been autogenerated by the
+// LLVM build. Do not edit!
+//
+//===----------------------------------------------------------------------===//
+
+#define LLVM_SRC_ROOT "@LLVM_SRC_ROOT@"
+#define LLVM_OBJ_ROOT "@LLVM_OBJ_ROOT@"
+#define LLVM_CPPFLAGS "@LLVM_CPPFLAGS@"
+#define LLVM_CFLAGS "@LLVM_CFLAGS@"
+#define LLVM_LDFLAGS "@LLVM_LDFLAGS@"
+#define LLVM_CXXFLAGS "@LLVM_CXXFLAGS@"
+#define LLVM_BUILDMODE "@LLVM_BUILDMODE@"
+#define LLVM_LIBDIR_SUFFIX "@LLVM_LIBDIR_SUFFIX@"
+#define LLVM_TARGETS_BUILT "@LLVM_TARGETS_BUILT@"
+#define LLVM_SYSTEM_LIBS "@LLVM_SYSTEM_LIBS@"
+#define LLVM_BUILD_SYSTEM "@LLVM_BUILD_SYSTEM@"
+#define LLVM_HAS_RTTI @LLVM_HAS_RTTI@
+#define LLVM_ENABLE_DYLIB @LLVM_BUILD_LLVM_DYLIB@
+#define LLVM_LINK_DYLIB @LLVM_LINK_LLVM_DYLIB@
+#define LLVM_ENABLE_SHARED @BUILD_SHARED_LIBS@
+#define LLVM_DYLIB_COMPONENTS "@LLVM_DYLIB_COMPONENTS@"
+#define LLVM_DYLIB_VERSION "@LLVM_DYLIB_VERSION@"
+#define LLVM_HAS_GLOBAL_ISEL @LLVM_HAS_GLOBAL_ISEL@
+#define LLVM_TOOLS_INSTALL_DIR "@LLVM_TOOLS_INSTALL_DIR@"
diff --git a/hpvm/tools/hpvm-config/CMakeLists.txt b/hpvm/tools/hpvm-config/CMakeLists.txt
new file mode 100644
index 0000000000000000000000000000000000000000..4b3730cc7ab07da1af569e80208fb9a3ddd60b46
--- /dev/null
+++ b/hpvm/tools/hpvm-config/CMakeLists.txt
@@ -0,0 +1,88 @@
+set(LLVM_LINK_COMPONENTS support)
+
+set(BUILDVARIABLES_SRCPATH ${CMAKE_CURRENT_SOURCE_DIR}/BuildVariables.inc.in)
+set(BUILDVARIABLES_OBJPATH ${CMAKE_CURRENT_BINARY_DIR}/BuildVariables.inc)
+
+# Add the hpvm-config tool.
+add_llvm_tool(hpvm-config
+ hpvm-config.cpp
+ )
+
+target_include_directories(
+ hpvm-config PRIVATE "${LLVM_BINARY_DIR}/tools/llvm-config"
+)
+
+
+# Compute the substitution values for various items.
+get_property(SUPPORT_SYSTEM_LIBS TARGET LLVMSupport PROPERTY LLVM_SYSTEM_LIBS)
+get_property(WINDOWSMANIFEST_SYSTEM_LIBS TARGET LLVMWindowsManifest PROPERTY LLVM_SYSTEM_LIBS)
+foreach(l ${SUPPORT_SYSTEM_LIBS} ${WINDOWSMANIFEST_SYSTEM_LIBS})
+ if(MSVC)
+ set(SYSTEM_LIBS ${SYSTEM_LIBS} "${l}.lib")
+ else()
+ if (l MATCHES "^-")
+ # If it's an option, pass it without changes.
+ set(SYSTEM_LIBS ${SYSTEM_LIBS} "${l}")
+ else()
+ # Otherwise assume it's a library name we need to link with.
+ set(SYSTEM_LIBS ${SYSTEM_LIBS} "-l${l}")
+ endif()
+ endif()
+endforeach()
+string(REPLACE ";" " " SYSTEM_LIBS "${SYSTEM_LIBS}")
+
+# Fetch target specific compile options, e.g. RTTI option
+get_property(COMPILE_FLAGS TARGET hpvm-config PROPERTY COMPILE_FLAGS)
+
+# The language standard potentially affects the ABI/API of LLVM, so we want
+# to make sure it is reported by hpvm-config.
+# NOTE: We don't want to start extracting any random C/CXX flags that the
+# user may add that could affect the ABI. We only want to extract flags
+# that have been added by the LLVM build system.
+string(REGEX MATCH "-std=[^ ]\+" LLVM_CXX_STD_FLAG ${CMAKE_CXX_FLAGS})
+string(REGEX MATCH "-stdlib=[^ ]\+" LLVM_CXX_STDLIB_FLAG ${CMAKE_CXX_FLAGS})
+string(REGEX MATCH "-std=[^ ]\+" LLVM_C_STD_FLAG ${CMAKE_C_FLAGS})
+
+# Use configure_file to create BuildVariables.inc.
+set(LLVM_SRC_ROOT ${LLVM_MAIN_SRC_DIR})
+set(LLVM_OBJ_ROOT ${LLVM_BINARY_DIR})
+set(LLVM_CPPFLAGS "${LLVM_DEFINITIONS}")
+set(LLVM_CFLAGS "${LLVM_C_STD_FLAG} ${LLVM_DEFINITIONS}")
+set(LLVM_CXXFLAGS "${LLVM_CXX_STD_FLAG} ${LLVM_CXX_STDLIB_FLAG} ${COMPILE_FLAGS} ${LLVM_DEFINITIONS}")
+set(LLVM_BUILD_SYSTEM cmake)
+set(LLVM_HAS_RTTI ${LLVM_CONFIG_HAS_RTTI})
+set(LLVM_DYLIB_VERSION "${LLVM_VERSION_MAJOR}${LLVM_VERSION_SUFFIX}")
+set(LLVM_HAS_GLOBAL_ISEL "ON")
+
+# Use the C++ link flags, since they should be a superset of C link flags.
+set(LLVM_LDFLAGS "${CMAKE_CXX_LINK_FLAGS}")
+set(LLVM_BUILDMODE ${CMAKE_BUILD_TYPE})
+set(LLVM_SYSTEM_LIBS ${SYSTEM_LIBS})
+string(REPLACE ";" " " LLVM_TARGETS_BUILT "${LLVM_TARGETS_TO_BUILD}")
+llvm_canonicalize_cmake_booleans(
+ LLVM_BUILD_LLVM_DYLIB
+ LLVM_LINK_LLVM_DYLIB
+ LLVM_HAS_RTTI
+ LLVM_HAS_GLOBAL_ISEL
+ BUILD_SHARED_LIBS)
+configure_file(${BUILDVARIABLES_SRCPATH} ${BUILDVARIABLES_OBJPATH} @ONLY)
+
+# Set build-time environment(s).
+add_definitions(-DCMAKE_CFG_INTDIR="${CMAKE_CFG_INTDIR}")
+
+if(LLVM_ENABLE_MODULES)
+ target_compile_options(hpvm-config PUBLIC
+ "-fmodules-ignore-macro=CMAKE_CFG_INTDIR"
+ )
+endif()
+
+# Add the dependency on the generation step.
+add_file_dependencies(${CMAKE_CURRENT_SOURCE_DIR}/hpvm-config.cpp ${BUILDVARIABLES_OBJPATH})
+
+if(CMAKE_CROSSCOMPILING AND NOT HPVM_CONFIG_PATH)
+ build_native_tool(hpvm-config HPVM_CONFIG_PATH)
+ set(HPVM_CONFIG_PATH "${LLVM_CONFIG_PATH}" CACHE STRING "")
+
+ add_custom_target(NativeHPVMConfig DEPENDS ${HPVM_CONFIG_PATH})
+ add_dependencies(hpvm-config NativeHPVMConfig)
+endif()
diff --git a/hpvm/tools/hpvm-config/LibraryDependencies.inc b/hpvm/tools/hpvm-config/LibraryDependencies.inc
new file mode 100644
index 0000000000000000000000000000000000000000..93cb24852b7c49249071d7dcc41ad1baaf0e8f8d
--- /dev/null
+++ b/hpvm/tools/hpvm-config/LibraryDependencies.inc
@@ -0,0 +1,95 @@
+//===- llvm-build generated file --------------------------------*- C++ -*-===//
+//
+// Component Library Dependency Table
+//
+// Automatically generated file, do not edit!
+//
+//===----------------------------------------------------------------------===//
+
+struct AvailableComponent {
+ /// The name of the component.
+ const char *Name;
+
+ /// The name of the library for this component (or NULL).
+ const char *Library;
+
+ /// Whether the component is installed.
+ bool IsInstalled;
+
+ /// The list of libraries required when linking this component.
+ const char *RequiredLibraries[24];
+} AvailableComponents[73] = {
+ { "aggressiveinstcombine", "LLVMAggressiveInstCombine", true, { "analysis", "core", "support", "transformutils" } },
+ { "all", nullptr, true, { "all-targets", "coroutines", "coverage", "debuginfogsym", "dlltooldriver", "engine", "fuzzmutate", "gtest_main", "interpreter", "libdriver", "lineeditor", "lto", "mca", "mirparser", "nativecodegen", "objectyaml", "orcjit", "symbolize", "tablegen", "testingsupport", "textapi", "windowsmanifest", "xray" } },
+ { "all-targets", nullptr, true, { "x86" } },
+ { "analysis", "LLVMAnalysis", true, { "binaryformat", "core", "object", "profiledata", "support" } },
+ { "asmparser", "LLVMAsmParser", true, { "binaryformat", "core", "support" } },
+ { "asmprinter", "LLVMAsmPrinter", true, { "analysis", "binaryformat", "codegen", "core", "debuginfocodeview", "debuginfodwarf", "debuginfomsf", "mc", "mcparser", "remarks", "support", "target" } },
+ { "binaryformat", "LLVMBinaryFormat", true, { "support" } },
+ { "bitreader", "LLVMBitReader", true, { "bitstreamreader", "core", "support" } },
+ { "bitstreamreader", "LLVMBitstreamReader", true, { "support" } },
+ { "bitwriter", "LLVMBitWriter", true, { "analysis", "core", "mc", "object", "support" } },
+ { "codegen", "LLVMCodeGen", true, { "analysis", "bitreader", "bitwriter", "core", "mc", "profiledata", "scalaropts", "support", "target", "transformutils" } },
+ { "core", "LLVMCore", true, { "binaryformat", "remarks", "support" } },
+ { "coroutines", "LLVMCoroutines", true, { "analysis", "core", "ipo", "scalaropts", "support", "transformutils" } },
+ { "coverage", "LLVMCoverage", true, { "core", "object", "profiledata", "support" } },
+ { "debuginfocodeview", "LLVMDebugInfoCodeView", true, { "support", "debuginfomsf" } },
+ { "debuginfodwarf", "LLVMDebugInfoDWARF", true, { "binaryformat", "object", "mc", "support" } },
+ { "debuginfogsym", "LLVMDebugInfoGSYM", true, { "support" } },
+ { "debuginfomsf", "LLVMDebugInfoMSF", true, { "support" } },
+ { "debuginfopdb", "LLVMDebugInfoPDB", true, { "object", "support", "debuginfocodeview", "debuginfomsf" } },
+ { "demangle", "LLVMDemangle", true, { } },
+ { "dlltooldriver", "LLVMDlltoolDriver", true, { "object", "option", "support" } },
+ { "engine", nullptr, true, { "mcjit", "native" } },
+ { "executionengine", "LLVMExecutionEngine", true, { "core", "mc", "object", "runtimedyld", "support", "target" } },
+ { "fuzzmutate", "LLVMFuzzMutate", true, { "analysis", "bitreader", "bitwriter", "core", "scalaropts", "support", "target" } },
+ { "globalisel", "LLVMGlobalISel", true, { "analysis", "codegen", "core", "mc", "selectiondag", "support", "target", "transformutils" } },
+ { "gtest", "gtest", false, { "support" } },
+ { "gtest_main", "gtest_main", false, { "gtest" } },
+ { "instcombine", "LLVMInstCombine", true, { "analysis", "core", "support", "transformutils" } },
+ { "instrumentation", "LLVMInstrumentation", true, { "analysis", "core", "mc", "support", "transformutils", "profiledata" } },
+ { "interpreter", "LLVMInterpreter", true, { "codegen", "core", "executionengine", "support" } },
+ { "ipo", "LLVMipo", true, { "aggressiveinstcombine", "analysis", "bitreader", "bitwriter", "core", "instcombine", "irreader", "linker", "object", "profiledata", "scalaropts", "support", "transformutils", "vectorize", "instrumentation" } },
+ { "irreader", "LLVMIRReader", true, { "asmparser", "bitreader", "core", "support" } },
+ { "jitlink", "LLVMJITLink", true, { "binaryformat", "object", "support" } },
+ { "libdriver", "LLVMLibDriver", true, { "binaryformat", "bitreader", "object", "option", "support" } },
+ { "lineeditor", "LLVMLineEditor", true, { "support" } },
+ { "linker", "LLVMLinker", true, { "core", "support", "transformutils" } },
+ { "lto", "LLVMLTO", true, { "aggressiveinstcombine", "analysis", "bitreader", "bitwriter", "codegen", "core", "ipo", "instcombine", "linker", "mc", "objcarcopts", "object", "passes", "remarks", "scalaropts", "support", "target", "transformutils" } },
+ { "mc", "LLVMMC", true, { "support", "binaryformat", "debuginfocodeview" } },
+ { "mca", "LLVMMCA", true, { "mc", "support" } },
+ { "mcdisassembler", "LLVMMCDisassembler", true, { "mc", "support" } },
+ { "mcjit", "LLVMMCJIT", true, { "core", "executionengine", "object", "runtimedyld", "support", "target" } },
+ { "mcparser", "LLVMMCParser", true, { "mc", "support" } },
+ { "mirparser", "LLVMMIRParser", true, { "asmparser", "binaryformat", "codegen", "core", "mc", "support", "target" } },
+ { "native", nullptr, true, { "x86" } },
+ { "nativecodegen", nullptr, true, { "x86codegen" } },
+ { "objcarcopts", "LLVMObjCARCOpts", true, { "analysis", "core", "support", "transformutils" } },
+ { "object", "LLVMObject", true, { "bitreader", "core", "mc", "binaryformat", "mcparser", "support" } },
+ { "objectyaml", "LLVMObjectYAML", true, { "object", "support", "debuginfocodeview" } },
+ { "option", "LLVMOption", true, { "support" } },
+ { "orcjit", "LLVMOrcJIT", true, { "core", "executionengine", "jitlink", "object", "mc", "runtimedyld", "support", "target", "transformutils" } },
+ { "passes", "LLVMPasses", true, { "aggressiveinstcombine", "analysis", "codegen", "core", "ipo", "instcombine", "scalaropts", "support", "target", "transformutils", "vectorize", "instrumentation" } },
+ { "profiledata", "LLVMProfileData", true, { "core", "support" } },
+ { "remarks", "LLVMRemarks", true, { "support" } },
+ { "runtimedyld", "LLVMRuntimeDyld", true, { "mc", "object", "support" } },
+ { "scalaropts", "LLVMScalarOpts", true, { "aggressiveinstcombine", "analysis", "core", "instcombine", "support", "transformutils" } },
+ { "selectiondag", "LLVMSelectionDAG", true, { "analysis", "codegen", "core", "mc", "support", "target", "transformutils" } },
+ { "support", "LLVMSupport", true, { "demangle" } },
+ { "symbolize", "LLVMSymbolize", true, { "debuginfodwarf", "debuginfopdb", "object", "support", "demangle" } },
+ { "tablegen", "LLVMTableGen", true, { "support" } },
+ { "target", "LLVMTarget", true, { "analysis", "core", "mc", "support" } },
+ { "testingsupport", "LLVMTestingSupport", false, { "support" } },
+ { "textapi", "LLVMTextAPI", true, { "support", "binaryformat" } },
+ { "transformutils", "LLVMTransformUtils", true, { "analysis", "core", "support" } },
+ { "vectorize", "LLVMVectorize", true, { "analysis", "core", "support", "transformutils" } },
+ { "windowsmanifest", "LLVMWindowsManifest", true, { "support" } },
+ { "x86", nullptr, true, { "x86info", "x86utils", "x86desc", "x86codegen", "x86asmparser", "x86disassembler" } },
+ { "x86asmparser", "LLVMX86AsmParser", true, { "mc", "mcparser", "support", "x86desc", "x86info" } },
+ { "x86codegen", "LLVMX86CodeGen", true, { "analysis", "asmprinter", "codegen", "core", "mc", "selectiondag", "support", "target", "x86desc", "x86info", "x86utils", "globalisel", "profiledata" } },
+ { "x86desc", "LLVMX86Desc", true, { "mc", "mcdisassembler", "object", "support", "x86info", "x86utils" } },
+ { "x86disassembler", "LLVMX86Disassembler", true, { "mcdisassembler", "support", "x86info" } },
+ { "x86info", "LLVMX86Info", true, { "support" } },
+ { "x86utils", "LLVMX86Utils", true, { "support" } },
+ { "xray", "LLVMXRay", true, { "support", "object" } },
+};
diff --git a/hpvm/tools/hpvm-config/hpvm-config.cpp b/hpvm/tools/hpvm-config/hpvm-config.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..88f7a0ae9b4ab051927435b2f5dff0dfae13a41c
--- /dev/null
+++ b/hpvm/tools/hpvm-config/hpvm-config.cpp
@@ -0,0 +1,725 @@
+//===-- hpvm-config.cpp - HPVM project configuration utility --------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This tool encapsulates information about an HPVM project configuration for
+// use by other project's build environments (to determine installed path,
+// available features, required libraries, etc.).
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/Config/llvm-config.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/StringMap.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/ADT/Triple.h"
+#include "llvm/ADT/Twine.h"
+#include "llvm/Config/config.h"
+#include "llvm/Support/FileSystem.h"
+#include "llvm/Support/Path.h"
+#include "llvm/Support/WithColor.h"
+#include "llvm/Support/raw_ostream.h"
+#include <cstdlib>
+#include <set>
+#include <unordered_set>
+#include <vector>
+
+using namespace llvm;
+
+#define HPVM_PACKAGE_VERSION "0.5"
+
+
+// Include the build time variables we can report to the user. This is generated
+// at build time from the BuildVariables.inc.in file by the build system.
+#include "BuildVariables.inc"
+
+// Include the component table. This creates an array of struct
+// AvailableComponent entries, which record the component name, library name,
+// and required components for all of the available libraries.
+//
+// Not all components define a library, we also use "library groups" as a way to
+// create entries for pseudo groups like x86 or all-targets.
+#include "LibraryDependencies.inc"
+
+// LinkMode determines what libraries and flags are returned by llvm-config.
+enum LinkMode {
+ // LinkModeAuto will link with the default link mode for the installation,
+ // which is dependent on the value of LLVM_LINK_LLVM_DYLIB, and fall back
+ // to the alternative if the required libraries are not available.
+ LinkModeAuto = 0,
+
+ // LinkModeShared will link with the dynamic component libraries if they
+ // exist, and return an error otherwise.
+ LinkModeShared = 1,
+
+ // LinkModeStatic will link with the static component libraries if they
+ // exist, and return an error otherwise.
+ LinkModeStatic = 2,
+};
+
+/// Traverse a single component adding to the topological ordering in
+/// \arg RequiredLibs.
+///
+/// \param Name - The component to traverse.
+/// \param ComponentMap - A prebuilt map of component names to descriptors.
+/// \param VisitedComponents [in] [out] - The set of already visited components.
+/// \param RequiredLibs [out] - The ordered list of required
+/// libraries.
+/// \param GetComponentNames - Get the component names instead of the
+/// library name.
+static void VisitComponent(const std::string &Name,
+ const StringMap<AvailableComponent *> &ComponentMap,
+ std::set<AvailableComponent *> &VisitedComponents,
+ std::vector<std::string> &RequiredLibs,
+ bool IncludeNonInstalled, bool GetComponentNames,
+ const std::function<std::string(const StringRef &)>
+ *GetComponentLibraryPath,
+ std::vector<std::string> *Missing,
+ const std::string &DirSep) {
+ // Lookup the component.
+ AvailableComponent *AC = ComponentMap.lookup(Name);
+ if (!AC) {
+ errs() << "Can't find component: '" << Name << "' in the map. Available components are: ";
+ for (const auto &Component : ComponentMap) {
+ errs() << "'" << Component.first() << "' ";
+ }
+ errs() << "\n";
+ report_fatal_error("abort");
+ }
+ assert(AC && "Invalid component name!");
+
+ // Add to the visited table.
+ if (!VisitedComponents.insert(AC).second) {
+ // We are done if the component has already been visited.
+ return;
+ }
+
+ // Only include non-installed components if requested.
+ if (!AC->IsInstalled && !IncludeNonInstalled)
+ return;
+
+ // Otherwise, visit all the dependencies.
+ for (unsigned i = 0; AC->RequiredLibraries[i]; ++i) {
+ VisitComponent(AC->RequiredLibraries[i], ComponentMap, VisitedComponents,
+ RequiredLibs, IncludeNonInstalled, GetComponentNames,
+ GetComponentLibraryPath, Missing, DirSep);
+ }
+
+ if (GetComponentNames) {
+ RequiredLibs.push_back(Name);
+ return;
+ }
+
+ // Add to the required library list.
+ if (AC->Library) {
+ if (Missing && GetComponentLibraryPath) {
+ std::string path = (*GetComponentLibraryPath)(AC->Library);
+ if (DirSep == "\\") {
+ std::replace(path.begin(), path.end(), '/', '\\');
+ }
+ if (!sys::fs::exists(path))
+ Missing->push_back(path);
+ }
+ RequiredLibs.push_back(AC->Library);
+ }
+}
+
+/// Compute the list of required libraries for a given list of
+/// components, in an order suitable for passing to a linker (that is, libraries
+/// appear prior to their dependencies).
+///
+/// \param Components - The names of the components to find libraries for.
+/// \param IncludeNonInstalled - Whether non-installed components should be
+/// reported.
+/// \param GetComponentNames - True if one would prefer the component names.
+static std::vector<std::string> ComputeLibsForComponents(
+ const std::vector<StringRef> &Components, bool IncludeNonInstalled,
+ bool GetComponentNames, const std::function<std::string(const StringRef &)>
+ *GetComponentLibraryPath,
+ std::vector<std::string> *Missing, const std::string &DirSep) {
+ std::vector<std::string> RequiredLibs;
+ std::set<AvailableComponent *> VisitedComponents;
+
+ // Build a map of component names to information.
+ StringMap<AvailableComponent *> ComponentMap;
+ for (unsigned i = 0; i != array_lengthof(AvailableComponents); ++i) {
+ AvailableComponent *AC = &AvailableComponents[i];
+ ComponentMap[AC->Name] = AC;
+ }
+
+ // Visit the components.
+ for (unsigned i = 0, e = Components.size(); i != e; ++i) {
+ // Users are allowed to provide mixed case component names.
+ std::string ComponentLower = Components[i].lower();
+
+ // Validate that the user supplied a valid component name.
+ if (!ComponentMap.count(ComponentLower)) {
+ llvm::errs() << "llvm-config: unknown component name: " << Components[i]
+ << "\n";
+ exit(1);
+ }
+
+ VisitComponent(ComponentLower, ComponentMap, VisitedComponents,
+ RequiredLibs, IncludeNonInstalled, GetComponentNames,
+ GetComponentLibraryPath, Missing, DirSep);
+ }
+
+ // The list is now ordered with leafs first, we want the libraries to printed
+ // in the reverse order of dependency.
+ std::reverse(RequiredLibs.begin(), RequiredLibs.end());
+
+ return RequiredLibs;
+}
+
+/* *** */
+
+static void usage() {
+ errs() << "\
+usage: llvm-config <OPTION>... [<COMPONENT>...]\n\
+\n\
+Get various configuration information needed to compile programs which use\n\
+LLVM. Typically called from 'configure' scripts. Examples:\n\
+ llvm-config --cxxflags\n\
+ llvm-config --ldflags\n\
+ llvm-config --libs engine bcreader scalaropts\n\
+\n\
+Options:\n\
+ --llvm-version Print LLVM version.\n\
+ --hpvm-version Print HPVM version.\n\
+ --version Print HPVM version.\n\
+ --prefix Print the installation prefix.\n\
+ --src-root Print the source root LLVM was built from.\n\
+ --obj-root Print the object root used to build LLVM.\n\
+ --bindir Directory containing LLVM executables.\n\
+ --includedir Directory containing LLVM headers.\n\
+ --libdir Directory containing LLVM libraries.\n\
+ --cmakedir Directory containing LLVM cmake modules.\n\
+ --cppflags C preprocessor flags for files that include LLVM headers.\n\
+ --cflags C compiler flags for files that include LLVM headers.\n\
+ --cxxflags C++ compiler flags for files that include LLVM headers.\n\
+ --ldflags Print Linker flags.\n\
+ --system-libs System Libraries needed to link against LLVM components.\n\
+ --libs Libraries needed to link against LLVM components.\n\
+ --libnames Bare library names for in-tree builds.\n\
+ --libfiles Fully qualified library filenames for makefile depends.\n\
+ --components List of all possible components.\n\
+ --targets-built List of all targets currently built.\n\
+ --host-target Target triple used to configure LLVM.\n\
+ --build-mode Print build mode of LLVM tree (e.g. Debug or Release).\n\
+ --assertion-mode Print assertion mode of LLVM tree (ON or OFF).\n\
+ --build-system Print the build system used to build LLVM (always cmake).\n\
+ --has-rtti Print whether or not LLVM was built with rtti (YES or NO).\n\
+ --has-global-isel Print whether or not LLVM was built with global-isel support (ON or OFF).\n\
+ --shared-mode Print how the provided components can be collectively linked (`shared` or `static`).\n\
+ --link-shared Link the components as shared libraries.\n\
+ --link-static Link the component libraries statically.\n\
+ --ignore-libllvm Ignore libLLVM and link component libraries instead.\n\
+ --genhpvm Generate textual representation for HPVM IR from LLVM IR.\n\
+ --dfg2llvm-cpu Generate code for CPU.\n\
+ --dfg2llvm-opencl Generate kernel code in OpenCL.\n\
+ --clearDFG Clean up HPVM dataflow graph and extraneous HPVM-specific instructions.\n\
+Typical components:\n\
+ all All LLVM libraries (default).\n\
+ engine Either a native JIT or a bitcode interpreter.\n";
+ exit(1);
+}
+
+/// Compute the path to the main executable.
+std::string GetExecutablePath(const char *Argv0) {
+ // This just needs to be some symbol in the binary; C++ doesn't
+ // allow taking the address of ::main however.
+ void *P = (void *)(intptr_t)GetExecutablePath;
+ return llvm::sys::fs::getMainExecutable(Argv0, P);
+}
+
+/// Expand the semi-colon delimited LLVM_DYLIB_COMPONENTS into
+/// the full list of components.
+std::vector<std::string> GetAllDyLibComponents(const bool IsInDevelopmentTree,
+ const bool GetComponentNames,
+ const std::string &DirSep) {
+ std::vector<StringRef> DyLibComponents;
+
+ StringRef DyLibComponentsStr(LLVM_DYLIB_COMPONENTS);
+ size_t Offset = 0;
+ while (true) {
+ const size_t NextOffset = DyLibComponentsStr.find(';', Offset);
+ DyLibComponents.push_back(DyLibComponentsStr.substr(Offset, NextOffset-Offset));
+ if (NextOffset == std::string::npos) {
+ break;
+ }
+ Offset = NextOffset + 1;
+ }
+
+ assert(!DyLibComponents.empty());
+
+ return ComputeLibsForComponents(DyLibComponents,
+ /*IncludeNonInstalled=*/IsInDevelopmentTree,
+ GetComponentNames, nullptr, nullptr, DirSep);
+}
+
+int main(int argc, char **argv) {
+ std::vector<StringRef> Components;
+ bool PrintLibs = false, PrintLibNames = false, PrintLibFiles = false;
+ bool PrintSystemLibs = false, PrintSharedMode = false;
+ bool HasAnyOption = false;
+
+ // llvm-config is designed to support being run both from a development tree
+ // and from an installed path. We try and auto-detect which case we are in so
+ // that we can report the correct information when run from a development
+ // tree.
+ bool IsInDevelopmentTree;
+ enum { CMakeStyle, CMakeBuildModeStyle } DevelopmentTreeLayout;
+ llvm::SmallString<256> CurrentPath(GetExecutablePath(argv[0]));
+ std::string CurrentExecPrefix;
+ std::string ActiveObjRoot;
+
+ // If CMAKE_CFG_INTDIR is given, honor it as build mode.
+ char const *build_mode = LLVM_BUILDMODE;
+#if defined(CMAKE_CFG_INTDIR)
+ if (!(CMAKE_CFG_INTDIR[0] == '.' && CMAKE_CFG_INTDIR[1] == '\0'))
+ build_mode = CMAKE_CFG_INTDIR;
+#endif
+
+ // Create an absolute path, and pop up one directory (we expect to be inside a
+ // bin dir).
+ sys::fs::make_absolute(CurrentPath);
+ CurrentExecPrefix =
+ sys::path::parent_path(sys::path::parent_path(CurrentPath)).str();
+
+ // Check to see if we are inside a development tree by comparing to possible
+ // locations (prefix style or CMake style).
+ if (sys::fs::equivalent(CurrentExecPrefix, LLVM_OBJ_ROOT)) {
+ IsInDevelopmentTree = true;
+ DevelopmentTreeLayout = CMakeStyle;
+ ActiveObjRoot = LLVM_OBJ_ROOT;
+ } else if (sys::fs::equivalent(CurrentExecPrefix,
+ Twine(LLVM_OBJ_ROOT) + "/bin")) {
+ IsInDevelopmentTree = true;
+ DevelopmentTreeLayout = CMakeBuildModeStyle;
+ ActiveObjRoot = LLVM_OBJ_ROOT;
+ } else {
+ IsInDevelopmentTree = false;
+ DevelopmentTreeLayout = CMakeStyle; // Initialized to avoid warnings.
+ }
+
+ // Compute various directory locations based on the derived location
+ // information.
+ std::string ActivePrefix, ActiveBinDir, ActiveIncludeDir, ActiveLibDir,
+ ActiveCMakeDir;
+ std::string ActiveIncludeOption;
+ if (IsInDevelopmentTree) {
+ ActiveIncludeDir = std::string(LLVM_SRC_ROOT) + "/include";
+ ActivePrefix = CurrentExecPrefix;
+
+ // CMake organizes the products differently than a normal prefix style
+ // layout.
+ switch (DevelopmentTreeLayout) {
+ case CMakeStyle:
+ ActiveBinDir = ActiveObjRoot + "/bin";
+ ActiveLibDir = ActiveObjRoot + "/lib" + LLVM_LIBDIR_SUFFIX;
+ ActiveCMakeDir = ActiveLibDir + "/cmake/llvm";
+ break;
+ case CMakeBuildModeStyle:
+ ActivePrefix = ActiveObjRoot;
+ ActiveBinDir = ActiveObjRoot + "/bin/" + build_mode;
+ ActiveLibDir =
+ ActiveObjRoot + "/lib" + LLVM_LIBDIR_SUFFIX + "/" + build_mode;
+ ActiveCMakeDir = ActiveLibDir + "/cmake/llvm";
+ break;
+ }
+
+ // We need to include files from both the source and object trees.
+ ActiveIncludeOption =
+ ("-I" + ActiveIncludeDir + " " + "-I" + ActiveObjRoot + "/include");
+ } else {
+ ActivePrefix = CurrentExecPrefix;
+ ActiveIncludeDir = ActivePrefix + "/include";
+ SmallString<256> path(StringRef(LLVM_TOOLS_INSTALL_DIR));
+ sys::fs::make_absolute(ActivePrefix, path);
+ ActiveBinDir = path.str();
+ ActiveLibDir = ActivePrefix + "/lib" + LLVM_LIBDIR_SUFFIX;
+ ActiveCMakeDir = ActiveLibDir + "/cmake/llvm";
+ ActiveIncludeOption = "-I" + ActiveIncludeDir;
+ }
+
+ /// We only use `shared library` mode in cases where the static library form
+ /// of the components provided are not available; note however that this is
+ /// skipped if we're run from within the build dir. However, once installed,
+ /// we still need to provide correct output when the static archives are
+ /// removed or, as in the case of CMake's `BUILD_SHARED_LIBS`, never present
+ /// in the first place. This can't be done at configure/build time.
+
+ StringRef SharedExt, SharedVersionedExt, SharedDir, SharedPrefix, StaticExt,
+ StaticPrefix, StaticDir = "lib", DirSep = "/";
+ const Triple HostTriple(Triple::normalize(LLVM_HOST_TRIPLE));
+ if (HostTriple.isOSWindows()) {
+ SharedExt = "dll";
+ SharedVersionedExt = LLVM_DYLIB_VERSION ".dll";
+ if (HostTriple.isOSCygMing()) {
+ StaticExt = "a";
+ StaticPrefix = "lib";
+ } else {
+ StaticExt = "lib";
+ DirSep = "\\";
+ std::replace(ActiveObjRoot.begin(), ActiveObjRoot.end(), '/', '\\');
+ std::replace(ActivePrefix.begin(), ActivePrefix.end(), '/', '\\');
+ std::replace(ActiveBinDir.begin(), ActiveBinDir.end(), '/', '\\');
+ std::replace(ActiveLibDir.begin(), ActiveLibDir.end(), '/', '\\');
+ std::replace(ActiveCMakeDir.begin(), ActiveCMakeDir.end(), '/', '\\');
+ std::replace(ActiveIncludeOption.begin(), ActiveIncludeOption.end(), '/',
+ '\\');
+ }
+ SharedDir = ActiveBinDir;
+ StaticDir = ActiveLibDir;
+ } else if (HostTriple.isOSDarwin()) {
+ SharedExt = "dylib";
+ SharedVersionedExt = LLVM_DYLIB_VERSION ".dylib";
+ StaticExt = "a";
+ StaticDir = SharedDir = ActiveLibDir;
+ StaticPrefix = SharedPrefix = "lib";
+ } else {
+ // default to the unix values:
+ SharedExt = "so";
+ SharedVersionedExt = LLVM_DYLIB_VERSION ".so";
+ StaticExt = "a";
+ StaticDir = SharedDir = ActiveLibDir;
+ StaticPrefix = SharedPrefix = "lib";
+ }
+
+ const bool BuiltDyLib = !!LLVM_ENABLE_DYLIB;
+
+ /// CMake style shared libs, ie each component is in a shared library.
+ const bool BuiltSharedLibs = !!LLVM_ENABLE_SHARED;
+
+ bool DyLibExists = false;
+ const std::string DyLibName =
+ (SharedPrefix + "LLVM-" + SharedVersionedExt).str();
+
+ // If LLVM_LINK_DYLIB is ON, the single shared library will be returned
+ // for "--libs", etc, if they exist. This behaviour can be overridden with
+ // --link-static or --link-shared.
+ bool LinkDyLib = !!LLVM_LINK_DYLIB;
+
+ if (BuiltDyLib) {
+ std::string path((SharedDir + DirSep + DyLibName).str());
+ if (DirSep == "\\") {
+ std::replace(path.begin(), path.end(), '/', '\\');
+ }
+ DyLibExists = sys::fs::exists(path);
+ if (!DyLibExists) {
+ // The shared library does not exist: don't error unless the user
+ // explicitly passes --link-shared.
+ LinkDyLib = false;
+ }
+ }
+ LinkMode LinkMode =
+ (LinkDyLib || BuiltSharedLibs) ? LinkModeShared : LinkModeAuto;
+
+ /// Get the component's library name without the lib prefix and the
+ /// extension. Returns true if Lib is in a recognized format.
+ auto GetComponentLibraryNameSlice = [&](const StringRef &Lib,
+ StringRef &Out) {
+ if (Lib.startswith("lib")) {
+ unsigned FromEnd;
+ if (Lib.endswith(StaticExt)) {
+ FromEnd = StaticExt.size() + 1;
+ } else if (Lib.endswith(SharedExt)) {
+ FromEnd = SharedExt.size() + 1;
+ } else {
+ FromEnd = 0;
+ }
+
+ if (FromEnd != 0) {
+ Out = Lib.slice(3, Lib.size() - FromEnd);
+ return true;
+ }
+ }
+
+ return false;
+ };
+ /// Maps Unixizms to the host platform.
+ auto GetComponentLibraryFileName = [&](const StringRef &Lib,
+ const bool Shared) {
+ std::string LibFileName;
+ if (Shared) {
+ if (Lib == DyLibName) {
+ // Treat the DyLibName specially. It is not a component library and
+ // already has the necessary prefix and suffix (e.g. `.so`) added so
+ // just return it unmodified.
+ assert(Lib.endswith(SharedExt) && "DyLib is missing suffix");
+ LibFileName = Lib;
+ } else {
+ LibFileName = (SharedPrefix + Lib + "." + SharedExt).str();
+ }
+ } else {
+ // default to static
+ LibFileName = (StaticPrefix + Lib + "." + StaticExt).str();
+ }
+
+ return LibFileName;
+ };
+ /// Get the full path for a possibly shared component library.
+ auto GetComponentLibraryPath = [&](const StringRef &Name, const bool Shared) {
+ auto LibFileName = GetComponentLibraryFileName(Name, Shared);
+ if (Shared) {
+ return (SharedDir + DirSep + LibFileName).str();
+ } else {
+ return (StaticDir + DirSep + LibFileName).str();
+ }
+ };
+
+ raw_ostream &OS = outs();
+ for (int i = 1; i != argc; ++i) {
+ StringRef Arg = argv[i];
+
+ if (Arg.startswith("-")) {
+ HasAnyOption = true;
+ if(Arg == "--llvm-version") {
+ OS << PACKAGE_VERSION << '\n';
+ } else if(Arg == "--hpvm-version") {
+ OS << HPVM_PACKAGE_VERSION << '\n';
+ } else if (Arg == "--version") {
+ OS << HPVM_PACKAGE_VERSION << '\n';
+ } else if (Arg == "--prefix") {
+ OS << ActivePrefix << '\n';
+ } else if (Arg == "--bindir") {
+ OS << ActiveBinDir << '\n';
+ } else if (Arg == "--includedir") {
+ OS << ActiveIncludeDir << '\n';
+ } else if (Arg == "--libdir") {
+ OS << ActiveLibDir << '\n';
+ } else if (Arg == "--cmakedir") {
+ OS << ActiveCMakeDir << '\n';
+ } else if (Arg == "--cppflags") {
+ OS << ActiveIncludeOption << ' ' << LLVM_CPPFLAGS << '\n';
+ } else if (Arg == "--cflags") {
+ OS << ActiveIncludeOption << ' ' << LLVM_CFLAGS << '\n';
+ } else if (Arg == "--cxxflags") {
+ OS << ActiveIncludeOption << ' ' << LLVM_CXXFLAGS << '\n';
+ } else if (Arg == "--ldflags") {
+ OS << ((HostTriple.isWindowsMSVCEnvironment()) ? "-LIBPATH:" : "-L")
+ << ActiveLibDir << ' ' << LLVM_LDFLAGS << '\n';
+ } else if (Arg == "--system-libs") {
+ PrintSystemLibs = true;
+ } else if (Arg == "--libs") {
+ PrintLibs = true;
+ } else if (Arg == "--libnames") {
+ PrintLibNames = true;
+ } else if (Arg == "--libfiles") {
+ PrintLibFiles = true;
+ } else if (Arg == "--components") {
+ /// If there are missing static archives and a dylib was
+ /// built, print LLVM_DYLIB_COMPONENTS instead of everything
+ /// in the manifest.
+ std::vector<std::string> Components;
+ for (unsigned j = 0; j != array_lengthof(AvailableComponents); ++j) {
+ // Only include non-installed components when in a development tree.
+ if (!AvailableComponents[j].IsInstalled && !IsInDevelopmentTree)
+ continue;
+
+ Components.push_back(AvailableComponents[j].Name);
+ if (AvailableComponents[j].Library && !IsInDevelopmentTree) {
+ std::string path(
+ GetComponentLibraryPath(AvailableComponents[j].Library, false));
+ if (DirSep == "\\") {
+ std::replace(path.begin(), path.end(), '/', '\\');
+ }
+ if (DyLibExists && !sys::fs::exists(path)) {
+ Components =
+ GetAllDyLibComponents(IsInDevelopmentTree, true, DirSep);
+ llvm::sort(Components);
+ break;
+ }
+ }
+ }
+
+ for (unsigned I = 0; I < Components.size(); ++I) {
+ if (I) {
+ OS << ' ';
+ }
+
+ OS << Components[I];
+ }
+ OS << '\n';
+ } else if (Arg == "--targets-built") {
+ OS << LLVM_TARGETS_BUILT << '\n';
+ } else if (Arg == "--host-target") {
+ OS << Triple::normalize(LLVM_DEFAULT_TARGET_TRIPLE) << '\n';
+ } else if (Arg == "--build-mode") {
+ OS << build_mode << '\n';
+ } else if (Arg == "--assertion-mode") {
+#if defined(NDEBUG)
+ OS << "OFF\n";
+#else
+ OS << "ON\n";
+#endif
+ } else if (Arg == "--build-system") {
+ OS << LLVM_BUILD_SYSTEM << '\n';
+ } else if (Arg == "--has-rtti") {
+ OS << (LLVM_HAS_RTTI ? "YES" : "NO") << '\n';
+ } else if (Arg == "--has-global-isel") {
+ OS << (LLVM_HAS_GLOBAL_ISEL ? "ON" : "OFF") << '\n';
+ } else if (Arg == "--shared-mode") {
+ PrintSharedMode = true;
+ } else if (Arg == "--obj-root") {
+ OS << ActivePrefix << '\n';
+ } else if (Arg == "--src-root") {
+ OS << LLVM_SRC_ROOT << '\n';
+ } else if (Arg == "--ignore-libllvm") {
+ LinkDyLib = false;
+ LinkMode = BuiltSharedLibs ? LinkModeShared : LinkModeAuto;
+ } else if (Arg == "--link-shared") {
+ LinkMode = LinkModeShared;
+ } else if (Arg == "--link-static") {
+ LinkMode = LinkModeStatic;
+ } else {
+ usage();
+ }
+ } else {
+ Components.push_back(Arg);
+ }
+ }
+
+ if (!HasAnyOption)
+ usage();
+
+ if (LinkMode == LinkModeShared && !DyLibExists && !BuiltSharedLibs) {
+ WithColor::error(errs(), "llvm-config") << DyLibName << " is missing\n";
+ return 1;
+ }
+
+ if (PrintLibs || PrintLibNames || PrintLibFiles || PrintSystemLibs ||
+ PrintSharedMode) {
+
+ if (PrintSharedMode && BuiltSharedLibs) {
+ OS << "shared\n";
+ return 0;
+ }
+
+ // If no components were specified, default to "all".
+ if (Components.empty())
+ Components.push_back("all");
+
+ // Construct the list of all the required libraries.
+ std::function<std::string(const StringRef &)>
+ GetComponentLibraryPathFunction = [&](const StringRef &Name) {
+ return GetComponentLibraryPath(Name, LinkMode == LinkModeShared);
+ };
+ std::vector<std::string> MissingLibs;
+ std::vector<std::string> RequiredLibs = ComputeLibsForComponents(
+ Components,
+ /*IncludeNonInstalled=*/IsInDevelopmentTree, false,
+ &GetComponentLibraryPathFunction, &MissingLibs, DirSep);
+ if (!MissingLibs.empty()) {
+ switch (LinkMode) {
+ case LinkModeShared:
+ if (LinkDyLib && !BuiltSharedLibs)
+ break;
+ // Using component shared libraries.
+ for (auto &Lib : MissingLibs)
+ WithColor::error(errs(), "llvm-config") << "missing: " << Lib << "\n";
+ return 1;
+ case LinkModeAuto:
+ if (DyLibExists) {
+ LinkMode = LinkModeShared;
+ break;
+ }
+ WithColor::error(errs(), "llvm-config")
+ << "component libraries and shared library\n\n";
+ LLVM_FALLTHROUGH;
+ case LinkModeStatic:
+ for (auto &Lib : MissingLibs)
+ WithColor::error(errs(), "llvm-config") << "missing: " << Lib << "\n";
+ return 1;
+ }
+ } else if (LinkMode == LinkModeAuto) {
+ LinkMode = LinkModeStatic;
+ }
+
+ if (PrintSharedMode) {
+ std::unordered_set<std::string> FullDyLibComponents;
+ std::vector<std::string> DyLibComponents =
+ GetAllDyLibComponents(IsInDevelopmentTree, false, DirSep);
+
+ for (auto &Component : DyLibComponents) {
+ FullDyLibComponents.insert(Component);
+ }
+ DyLibComponents.clear();
+
+ for (auto &Lib : RequiredLibs) {
+ if (!FullDyLibComponents.count(Lib)) {
+ OS << "static\n";
+ return 0;
+ }
+ }
+ FullDyLibComponents.clear();
+
+ if (LinkMode == LinkModeShared) {
+ OS << "shared\n";
+ return 0;
+ } else {
+ OS << "static\n";
+ return 0;
+ }
+ }
+
+ if (PrintLibs || PrintLibNames || PrintLibFiles) {
+
+ auto PrintForLib = [&](const StringRef &Lib) {
+ const bool Shared = LinkMode == LinkModeShared;
+ if (PrintLibNames) {
+ OS << GetComponentLibraryFileName(Lib, Shared);
+ } else if (PrintLibFiles) {
+ OS << GetComponentLibraryPath(Lib, Shared);
+ } else if (PrintLibs) {
+ // On Windows, output full path to library without parameters.
+ // Elsewhere, if this is a typical library name, include it using -l.
+ if (HostTriple.isWindowsMSVCEnvironment()) {
+ OS << GetComponentLibraryPath(Lib, Shared);
+ } else {
+ StringRef LibName;
+ if (GetComponentLibraryNameSlice(Lib, LibName)) {
+ // Extract library name (remove prefix and suffix).
+ OS << "-l" << LibName;
+ } else {
+ // Lib is already a library name without prefix and suffix.
+ OS << "-l" << Lib;
+ }
+ }
+ }
+ };
+
+ if (LinkMode == LinkModeShared && LinkDyLib) {
+ PrintForLib(DyLibName);
+ } else {
+ for (unsigned i = 0, e = RequiredLibs.size(); i != e; ++i) {
+ auto Lib = RequiredLibs[i];
+ if (i)
+ OS << ' ';
+
+ PrintForLib(Lib);
+ }
+ }
+ OS << '\n';
+ }
+
+ // Print SYSTEM_LIBS after --libs.
+ // FIXME: Each LLVM component may have its dependent system libs.
+ if (PrintSystemLibs) {
+ // Output system libraries only if linking against a static
+ // library (since the shared library links to all system libs
+ // already)
+ OS << (LinkMode == LinkModeStatic ? LLVM_SYSTEM_LIBS : "") << '\n';
+ }
+ } else if (!Components.empty()) {
+ WithColor::error(errs(), "llvm-config")
+ << "components given, but unused\n\n";
+ usage();
+ }
+
+ return 0;
+}