diff --git a/hpvm/lib/Transforms/DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp b/hpvm/lib/Transforms/DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp
index 5a778566043ff59e1135c93447ae48fef051ec3c..d9692b691a853202ca22b9a05258099198c45528 100644
--- a/hpvm/lib/Transforms/DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp
+++ b/hpvm/lib/Transforms/DFG2LLVM_NVPTX/DFG2LLVM_NVPTX.cpp
@@ -202,7 +202,7 @@ private:
public:
// Constructor
- CGT_NVPTX(Module &_M, BuildDFG &_DFG) : CodeGenTraversal(_M, _DFG), KernelM(CloneModule(&_M)) {
+ CGT_NVPTX(Module &_M, BuildDFG &_DFG) : CodeGenTraversal(_M, _DFG), KernelM(CloneModule(_M)) {
init();
initRuntimeAPI();
errs() << "Old module pointer: " << &_M << "\n";
@@ -527,7 +527,7 @@ void CGT_NVPTX::insertRuntimeCalls(DFInternalNode* N, Kernel* K, const Twine& Fi
// Scalar Input
// Store the scalar value on stack and then pass the pointer to its
// location
- AllocaInst* inputValPtr = new AllocaInst(inputVal->getType(), inputVal->getName()+".ptr", RI);
+ AllocaInst* inputValPtr = new AllocaInst(inputVal->getType(), 0, inputVal->getName()+".ptr", RI);
StoreInst* SI = new StoreInst(inputVal, inputValPtr, RI);
Value* inputValI8Ptr = CastInst::CreatePointerCast(inputValPtr,
@@ -580,7 +580,7 @@ void CGT_NVPTX::insertRuntimeCalls(DFInternalNode* N, Kernel* K, const Twine& Fi
switchToTimer(visc_TimerID_COPY_SCALAR, RI);
// Store the scalar value on stack and then pass the pointer to its
// location
- AllocaInst* allocSizePtr = new AllocaInst(allocSize->getType(),
+ AllocaInst* allocSizePtr = new AllocaInst(allocSize->getType(), 0,
allocSize->getName()+".sharedMem.ptr", RI);
StoreInst* SI = new StoreInst(allocSize, allocSizePtr, RI);
@@ -638,7 +638,7 @@ void CGT_NVPTX::insertRuntimeCalls(DFInternalNode* N, Kernel* K, const Twine& Fi
switchToTimer(visc_TimerID_COPY_SCALAR, RI);
// Store the scalar value on stack and then pass the pointer to its
// location
- AllocaInst* allocSizePtr = new AllocaInst(allocSize->getType(),
+ AllocaInst* allocSizePtr = new AllocaInst(allocSize->getType(), 0,
allocSize->getName()+".sharedMem.ptr", RI);
StoreInst* SI = new StoreInst(allocSize, allocSizePtr, RI);
@@ -1032,8 +1032,8 @@ void CGT_NVPTX::codeGen(DFLeafNode* N) {
N->addGenFunc(F_nvptx, visc::GPU_TARGET, false);
DEBUG(errs() << "Removing all attributes from Kernel Function and adding nounwind\n");
- F_nvptx->removeAttributes(AttributeSet::FunctionIndex, F_nvptx->getAttributes().getFnAttributes());
- F_nvptx->addAttribute(AttributeSet::FunctionIndex, Attribute::NoUnwind);
+ F_nvptx->removeAttributes(AttributeList::FunctionIndex, F_nvptx->getAttributes().getFnAttributes());
+ F_nvptx->addAttribute(AttributeList::FunctionIndex, Attribute::NoUnwind);
//FIXME: For now, assume only one allocation node
kernel->AllocationNode = NULL;
@@ -1230,19 +1230,19 @@ void CGT_NVPTX::codeGen(DFLeafNode* N) {
DEBUG(errs() << "Substitute with get_global_id()\n");
DEBUG(errs() << *II << "\n");
OpenCLFunction = cast<Function>
- (KernelM->getOrInsertFunction(StringRef("get_global_id"), FT));
+ ((KernelM->getOrInsertFunction(StringRef("get_global_id"), FT)).getCallee());
} else if (Leaf_HandleToDFNodeMap[ArgII] == N) {
//DEBUG(errs() << "Here inside cond 2\n");
// We are asking for this node's id with respect to its parent
// this is a local id call
OpenCLFunction = cast<Function>
- (KernelM->getOrInsertFunction(StringRef("get_local_id"), FT));
+ ((KernelM->getOrInsertFunction(StringRef("get_local_id"), FT)).getCallee());
//DEBUG(errs() << "exiting condition 2\n");
} else if (Leaf_HandleToDFNodeMap[ArgII] == N->getParent()) {
// We are asking for this node's parent's id with respect to its
// parent: this is a group id call
OpenCLFunction = cast<Function>
- (KernelM->getOrInsertFunction(StringRef("get_group_id"), FT));
+ ((KernelM->getOrInsertFunction(StringRef("get_group_id"), FT)).getCallee());
} else {
errs() << N->getFuncPointer()->getName() << "\n";
errs() << N->getParent()->getFuncPointer()->getName() << "\n";
@@ -1308,17 +1308,17 @@ void CGT_NVPTX::codeGen(DFLeafNode* N) {
// replicated. This indicates that the parent node is the kernel
// launch, so the instances are global_size (gridDim x blockDim)
OpenCLFunction = cast<Function>
- (KernelM->getOrInsertFunction(StringRef("get_global_size"), FT));
+ ((KernelM->getOrInsertFunction(StringRef("get_global_size"), FT)).getCallee());
} else if (Leaf_HandleToDFNodeMap[ArgII] == N) {
// We are asking for this node's instances
// this is a local size (block dim) call
OpenCLFunction = cast<Function>
- (KernelM->getOrInsertFunction(StringRef("get_local_size"), FT));
+ ((KernelM->getOrInsertFunction(StringRef("get_local_size"), FT)).getCallee());
} else if (Leaf_HandleToDFNodeMap[ArgII] == N->getParent()) {
// We are asking for this node's parent's instances
// this is a (global_size/local_size) (grid dim) call
OpenCLFunction = cast<Function>
- (KernelM->getOrInsertFunction(StringRef("get_num_groups"), FT));
+ ((KernelM->getOrInsertFunction(StringRef("get_num_groups"), FT)).getCallee());
} else {
assert(false && "Unable to translate getNumNodeInstances intrinsic");
}
@@ -1340,7 +1340,7 @@ void CGT_NVPTX::codeGen(DFLeafNode* N) {
std::vector<Type*>(1, Type::getInt32Ty(KernelM->getContext())),
false);
Function* OpenCLFunction = cast<Function>
- (KernelM->getOrInsertFunction(StringRef("barrier"), FT));
+ ((KernelM->getOrInsertFunction(StringRef("barrier"), FT)).getCallee());
CallInst* CI = CallInst::Create(OpenCLFunction,
ArrayRef<Value*>(ConstantInt::get(Type::getInt32Ty(KernelM->getContext()), 1)),
"", II);
@@ -1891,7 +1891,7 @@ static Value* genWorkGroupPtr(Module &M, std::vector<Value*> WGSize, ValueToValu
Type* WGTy = ArrayType::get(Int64Ty, WGSize.size());
// Allocate space of Global work group data on stack and get pointer to
// first element.
- AllocaInst* WG = new AllocaInst(WGTy, WGName, IB);
+ AllocaInst* WG = new AllocaInst(WGTy, 0, WGName, IB);
WGPtr = BitCastInst::CreatePointerCast(WG, Int64Ty->getPointerTo(), WG->getName()+".0", IB);
Value* nextDim = WGPtr;
DEBUG(errs() << *WGPtr << "\n");