Made changes to GenVISC pass to correctly change the parameter list and return

type of the kernel function

llvm/include/llvm/GenVISC/GenVISC.h

@@ -34,8 +34,9 @@ public:
   // Functions
   virtual bool runOnModule(Module &M);
 
-  void genKernel(Function* KernelFunction, CallInst* CI);
-  void genHost(CallInst* CI);
+  void generateTest(CallInst* CI);
+  void genKernel(Function* KernelFunction, CallInst* CI, StructType* RetTy);
+  void genHost(CallInst*, unsigned, unsigned, unsigned, StructType*);
 };
 
 } // End of namespace

llvm/lib/Transforms/GenVISC/GenVISC.cpp

@@ -186,7 +186,7 @@ static Function* genRootNode(Function* KernelF, unsigned numArgs, unsigned numDi
         createNodeXD = Intrinsic::visc_createNode3D;
         break;
     default:
-        assert(false && "Invalid number of dimensions!");
+        llvm_unreachable("Invalid number of dimensions!");
         break;
     };
     
@@ -237,8 +237,7 @@ bool GenVISC::runOnModule(Module &M) {
                 assert(CI->getNumArgOperands() >= 4
                        && "__visc__node call should have atleast 4 arguments!");
                 errs() << "Kernel Function = " << *CI->getArgOperand(1) << "\n";
-                genKernel(cast<Function>(CI->getArgOperand(0)), CI);
-                genHost(CI);
+                generateTest(CI);
                 toBeErased.push_back(CI);
             }
         }
@@ -255,15 +254,29 @@ bool GenVISC::runOnModule(Module &M) {
     return false; //TODO: What does returning "false" mean?
 }
 
-void GenVISC::genKernel(Function* KernelF, CallInst* CI) {
-    // Make changes to kernel here
-    errs() << "Modifying Node Function: " << KernelF->getName() << "\n";
+static void shiftAttributes(Function::arg_iterator from, unsigned shift, Function* F) {
+  Function::arg_iterator ai = --F->arg_end(), ae = --from;
+  errs() << "Total arguments: " << F->getArgumentList().size() << "\n";
+  errs() << "Total slots = " << F->getAttributes().getNumSlots() << "\n";
+  for(; ai!=ae; ai--) {
+    if(ai->getArgNo()+1+shift > F->getArgumentList().size()) {
+      errs() << "Skipping argument " << ai->getArgNo() << "\n";
+      continue;
+    }
 
+    AttributeSet AS = F->getAttributes();
+    errs() << "Moving attributes from " << ai->getArgNo() << " to " << ai->getArgNo()+shift << "\n";
+    AttrBuilder AB(AS, ai->getArgNo()+1);
+    AttributeSet argAS = AttributeSet::get(F->getContext(), ai->getArgNo()+1+shift, AB);
+    F->addAttributes(ai->getArgNo()+1+shift, argAS);
+    F->removeAttributes(ai->getArgNo()+1,AS.getParamAttributes(ai->getArgNo()+1));
+    AS = F->getAttributes();
+    errs() << "Argument: " << *ai << "\n";
+  }
 }
 
-void GenVISC::genHost(CallInst* CI) {
-    // Make host code changes here
-    errs() << "Modifying Host code for __visc__node call site: " << *CI << "\n";
+void GenVISC::generateTest(CallInst* CI) {
+    // Parse the dummy function call here
     LLVMContext& Ctx = CI->getParent()->getContext();
 
     unsigned offset = 1; // argument at offset 1 is the number of dimensions
@@ -292,9 +305,66 @@ void GenVISC::genHost(CallInst* CI) {
     assert(numOutputs == 0 && "Not handled case where number of outputs is non-zero!");
     StructType* RetTy = StructType::create(Ctx, None, "rtype");
 
+    genKernel(cast<Function>(CI->getArgOperand(0)), CI, RetTy);
+    genHost(CI, numDims, numArgs, numOutputs, RetTy);
+}
+
+void GenVISC::genKernel(Function* KernelF, CallInst* CI, StructType* RetTy) {
+    // Make changes to kernel here
+    errs() << "Modifying Node Function: " << KernelF->getName() << "\n";
+    // Change arguments and types
+    // Create the argument type list with added argument types
+    Function::ArgumentListType& argList = KernelF->getArgumentList();
+    std::vector<Type*> argTypes;
+    bool shiftAttr = false;
+    for(Function::arg_iterator ai = KernelF->arg_begin(), ae = KernelF->arg_end();
+            ai != ae; ++ai) {
+    
+      argTypes.push_back(ai->getType());  
+      if(shiftAttr) {
+        shiftAttr = false;
+        shiftAttributes(ai, 1, KernelF);
+
+      }
+      if(ai->getType()->isPointerTy()) {
+        Argument* arg = new Argument(Type::getInt32Ty(KernelF->getContext()),
+            "bytes_"+ai->getName());
+        argList.insertAfter(ai, arg);
+        shiftAttr = true;
+      }
+      
+    }
+    // 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
+    // Find return struct type
+    FunctionType* FTy = FunctionType::get(RetTy, argTypes, KernelF->isVarArg());
+    PointerType* PTy = PointerType::get(FTy, cast<PointerType>(KernelF->getType())->getAddressSpace());
+
+    // Change the function type
+    KernelF->mutateType(PTy);
+    errs() << *KernelF << "\n";
+    
+    // Replace ret void with ret %RetTy undef
+    for(auto& BB: KernelF->getBasicBlockList()) {
+      if(ReturnInst* RI = dyn_cast<ReturnInst>(BB.getTerminator())) {
+        errs() << "Found return inst: "<< *RI << "\n";
+        ReturnInst* newRI = ReturnInst::Create(KernelF->getContext(), UndefValue::get(RetTy));
+        ReplaceInstWithInst(RI, newRI);
+      }
+    }
+    // Replace opencl query intrinsics with visc query intrinsics
+}
+
+
+void GenVISC::genHost(CallInst* CI, unsigned numDims, unsigned numArgs, unsigned numOutputs, StructType* RetTy) {
+    // Make host code changes here
+    errs() << "Modifying Host code for __visc__node call site: " << *CI << "\n";
+    LLVMContext& Ctx = CI->getParent()->getContext();
+    
     // Generate argument struct type (All arguments followed by return struct type)
     std::vector<Type*> ArgList;
-    offset = numDims + 2 + 1;
+    unsigned offset = numDims + 2 + 1;
     for(unsigned i=0; i<numArgs; i++) {
         Type* Ty = CI->getArgOperand(i + offset)->getType();
         ArgList.push_back(Ty);
@@ -347,6 +417,8 @@ void GenVISC::genHost(CallInst* CI) {
     // Get result (optional)
 }
 
+
+
 char GenVISC::ID = 0;
 static RegisterPass<GenVISC> X("genvisc", "Pass to generate VISC IR from LLVM IR (with dummy function calls)", false, false);
 

llvm/test/VISC/MatrixMultiplication/auto_gemm.c

@@ -60,8 +60,8 @@ __attribute__ ((noinline)) int checkResults(float* A, float* B, float* C) {
 // Dummy visc node execution call
 //void __visc__node(void kernel (float*, float*, float*, unsigned, unsigned), int numDims, void* dims, int numInputs, void* inputs, int numOutputs, void* outputs);
 
-//void matrixMul(float* A, float* B, float* C, unsigned k, unsigned n) {
-void matrixMul(float* A, unsigned bytesA, float* B, unsigned bytesB, float* C, unsigned bytesC, unsigned k, unsigned n) {
+void matrixMul(float* A, float* B, float* C, unsigned k, unsigned n) {
+//void matrixMul(float* A, unsigned bytesA, float* B, unsigned bytesB, float* C, unsigned bytesC, unsigned k, unsigned n) {
 
   printf("Entered function\n");
   int tx = get_global_id(0); //2D Global Thread ID x