diff --git a/llvm/include/llvm/IR/DFGraph.h b/llvm/include/llvm/IR/DFGraph.h index 06e8e499b2de306c849643795efecc0756dbff50..b6471a4cc1420458be65467df814780ea28196e4 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 9bfc0ba9abfcc4bbebae10fd7f982dac9ebbfa0d..7186008fe895e8f025603bc9949ee5e46f7bcb48 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 1023c6d35c00f711d3791a42cd4bb5839c23967c..0d50832387b83da17a25e9e06a113ccee8e37f02 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 b7a91a5abbe395674706c0b600352f719d09549e..9c66c69124646b0c0eabdba81f4bda614f408808 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 89c0e376b5f19e6fdef531be572f07245664638b..d844cabfddf6acc5603f7b52fde489f50db7162a 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 b33ffaad88afb09a88546c083555808626a80506..452762267e2b56c738ecc8e6cae0b67d5aa8365b 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 2e6b21b90833db4c3a3597899ac9237cb6fb27be..e21a2c9ff7159c6217b4a7dca672c867238c59de 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; }