(1) Updated test cases to make mri-q work. Still not done ...

(2) Changes to DFGraph.h and visc-rt.cpp were not committed before.

M    llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/kernels.cl
M    llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/main.c
M    llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/macros.h
M    llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/ocl.c
M    llvm/test/VISC/parboil/benchmarks/mri-q/Makefile
M    llvm/include/llvm/IR/DFGraph.h
M    llvm/projects/visc-rt/visc-rt.cpp

llvm/include/llvm/IR/DFGraph.h

@@ -335,7 +335,7 @@ public:
   bool isEntryNode();
   bool isExitNode();
   DFEdge* getInDFEdgeAt(unsigned inPort);
-  std::vector<unsigned>* getInArgMap();
+  std::vector<unsigned> getInArgMap();
   int getAncestorHops(DFNode* N);
 
   virtual void applyDFNodeVisitor(DFNodeVisitor &V) = 0;
@@ -567,8 +567,8 @@ DFEdge* DFNode::getInDFEdgeAt(unsigned inPort) {
   return NULL;
 }
 
-std::vector<unsigned>* DFNode::getInArgMap() {
-  std::vector<unsigned>* map = new std::vector<unsigned>(InDFEdges.size());
+std::vector<unsigned> DFNode::getInArgMap() {
+  std::vector<unsigned> map(InDFEdges.size());
   for (unsigned i = 0; i < InDFEdges.size(); i++) {
     DFEdge* E = getInDFEdgeAt(i);
     unsigned pos = E->getSourcePosition();

llvm/projects/visc-rt/visc-rt.cpp

@@ -182,7 +182,7 @@ void* llvm_visc_ptx_executeNode(void* graphID, unsigned workDim , const size_t*
   }
 
   cl_int errcode = clEnqueueNDRangeKernel(Context->clCommandQue,
-      Context->clKernel, workDim, NULL, GlobalWG, NULL, 0, NULL, NULL);
+      Context->clKernel, workDim, NULL, GlobalWG, LocalWG, 0, NULL, NULL);
   DEBUG(cout << "Enqueued kernel\n");
   checkErr(errcode, CL_SUCCESS, "Failure to enqueue kernel");
   return event;

llvm/test/VISC/parboil/benchmarks/mri-q/Makefile

@@ -20,12 +20,12 @@ DATASET_DIR = $(PARBOIL_ROOT)/datasets/$(APP)
 ifeq ($(TEST),small)
   INPUT = $(DATASET_DIR)/small/input/32_32_32_dataset.bin
   REF_OUTPUT = $(DATASET_DIR)/small/output/32_32_32_dataset.out
-  RUNDIR = run/small
+  RUNDIR = run/$(VERSION)/small
   OUTPUT = $(RUNDIR)/32_32_32_dataset.out
 else
   INPUT = $(DATASET_DIR)/large/input/64_64_64_dataset.bin
   REF_OUTPUT = $(DATASET_DIR)/large/output/64_64_64_dataset.out
-  RUNDIR = run/large
+  RUNDIR = run/$(VERSION)/large
   OUTPUT = $(RUNDIR)/64_64_64_dataset.out
 endif
 

llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/kernels.cl

 #include "macros.h"
 
 #define NC  4
-#define KERNEL_Q_K_ELEMS_PER_GRID 256
+#define KERNEL_Q_K_ELEMS_PER_GRID 1024
 #define COARSE_GENERAL
 // #define COARSE_SPEC NC
 

llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/macros.h

-#ifndef __MACROS__
+//#ifndef __MACROS__
 //#define __MACROS__
 
 #define PI   3.1415926535897932384626433832795029f
@@ -9,7 +9,7 @@
 
 #define KERNEL_PHI_MAG_THREADS_PER_BLOCK 256 /* 512 */
 #define KERNEL_Q_THREADS_PER_BLOCK 256
-#define KERNEL_Q_K_ELEMS_PER_GRID 256
+#define KERNEL_Q_K_ELEMS_PER_GRID 1024
 
 struct kValues {
   float Kx;
@@ -18,4 +18,4 @@ struct kValues {
   float PhiMag;
 };
 
-#endif
+//#endif

llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/main.c

@@ -136,7 +136,7 @@ main (int argc, char *argv[]) {
   CHECK_ERROR("clCreateProgramWithSource")
 
   char options[50];
-  sprintf(options,"-I src/opencl_nvidia");
+  sprintf(options,"-I src/opencl");
   clStatus = clBuildProgram(clProgram,0,NULL,options,NULL,NULL);
   if (clStatus != CL_SUCCESS) {
     char buf[4096];

llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/ocl.c

@@ -35,6 +35,7 @@ char* readFile(const char* fileName)
 
 	buffer[size] = 0;
         fclose(fp);
+        printf("Kernel file:\n%s\n", buffer);
         return buffer;
 }