diff --git a/hpvm/test/parboil/benchmarks/pipeline/src/visc_parallel/main.cc b/hpvm/test/parboil/benchmarks/pipeline/src/visc_parallel/main.cc
index 8438b9e31da2ffc9061106bcc08abdc01b7e04a2..b516f860a8c6b8317b11860127b8dc210c66b246 100644
--- a/hpvm/test/parboil/benchmarks/pipeline/src/visc_parallel/main.cc
+++ b/hpvm/test/parboil/benchmarks/pipeline/src/visc_parallel/main.cc
@@ -11,7 +11,7 @@
*/
#include "opencv2/opencv.hpp"
-#include "opencv2/ocl/ocl.hpp"
+#include "opencv2/core/ocl.hpp"
#include <stdio.h>
#include <math.h>
#include <stdlib.h>
@@ -72,10 +72,6 @@ std::string output_window = "GPU Pipeline - Edge Mapping";
extern "C" {
-struct RetStruct {
- size_t bytesRet;
-};
-
struct __attribute__((__packed__)) InStruct {
float* I ;
size_t bytesI;
@@ -247,14 +243,9 @@ void laplacianEstimate(float *Is, size_t bytesIs,
// 3x3 image area
float imageArea[SZB*SZB];
- //int gx = get_global_id(0);
- //int gy = get_global_id(1);
void* thisNode = __visc__getNode();
long gx = __visc__getNodeInstanceID_x(thisNode);
long gy = __visc__getNodeInstanceID_y(thisNode);
- //if(gx == 0 && gy == 0)
- //std::cout << "Entered laplacian\n";
-
int i, j;
if ((gx < n) && (gy < m)) {
@@ -313,9 +304,6 @@ void laplacianEstimate(float *Is, size_t bytesIs,
float laplacian = dilatedPixel + erodedPixel - 2 * imageArea[1 * SZB +1];
L[gy*n+gx] = laplacian;
}
- //OutStruct output = {bytesB, bytesL};
- //if(gx == m-1 && gy == n-1)
- //std::cout << "Exit laplacian\n";
__visc__return(1, bytesL);
}
@@ -358,15 +346,11 @@ void computeZeroCrossings(float *L, size_t bytesL,
// 3x3 image area
float imageArea[SZB][SZB];
- //int gx = get_global_id(0);
- //int gy = get_global_id(1);
void* thisNode = __visc__getNode();
long gx = __visc__getNodeInstanceID_x(thisNode);
long gy = __visc__getNodeInstanceID_y(thisNode);
int i, j;
- //if(gx == 0 && gy == 0)
- //std::cout << "Entered ZC\n";
if ((gx < n) && (gy < m)) {
// Data copy for dilation filter
imageArea[1][1] = L[gy * n + gx] > MIN_BR? MAX_BR : MIN_BR;
@@ -435,9 +419,6 @@ void computeZeroCrossings(float *L, size_t bytesL,
float pixelSign = dilatedPixel - erodedPixel;
S[gy*n+gx] = pixelSign;
}
- //OutStruct output = {bytesB, bytesS};
- //if(gx == n-1 && gy == n-1)
- //std::cout << "Exit ZC\n";
__visc__return(1, bytesS);
}
@@ -523,7 +504,6 @@ void computeGradient(float *Is, size_t bytesIs,
}
G[gloc] = sqrt(Gx*Gx + Gy*Gy);
- //G[gloc] = Gx*Gx + Gy*Gy;
}
__visc__return(1, bytesG);
}
@@ -562,7 +542,6 @@ void computeMaxGradientLeaf(float *G, size_t bytesG,
long m, long n) {
__visc__hint(visc::CPU_TARGET);
- //__visc__hint(visc::CPU_TARGET);
__visc__attributes(1, G, 1, maxG);
void* thisNode = __visc__getNode();
@@ -570,7 +549,6 @@ void computeMaxGradientLeaf(float *G, size_t bytesG,
long lx = __visc__getNodeInstanceID_x(thisNode); // threadIdx.x
long dimx = __visc__getNumNodeInstances_x(thisNode); // blockDim.x
-
// Assume a single thread block
// Thread block iterates over all elements
for (int i = lx + dimx; i < m*n; i+= dimx) {
@@ -591,54 +569,10 @@ void computeMaxGradientLeaf(float *G, size_t bytesG,
__visc__return(1, bytesMaxG);
}
-/*
- * Reduction
- * G : input
- * maxG: output
- * Each static node processes 2*nodeDim elements
- * Need 1D grid, a thread per 2 pixels
- */
-//void computeMaxGradientLeaf(float *G, size_t bytesG,
- //float *maxG, size_t bytesMaxG,
- //int m, int n) {
-
- //__visc__hint(visc::DEVICE);
- //TODO: maxG should be initialized to zero (MIN_BR) every time
- //__visc__attributes(2, G, maxG, 1, maxG);
-
- //void* thisNode = __visc__getNode();
- //void* parentNode = __visc__getParentNode(thisNode);
-
- //int lx = __visc__getNodeInstanceID_x(thisNode);
- //int px = __visc__getNodeInstanceID_x(parentNode);
- //int dimx = __visc__getNumNodeInstances_x(thisNode);
-
- //int gid = lx + 2*px*dimx;
-
- //for (unsigned stride = dimx; stride > 32; stride >>= 1) {
- //if ((gid + stride < m*n) && (lx < stride))
- //if (G[gid + stride] > G[gid])
- //G[gid] = G[gid + stride];
- //__visc__barrier();
- //}
-
- //for (unsigned stride = 32; stride >= 1; stride >>= 1) {
- //if ((gid + stride < m*n) && (lx < stride))
- //if (G[gid + stride] > G[gid])
- //G[gid] = G[gid + stride];
- //}
-
- //if (lx == 0)
- //__visc__atomic_max(maxG,G[gid]);
-
- //__visc__return(1, m);
-//}
-
void computeMaxGradientTB(float *G, size_t bytesG,
float *maxG, size_t bytesMaxG,
long m, long n,
long block_x) {
- //__visc__hint(visc::DEVICE);
__visc__hint(visc::CPU_TARGET);
__visc__attributes(2, G, maxG, 1, maxG);
void* CMGLeafNode = __visc__createNodeND(1, computeMaxGradientLeaf, block_x);
@@ -686,7 +620,6 @@ void rejectZeroCrossings(float *S, size_t bytesS,
float *E, size_t bytesE,
long m, long n) {
__visc__hint(visc::DEVICE);
- //__visc__hint(visc::CPU_TARGET);
__visc__attributes(3, S, G, maxG, 1, E);
void* thisNode = __visc__getNode();
@@ -694,7 +627,6 @@ void rejectZeroCrossings(float *S, size_t bytesS,
int gy = __visc__getNodeInstanceID_y(thisNode);
float mG = *maxG;
- //float mG = 1.39203;
if ((gx < n) && (gy < m)) {
E[gy*n+gx] = ((S[gy*n+gx] > 0.0) && (G[gy*n+gx] > THETA*mG)) ? 1.0 : 0.0 ;
}
@@ -829,18 +761,8 @@ void getNextFrame(VideoCapture& VC, Mat& F) {
cvtColor( F, F, CV_BGR2GRAY );
F.convertTo(F, CV_32F, 1.0/255.0);
-
}
-//void showInOut(Mat& Input, Mat& Output) {
- //Mat in, out;
- //resize(Input, in, Size(512, 768));
- //resize(Output, out, Size(512, 768));
- //imshow(input_window, in);
- //imshow(output_window, out);
-//}
-
-
int main (int argc, char *argv[]) {
struct pb_Parameters *params;
@@ -873,11 +795,6 @@ int main (int argc, char *argv[]) {
NUM_FRAMES = 600;
std::cout << "Number of frames = " << NUM_FRAMES << "\n";
- // Used to store time after each frame computation is completed
- //timeval *timeStamps = (timeval *)malloc(NUM_RUNS*NUM_FRAMES*sizeof(timeval));
- //struct timeval tv_start;
- //unsigned long cnt = 0; // Counter, for the video frames
-
namedWindow(input_window, CV_WINDOW_AUTOSIZE);
namedWindow(output_window, CV_WINDOW_AUTOSIZE);
moveWindow(input_window, POSX_IN, POSY_IN);
@@ -905,7 +822,6 @@ int main (int argc, char *argv[]) {
pb_InitializeTimerSet(&timers);
__visc__init();
- //pb_SwitchToTimer( &timers, pb_TimerID_COMPUTE );
// copy A to device memory
I_sz = src.size[0]*src.size[1]*sizeof(float);
@@ -916,8 +832,6 @@ int main (int argc, char *argv[]) {
1, 1, 1,
1, 1, 1 };
size_t bytesB = 9*sizeof(float);
- //Sx = [-1 0 1 ; -2 0 2 ; -1 0 1 ]
- //Sy = [-1 -2 -1 ; 0 0 0 ; 1 2 1 ]
float Sx[] = { -1, 0, 1,
-2, 0, 2,
-1, 0, 1 };
@@ -941,10 +855,6 @@ int main (int argc, char *argv[]) {
// grid_x should be equal to the number of SMs on GPU. FTX 680 has 8 SMs
grid_x = 1;
- // Copy A and B^T into device memory
- //pb_SwitchToTimer( &timers, pb_TimerID_COMPUTE );
-
- //showInOut(src, E);
Mat in, out;
resize(src, in, Size(HEIGHT, WIDTH));
resize(E, out, Size(HEIGHT, WIDTH));
@@ -952,7 +862,6 @@ int main (int argc, char *argv[]) {
imshow(output_window, out);
waitKey(0);
- //NUM_FRAMES = 20;
pb_SwitchToTimer( &timers, visc_TimerID_COMPUTATION );
struct InStruct* args = (struct InStruct*)malloc (sizeof(InStruct));
packData(args, (float*)src.data, I_sz,
@@ -972,13 +881,6 @@ int main (int argc, char *argv[]) {
// Check if the total elements is a multiple of block size
assert(src.size[0]*src.size[1] % block_x == 0);
- //imshow(input_window, src);
- //imshow(output_window, E);
- //waitKey(0);
-
- // Get the time just before computation starts
- //gettimeofday(&tv_start,NULL);
-
for(unsigned j=0; j<NUM_RUNS; j++) {
std::cout << "Run: " << j << "\n";
void* DFG = __visc__launch(1, edgeDetection, (void*)args);
@@ -986,11 +888,7 @@ int main (int argc, char *argv[]) {
cap = VideoCapture(params->inpFiles[0]);
getNextFrame(cap, src);
- //packData(args, A.data, BlockSize, &matB[i], BlockSize, &matC[i], BlockSize, BlockElements);
-
if(NUM_FRAMES >=2) {
- //__visc__push(DFG, args);
- //__visc__push(DFG, args);
for(int i=0; i<NUM_FRAMES; i++) {
std::cout << "Frame " << i << "\n";
args->I = (float*) src.data;
@@ -1014,51 +912,15 @@ int main (int argc, char *argv[]) {
std::cout << "Returned size: " << *(size_t *)ret
<< " expected " << I_sz << '\n';
- //llvm_visc_request_mem(E.data, I_sz);
- //std::cout << "Show E" << "\n";
- //imshow(window_name, E);
- //waitKey(0);
- //llvm_visc_request_mem(src.data, I_sz);
- //llvm_visc_request_mem(Is.data, I_sz);
- //llvm_visc_request_mem(L.data, I_sz);
- //llvm_visc_request_mem(S.data, I_sz);
- //llvm_visc_request_mem(G.data, I_sz);
llvm_visc_request_mem(maxG, bytesMaxG);
llvm_visc_request_mem(E.data, I_sz);
- //std::cout << "src.data = " << (float*)src.data << "\n";
- //std::cout << "Is.data = " << (float*)Is.data << "\n";
- //std::cout << "L.data = " << (float*)L.data << "\n";
- //std::cout << "S.data = " << (float*)S.data << "\n";
- //std::cout << "G.data = " << (float*)G.data << "\n";
- //std::cout << "E.data = " << (float*)E.data << "\n";
- //std::cout << "Max G = " << *maxG << "\n";
-
- //gettimeofday(&timeStamps[cnt], NULL);
- //cnt++;
+
Mat in, out;
resize(src, in, Size(HEIGHT, WIDTH));
- //std::cout << "Show E\n";
resize(E, out, Size(HEIGHT, WIDTH));
imshow(output_window, out);
imshow(input_window, in);
waitKey(1);
- //waitKey(0);
- //std::cout << "Show Is\n";
- //resize(Is, out, Size(HEIGHT, WIDTH));
- //imshow(output_window, out);
- //waitKey(0);
- //std::cout << "Show L\n";
- //resize(L, out, Size(HEIGHT, WIDTH));
- //imshow(output_window, out);
- //waitKey(0);
- //std::cout << "Show S\n";
- //resize(S, out, Size(HEIGHT, WIDTH));
- //imshow(output_window, out);
- //waitKey(0);
- //std::cout << "Show G\n";
- //resize(G, out, Size(HEIGHT, WIDTH));
- //imshow(output_window, out);
- //waitKey(0);
llvm_visc_untrack_mem(src.data);
llvm_visc_untrack_mem(Is.data);
@@ -1073,72 +935,17 @@ int main (int argc, char *argv[]) {
llvm_visc_untrack_mem(Sy);
getNextFrame(cap, src);
-
}
- //__visc__pop(DFG);
- //__visc__pop(DFG);
}
else {
__visc__push(DFG, args);
__visc__pop(DFG);
}
-
-
__visc__wait(DFG);
}
-
pb_SwitchToTimer(&timers, pb_TimerID_NONE);
-
-
pb_PrintTimerSet(&timers);
-
- //const char *fn = "timestamps.txt";
- //std::ofstream outfile;
- //outfile.open(fn);
- //if (!outfile.is_open()) {
- //std::cout << "Failed to open " << fn << " for writing\n";
- //}
- //for (unsigned long i = 0; i < cnt; i++) {
- //double elapsed = (timeStamps[i].tv_sec - tv_start.tv_sec) +
- //((timeStamps[i].tv_usec - tv_start.tv_usec)/1000000.0);
- //outfile << elapsed << "\n";
- //}
- //free(timeStamps);
- //outfile.close();
__visc__cleanup();
-
- //if (params->outFile) {
-
- /* Write C to file */
- //pb_SwitchToTimer(&timers, pb_TimerID_IO);
- //writeColMajorMatrixFile(params->outFile,
- //src.size[0], src.size[1], matE);
- //}
-
- //std::cout << "Show Is" << "\n";
- //Mat output(src.size[0], src.size[1], CV_32F);
- //imshow(output_window, Is);
- //waitKey(0);
-
- //std::cout << "Show G" << "\n";
- //imshow(output_window, L);
- //waitKey(0);
-
- //std::cout << "Show L" << "\n";
- //imshow(output_window, S);
- //waitKey(0);
-
- //std::cout << "Show S" << "\n";
- //imshow(output_window, G);
- //waitKey(0);
-
- //std::cout << "Show E" << "\n";
- //imshow(output_window, E);
- //waitKey(0);
-
- //double GPUtime = pb_GetElapsedTime(&(timers.timers[pb_TimerID_KERNEL]));
- //std::cout<< "GFLOPs = " << 2.* src.size[0] * src.size[1] * src.size[1]/GPUtime/1e9 << std::endl;
pb_FreeParameters(params);
-
return 0;
}