Update DFG2LLVM for NVPTX with changes to allocas and API for getting and insertion of functions

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");