Fixing a slew of compiler warnings

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

9 files changed, 44 insertions, 20 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 7535fa3..dad7669 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -558,9 +558,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 );
@@ -1616,9 +1614,9 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 		// PTX/SASS code for the launched kernel function.  
 		// This allows us to work around the fact that cuobjdump only outputs the
 		// file name associated with each section. 
-		unsigned fat_cubin_handle = next_fat_bin_handle;
+		unsigned long long fat_cubin_handle = next_fat_bin_handle;
 		next_fat_bin_handle++;
-		printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %u, filename=%s\n", fat_cubin_handle, filename);
+		printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %llu, filename=%s\n", fat_cubin_handle, filename);
 		/*!
 		 * This function extracts all data from all files in first call
 		 * then for next calls, only returns the appropriate number
@@ -1630,7 +1628,7 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 		return (void**)fat_cubin_handle;
 	} else {
 		static unsigned source_num=1;
-		unsigned fat_cubin_handle = next_fat_bin_handle++;
+		unsigned long long fat_cubin_handle = next_fat_bin_handle++;
 		__cudaFatCudaBinary *info =   (__cudaFatCudaBinary *)fatCubin;
 		assert( info->version >= 3 );
 		unsigned num_ptx_versions=0;
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index ad32a18..630c491 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -609,7 +609,7 @@ char *_cl_program::get_ptx()
    }
    size_t buffer_length= get_ptx_size();
    char *tmp = (char*)calloc(buffer_length + 1,1);
-   tmp[ buffer_length ] = NULL;
+   tmp[ buffer_length ] = '\0';
    unsigned n=0;
    std::map<cl_uint,pgm_info>::iterator p;
    for( p=m_pgm.begin(); p != m_pgm.end(); p++ ) {
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 1e0ed6f..cf64e82 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -441,7 +441,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 );
@@ -449,7 +448,7 @@ void ptx_print_insn( address_type pc, FILE *fp )
    }
    function_info *finfo = f->second;
    assert( finfo );
-   size = finfo->print_insn(pc,fp);
+   finfo->print_insn(pc,fp);
 }
 void ptx_instruction::set_fp_or_int_archop(){
 	op2=UN_OP;
@@ -1207,8 +1206,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(),
@@ -1587,7 +1584,7 @@ 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";
+      mem_name = "const";
       break;
    case global_space:
       mem = gpu->get_global_memory();
@@ -1596,8 +1593,8 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
    default:
       abort();
    }
-   printf("GPGPU-Sim PTX: gpgpu_ptx_sim_memcpy_symbol: copying %zu bytes %s symbol %s+%zu @0x%x ...\n", 
-          count, (to?" to ":"from"), sym_name.c_str(), offset, dst );
+   printf("GPGPU-Sim PTX: gpgpu_ptx_sim_memcpy_symbol: copying %s memory %zu bytes %s symbol %s+%zu @0x%x ...\n", 
+          mem_name, count, (to?" to ":"from"), sym_name.c_str(), offset, dst );
    for ( unsigned n=0; n < count; n++ ) {
       if( to ) mem->write(dst+n,1,((char*)src)+n,NULL,NULL); 
       else mem->read(dst+n,1,((char*)src)+n); 
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 99d7399..99e81e5 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -1168,7 +1168,12 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
 
    // Write operation result into  memory
    // (i.e. copy src1_data to dst)
-   mem->write(effective_address,size/8,&op_result.s64,thread,pI);
+   if ( data_ready ) {
+      mem->write(effective_address,size/8,&op_result.s64,thread,pI);
+   } else {
+      printf("Execution error: data_ready not set\n");
+      assert(0);
+   }
 }
 
 // atom_impl will now result in a callback being called in mem_ctrl_pop (gpu-sim.c)
diff --git a/src/gpgpu-sim/dram_sched.cc b/src/gpgpu-sim/dram_sched.cc
index e30bd24..2f991c1 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 3b25565..9317c18 100644
--- a/src/gpgpu-sim/l2cache.h
+++ b/src/gpgpu-sim/l2cache.h
@@ -142,6 +142,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 72a1936..8573344 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -830,6 +830,20 @@ void TwoLevelScheduler::cycle() {
 			break;
 		}
 	}
+
+    // tgrogers - fixing a warning about unused variables at the top of this fucntion...
+    // This scheduler has A LOT of copied code from the LRR scheduler.
+    // TODO - this thing really needs to be re-written in a modular way.
+    // For now, to get rid of the warnings I am including the stats epilogue
+    // from the original scheduler function that the author of this function forgot to copy/paste
+    //
+    // issue stall statistics:
+    if( !valid_inst ) 
+        m_stats->shader_cycle_distro[0]++; // idle or control hazard
+    else if( !ready_inst ) 
+        m_stats->shader_cycle_distro[1]++; // waiting for RAW hazards (possibly due to memory) 
+    else if( !issued_inst ) 
+        m_stats->shader_cycle_distro[2]++; // pipeline stalled
 }
 
 void shader_core_ctx::read_operands()
diff --git a/src/gpuwattch/XML_Parse.cc b/src/gpuwattch/XML_Parse.cc
index 28b0efe..8e33f3f 100644
--- a/src/gpuwattch/XML_Parse.cc
+++ b/src/gpuwattch/XML_Parse.cc
@@ -43,6 +43,11 @@
 

 using namespace std;

 

+const char * perf_count_label[] = {"TOT_INST,", "FP_INT,", "IC_H,", "IC_M,", "DC_RH,", "DC_RM,", "DC_WH,", "DC_WM,",

+                                "TC_H,", "TC_M,", "CC_H,", "CC_M,", "SHRD_ACC,", "REG_RD,", "REG_WR,", "NON_REG_OPs,",

+                                "SP_ACC,", "SFU_ACC,", "FPU_ACC,", "MEM_RD,","MEM_WR,", "MEM_PRE,", "L2_RH,", "L2_RM,", "L2_WH,",

+                                "L2_WM,", "NOC_A,", "PIPE_A,", "IDLE_CORE_N,", "CONST_DYNAMICN"};

+

 void ParseXML::parse(char* filepath)

 {

 	unsigned int i,j,k,m,n;

diff --git a/src/gpuwattch/XML_Parse.h b/src/gpuwattch/XML_Parse.h
index 53e408b..698445e 100644
--- a/src/gpuwattch/XML_Parse.h
+++ b/src/gpuwattch/XML_Parse.h
@@ -56,10 +56,16 @@ ToXMLStringTool tx,tx2;
 */

 //all subnodes at the level of system.core(0-n)

 //cache_policy is added into cache property arrays;//0 no write or write-though with non-write allocate;1 write-back with write-allocate

-static const char * perf_count_label[] = {"TOT_INST,", "FP_INT,", "IC_H,", "IC_M,", "DC_RH,", "DC_RM,", "DC_WH,", "DC_WM,",

-                                "TC_H,", "TC_M,", "CC_H,", "CC_M,", "SHRD_ACC,", "REG_RD,", "REG_WR,", "NON_REG_OPs,",

-                                "SP_ACC,", "SFU_ACC,", "FPU_ACC,", "MEM_RD,","MEM_WR,", "MEM_PRE,", "L2_RH,", "L2_RM,", "L2_WH,",

-                                "L2_WM,", "NOC_A,", "PIPE_A,", "IDLE_CORE_N,", "CONST_DYNAMICN"};

+//
+//tgrogers - This was a static array declared in the header...
+//           Not too sure why the authors did this, maybe they didn't understand the context
+//           of the "static" keyword outside a class declaration.  As it was written, each object file
+//           had it's own copy of the string list - which is okay, since the string is constant but
+//           wastes space and causes the compiler to complain about unused vars in files where this header
+//           is included but they don't use the variable.  Now this is extern'd here and the storage/definition
+//           is in the XML_Parse.cc file
+extern const char * perf_count_label[];
+
 enum perf_count_t {

    TOT_INST=0,

    FP_INT,