Checkpoint commit in implementing ato test generator

llvm/include/llvm/GenVISC/GenVISC.h

@@ -28,14 +28,12 @@ struct GenVISC : public ModulePass {
 private:
   // Member variables
 
-
   // Functions
 
 public:
   // Functions
   virtual bool runOnModule(Module &M);
 
-  static bool isVISCNodeCall(Instruction* I);
   void genKernel(Function* KernelFunction, CallInst* CI);
   void genHost(CallInst* CI);
 };

llvm/lib/Transforms/GenVISC/GenVISC.cpp

@@ -21,6 +21,163 @@ STATISTIC(viscCallCtr, "Counts number of visc launch node calls greeted");
 
 namespace genvisc {
 
+// Helper Functions
+static unsigned getNumericValue(Value* V) {
+  assert(isa<ConstantInt>(V)
+      && "Number of arguments should be a constant");
+  return cast<ConstantInt>(V)->getZExtValue();
+}
+
+static void addArgs(Function* F, unsigned numArgs, std::string names[]) {
+  if(numArgs == 0) return;
+  for (unsigned i = 0; i < numArgs; ++i) {
+    new Argument(Type::getInt32Ty(F->getContext()), names[i], 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);
+}
+
+
+static void marshallArguments(unsigned numArgs, unsigned argOffset, Value* InStruct, CallInst* CI) {
+  assert(isa<Function>(CI->getArgOperand(0))
+      && "InStruct should be of type struct");
+  Function* KernelF = cast<Function>(CI->getArgOperand(0));
+  errs() << "Kernel F = " << KernelF->getName() << "\n"; 
+  LLVMContext& Ctx = CI->getParent()->getContext();
+  Constant* IntZero = ConstantInt::get(Type::getInt32Ty(Ctx), 0);
+  Function::arg_iterator ai = KernelF->arg_begin();
+  Function::arg_iterator ae = KernelF->arg_end();
+
+  for(unsigned i = 0; i < numArgs && ai != ae; i++, ai++) {
+    errs() << "Argument: " << ai->getName() << "\n";
+    Value* arg = CI->getArgOperand(i+argOffset);
+    // Create constant int (i)
+    Constant* Int_i = ConstantInt::get(Type::getInt32Ty(Ctx), i);
+    // Get Element pointer instruction
+    Value* GEPIndices[] = { IntZero, Int_i };
+    GetElementPtrInst* GEP = GetElementPtrInst::Create(InStruct, 
+                                                      ArrayRef<Value*>(GEPIndices, 2),
+                                                      InStruct->getName()+"."+ai->getName(),
+                                                      CI);
+    // Store instruction
+    new StoreInst(arg, GEP, CI);
+  }
+}
+
+static Function* genRootNode(Function* KernelF, unsigned numDims, unsigned dimOffset, CallInst* CI) {
+  // Create new function with the same type
+  Module* module = KernelF->getParent();
+  Function* Root = Function::Create(KernelF->getFunctionType(),
+      KernelF->getLinkage(), KernelF->getName()+"Root", module);
+
+  // Loop over the arguments, copying the names of arguments over.
+  Function::arg_iterator dest_iterator = Root->arg_begin();
+  for (Function::const_arg_iterator i = KernelF->arg_begin(), e = KernelF->arg_end();
+       i != e; ++i) {
+    dest_iterator->setName(i->getName()); // Copy the name over...
+    // Add mapping to VMap and increment dest iterator
+    // VMap[i] = dest_iterator++;
+  }
+
+  // Add extra dimesnion arguments
+  std::string dimNames[] = {"dim_x", "dim_y", "dim_z"};
+  addArgs(Root, numDims, dimNames);
+  // Create a basic block in this function
+  BasicBlock *BB = BasicBlock::Create(Root->getContext(), "entry", Root);
+  ReturnInst* RI = ReturnInst::Create(Root->getContext(),
+                                             UndefValue::get(Root->getReturnType()), BB);
+
+  // Insert createNode intrinsic
+  std::vector<Type*> argTy;
+  std::vector<Value*> args;
+
+  Value* NodeF = ConstantExpr::getPointerCast(KernelF, Type::getInt8PtrTy(module->getContext()));
+  argTy.push_back(NodeF->getType());
+  args.push_back(NodeF);
+  Intrinsic::ID createNodeXD;
+  StringRef createNodeX;
+  Function::arg_iterator ai = Root->arg_end();
+  for(unsigned i=0; i<numDims; i++, ai--);
+  errs() << "Iterator at: " << *ai << "\n";
+  for(unsigned i=0; i < numDims; i++) {
+    argTy.push_back(ai->getType());
+    args.push_back(ai);
+    ai++;
+  }
+  errs() << "Number of dims = " << numDims << "\n";
+  switch(numDims) {
+    case 0:
+      createNodeXD = Intrinsic::visc_createNode;
+      createNodeX = "llvm.visc.createNode";
+      break;
+    case 1:
+      createNodeXD = Intrinsic::visc_createNode1D;
+      createNodeX = "llvm.visc.createNode1D";
+      break;
+    case 2:
+      createNodeXD = Intrinsic::visc_createNode2D;
+      createNodeX = "llvm.visc.createNode2D";
+      break;
+    case 3:
+      createNodeXD = Intrinsic::visc_createNode3D;
+      createNodeX = "llvm.visc.createNode3D";
+      break;
+    default:
+      assert(false && "Invalid number of dimensions!");
+      break;
+  };
+  errs() << "*** Exited genRoot once!! ***\n";
+  FunctionType* CreateNodeFTy = FunctionType::get(Type::getInt8PtrTy(module->getContext()),
+                                                  argTy,
+                                                  false);
+  Function* CreateNodeF = cast<Function>(module->getOrInsertFunction(createNodeX,
+                                         CreateNodeFTy));
+  errs() << "Function chosen:\n" << *CreateNodeF << "\n";
+  CallInst *CreateNodeCall = CallInst::Create(CreateNodeF, args, KernelF->getName()+".node", RI);
+  errs() << "Generate call: " << *CreateNodeCall << "\n";
+  // Bind intrinsics
+  return Root; 
+}
+
+
+static bool isVISCNodeCall(Instruction* I) {
+  if(!isa<CallInst>(I))
+    return false;
+  CallInst* CI = cast<CallInst>(I);
+  DEBUG(errs() << *I << "\n");
+  return (CI->getCalledValue()->stripPointerCasts()->getName()).equals("__visc__node");
+}
+
+static bool isVISCLaunchCall(Instruction* I) {
+  if(!isa<CallInst>(I))
+    return false;
+  CallInst* CI = cast<CallInst>(I);
+  DEBUG(errs() << *I << "\n");
+  return (CI->getCalledValue()->stripPointerCasts()->getName()).equals("__visc__launch");
+}
+
+static bool isVISCEdgeCall(Instruction* I) {
+  if(!isa<CallInst>(I))
+    return false;
+  CallInst* CI = cast<CallInst>(I);
+  DEBUG(errs() << *I << "\n");
+  return (CI->getCalledValue()->stripPointerCasts()->getName()).equals("__visc__edge");
+}
+
+// Public Functions of GenVISC pass
 bool GenVISC::runOnModule(Module &M) {
 
   errs() << "-------- Searching for launch sites ----------\n";
@@ -36,7 +193,8 @@ bool GenVISC::runOnModule(Module &M) {
         CallInst* CI = cast<CallInst>(I);
         errs() << "Found visc node call\n";
         assert(CI->getNumArgOperands() >= 4
-            && "__visc_node call should have atleast 4 arguments!");
+            && "__visc__node call should have atleast 4 arguments!");
+        errs() << "Kernel Function = " << *CI->getArgOperand(1) << "\n";
         genKernel(cast<Function>(CI->getArgOperand(0)), CI);
         genHost(CI);
       }
@@ -46,22 +204,6 @@ bool GenVISC::runOnModule(Module &M) {
   return false; //TODO: What does returning "false" mean?
 }
 
-// Helper Functions
-static unsigned getNumericValue(Value* V) {
-  assert(isa<ConstantInt>(V)
-      && "Number of arguments should be a constant");
-  return cast<ConstantInt>(V)->getZExtValue();
-}
-
-// Public Functions of GenVISC pass
-bool GenVISC::isVISCNodeCall(Instruction* I) {
-  if(!isa<CallInst>(I))
-    return false;
-  CallInst* CI = cast<CallInst>(I);
-  DEBUG(errs() << *I << "\n");
-  return (CI->getCalledValue()->stripPointerCasts()->getName()).equals("__visc__node");
-}
-
 void GenVISC::genKernel(Function* KernelF, CallInst* CI) {
   // Make changes to kernel here
   errs() << "Modifying Node Function: " << KernelF->getName() << "\n";
@@ -70,34 +212,62 @@ void GenVISC::genKernel(Function* KernelF, CallInst* CI) {
 void GenVISC::genHost(CallInst* CI) {
   // Make host code changes here
   errs() << "Modifying Host code for __visc__node call site: " << *CI << "\n";
+  LLVMContext& Ctx = CI->getParent()->getContext();
 
-  // Find number of dimensions
   unsigned offset = 1; // argument at offset 1 is the number of dimensions
+  // Find number of dimensions
   assert(CI->getNumArgOperands() > offset
-      && "Too few arguments for __visc_node call!");
+      && "Too few arguments for __visc__node call!");
   unsigned numDims = getNumericValue(CI->getOperand(offset));
   errs () << "Num of dimensions = " << numDims << "\n";
+  
 
   // Find number of arguments
   offset += numDims + 1; // skip the dimesnions
   assert(CI->getNumArgOperands() > offset
-      && "Too few arguments for __visc_node call!");
+      && "Too few arguments for __visc__node call!");
   unsigned numArgs = getNumericValue(CI->getArgOperand(offset));
   errs () << "Num of kernel arguments = " << numArgs << "\n";
 
   // Find number of outputs
   offset += numArgs + 1; // skip the kernel arguments
   assert(CI->getNumArgOperands() > offset
-      && "Too few arguments for __visc_node call!");
+      && "Too few arguments for __visc__node call!");
   unsigned numOutputs = getNumericValue(CI->getArgOperand(offset));
   errs () << "Num of kernel outputs = " << numOutputs << "\n";
 
   // Find return struct type
+  assert(numOutputs == 0 && "Not handled case where number of outputs is non-zero!");
+  StructType* RetTy = StructType::create(Ctx, None, "rtype");
+
   // Generate argument struct type (All arguments followed by return struct type)
+  std::vector<Type*> ArgList;
+  offset = numDims + 2 + 1;
+  for(unsigned i=0; i<numArgs; i++) {
+    Type* Ty = CI->getArgOperand(i + offset)->getType();
+    ArgList.push_back(Ty);
+    /* Not required as the source code visc node call has all the correct
+     * arguments
+     * if(Ty->isPointerTy()) {
+      ArgList.push_back(IntegerType::getInt64Ty(Ctx));
+    }*/
+  }
+  ArgList.push_back(RetTy);
+  StructType* ArgStructTy = StructType::create(ArgList, "struct.arg", true);
+
   // Insert alloca inst for this argument struct type
+  AllocaInst* AI = new AllocaInst(ArgStructTy, "in.addr", CI);
   // Marshall all input arguments into argument struct type
+  marshallArguments(numArgs, offset, AI, CI);
   // Type cast argument struct to i8*
-  // Replace CI with launch call
+  CastInst* BI = BitCastInst::CreatePointerCast(AI,
+                                                Type::getInt8PtrTy(Ctx),
+                                                "args",
+                                                CI);
+  // Create a root funtion which has this as internal node
+  Function* Root = genRootNode(cast<Function>(CI->getArgOperand(0)), numDims, 2, CI);
+  // Replace CI with launch call to a Root function
+  
   // Add wait call
   // Get result (optional)
 }

llvm/test/VISC/MatrixMultiplication/auto_gemm.c

@@ -60,7 +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, 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
@@ -85,7 +86,6 @@ void matrixMul(float* A, float* B, float* C, unsigned k, unsigned n) {
   printf("Result written to C\n");
 }
 
-
 // Main
 int main(int argc, char** argv) {
 
@@ -132,7 +132,7 @@ int main(int argc, char** argv) {
    // Compute using OpenCL
   //matrixMul(h_A, h_B, h_C, WA, WB);
   //__visc__node(matrixMul, 2, WB, HA, 3, h_A, h_B, h_C, 0);
-  __visc__node(matrixMul, 2, WB, HA, 5, h_A, h_B, h_C, WA, WB, 0);
+  __visc__node(matrixMul, 2, WB, HA, 8, h_A, bytes_A, h_B, bytes_B, h_C, bytes_C, WA, WB, 0);
   if(checkResults(h_A, h_B, h_C))
     printf("\nPass!\n");
   else