From cd8e1fc54413b17e193ce76cc8cba5a8a76ddab9 Mon Sep 17 00:00:00 2001
From: Prakalp Srivastava <psrivas2@illinois.edu>
Date: Thu, 27 Nov 2014 22:08:08 +0000
Subject: [PATCH] (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 | 6 +++---
llvm/projects/visc-rt/visc-rt.cpp | 2 +-
llvm/test/VISC/parboil/benchmarks/mri-q/Makefile | 4 ++--
.../VISC/parboil/benchmarks/mri-q/src/opencl/kernels.cl | 2 +-
llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/macros.h | 6 +++---
llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/main.c | 2 +-
llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/ocl.c | 1 +
7 files changed, 12 insertions(+), 11 deletions(-)
diff --git a/llvm/include/llvm/IR/DFGraph.h b/llvm/include/llvm/IR/DFGraph.h
index 06e8e499b2..b6471a4cc1 100644
--- a/llvm/include/llvm/IR/DFGraph.h
+++ b/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();
diff --git a/llvm/projects/visc-rt/visc-rt.cpp b/llvm/projects/visc-rt/visc-rt.cpp
index 9bfc0ba9ab..7186008fe8 100644
--- a/llvm/projects/visc-rt/visc-rt.cpp
+++ b/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;
diff --git a/llvm/test/VISC/parboil/benchmarks/mri-q/Makefile b/llvm/test/VISC/parboil/benchmarks/mri-q/Makefile
index 1023c6d35c..0d50832387 100644
--- a/llvm/test/VISC/parboil/benchmarks/mri-q/Makefile
+++ b/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
diff --git a/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/kernels.cl b/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/kernels.cl
index b7a91a5abb..9c66c69124 100644
--- a/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/kernels.cl
+++ b/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/kernels.cl
@@ -1,7 +1,7 @@
#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
diff --git a/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/macros.h b/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/macros.h
index 89c0e376b5..d844cabfdd 100644
--- a/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/macros.h
+++ b/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/macros.h
@@ -1,4 +1,4 @@
-#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
diff --git a/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/main.c b/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/main.c
index b33ffaad88..452762267e 100644
--- a/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/main.c
+++ b/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];
diff --git a/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/ocl.c b/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/ocl.c
index 2e6b21b908..e21a2c9ff7 100644
--- a/llvm/test/VISC/parboil/benchmarks/mri-q/src/opencl/ocl.c
+++ b/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;
}
--
GitLab