Removing compiler warnings

Fixing failed compilation due to double definition of parsing functions

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 13683]

9 files changed, 9 insertions, 47 deletions

diff --git a/CHANGES b/CHANGES
index 2e2db44..3a19891 100644
--- a/CHANGES
+++ b/CHANGES
@@ -1,5 +1,10 @@
 CHANGE LOG:
 
+Version 3.1.2 versus 3.1.1
+- Bug Fixes:
+    - Fixed a minor compilation bug with newer gcc/g++ versions
+    - Removed compiler warnings with gcc 4.7.1
+
 Version 3.1.1 versus 3.1.0
 - Add checks to top level makefile to ensure setup_environment is run and checks to
   setup_environment to ensure a compatible version of CUDA is installed.
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index e189403..cdfd99c 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -551,9 +551,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void
 {
 	CUctx_st *context = GPGPUSim_Context();
 	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	struct cudaArray *cuArray_ptr;
 	size_t size = spitch*height;
-	cuArray_ptr = (cudaArray*)dst;
 	gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" );
 	if( kind == cudaMemcpyHostToDevice )
 		gpu->memcpy_to_gpu( (size_t)dst, src, size );
@@ -1801,7 +1799,7 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
 
 ////////
 
-extern "C" int ptx_parse();
+//extern "C" int ptx_parse();
 extern "C" int ptx__scan_string(const char*);
 extern "C" FILE *ptx_in;
 
diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y
index 914a3d6..b07c754 100644
--- a/libcuda/cuobjdump.y
+++ b/libcuda/cuobjdump.y
@@ -103,7 +103,7 @@ ptxcode :	ptxcode PTXLINE {fprintf(ptxfile, "%s", $2);}
 elfcode :	elfcode ELFLINE {fprintf(elffile, "%s", $2);}
 		|	;
 
-sasscode :	sasscode SASSLINE {fprintf(sassfile, "%s", $2)}
+sasscode :	sasscode SASSLINE {fprintf(sassfile, "%s", $2);}
 		 |	;
 
 
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 02c7dfd..89c240f 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -438,7 +438,6 @@ void gpgpu_t::gpu_memset( size_t dst_start_addr, int c, size_t count )
 
 void ptx_print_insn( address_type pc, FILE *fp )
 {
-   static unsigned size=1;
    std::map<unsigned,function_info*>::iterator f = g_pc_to_finfo.find(pc);
    if( f == g_pc_to_finfo.end() ) {
        fprintf(fp,"<no instruction at address 0x%x>", pc );
@@ -446,7 +445,6 @@ void ptx_print_insn( address_type pc, FILE *fp )
    }
    function_info *finfo = f->second;
    assert( finfo );
-   size = finfo->print_insn(pc,fp);
 }
 
 void ptx_instruction::set_opcode_and_latency()
@@ -1090,8 +1088,6 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
    // Output instruction information to file and stdout
    if( config.get_ptx_inst_debug_to_file() != 0 && 
         (config.get_ptx_inst_debug_thread_uid() == 0 || config.get_ptx_inst_debug_thread_uid() == get_uid()) ) {
-      dim3 ctaid = get_ctaid();
-      dim3 tid = get_tid();
       fprintf(m_gpu->get_ptx_inst_debug_file(),
              "[thd=%u] : (%s:%u - %s)\n",
              get_uid(),
@@ -1457,7 +1453,6 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
       printf("Execution error: No information for PTX symbol w/ hostVar=0x%Lx\n", (unsigned long long)hostVar );
       abort();
    } else printf("GPGPU-Sim PTX: gpgpu_ptx_sim_memcpy_symbol: Found PTX symbol w/ hostVar=0x%Lx\n", (unsigned long long)hostVar ); 
-   const char *mem_name = NULL;
    memory_space *mem = NULL;
 
    std::map<std::string,symbol_table*>::iterator st = g_sym_name_to_symbol_table.find(sym_name.c_str());
@@ -1470,11 +1465,9 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
    switch (mem_region.get_type()) {
    case const_space:
       mem = gpu->get_global_memory();
-      mem_name = "global";
       break;
    case global_space:
       mem = gpu->get_global_memory();
-      mem_name = "global";
       break;
    default:
       abort();
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 6976ba0..8a737bc 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -882,8 +882,6 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
    src2_data,   // b
    op_result;   // temp variable to hold operation result
 
-   bool data_ready = false;
-
    // Get operand info of sources and destination
    const operand_info &dst  = pI->dst();     // d
    const operand_info &src1 = pI->src1();    // a
@@ -918,11 +916,9 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          case B32_TYPE:
          case U32_TYPE:
             op_result.u32 = data.u32 & src2_data.u32;
-            data_ready = true;
             break;
          case S32_TYPE:
             op_result.s32 = data.s32 & src2_data.s32;
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch (%x) with instruction\natom.AND only accepts b32\n", to_type);
@@ -940,11 +936,9 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          case B32_TYPE:
          case U32_TYPE:
             op_result.u32 = data.u32 | src2_data.u32;
-            data_ready = true;
             break;
          case S32_TYPE:
             op_result.s32 = data.s32 | src2_data.s32;
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch (%x) with instruction\natom.OR only accepts b32\n", to_type);
@@ -962,11 +956,9 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          case B32_TYPE:
          case U32_TYPE:
             op_result.u32 = data.u32 ^ src2_data.u32;
-            data_ready = true;
             break;
          case S32_TYPE:
             op_result.s32 = data.s32 ^ src2_data.s32;
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch (%x) with instruction\natom.XOR only accepts b32\n", to_type);
@@ -988,16 +980,13 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          case B32_TYPE:
          case U32_TYPE:
             op_result.u32 = MY_CAS_I(data.u32, src2_data.u32, src3_data.u32);
-            data_ready = true;
             break;
          case B64_TYPE:
          case U64_TYPE:
             op_result.u64 = MY_CAS_I(data.u64, src2_data.u64, src3_data.u64);
-            data_ready = true;
             break;
          case S32_TYPE:
             op_result.s32 = MY_CAS_I(data.s32, src2_data.s32, src3_data.s32);
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch (%x) with instruction\natom.CAS only accepts b32 and b64\n", to_type);
@@ -1014,16 +1003,13 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          case B32_TYPE:
          case U32_TYPE:
             op_result.u32 = MY_EXCH(data.u32, src2_data.u32);
-            data_ready = true;
             break;
          case B64_TYPE:
          case U64_TYPE:
             op_result.u64 = MY_EXCH(data.u64, src2_data.u64);
-            data_ready = true;
             break;
          case S32_TYPE:
             op_result.s32 = MY_EXCH(data.s32, src2_data.s32);
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch (%x) with instruction\natom.EXCH only accepts b32\n", to_type);
@@ -1040,19 +1026,15 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          switch ( to_type ) {
          case U32_TYPE:
             op_result.u32 = data.u32 + src2_data.u32;
-            data_ready = true;
             break;
          case S32_TYPE:
             op_result.s32 = data.s32 + src2_data.s32;
-            data_ready = true;
             break;
          case U64_TYPE: 
             op_result.u64 = data.u64 + src2_data.u64; 
-            data_ready = true; 
             break; 
          case F32_TYPE: 
             op_result.f32 = data.f32 + src2_data.f32; 
-            data_ready = true; 
             break; 
          default:
             printf("Execution error: type mismatch with instruction\natom.ADD only accepts u32, s32, u64, and f32\n");
@@ -1068,7 +1050,6 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          switch ( to_type ) {
          case U32_TYPE: 
             op_result.u32 = MY_INC_I(data.u32, src2_data.u32);
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch with instruction\natom.INC only accepts u32 and s32\n");
@@ -1084,7 +1065,6 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          switch ( to_type ) {
          case U32_TYPE: 
             op_result.u32 = MY_DEC_I(data.u32, src2_data.u32);
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch with instruction\natom.DEC only accepts u32 and s32\n");
@@ -1100,11 +1080,9 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          switch ( to_type ) {
          case U32_TYPE: 
             op_result.u32 = MY_MIN_I(data.u32, src2_data.u32);
-            data_ready = true;
             break;
          case S32_TYPE:
             op_result.s32 = MY_MIN_I(data.s32, src2_data.s32);
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch with instruction\natom.MIN only accepts u32 and s32\n");
@@ -1120,11 +1098,9 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
          switch ( to_type ) {
          case U32_TYPE:
             op_result.u32 = MY_MAX_I(data.u32, src2_data.u32);
-            data_ready = true;
             break;
          case S32_TYPE:
             op_result.s32 = MY_MAX_I(data.s32, src2_data.s32);
-            data_ready = true;
             break;
          default:
             printf("Execution error: type mismatch with instruction\natom.MAX only accepts u32 and s32\n");
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 56b9449..b5103df 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -43,7 +43,7 @@ bool g_override_embedded_ptx = false;
 
 /// extern prototypes
 
-extern "C" int ptx_parse();
+//extern "C" int ptx_parse();
 extern "C" int ptx__scan_string(const char*);
 
 const char *g_ptxinfo_filename;
diff --git a/src/gpgpu-sim/dram_sched.cc b/src/gpgpu-sim/dram_sched.cc
index e556edc..48e96ee 100644
--- a/src/gpgpu-sim/dram_sched.cc
+++ b/src/gpgpu-sim/dram_sched.cc
@@ -78,7 +78,6 @@ void frfcfs_scheduler::data_collection(unsigned int bank)
 
 dram_req_t *frfcfs_scheduler::schedule( unsigned bank, unsigned curr_row )
 {
-   int row_hit = 0;
    if ( m_last_row[bank] == NULL ) {
       if ( m_queue[bank].empty() )
          return NULL;
@@ -95,7 +94,6 @@ dram_req_t *frfcfs_scheduler::schedule( unsigned bank, unsigned curr_row )
 
       }
    }
-   row_hit=1;
    std::list<dram_req_t*>::iterator next = m_last_row[bank]->back();
    dram_req_t *req = (*next);
 
diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h
index 22640b8..15c3806 100644
--- a/src/gpgpu-sim/l2cache.h
+++ b/src/gpgpu-sim/l2cache.h
@@ -121,6 +121,7 @@ private:
 class L2interface : public mem_fetch_interface {
 public:
     L2interface( memory_partition_unit *unit ) { m_unit=unit; }
+    virtual ~L2interface(){}
     virtual bool full( unsigned size, bool write) const 
     {
         // assume read and write packets all same size
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index ddb544a..a7daeaa 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -672,10 +672,6 @@ void TwoLevelScheduler::cycle() {
 	//Do the scheduling only from activeWarps
 	//If you schedule an instruction, move it to the end of the list
 
-	bool valid_inst = false;  // there was one warp with a valid instruction to issue (didn't require flush due to control hazard)
-	bool ready_inst = false;  // of the valid instructions, there was one not waiting for pending register writes
-	bool issued_inst = false; // of these we issued one
-
 	for (	std::list<int>::iterator warp_id = activeWarps.begin();
 			warp_id != activeWarps.end();
 			warp_id++) {
@@ -695,16 +691,13 @@ void TwoLevelScheduler::cycle() {
 					warp(*warp_id).set_next_pc(pc);
 					warp(*warp_id).ibuffer_flush();
 				} else {
-					valid_inst = true;
 					if ( !m_scoreboard->checkCollision(*warp_id, pI) ) {
-						ready_inst = true;
 						const active_mask_t &active_mask = m_simt_stack[*warp_id]->get_active_mask();
 						assert( warp(*warp_id).inst_in_pipeline() );
 						if ( (pI->op == LOAD_OP) || (pI->op == STORE_OP) || (pI->op == MEMORY_BARRIER_OP) ) {
 							if( m_mem_out->has_free() ) {
 								m_shader->issue_warp(*m_mem_out,pI,active_mask,*warp_id);
 								issued++;
-								issued_inst=true;
 								warp_inst_issued = true;
 								// Move it to pendingWarps
 								unsigned currwarp = *warp_id;
@@ -718,7 +711,6 @@ void TwoLevelScheduler::cycle() {
 								// always prefer SP pipe for operations that can use both SP and SFU pipelines
 								m_shader->issue_warp(*m_sp_out,pI,active_mask,*warp_id);
 								issued++;
-								issued_inst=true;
 								warp_inst_issued = true;
 								//Move it to end of the activeWarps
 								unsigned currwarp = *warp_id;
@@ -728,7 +720,6 @@ void TwoLevelScheduler::cycle() {
 								if( sfu_pipe_avail ) {
 									m_shader->issue_warp(*m_sfu_out,pI,active_mask,*warp_id);
 									issued++;
-									issued_inst=true;
 									warp_inst_issued = true;
 									//Move it to end of the activeWarps
 									unsigned currwarp = *warp_id;