Back out changelist 13683

There is a problem with the linkage on my machine. Before this changelist the code didn't build on my machine. After it it builds but fails to run due to missing dynamic linkage. And obviously it breaks the jenkins build.

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

9 files changed, 47 insertions, 9 deletions

diff --git a/CHANGES b/CHANGES
index 3a19891..2e2db44 100644
--- a/CHANGES
+++ b/CHANGES
@@ -1,10 +1,5 @@
 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 cdfd99c..e189403 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -551,7 +551,9 @@ __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 );
@@ -1799,7 +1801,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 b07c754..914a3d6 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 89c240f..02c7dfd 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -438,6 +438,7 @@ 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 );
@@ -445,6 +446,7 @@ 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()
@@ -1088,6 +1090,8 @@ 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(),
@@ -1453,6 +1457,7 @@ 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());
@@ -1465,9 +1470,11 @@ 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 8a737bc..6976ba0 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -882,6 +882,8 @@ 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
@@ -916,9 +918,11 @@ 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);
@@ -936,9 +940,11 @@ 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);
@@ -956,9 +962,11 @@ 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);
@@ -980,13 +988,16 @@ 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);
@@ -1003,13 +1014,16 @@ 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);
@@ -1026,15 +1040,19 @@ 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");
@@ -1050,6 +1068,7 @@ 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");
@@ -1065,6 +1084,7 @@ 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");
@@ -1080,9 +1100,11 @@ 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");
@@ -1098,9 +1120,11 @@ 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 b5103df..56b9449 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 48e96ee..e556edc 100644
--- a/src/gpgpu-sim/dram_sched.cc
+++ b/src/gpgpu-sim/dram_sched.cc
@@ -78,6 +78,7 @@ 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;
@@ -94,6 +95,7 @@ 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 15c3806..22640b8 100644
--- a/src/gpgpu-sim/l2cache.h
+++ b/src/gpgpu-sim/l2cache.h
@@ -121,7 +121,6 @@ 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 a7daeaa..ddb544a 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -672,6 +672,10 @@ 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++) {
@@ -691,13 +695,16 @@ 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;
@@ -711,6 +718,7 @@ 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;
@@ -720,6 +728,7 @@ 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;