Support for named bariers + bar.red + bar.arrive instructions

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

8 files changed, 198 insertions, 10 deletions

diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index dd76e4b..1ccc156 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -541,6 +541,40 @@ void ptx_instruction::set_mul_div_or_other_archop(){
 	}
 
 }
+
+
+
+void ptx_instruction::set_bar_type()
+{
+	   if(m_opcode==BAR_OP) {
+		   switch(m_barrier_op){
+		   	   case SYNC_OPTION:
+		   		   bar_type = SYNC;
+		   		   break;
+		   	   case ARRIVE_OPTION:
+		   		   bar_type = ARRIVE;
+		   		   break;
+		   	   case RED_OPTION:
+		   		   bar_type = RED;
+		   		   switch(m_reduction_op){
+		   		   	   case POPC_REDUCTION:
+				   		   red_type = POPC_RED;
+				   		   break;
+		   		   	   case AND_REDUCTION:
+				   		   red_type = AND_RED;
+				   		   break;
+		   		   	   case OR_REDUCTION:
+				   		   red_type = OR_RED;
+				   		   break;
+		   		   }
+		   		break;
+		   	   default:
+		   		   abort();
+		   }
+	   }
+}
+
+
 void ptx_instruction::set_opcode_and_latency()
 {
 	unsigned int_latency[5];
@@ -837,7 +871,7 @@ void ptx_instruction::pre_decode()
    }
 
    set_opcode_and_latency();
-
+   set_bar_type();
    // Get register operands
    int n=0,m=0;
    ptx_instruction::const_iterator opr=op_iter_begin();
@@ -1257,11 +1291,15 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
       insn_data_size = datatype2size(to_type);
       insn_memory_op = pI->has_memory_read() ? memory_load : memory_store;
    }
-   
+
+   if ( pI->get_opcode() == BAR_OP && pI->barrier_op() == RED_OPTION) {
+	   inst.add_callback( lane_id, last_callback().function, last_callback().instruction, this,false /*not atomic*/);
+   }
+
    if ( pI->get_opcode() == ATOM_OP ) {
       insn_memaddr = last_eaddr();
       insn_space = last_space();
-      inst.add_callback( lane_id, last_callback().function, last_callback().instruction, this );
+      inst.add_callback( lane_id, last_callback().function, last_callback().instruction, this,true /*atomic*/);
       unsigned to_type = pI->get_type();
       insn_data_size = datatype2size(to_type);
    }
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index ce39ca4..5d17f95 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -865,7 +865,19 @@ void andn_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    thread->set_operand_value(dst,data, i_type, thread, pI);
 }
 
-void atom_callback( const inst_t* inst, ptx_thread_info* thread )
+void bar_callback( const inst_t* inst, ptx_thread_info* thread)
+{
+	unsigned ctaid = thread->get_cta_uid();
+	unsigned barid = inst->bar_id;
+	unsigned value = thread->get_reduction_value(ctaid,barid);
+	const ptx_instruction *pI = dynamic_cast<const ptx_instruction*>(inst);
+	const operand_info &dst  = pI->dst();
+	ptx_reg_t data;
+	data.u32 = value;
+	thread->set_operand_value(dst,value, U32_TYPE, thread, pI);
+}
+
+void atom_callback( const inst_t* inst, ptx_thread_info* thread)
 {
    const ptx_instruction *pI = dynamic_cast<const ptx_instruction*>(inst);
 
@@ -1221,11 +1233,107 @@ void atom_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    thread->m_last_dram_callback.instruction = pI; 
 }
 
-void bar_sync_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
+void bar_impl( const ptx_instruction *pIin, ptx_thread_info *thread )
 { 
-   const operand_info &dst  = pI->dst();
-   ptx_reg_t b = thread->get_operand_value(dst, dst, U32_TYPE, thread, 1);
-   assert( b.u32 == 0 ); // support for bar.sync a{,b}; where a != 0 not yet implemented
+   ptx_instruction * pI = const_cast<ptx_instruction *>(pIin);
+   unsigned bar_op = pI->barrier_op();
+   unsigned red_op = pI->reduction_op();
+   unsigned ctaid = thread->get_cta_uid();
+
+   switch(bar_op){
+   	   case SYNC_OPTION:
+   	   {
+   		   if(pI->get_num_operands()>1){
+   			   const operand_info &op0 = pI->dst();
+   			   const operand_info &op1 = pI->src1();
+   			   ptx_reg_t op0_data;
+   			   ptx_reg_t op1_data;
+   			   op0_data = thread->get_operand_value(op0, op0, U32_TYPE, thread, 1);
+   			   op1_data = thread->get_operand_value(op1, op1, U32_TYPE, thread, 1);
+   			   pI->set_bar_id(op0_data.u32);
+   			   pI->set_bar_count(op1_data.u32);
+   		   }else{
+   			   const operand_info &op0 = pI->dst();
+   			   ptx_reg_t op0_data;
+   			   op0_data = thread->get_operand_value(op0, op0, U32_TYPE, thread, 1);
+   			   pI->set_bar_id(op0_data.u32);
+   		   }
+   		   break;
+   	   }
+   	   case ARRIVE_OPTION:
+   	   {
+			   const operand_info &op0 = pI->dst();
+			   const operand_info &op1 = pI->src1();
+			   ptx_reg_t op0_data;
+			   ptx_reg_t op1_data;
+			   op0_data = thread->get_operand_value(op0, op0, U32_TYPE, thread, 1);
+			   op1_data = thread->get_operand_value(op1, op1, U32_TYPE, thread, 1);
+   			   pI->set_bar_id(op0_data.u32);
+   			   pI->set_bar_count(op1_data.u32);
+   		   break;
+   	   }
+   	   case RED_OPTION:
+   	   {
+   		   if(pI->get_num_operands()>3){
+   			   const operand_info &op1 = pI->src1();
+   			   const operand_info &op2 = pI->src2();
+   			   const operand_info &op3 = pI->src3();
+   			   ptx_reg_t op1_data;
+   			   ptx_reg_t op2_data;
+   			   ptx_reg_t op3_data;
+   			   op1_data = thread->get_operand_value(op1, op1, U32_TYPE, thread, 1);
+   			   op2_data = thread->get_operand_value(op2, op2, U32_TYPE, thread, 1);
+   			   op3_data = thread->get_operand_value(op3, op3, PRED_TYPE, thread, 1);
+   			   op3_data.u32=!(op3_data.pred & 0x0001);
+   			   pI->set_bar_id(op1_data.u32);
+   			   pI->set_bar_count(op2_data.u32);
+   			   switch(red_op){
+   			   	   case POPC_REDUCTION:
+   			   		   thread->popc_reduction(ctaid,op1_data.u32,op3_data.u32);
+   			   		   break;
+   			   	   case AND_REDUCTION:
+   			   		   thread->and_reduction(ctaid,op1_data.u32,op3_data.u32);
+   			   		   break;
+   			   	   case OR_REDUCTION:
+   			   		   thread->or_reduction(ctaid,op1_data.u32,op3_data.u32);
+   			   		   break;
+   			   	   default:
+   			   		   abort();
+   			   		   break;
+   			   }
+   		   }else{
+   			   const operand_info &op1 = pI->src1();
+   			   const operand_info &op2 = pI->src2();
+   			   ptx_reg_t op1_data;
+   			   ptx_reg_t op2_data;
+   			   op1_data = thread->get_operand_value(op1, op1, U32_TYPE, thread, 1);
+   			   op2_data = thread->get_operand_value(op2, op2, PRED_TYPE, thread, 1);
+               op2_data.u32=!(op2_data.pred & 0x0001);
+   			   pI->set_bar_id(op1_data.u32);
+   			   switch(red_op){
+   			   	   case POPC_REDUCTION:
+   			   		   thread->popc_reduction(ctaid,op1_data.u32,op2_data.u32);
+   			   		   break;
+   			   	   case AND_REDUCTION:
+   			   		   thread->and_reduction(ctaid,op1_data.u32,op2_data.u32);
+   			   		   break;
+   			   	   case OR_REDUCTION:
+   			   		   thread->or_reduction(ctaid,op1_data.u32,op2_data.u32);
+   			   		   break;
+   			   	   default:
+   			   		   abort();
+   			   		   break;
+   			   }
+   		   }
+   		   break;
+   	   }
+   	   default:
+   		   abort();
+   		   break;
+   }
+
+   thread->m_last_dram_callback.function = bar_callback;
+   thread->m_last_dram_callback.instruction = pIin;
 }
 
 void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def
index b4912c0..7aaa14f 100644
--- a/src/cuda-sim/opcodes.def
+++ b/src/cuda-sim/opcodes.def
@@ -44,7 +44,7 @@ OP_DEF(ADDC_OP,addc_impl,"addc",1,1)
 OP_DEF(AND_OP,and_impl,"and",1,1)
 OP_DEF(ANDN_OP,andn_impl,"andn",1,1)
 OP_DEF(ATOM_OP,atom_impl,"atom",1,3)
-OP_DEF(BAR_OP,bar_sync_impl,"bar.sync",1,3)
+OP_DEF(BAR_OP,bar_impl,"bar",1,3)
 OP_DEF(BFE_OP,bfe_impl,"bfe",1,1)
 OP_DEF(BFI_OP,bfi_impl,"bfi",1,1)
 OP_DEF(BFIND_OP,bfind_impl,"bfind",1,1)
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index dd45a67..ebe07c7 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -60,7 +60,7 @@ addc    TC; ptx_lval.int_value = ADDC_OP; return OPCODE;
 and	TC; ptx_lval.int_value = AND_OP; return OPCODE;
 andn	TC; ptx_lval.int_value = ANDN_OP; return OPCODE;
 atom	TC; ptx_lval.int_value = ATOM_OP; return OPCODE;
-bar.sync TC; ptx_lval.int_value = BAR_OP; return OPCODE;
+bar TC; ptx_lval.int_value = BAR_OP; return OPCODE;
 bfe     TC; ptx_lval.int_value = BFE_OP; return OPCODE;
 bfi     TC; ptx_lval.int_value = BFI_OP; return OPCODE;
 bfind   TC; ptx_lval.int_value = BFIND_OP; return OPCODE;
@@ -294,6 +294,14 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.wide  TC; return WIDE_OPTION;
 \.uni	TC; return UNI_OPTION;
 
+\.sync TC; return SYNC_OPTION;
+\.arrive TC; return ARRIVE_OPTION;
+\.red TC; return RED_OPTION;
+
+\.popc	TC; return POPC_REDUCTION;
+\.and	TC;	return AND_REDUCTION;
+\.or	TC; return OR_REDUCTION;
+
 \.approx TC; return APPROX_OPTION;
 \.full  TC; return FULL_OPTION;
 
@@ -330,6 +338,8 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.min	TC; return ATOMIC_MIN;
 \.max	TC; return ATOMIC_MAX;
 
+
+
 \.1d	TC; return GEOM_MODIFIER_1D;
 \.2d	TC; return GEOM_MODIFIER_2D;
 \.3d	TC; return GEOM_MODIFIER_3D;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 15662aa..6c614bf 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -155,6 +155,12 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %token SAT_OPTION
 %token FTZ_OPTION
 %token NEG_OPTION
+%token SYNC_OPTION
+%token RED_OPTION
+%token POPC_REDUCTION
+%token AND_REDUCTION
+%token OR_REDUCTION
+%token ARRIVE_OPTION
 %token ATOMIC_AND
 %token ATOMIC_OR
 %token ATOMIC_XOR
@@ -408,6 +414,12 @@ option: type_spec
 	| compare_spec
 	| addressable_spec
 	| rounding_mode
+	| SYNC_OPTION { add_option(SYNC_OPTION); }	
+	| ARRIVE_OPTION { add_option(ARRIVE_OPTION); }
+	| RED_OPTION { add_option(RED_OPTION); }					
+	| POPC_REDUCTION { add_option(POPC_REDUCTION); }
+	| AND_REDUCTION { add_option(AND_REDUCTION); }
+	| OR_REDUCTION { add_option(OR_REDUCTION); }	
 	| UNI_OPTION { add_option(UNI_OPTION); }
 	| WIDE_OPTION { add_option(WIDE_OPTION); }
 	| ANY_OPTION { add_option(ANY_OPTION); }
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index a442e2e..b6c1b9a 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1044,6 +1044,16 @@ ptx_instruction::ptx_instruction( int opcode,
    for ( i=options.begin(); i!= options.end(); i++, n++ ) {
       int last_ptx_inst_option = *i;
       switch ( last_ptx_inst_option ) {
+      case SYNC_OPTION:
+      case ARRIVE_OPTION:
+      case RED_OPTION:
+          m_barrier_op = last_ptx_inst_option;
+          break;
+      case OR_REDUCTION:
+      case AND_REDUCTION:
+      case POPC_REDUCTION:
+          m_reduction_op = last_ptx_inst_option;
+          break;
       case EQU_OPTION:
       case NEU_OPTION:
       case LTU_OPTION:
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 5efa3a8..1dd851a 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -992,6 +992,8 @@ public:
    unsigned rounding_mode() const { return m_rounding_mode;}
    unsigned saturation_mode() const { return m_saturation_mode;}
    unsigned dimension() const { return m_geom_spec;}
+   unsigned barrier_op() const {return m_barrier_op;}
+   unsigned reduction_op() const {return m_reduction_op;}
    enum vote_mode_t { vote_any, vote_all, vote_uni, vote_ballot };
    enum vote_mode_t vote_mode() const { return m_vote_mode; }
 
@@ -1021,8 +1023,10 @@ public:
       return false;
    }
 
+
 private:
    void set_opcode_and_latency();
+   void set_bar_type();
    void set_fp_or_int_archop();
    void set_mul_div_or_other_archop();
 
@@ -1054,6 +1058,8 @@ private:
    unsigned            m_rounding_mode;
    unsigned            m_compare_op;
    unsigned            m_saturation_mode;
+   unsigned 		   m_barrier_op;
+   unsigned			   m_reduction_op;
 
    std::list<int>          m_scalar_type;
    memory_space_t m_space_spec;
diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h
index 1d30eb6..610efa2 100644
--- a/src/cuda-sim/ptx_sim.h
+++ b/src/cuda-sim/ptx_sim.h
@@ -413,6 +413,10 @@ public:
    }
    
    void registerExit(){m_cta_info->register_thread_exit(this);}
+   unsigned get_reduction_value(unsigned ctaid, unsigned barid) {return m_core->get_reduction_value(ctaid,barid);}
+   void and_reduction(unsigned ctaid, unsigned barid, bool value) {m_core->and_reduction(ctaid,barid,value);}
+   void or_reduction(unsigned ctaid, unsigned barid, bool value) {m_core->or_reduction(ctaid,barid,value);}
+   void popc_reduction(unsigned ctaid, unsigned barid, bool value) {m_core->popc_reduction(ctaid,barid,value);}
 
 public:
    addr_t         m_last_effective_address;