diff options
| author | sspenst <[email protected]> | 2016-07-11 11:16:46 -0700 |
|---|---|---|
| committer | sspenst <[email protected]> | 2016-07-11 11:16:46 -0700 |
| commit | adc311951d67b0685ebf2fab4ce6410f96f0039a (patch) | |
| tree | e16dbeddf663cfc11bbefaa517101cd2cb5b48fe /src/cuda-sim | |
| parent | 6c1fb702e17b00fd7de72ac7dd4a31584d5978b9 (diff) | |
Reverted the previous commit to add a cleaner way of getting NUM_THREADS. Now, sst_impl doesn't functionally execute on the last indexed element of an array, but instead on the actual last thread that executes
Diffstat (limited to 'src/cuda-sim')
| -rw-r--r-- | src/cuda-sim/instructions.cc | 13 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_sim.cc | 16 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_sim.h | 4 |
3 files changed, 28 insertions, 5 deletions
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; |
