Fixed compiling of spmv parboil

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

@@ -29,7 +29,7 @@ cl_context globalGPUContext;
 
 static inline void checkErr(cl_int err, cl_int success, const char * name) {
   if (err != success) {
-    cout << "ERROR:" << name << "\n";
+  cout << "ERROR:" << name << "\n";
     exit(EXIT_FAILURE);
   }
 }
@@ -39,17 +39,17 @@ void* llvm_visc_x86_launch(void* (*rootFunc)(void*), void* arguments) {
   DFNodeContext_X86 *Context = (DFNodeContext_X86 *) malloc(sizeof(DFNodeContext_X86));
   int err;
   if((err = pthread_create(&Context->threadID, NULL, rootFunc, arguments)) != 0)
-    cout << "Failed to create pthread. Error code = " << err << "\n";
+  cout << "Failed to create pthread. Error code = " << err << "\n";
   return Context;
 }
 
 extern "C"
 void llvm_visc_x86_wait(void* graphID) {
-  cout << "Waiting for pthread to finish ...\n";
+  DEBUG(cout << "Waiting for pthread to finish ...\n");
   DFNodeContext_X86* Context = (DFNodeContext_X86*) graphID;
   pthread_join(Context->threadID, NULL);
   free(graphID);
-  cout << "\t... pthread Done!\n";
+  DEBUG(cout << "\t... pthread Done!\n");
 }
 
 extern "C"
@@ -85,7 +85,7 @@ void* llvm_visc_ptx_initContext() {
   globalGPUContext = clCreateContextFromType(properties, CL_DEVICE_TYPE_GPU,
                                          NULL, NULL, &errcode);
   free(platforms);
-  cout << "\tContext " << globalGPUContext << "\n";
+DEBUG(cout << "\tContext " << globalGPUContext << "\n");
   checkErr(errcode, CL_SUCCESS, "Failure to create GPU context");
 
   return globalGPUContext;

llvm/test/VISC/parboil/benchmarks/lbm/src/visc/main.c

@@ -314,13 +314,6 @@ void MAIN_initialize( const MAIN_Param* param ) {
 	LBM_initializeSpecialCellsForLDC( srcGrid );
 	LBM_initializeSpecialCellsForLDC( dstGrid );
 	
-        pb_SwitchToTimer(&timers, pb_TimerID_COPY);
-	
-	//Setup DEVICE datastructures
-	
-	//Initialize DEVICE datastructures
-	
-        pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
 	LBM_showGridStatistics( srcGrid );
 
 	//LBM_freeGrid( (float**) &srcGrid );

llvm/test/VISC/parboil/benchmarks/spmv/Makefile

-PARBOIL_ROOT = /home/kotsifa2/visc/llvm/test/VISC/parboil
+PARBOIL_ROOT = $(LLVM_SRC_ROOT)/test/VISC/parboil
 APP = spmv
 
 # Default compile visc
@@ -37,7 +37,7 @@ else
   OUTPUT = $(RUNDIR)/Dubcova3.mtx.out
 endif
 
-ARGS = -i $(INPUT1) $(INPUT2) -o $(OUTPUT)
+ARGS = -i $(INPUT1),$(INPUT2) -o $(OUTPUT)
 TOOL = tools/compare-output
 
 include $(PARBOIL_ROOT)/common/mk/Makefile

llvm/test/VISC/parboil/benchmarks/spmv/common_src/convert-dataset/commontools.mk

 
-SRCDIR_OBJS += ../../common_src/convert-dataset/convert_dataset.o ../../common_src/convert-dataset/mmio.o
+SRCDIR_OBJS += ../../common_src/convert-dataset/convert_dataset.ll ../../common_src/convert-dataset/mmio.ll

llvm/test/VISC/parboil/benchmarks/spmv/src/visc/Makefile

@@ -2,9 +2,10 @@
 
 LANGUAGE=visc
 TOOLS_SRC=common_src/convert-dataset
-SRCDIR_OBJS=main.visc.ll gpu_info.ll file.ll
+SRCDIR_OBJS=gpu_info.ll file.ll
+VISC_OBJS=main.visc.ll
 APP_CUDALDFLAGS=-lm
-APP_CFLAGS=-ffast-math -g3 -O3 -I$(TOOLS_SRC)
-APP_CXXFLAGS=-ffast-math -g3 -O3
+APP_CFLAGS=-ffast-math -O3 -I$(TOOLS_SRC)
+APP_CXXFLAGS=-ffast-math -O3
 
 include $(TOOLS_SRC)/commontools.mk

llvm/test/VISC/parboil/benchmarks/spmv/src/visc/main.c

@@ -17,10 +17,7 @@
 #include "gpu_info.h"
 #include "convert_dataset.h"
 
-extern "C" void __visc__attributes(...);
-extern "C" void __visc__wait(...);
-extern "C" unsigned __visc__node(...);
-extern "C" int get_global_id(int);
+#define WARP_BITS 5
 
 static int generate_vector(float *x_vector, int dim) 
 {	
@@ -33,7 +30,7 @@ static int generate_vector(float *x_vector, int dim)
 	return 0;
 }
 
-__kernel void spmv_jds(float *dst_vector, float *d_data,
+void spmv_jds(float *dst_vector, float *d_data,
 		               int *d_index, int *d_perm,
 		               float *x_vec, int dim, 
 		               int *jds_ptr_int,
@@ -178,7 +175,7 @@ int main(int argc, char** argv) {
 	int i;
 	for(i=0; i<50; i++)
 	{
-	unsigned spmvDFG = __visc__node(spmv_jds, 2, 1, block, grid,
+	unsigned spmvDFG = __visc__node(spmv_jds, 2, 1, (unsigned)block, (unsigned)grid,
                                     15,
                                     h_Ax_vector,
                                     dim*sizeof(float),

llvm/test/VISC/parboil/benchmarks/spmv/tools/compare-output

@@ -3,8 +3,9 @@
 # (c) Copyright 2007 The Board of Trustees of the University of Illinois.
 
 import sys
-import itertools
+sys.path.insert(0, '../../common/python')
 
+import itertools
 import filecompare as fc
 import binaryfilecompare as bfc