Skip to content
Snippets Groups Projects
Commit 2ae821f7 authored by Yifan Zhao's avatar Yifan Zhao
Browse files

Cleaned up unused code

parent de048533
No related branches found
No related tags found
No related merge requests found
Hide whitespace changes
Inline Side-by-side
hpvm/test/parboil/benchmarks/pipeline/src/visc_parallel/main.cc
+ 2
− 195
View file @ 2ae821f7
...@@ -11,7 +11,7 @@ ...@@ -11,7 +11,7 @@
*/ */
#include "opencv2/opencv.hpp" #include "opencv2/opencv.hpp"
#include "opencv2/ocl/ocl.hpp" #include "opencv2/core/ocl.hpp"
#include <stdio.h> #include <stdio.h>
#include <math.h> #include <math.h>
#include <stdlib.h> #include <stdlib.h>
...@@ -72,10 +72,6 @@ std::string output_window = "GPU Pipeline - Edge Mapping"; ...@@ -72,10 +72,6 @@ std::string output_window = "GPU Pipeline - Edge Mapping";
extern "C" { extern "C" {
struct RetStruct {
size_t bytesRet;
};
struct __attribute__((__packed__)) InStruct { struct __attribute__((__packed__)) InStruct {
float* I ; float* I ;
size_t bytesI; size_t bytesI;
...@@ -247,14 +243,9 @@ void laplacianEstimate(float *Is, size_t bytesIs, ...@@ -247,14 +243,9 @@ void laplacianEstimate(float *Is, size_t bytesIs,
// 3x3 image area // 3x3 image area
float imageArea[SZB*SZB]; float imageArea[SZB*SZB];
//int gx = get_global_id(0);
//int gy = get_global_id(1);
void* thisNode = __visc__getNode(); void* thisNode = __visc__getNode();
long gx = __visc__getNodeInstanceID_x(thisNode); long gx = __visc__getNodeInstanceID_x(thisNode);
long gy = __visc__getNodeInstanceID_y(thisNode); long gy = __visc__getNodeInstanceID_y(thisNode);
//if(gx == 0 && gy == 0)
//std::cout << "Entered laplacian\n";
int i, j; int i, j;
if ((gx < n) && (gy < m)) { if ((gx < n) && (gy < m)) {
...@@ -313,9 +304,6 @@ void laplacianEstimate(float *Is, size_t bytesIs, ...@@ -313,9 +304,6 @@ void laplacianEstimate(float *Is, size_t bytesIs,
float laplacian = dilatedPixel + erodedPixel - 2 * imageArea[1 * SZB +1]; float laplacian = dilatedPixel + erodedPixel - 2 * imageArea[1 * SZB +1];
L[gy*n+gx] = laplacian; 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); __visc__return(1, bytesL);
} }
...@@ -358,15 +346,11 @@ void computeZeroCrossings(float *L, size_t bytesL, ...@@ -358,15 +346,11 @@ void computeZeroCrossings(float *L, size_t bytesL,
// 3x3 image area // 3x3 image area
float imageArea[SZB][SZB]; float imageArea[SZB][SZB];
//int gx = get_global_id(0);
//int gy = get_global_id(1);
void* thisNode = __visc__getNode(); void* thisNode = __visc__getNode();
long gx = __visc__getNodeInstanceID_x(thisNode); long gx = __visc__getNodeInstanceID_x(thisNode);
long gy = __visc__getNodeInstanceID_y(thisNode); long gy = __visc__getNodeInstanceID_y(thisNode);
int i, j; int i, j;
//if(gx == 0 && gy == 0)
//std::cout << "Entered ZC\n";
if ((gx < n) && (gy < m)) { if ((gx < n) && (gy < m)) {
// Data copy for dilation filter // Data copy for dilation filter
imageArea[1][1] = L[gy * n + gx] > MIN_BR? MAX_BR : MIN_BR; imageArea[1][1] = L[gy * n + gx] > MIN_BR? MAX_BR : MIN_BR;
...@@ -435,9 +419,6 @@ void computeZeroCrossings(float *L, size_t bytesL, ...@@ -435,9 +419,6 @@ void computeZeroCrossings(float *L, size_t bytesL,
float pixelSign = dilatedPixel - erodedPixel; float pixelSign = dilatedPixel - erodedPixel;
S[gy*n+gx] = pixelSign; 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); __visc__return(1, bytesS);
} }
...@@ -523,7 +504,6 @@ void computeGradient(float *Is, size_t bytesIs, ...@@ -523,7 +504,6 @@ void computeGradient(float *Is, size_t bytesIs,
} }
G[gloc] = sqrt(Gx*Gx + Gy*Gy); G[gloc] = sqrt(Gx*Gx + Gy*Gy);
//G[gloc] = Gx*Gx + Gy*Gy;
} }
__visc__return(1, bytesG); __visc__return(1, bytesG);
} }
...@@ -562,7 +542,6 @@ void computeMaxGradientLeaf(float *G, size_t bytesG, ...@@ -562,7 +542,6 @@ void computeMaxGradientLeaf(float *G, size_t bytesG,
long m, long n) { long m, long n) {
__visc__hint(visc::CPU_TARGET); __visc__hint(visc::CPU_TARGET);
//__visc__hint(visc::CPU_TARGET);
__visc__attributes(1, G, 1, maxG); __visc__attributes(1, G, 1, maxG);
void* thisNode = __visc__getNode(); void* thisNode = __visc__getNode();
...@@ -570,7 +549,6 @@ void computeMaxGradientLeaf(float *G, size_t bytesG, ...@@ -570,7 +549,6 @@ void computeMaxGradientLeaf(float *G, size_t bytesG,
long lx = __visc__getNodeInstanceID_x(thisNode); // threadIdx.x long lx = __visc__getNodeInstanceID_x(thisNode); // threadIdx.x
long dimx = __visc__getNumNodeInstances_x(thisNode); // blockDim.x long dimx = __visc__getNumNodeInstances_x(thisNode); // blockDim.x
// Assume a single thread block // Assume a single thread block
// Thread block iterates over all elements // Thread block iterates over all elements
for (int i = lx + dimx; i < m*n; i+= dimx) { for (int i = lx + dimx; i < m*n; i+= dimx) {
...@@ -591,54 +569,10 @@ void computeMaxGradientLeaf(float *G, size_t bytesG, ...@@ -591,54 +569,10 @@ void computeMaxGradientLeaf(float *G, size_t bytesG,
__visc__return(1, bytesMaxG); __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, void computeMaxGradientTB(float *G, size_t bytesG,
float *maxG, size_t bytesMaxG, float *maxG, size_t bytesMaxG,
long m, long n, long m, long n,
long block_x) { long block_x) {
//__visc__hint(visc::DEVICE);
__visc__hint(visc::CPU_TARGET); __visc__hint(visc::CPU_TARGET);
__visc__attributes(2, G, maxG, 1, maxG); __visc__attributes(2, G, maxG, 1, maxG);
void* CMGLeafNode = __visc__createNodeND(1, computeMaxGradientLeaf, block_x); void* CMGLeafNode = __visc__createNodeND(1, computeMaxGradientLeaf, block_x);
...@@ -686,7 +620,6 @@ void rejectZeroCrossings(float *S, size_t bytesS, ...@@ -686,7 +620,6 @@ void rejectZeroCrossings(float *S, size_t bytesS,
float *E, size_t bytesE, float *E, size_t bytesE,
long m, long n) { long m, long n) {
__visc__hint(visc::DEVICE); __visc__hint(visc::DEVICE);
//__visc__hint(visc::CPU_TARGET);
__visc__attributes(3, S, G, maxG, 1, E); __visc__attributes(3, S, G, maxG, 1, E);
void* thisNode = __visc__getNode(); void* thisNode = __visc__getNode();
...@@ -694,7 +627,6 @@ void rejectZeroCrossings(float *S, size_t bytesS, ...@@ -694,7 +627,6 @@ void rejectZeroCrossings(float *S, size_t bytesS,
int gy = __visc__getNodeInstanceID_y(thisNode); int gy = __visc__getNodeInstanceID_y(thisNode);
float mG = *maxG; float mG = *maxG;
//float mG = 1.39203;
if ((gx < n) && (gy < m)) { 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 ; 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) { ...@@ -829,18 +761,8 @@ void getNextFrame(VideoCapture& VC, Mat& F) {
cvtColor( F, F, CV_BGR2GRAY ); cvtColor( F, F, CV_BGR2GRAY );
F.convertTo(F, CV_32F, 1.0/255.0); 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[]) { int main (int argc, char *argv[]) {
struct pb_Parameters *params; struct pb_Parameters *params;
...@@ -873,11 +795,6 @@ int main (int argc, char *argv[]) { ...@@ -873,11 +795,6 @@ int main (int argc, char *argv[]) {
NUM_FRAMES = 600; NUM_FRAMES = 600;
std::cout << "Number of frames = " << NUM_FRAMES << "\n"; 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(input_window, CV_WINDOW_AUTOSIZE);
namedWindow(output_window, CV_WINDOW_AUTOSIZE); namedWindow(output_window, CV_WINDOW_AUTOSIZE);
moveWindow(input_window, POSX_IN, POSY_IN); moveWindow(input_window, POSX_IN, POSY_IN);
...@@ -905,7 +822,6 @@ int main (int argc, char *argv[]) { ...@@ -905,7 +822,6 @@ int main (int argc, char *argv[]) {
pb_InitializeTimerSet(&timers); pb_InitializeTimerSet(&timers);
__visc__init(); __visc__init();
//pb_SwitchToTimer( &timers, pb_TimerID_COMPUTE );
// copy A to device memory // copy A to device memory
I_sz = src.size[0]*src.size[1]*sizeof(float); I_sz = src.size[0]*src.size[1]*sizeof(float);
...@@ -916,8 +832,6 @@ int main (int argc, char *argv[]) { ...@@ -916,8 +832,6 @@ int main (int argc, char *argv[]) {
1, 1, 1, 1, 1, 1,
1, 1, 1 }; 1, 1, 1 };
size_t bytesB = 9*sizeof(float); 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, float Sx[] = { -1, 0, 1,
-2, 0, 2, -2, 0, 2,
-1, 0, 1 }; -1, 0, 1 };
...@@ -941,10 +855,6 @@ int main (int argc, char *argv[]) { ...@@ -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 should be equal to the number of SMs on GPU. FTX 680 has 8 SMs
grid_x = 1; grid_x = 1;
// Copy A and B^T into device memory
//pb_SwitchToTimer( &timers, pb_TimerID_COMPUTE );
//showInOut(src, E);
Mat in, out; Mat in, out;
resize(src, in, Size(HEIGHT, WIDTH)); resize(src, in, Size(HEIGHT, WIDTH));
resize(E, out, Size(HEIGHT, WIDTH)); resize(E, out, Size(HEIGHT, WIDTH));
...@@ -952,7 +862,6 @@ int main (int argc, char *argv[]) { ...@@ -952,7 +862,6 @@ int main (int argc, char *argv[]) {
imshow(output_window, out); imshow(output_window, out);
waitKey(0); waitKey(0);
//NUM_FRAMES = 20;
pb_SwitchToTimer( &timers, visc_TimerID_COMPUTATION ); pb_SwitchToTimer( &timers, visc_TimerID_COMPUTATION );
struct InStruct* args = (struct InStruct*)malloc (sizeof(InStruct)); struct InStruct* args = (struct InStruct*)malloc (sizeof(InStruct));
packData(args, (float*)src.data, I_sz, packData(args, (float*)src.data, I_sz,
...@@ -972,13 +881,6 @@ int main (int argc, char *argv[]) { ...@@ -972,13 +881,6 @@ int main (int argc, char *argv[]) {
// Check if the total elements is a multiple of block size // Check if the total elements is a multiple of block size
assert(src.size[0]*src.size[1] % block_x == 0); 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++) { for(unsigned j=0; j<NUM_RUNS; j++) {
std::cout << "Run: " << j << "\n"; std::cout << "Run: " << j << "\n";
void* DFG = __visc__launch(1, edgeDetection, (void*)args); void* DFG = __visc__launch(1, edgeDetection, (void*)args);
...@@ -986,11 +888,7 @@ int main (int argc, char *argv[]) { ...@@ -986,11 +888,7 @@ int main (int argc, char *argv[]) {
cap = VideoCapture(params->inpFiles[0]); cap = VideoCapture(params->inpFiles[0]);
getNextFrame(cap, src); getNextFrame(cap, src);
//packData(args, A.data, BlockSize, &matB[i], BlockSize, &matC[i], BlockSize, BlockElements);
if(NUM_FRAMES >=2) { if(NUM_FRAMES >=2) {
//__visc__push(DFG, args);
//__visc__push(DFG, args);
for(int i=0; i<NUM_FRAMES; i++) { for(int i=0; i<NUM_FRAMES; i++) {
std::cout << "Frame " << i << "\n"; std::cout << "Frame " << i << "\n";
args->I = (float*) src.data; args->I = (float*) src.data;
...@@ -1014,51 +912,15 @@ int main (int argc, char *argv[]) { ...@@ -1014,51 +912,15 @@ int main (int argc, char *argv[]) {
std::cout << "Returned size: " << *(size_t *)ret std::cout << "Returned size: " << *(size_t *)ret
<< " expected " << I_sz << '\n'; << " 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(maxG, bytesMaxG);
llvm_visc_request_mem(E.data, I_sz); 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; Mat in, out;
resize(src, in, Size(HEIGHT, WIDTH)); resize(src, in, Size(HEIGHT, WIDTH));
//std::cout << "Show E\n";
resize(E, out, Size(HEIGHT, WIDTH)); resize(E, out, Size(HEIGHT, WIDTH));
imshow(output_window, out); imshow(output_window, out);
imshow(input_window, in); imshow(input_window, in);
waitKey(1); 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(src.data);
llvm_visc_untrack_mem(Is.data); llvm_visc_untrack_mem(Is.data);
...@@ -1073,72 +935,17 @@ int main (int argc, char *argv[]) { ...@@ -1073,72 +935,17 @@ int main (int argc, char *argv[]) {
llvm_visc_untrack_mem(Sy); llvm_visc_untrack_mem(Sy);
getNextFrame(cap, src); getNextFrame(cap, src);
} }
//__visc__pop(DFG);
//__visc__pop(DFG);
} }
else { else {
__visc__push(DFG, args); __visc__push(DFG, args);
__visc__pop(DFG); __visc__pop(DFG);
} }
__visc__wait(DFG); __visc__wait(DFG);
} }
pb_SwitchToTimer(&timers, pb_TimerID_NONE); pb_SwitchToTimer(&timers, pb_TimerID_NONE);
pb_PrintTimerSet(&timers); 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(); __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); pb_FreeParameters(params);
return 0; return 0;
} }
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment