diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/cuda-sim/instructions.cc | 85 |
1 files changed, 46 insertions, 39 deletions
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 4eb5ce3..47f7075 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3741,15 +3741,16 @@ void sqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) void sst_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { + // Step 1: store data in sstarr memory const operand_info &src1 = pI->src1(); - const operand_info &src3 = pI->src3(); //may be scalar or vector of regs + const operand_info &src2 = pI->src2(); + const operand_info &src3 = pI->src3(); unsigned type = pI->get_type(); - ptx_reg_t addr_reg = thread->get_operand_value(src1, src1, type, thread, 1); - ptx_reg_t src3_data; + ptx_reg_t src1_data = thread->get_operand_value(src1, src1, type, thread, 1); + ptx_reg_t src2_data, src3_data; memory_space_t space = pI->get_space(); - memory_space *mem = NULL; - addr_t addr = addr_reg.u32; + addr_t addr = src1_data.u32; decode_space(space,thread,src1,mem,addr); @@ -3757,51 +3758,57 @@ void sst_impl( const ptx_instruction *pI, ptx_thread_info *thread ) int t; type_info_key::type_decode(type,size,t); + src2_data = thread->get_operand_value(src2, src1, type, thread, 1); src3_data = thread->get_operand_value(src3, src1, type, thread, 1); mem->write(addr,size/8,&src3_data.s64,thread,pI); + thread->m_last_effective_address = addr; thread->m_last_memory_space = space; + // Step 2: __syncthreads() to make sure all data is stored in sstarr memory + // (function must be called with dst = 0 so that all threads execute bar.sync 0) + ptx_instruction * cpI = const_cast<ptx_instruction *>(pI); + const operand_info &dst = cpI->dst(); + ptx_reg_t dst_data; + dst_data = thread->get_operand_value(dst, dst, U32_TYPE, thread, 1); + cpI->set_bar_id(dst_data.u32); - printf("SST instruction found.\n"); + thread->m_last_dram_callback.function = bar_callback; + thread->m_last_dram_callback.instruction = pI; - /*const operand_info &dst = pI->dst(); - const operand_info &src1 = pI->src1(); - const operand_info &src2 = pI->src2(); - const operand_info &src3 = pI->src3(); + // Step 3: pick only one thread to load all of the data back from sstarr memory + // rearrange the data so that zeros are at the end of the array + // store this data back in the original array (each thread can maybe do this after another sync?) + int NUM_THREADS = 8; + if (src2_data.s64 == NUM_THREADS-1) { + addr -= (NUM_THREADS-1)*4; + unsigned offset = 0; + ptx_reg_t data; + // loop through all of the threads (how do you do this dynamically?) + for (int tid = 0; tid < NUM_THREADS; tid++) { + data.u64=0; + mem->read(addr+(tid*4),size/8,&data.s64); - unsigned type = pI->get_type(); - ptx_reg_t addr_reg = thread->get_operand_value(src1, src1, type, thread, 1); - memory_space_t space = pI->get_space(); - - memory_space *mem = NULL; - addr_t addr = addr_reg.u32; - - decode_space(space,thread,src1,mem,addr); - - size_t size; - int t; - type_info_key::type_decode(type,size,t); + // store nonzero entries + if (data.s64 != 0) { + mem->write(addr+(offset*4),size/8,&data.s64,thread,pI); + thread->m_last_effective_address = addr+(offset*4); + offset++; + } + } + // fill the rest of the array with zeros + while (offset < NUM_THREADS) { + mem->write(addr+(offset*4),size/8,&src2_data.s64,thread,pI); + thread->m_last_effective_address = addr+(offset*4); + offset++; + } - ptx_reg_t src2_data = thread->get_operand_value(src2, src1, type, thread, 1); - ptx_reg_t src3_data = thread->get_operand_value(src3, src1, type, thread, 1); - mem->write(addr,size/8,&src3_data.s64,thread,pI);*/ + // Step 4: load from sstarr memory and store data back into original array - /* - switch ( i_type ) { - case U32_TYPE: - data.u64 = (src1_data.u64 & 0xFFFFFFFF) + (src2_data.u64 & 0xFFFFFFFF); - carry = (data.u64 & 0x100000000)>>32; - break; - case U64_TYPE: - data.u64 = src1_data.u64 + src2_data.u64; - break; - default: assert(0); break; - }*/ + } - //thread->set_operand_value(dst, data, i_type, thread, pI, overflow, carry ); - //thread->m_last_effective_address = addr; - //thread->m_last_memory_space = space; + //if( type == S16_TYPE || type == S32_TYPE ) sign_extend(data,size,dst); + //thread->set_operand_value(dst,data, type, thread, pI); } void ssy_impl( const ptx_instruction *pI, ptx_thread_info *thread ) |