4 files changed, 28 insertions, 8 deletions

diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 46c3279..6ed9b8e 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -162,9 +162,6 @@ struct dim3 {
 };
 #endif
 
-extern dim3 gridDim;
-extern dim3 blockDim;
-
 void increment_x_then_y_then_z( dim3 &i, const dim3 &bound);
 
 class kernel_info_t {
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index fd3b1fa..b5a3db4 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -3754,6 +3754,7 @@ void sst_impl( const ptx_instruction *pI, ptx_thread_info *thread )
 	memory_space_t space = pI->get_space();
 	memory_space *mem = NULL;
 	addr_t addr = src2_data.u32 * 4; // this assumes sstarr memory starts at address 0
+	ptx_cta_info *cta_info = thread->m_cta_info;
 
 	decode_space(space,thread,src1,mem,addr);
 
@@ -3765,18 +3766,19 @@ void sst_impl( const ptx_instruction *pI, ptx_thread_info *thread )
 	mem->write(addr,size/8,&src3_data.s64,thread,pI);
 
 	// sync threads
-	cpI->set_bar_id(dst_data.u32);
+	cpI->set_bar_id(16); // use 16 for sst because bar uses an int from 0-15
 
 	thread->m_last_effective_address = addr;
 	thread->m_last_memory_space = space;
 	thread->m_last_dram_callback.function = bar_callback;
 	thread->m_last_dram_callback.instruction = cpI;
 
-	int NUM_THREADS = blockDim.x * blockDim.y * blockDim.z;
-	if (src2_data.s64 == NUM_THREADS-1) {
-		// pick only one thread to load all of the data back from sstarr memory
+	// the last thread that executes loads all of the data back from sstarr memory
+	int NUM_THREADS = cta_info->num_threads();
+	cta_info->inc_bar_threads();
+	if (NUM_THREADS == cta_info->get_bar_threads()) {
 		unsigned offset = 0;
-		addr -= (NUM_THREADS-1)*4;
+		addr = 0;
 		ptx_reg_t data;
 		float sstarr_fdata[NUM_THREADS];
 		signed long long sstarr_ldata[NUM_THREADS];
@@ -3812,6 +3814,7 @@ void sst_impl( const ptx_instruction *pI, ptx_thread_info *thread )
 			offset++;
 		}
 
+		cta_info->reset_bar_threads();
 		thread->m_last_effective_address = addr+(NUM_THREADS-1)*4;
 		thread->m_last_memory_space = space;
 	}
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 511e8d6..f48115b 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -44,6 +44,7 @@ ptx_cta_info::ptx_cta_info( unsigned sm_idx )
 
    m_sm_idx = sm_idx;
    m_uid = g_ptx_cta_info_uid++;
+   m_bar_threads = 0;
 }
 
 void ptx_cta_info::add_thread( ptx_thread_info *thd )
@@ -128,6 +129,21 @@ unsigned ptx_cta_info::get_sm_idx() const
    return m_sm_idx;
 }
 
+unsigned ptx_cta_info::get_bar_threads() const
+{
+   return m_bar_threads;
+}
+
+void ptx_cta_info::inc_bar_threads()
+{
+	m_bar_threads++;
+}
+
+void ptx_cta_info::reset_bar_threads()
+{
+	m_bar_threads = 0;
+}
+
 unsigned g_ptx_thread_info_uid_next=1;
 unsigned g_ptx_thread_info_delete_count=0;
 
diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h
index c66b68c..4e748e9 100644
--- a/src/cuda-sim/ptx_sim.h
+++ b/src/cuda-sim/ptx_sim.h
@@ -158,8 +158,12 @@ public:
    void register_thread_exit( ptx_thread_info *thd );
    void register_deleted_thread( ptx_thread_info *thd );
    unsigned get_sm_idx() const;
+   unsigned get_bar_threads() const;
+   void inc_bar_threads();
+   void reset_bar_threads();
 
 private:
+   unsigned           m_bar_threads;
    unsigned long long         m_uid;
    unsigned                m_sm_idx;
    std::set<ptx_thread_info*>    m_threads_in_cta;