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

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

13 files changed, 425 insertions, 68 deletions

diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index cd4e8d7..84d165c 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -152,6 +152,7 @@ void warp_inst_t::do_atomic(bool forceDo) {
     do_atomic( m_warp_active_mask,forceDo );
 }
 
+
 void warp_inst_t::do_atomic( const active_mask_t& access_mask,bool forceDo ) {
     assert( m_isatomic && (!m_empty||forceDo) );
     for( unsigned i=0; i < m_config->warp_size; i++ )
@@ -165,6 +166,20 @@ void warp_inst_t::do_atomic( const active_mask_t& access_mask,bool forceDo ) {
     }
 }
 
+void warp_inst_t::broadcast_barrier_reduction(const active_mask_t& access_mask)
+{
+	for( unsigned i=0; i < m_config->warp_size; i++ )
+    {
+        if( access_mask.test(i) )
+        {
+            dram_callback_t &cb = m_per_scalar_thread[i].callback;
+            if( cb.thread ){
+                cb.function(cb.instruction, cb.thread);
+            }
+        }
+    }
+}
+
 void warp_inst_t::generate_mem_accesses()
 {
     if( empty() || op == MEMORY_BARRIER_OP || m_mem_accesses_created ) 
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 09b36ff..936d208 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -28,11 +28,14 @@
 #ifndef ABSTRACT_HARDWARE_MODEL_INCLUDED
 #define ABSTRACT_HARDWARE_MODEL_INCLUDED
 
-
 // Forward declarations
 class gpgpu_sim;
 class kernel_info_t;
 
+//Set a hard limit of 32 CTAs per shader [cuda only has 8]
+#define MAX_CTA_PER_SHADER 32
+#define MAX_BARRIERS_PER_CTA 16
+
 enum _memory_space_t {
    undefined_space=0,
    reg_space,
@@ -84,6 +87,24 @@ enum uarch_op_t {
 };
 typedef enum uarch_op_t op_type;
 
+
+enum uarch_bar_t {
+   NOT_BAR=-1,
+   SYNC=1,
+   ARRIVE,
+   RED
+};
+typedef enum uarch_bar_t barrier_type;
+
+enum uarch_red_t {
+   NOT_RED=-1,
+   POPC_RED=1,
+   AND_RED,
+   OR_RED
+};
+typedef enum uarch_red_t reduction_type;
+
+
 enum uarch_operand_type_t {
 	UN_OP=-1,
     INT_OP,
@@ -680,6 +701,7 @@ public:
 struct dram_callback_t {
    dram_callback_t() { function=NULL; instruction=NULL; thread=NULL; }
    void (*function)(const class inst_t*, class ptx_thread_info*);
+
    const class inst_t* instruction;
    class ptx_thread_info *thread;
 };
@@ -691,7 +713,11 @@ public:
         m_decoded=false;
         pc=(address_type)-1;
         reconvergence_pc=(address_type)-1;
-        op=NO_OP; 
+        op=NO_OP;
+        bar_type=NOT_BAR;
+        red_type=NOT_RED;
+        bar_id=(unsigned)-1;
+        bar_count=(unsigned)-1;
         oprnd_type=UN_OP;
         sp_op=OTHER_OP;
         op_pipe=UNKOWN_OP;
@@ -723,10 +749,18 @@ public:
     unsigned get_num_regs() const {return num_regs;}
     void set_num_regs(unsigned num) {num_regs=num;}
     void set_num_operands(unsigned num) {num_operands=num;}
+    void set_bar_id(unsigned id) {bar_id=id;}
+    void set_bar_count(unsigned count) {bar_count=count;}
 
     address_type pc;        // program counter address of instruction
     unsigned isize;         // size of instruction in bytes 
     op_type op;             // opcode (uarch visible)
+
+    barrier_type bar_type;
+    reduction_type red_type;
+    unsigned bar_id;
+    unsigned bar_count;
+
     types_of_operands oprnd_type;     // code (uarch visible) identify if the operation is an interger or a floating point
     special_ops sp_op;           // code (uarch visible) identify if int_alu, fp_alu, int_mul ....
     operation_pipeline op_pipe;  // code (uarch visible) identify the pipeline of the operation (SP, SFU or MEM)
@@ -793,6 +827,7 @@ public:
     }
 
     // modifiers
+    void broadcast_barrier_reduction( const active_mask_t& access_mask);
     void do_atomic(bool forceDo=false);
     void do_atomic( const active_mask_t& access_mask, bool forceDo=false );
     void clear() 
@@ -816,6 +851,7 @@ public:
     	return m_warp_active_mask;
     }
     void completed( unsigned long long cycle ) const;  // stat collection: called when the instruction is completed  
+
     void set_addr( unsigned n, new_addr_type addr ) 
     {
         if( !m_per_scalar_thread_valid ) {
@@ -856,12 +892,13 @@ public:
     void add_callback( unsigned lane_id, 
                        void (*function)(const class inst_t*, class ptx_thread_info*),
                        const inst_t *inst, 
-                       class ptx_thread_info *thread )
+                       class ptx_thread_info *thread,
+                       bool atomic)
     {
         if( !m_per_scalar_thread_valid ) {
             m_per_scalar_thread.resize(m_config->warp_size);
             m_per_scalar_thread_valid=true;
-            m_isatomic=true;
+            if(atomic) m_isatomic=true;
         }
         m_per_scalar_thread[lane_id].callback.function = function;
         m_per_scalar_thread[lane_id].callback.instruction = inst;
@@ -927,6 +964,7 @@ public:
     void print( FILE *fout ) const;
     unsigned get_uid() const { return m_uid; }
 
+
 protected:
 
     unsigned m_uid;
@@ -989,6 +1027,13 @@ class core_t {
                      calloc( m_warp_count * m_warp_size,
                              sizeof( ptx_thread_info* ) );
             initilizeSIMTStack(m_warp_count,m_warp_size);
+
+            for(unsigned i=0; i<MAX_CTA_PER_SHADER; i++){
+            	for(unsigned j=0; j<MAX_BARRIERS_PER_CTA; j++){
+            		reduction_storage[i][j]=0;
+            	}
+            }
+
         }
         virtual ~core_t() { free(m_thread); }
         virtual void warp_exit( unsigned warp_id ) = 0;
@@ -1004,6 +1049,10 @@ class core_t {
         void get_pdom_stack_top_info( unsigned warpId, unsigned *pc, unsigned *rpc ) const;
         kernel_info_t * get_kernel_info(){ return m_kernel;}
         unsigned get_warp_size() const { return m_warp_size; }
+        void and_reduction(unsigned ctaid, unsigned barid, bool value) { reduction_storage[ctaid][barid] &= value; }
+        void or_reduction(unsigned ctaid, unsigned barid, bool value) { reduction_storage[ctaid][barid] |= value; }
+        void popc_reduction(unsigned ctaid, unsigned barid, bool value) { reduction_storage[ctaid][barid] += value;}
+        unsigned get_reduction_value(unsigned ctaid, unsigned barid) {return reduction_storage[ctaid][barid];}
     protected:
         class gpgpu_sim *m_gpu;
         kernel_info_t *m_kernel;
@@ -1011,6 +1060,7 @@ class core_t {
         class ptx_thread_info ** m_thread;
         unsigned m_warp_size;
         unsigned m_warp_count;
+        unsigned reduction_storage[MAX_CTA_PER_SHADER][MAX_BARRIERS_PER_CTA];
 };
 
 
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;
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 987942a..8b47c28 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -248,6 +248,9 @@ void shader_core_config::reg_options(class OptionParser * opp)
     option_parser_register(opp, "-gpgpu_shader_cta", OPT_UINT32, &max_cta_per_core, 
                  "Maximum number of concurrent CTAs in shader (default 8)",
                  "8");
+    option_parser_register(opp, "-gpgpu_num_cta_barriers", OPT_UINT32, &max_barriers_per_cta,
+                 "Maximum number of named barriers per CTA (default 16)",
+                 "16");
     option_parser_register(opp, "-gpgpu_n_clusters", OPT_UINT32, &n_simt_clusters, 
                  "number of processing clusters",
                  "10");
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 3692c01..2fddadd 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -73,7 +73,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
                                   const struct memory_config *mem_config,
                                   shader_core_stats *stats )
    : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ),
-     m_barriers( config->max_warps_per_shader, config->max_cta_per_core ),
+     m_barriers( this, config->max_warps_per_shader, config->max_cta_per_core, config->max_barriers_per_cta, config->warp_size ),
      m_dynamic_warp_id(0)
 {
     m_cluster = cluster;
@@ -687,10 +687,13 @@ void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t*
     (*pipe_reg)->issue( active_mask, warp_id, gpu_tot_sim_cycle + gpu_sim_cycle, m_warp[warp_id].get_dynamic_warp_id() ); // dynamic instruction information
     m_stats->shader_cycle_distro[2+(*pipe_reg)->active_count()]++;
     func_exec_inst( **pipe_reg );
-    if( next_inst->op == BARRIER_OP ) 
-        m_barriers.warp_reaches_barrier(m_warp[warp_id].get_cta_id(),warp_id);
-    else if( next_inst->op == MEMORY_BARRIER_OP ) 
+    if( next_inst->op == BARRIER_OP ){
+    	m_warp[warp_id].store_info_of_last_inst_at_barrier(*pipe_reg);
+        m_barriers.warp_reaches_barrier(m_warp[warp_id].get_cta_id(),warp_id,const_cast<warp_inst_t*> (next_inst));
+
+    }else if( next_inst->op == MEMORY_BARRIER_OP ){
         m_warp[warp_id].set_membar();
+    }
 
     updateSIMTStack(warp_id,*pipe_reg);
     m_scoreboard->reserveRegisters(*pipe_reg);
@@ -865,8 +868,7 @@ void scheduler_unit::cycle()
                                     issued_inst=true;
                                     warp_inst_issued = true;
                                 }
-                            } 
-                        }
+                            }                         }
                     } else {
                         SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) fails scoreboard\n",
                                        (*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id() );
@@ -1153,7 +1155,7 @@ void shader_core_ctx::execute()
         m_fu[n]->active_lanes_in_pipeline();
         enum pipeline_stage_name_t issue_port = m_issue_port[n];
         register_set& issue_inst = m_pipeline_reg[ issue_port ];
-	warp_inst_t** ready_reg = issue_inst.get_ready();
+        warp_inst_t** ready_reg = issue_inst.get_ready();
         if( issue_inst.has_ready() && m_fu[n]->can_issue( **ready_reg ) ) {
             bool schedule_wb_now = !m_fu[n]->stallable();
             int resbus = -1;
@@ -1234,7 +1236,7 @@ void shader_core_ctx::writeback()
 
     warp_inst_t** preg = m_pipeline_reg[EX_WB].get_ready();
     warp_inst_t* pipe_reg = (preg==NULL)? NULL:*preg;
-    while( preg and !pipe_reg->empty() ) {
+    while( preg and !pipe_reg->empty()) {
     	/*
     	 * Right now, the writeback stage drains all waiting instructions
     	 * assuming there are enough ports in the register file or the
@@ -1251,6 +1253,7 @@ void shader_core_ctx::writeback()
     	 * To handle this case, we ignore the return value (thus allowing
     	 * no stalling).
     	 */
+
         m_operand_collector.writeback(*pipe_reg);
         unsigned warp_id = pipe_reg->warp_id();
         m_scoreboard->releaseRegisters( pipe_reg );
@@ -1514,6 +1517,22 @@ pipelined_simd_unit::pipelined_simd_unit( register_set* result_port, const shade
     m_core=core;
 }
 
+void pipelined_simd_unit::cycle()
+{
+    if( !m_pipeline_reg[0]->empty() ){
+        m_result_port->move_in(m_pipeline_reg[0]);
+    }
+    for( unsigned stage=0; (stage+1)<m_pipeline_depth; stage++ )
+        move_warp(m_pipeline_reg[stage], m_pipeline_reg[stage+1]);
+    if( !m_dispatch_reg->empty() ) {
+        if( !m_dispatch_reg->dispatch_delay()){
+            int start_stage = m_dispatch_reg->latency - m_dispatch_reg->initiation_interval;
+            move_warp(m_pipeline_reg[start_stage],m_dispatch_reg);
+        }
+    }
+    occupied >>=1;
+}
+
 
 void pipelined_simd_unit::issue( register_set& source_reg )
 {
@@ -2520,17 +2539,28 @@ std::list<opndcoll_rfu_t::op_t> opndcoll_rfu_t::arbiter_t::allocate_reads()
    return result;
 }
 
-barrier_set_t::barrier_set_t( unsigned max_warps_per_core, unsigned max_cta_per_core )
+barrier_set_t::barrier_set_t(shader_core_ctx *shader,unsigned max_warps_per_core, unsigned max_cta_per_core, unsigned max_barriers_per_cta, unsigned warp_size)
 {
    m_max_warps_per_core = max_warps_per_core;
    m_max_cta_per_core = max_cta_per_core;
+   m_max_barriers_per_cta = max_barriers_per_cta;
+   m_warp_size = warp_size;
+   m_shader = shader;
    if( max_warps_per_core > WARP_PER_CTA_MAX ) {
       printf("ERROR ** increase WARP_PER_CTA_MAX in shader.h from %u to >= %u or warps per cta in gpgpusim.config\n",
              WARP_PER_CTA_MAX, max_warps_per_core );
       exit(1);
    }
+   if(max_barriers_per_cta > MAX_BARRIERS_PER_CTA){
+	   printf("ERROR ** increase MAX_BARRIERS_PER_CTA in abstract_hardware_model.h from %u to >= %u or barriers per cta in gpgpusim.config\n",
+			   MAX_BARRIERS_PER_CTA, max_barriers_per_cta );
+	   exit(1);
+   }
    m_warp_active.reset();
    m_warp_at_barrier.reset();
+   for(unsigned i=0; i<max_barriers_per_cta; i++){
+	   m_bar_id_to_warps[i].reset();
+   }
 }
 
 // during cta allocation
@@ -2544,6 +2574,10 @@ void barrier_set_t::allocate_barrier( unsigned cta_id, warp_set_t warps )
   
    m_warp_active |= warps;
    m_warp_at_barrier &= ~warps;
+   for(unsigned i=0; i<m_max_barriers_per_cta; i++){
+	   m_bar_id_to_warps[i] &=~warps;
+   }
+
 }
 
 // during cta deallocation
@@ -2559,12 +2593,23 @@ void barrier_set_t::deallocate_barrier( unsigned cta_id )
    assert( active.any() == false ); // no warps in CTA still running
    m_warp_active &= ~warps;
    m_warp_at_barrier &= ~warps;
+
+   for(unsigned i=0; i<m_max_barriers_per_cta; i++){
+	   warp_set_t at_a_specific_barrier = warps & m_bar_id_to_warps[i];
+	   assert( at_a_specific_barrier.any() == false ); // no warps stuck at barrier
+	   m_bar_id_to_warps[i] &=~warps;
+   }
    m_cta_to_warps.erase(w);
 }
 
 // individual warp hits barrier
-void barrier_set_t::warp_reaches_barrier( unsigned cta_id, unsigned warp_id )
+void barrier_set_t::warp_reaches_barrier(unsigned cta_id,unsigned warp_id,warp_inst_t* inst)
 {
+	barrier_type bar_type = inst->bar_type;
+	reduction_type red_type = inst->red_type;
+	unsigned bar_id = inst->bar_id;
+	unsigned bar_count = inst->bar_count;
+	assert(bar_id!=(unsigned)-1);
    cta_to_warp_t::iterator w=m_cta_to_warps.find(cta_id);
 
    if( w == m_cta_to_warps.end() ) { // cta is active
@@ -2574,24 +2619,42 @@ void barrier_set_t::warp_reaches_barrier( unsigned cta_id, unsigned warp_id )
    }
    assert( w->second.test(warp_id) == true ); // warp is in cta
 
-   m_warp_at_barrier.set(warp_id);
-
+   m_bar_id_to_warps[bar_id].set(warp_id);
+   if(bar_type==SYNC || bar_type==RED){
+	   m_warp_at_barrier.set(warp_id);
+   }
    warp_set_t warps_in_cta = w->second;
-   warp_set_t at_barrier = warps_in_cta & m_warp_at_barrier;
+   warp_set_t at_barrier = warps_in_cta & m_bar_id_to_warps[bar_id];
    warp_set_t active = warps_in_cta & m_warp_active;
+   if(bar_count==(unsigned)-1){
+	   if( at_barrier == active ) {
+		   printf("Barrier should be released\n");
+		   // all warps have reached barrier, so release waiting warps...
+		   m_bar_id_to_warps[bar_id] &= ~at_barrier;
+		   m_warp_at_barrier &= ~at_barrier;
+		   if(bar_type==RED){
+			   printf("broadcast_barrier\n");
+			   m_shader->broadcast_barrier_reduction(cta_id, bar_id,at_barrier);
+		   }
+		   printf("warps at barrier = %u\n",m_warp_at_barrier.count());
+	   }
+  }else{
+	  // TODO: check on the hardware if the count should include warp that exited
+	  if ((at_barrier.count() * m_warp_size) == bar_count){
+		   // required number of warps have reached barrier, so release waiting warps...
+		   m_bar_id_to_warps[bar_id] &= ~at_barrier;
+		   m_warp_at_barrier &= ~at_barrier;
+		   if(bar_type==RED){
+			   printf("REDUCTION CALLED-2\n");
+			   m_shader->broadcast_barrier_reduction(cta_id, bar_id,at_barrier);
+		   }
+	  }
+  }
 
-   if( at_barrier == active ) {
-      // all warps have reached barrier, so release waiting warps...
-      m_warp_at_barrier &= ~at_barrier;
-   }
-}
 
-// fetching a warp
-bool barrier_set_t::available_for_fetch( unsigned warp_id ) const
-{
-   return m_warp_active.test(warp_id) && m_warp_at_barrier.test(warp_id);
 }
 
+
 // warp reaches exit 
 void barrier_set_t::warp_exit( unsigned warp_id )
 {
@@ -2605,12 +2668,15 @@ void barrier_set_t::warp_exit( unsigned warp_id )
       if (w->second.test(warp_id) == true) break; 
    }
    warp_set_t warps_in_cta = w->second;
-   warp_set_t at_barrier = warps_in_cta & m_warp_at_barrier;
    warp_set_t active = warps_in_cta & m_warp_active;
 
-   if( at_barrier == active ) {
-      // all warps have reached barrier, so release waiting warps...
-      m_warp_at_barrier &= ~at_barrier;
+   for(unsigned i=0; i<m_max_barriers_per_cta; i++){
+	   warp_set_t at_a_specific_barrier = warps_in_cta & m_bar_id_to_warps[i];
+	   if( at_a_specific_barrier == active ) {
+	      // all warps have reached barrier, so release waiting warps...
+		   m_bar_id_to_warps[i] &= ~at_a_specific_barrier;
+		   m_warp_at_barrier &= ~at_a_specific_barrier;
+	   }
    }
 }
 
@@ -2620,11 +2686,12 @@ bool barrier_set_t::warp_waiting_at_barrier( unsigned warp_id ) const
    return m_warp_at_barrier.test(warp_id);
 }
 
-void barrier_set_t::dump() const
+void barrier_set_t::dump()
 {
    printf( "barrier set information\n");
    printf( "  m_max_cta_per_core = %u\n",  m_max_cta_per_core );
    printf( "  m_max_warps_per_core = %u\n", m_max_warps_per_core );
+   printf( " m_max_barriers_per_cta =%u\n", m_max_barriers_per_cta);
    printf( "  cta_to_warps:\n");
    
    cta_to_warp_t::const_iterator i;
@@ -2635,6 +2702,10 @@ void barrier_set_t::dump() const
    }
    printf("  warp_active: %s\n", m_warp_active.to_string().c_str() );
    printf("  warp_at_barrier: %s\n", m_warp_at_barrier.to_string().c_str() );
+   for( unsigned i=0; i<m_max_barriers_per_cta; i++){
+	   warp_set_t warps_reached_barrier = m_bar_id_to_warps[i];
+	   printf("  warp_at_barrier %u: %s\n", i, warps_reached_barrier.to_string().c_str() );
+   }
    fflush(stdout); 
 }
 
@@ -2658,6 +2729,18 @@ void shader_core_ctx::warp_exit( unsigned warp_id )
 		m_barriers.warp_exit( warp_id );
 }
 
+bool shader_core_ctx::check_if_non_released_reduction_barrier(warp_inst_t &inst)
+{
+	unsigned warp_id = inst.warp_id();
+	bool bar_red_op = (inst.op == BARRIER_OP) && (inst.bar_type == RED);
+    bool non_released_barrier_reduction = false;
+    bool warp_stucked_at_barrier = warp_waiting_at_barrier(warp_id);
+    bool single_inst_in_pipeline = (m_warp[warp_id].num_issued_inst_in_pipeline()==1);
+    non_released_barrier_reduction = single_inst_in_pipeline and warp_stucked_at_barrier and bar_red_op;
+    printf("non_released_barrier_reduction=%u\n",non_released_barrier_reduction);
+    return non_released_barrier_reduction;
+}
+
 bool shader_core_ctx::warp_waiting_at_barrier( unsigned warp_id ) const
 {
    return m_barriers.warp_waiting_at_barrier(warp_id);
@@ -2690,6 +2773,17 @@ void shader_core_ctx::decrement_atomic_count( unsigned wid, unsigned n )
    m_warp[wid].dec_n_atomic(n);
 }
 
+void shader_core_ctx::broadcast_barrier_reduction(unsigned cta_id,unsigned bar_id,warp_set_t warps)
+{
+	unsigned value = get_reduction_value(cta_id,bar_id);
+	for(unsigned i=0; i<m_config->max_warps_per_shader;i++){
+		if(warps.test(i)){
+			printf("braodcast for i=%u\n",i);
+			const warp_inst_t * inst = m_warp[i].restore_info_of_last_inst_at_barrier();
+			const_cast<warp_inst_t *> (inst)->broadcast_barrier_reduction(inst->get_active_mask());
+		}
+	}
+}
 
 bool shader_core_ctx::fetch_unit_response_buffer_full() const
 {
@@ -3340,7 +3434,7 @@ void simt_core_cluster::get_L1T_sub_stats(struct cache_sub_stats &css) const{
 
 void shader_core_ctx::checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t, unsigned tid)
 {
-    if( inst.has_callback(t) ) 
+    if(inst.isatomic())
            m_warp[inst.warp_id()].inc_n_atomic();
         if (inst.space.is_local() && (inst.is_load() || inst.is_store())) {
             new_addr_type localaddrs[MAX_ACCESSES_PER_INSN_PER_THREAD];
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index 78ac55e..cf7160f 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -70,8 +70,6 @@
 
 #define WRITE_MASK_SIZE 8
 
-//Set a hard limit of 32 CTAs per shader [cuda only has 8]
-#define MAX_CTA_PER_SHADER 32
 
 class thread_ctx_t {
 public:
@@ -109,6 +107,7 @@ public:
         m_done_exit=true;
         m_last_fetch=0;
         m_next=0;
+        m_inst_at_barrier=NULL;
     }
     void init( address_type start_pc,
                unsigned cta_id,
@@ -157,6 +156,9 @@ public:
     address_type get_pc() const { return m_next_pc; }
     void set_next_pc( address_type pc ) { m_next_pc = pc; }
 
+    void store_info_of_last_inst_at_barrier(const warp_inst_t *pI){ m_inst_at_barrier = pI;}
+    const warp_inst_t * restore_info_of_last_inst_at_barrier(){ return m_inst_at_barrier;}
+
     void ibuffer_fill( unsigned slot, const warp_inst_t *pI )
     {
        assert(slot < IBUFFER_SIZE );
@@ -201,6 +203,17 @@ public:
         m_stores_outstanding--;
     }
 
+    unsigned num_inst_in_buffer() const
+    {
+    	unsigned count=0;
+        for(unsigned i=0;i<IBUFFER_SIZE;i++) {
+            if( m_ibuffer[i].m_valid )
+            	count++;
+        }
+    	return count;
+    }
+    unsigned num_inst_in_pipeline() const { return m_inst_in_pipeline;}
+    unsigned num_issued_inst_in_pipeline() const {return (num_inst_in_pipeline()-num_inst_in_buffer());}
     bool inst_in_pipeline() const { return m_inst_in_pipeline > 0; }
     void inc_inst_in_pipeline() { m_inst_in_pipeline++; }
     void dec_inst_in_pipeline() 
@@ -233,6 +246,8 @@ private:
        const warp_inst_t *m_inst;
        bool m_valid;
     };
+
+    const warp_inst_t *m_inst_at_barrier;
     ibuffer_entry m_ibuffer[IBUFFER_SIZE]; 
     unsigned m_next;
                                    
@@ -879,7 +894,7 @@ private:
 
 class barrier_set_t {
 public:
-   barrier_set_t( unsigned max_warps_per_core, unsigned max_cta_per_core );
+   barrier_set_t(shader_core_ctx * shader, unsigned max_warps_per_core, unsigned max_cta_per_core, unsigned max_barriers_per_cta, unsigned warp_size);
 
    // during cta allocation
    void allocate_barrier( unsigned cta_id, warp_set_t warps );
@@ -888,12 +903,12 @@ public:
    void deallocate_barrier( unsigned cta_id );
 
    typedef std::map<unsigned, warp_set_t >  cta_to_warp_t;
+   typedef std::map<unsigned, warp_set_t >  bar_id_to_warp_t; /*set of warps reached a specific barrier id*/
+
 
    // individual warp hits barrier
-   void warp_reaches_barrier( unsigned cta_id, unsigned warp_id );
+   void warp_reaches_barrier( unsigned cta_id, unsigned warp_id, warp_inst_t* inst);
 
-   // fetching a warp
-   bool available_for_fetch( unsigned warp_id ) const;
 
    // warp reaches exit 
    void warp_exit( unsigned warp_id );
@@ -902,15 +917,19 @@ public:
    bool warp_waiting_at_barrier( unsigned warp_id ) const;
 
    // debug
-   void dump() const;
+   void dump();
 
 private:
    unsigned m_max_cta_per_core;
    unsigned m_max_warps_per_core;
-
-   cta_to_warp_t m_cta_to_warps; 
+   unsigned m_max_barriers_per_cta;
+   unsigned m_warp_size;
+   cta_to_warp_t m_cta_to_warps;
+   bar_id_to_warp_t m_bar_id_to_warps;
    warp_set_t m_warp_active;
    warp_set_t m_warp_at_barrier;
+   shader_core_ctx *m_shader;
+
 };
 
 struct insn_latency_info {
@@ -969,21 +988,7 @@ public:
     pipelined_simd_unit( register_set* result_port, const shader_core_config *config, unsigned max_latency, shader_core_ctx *core );
 
     //modifiers
-    virtual void cycle() 
-    {
-        if( !m_pipeline_reg[0]->empty() ){
-            m_result_port->move_in(m_pipeline_reg[0]);
-        }
-        for( unsigned stage=0; (stage+1)<m_pipeline_depth; stage++ ) 
-            move_warp(m_pipeline_reg[stage], m_pipeline_reg[stage+1]);
-        if( !m_dispatch_reg->empty() ) {
-            if( !m_dispatch_reg->dispatch_delay()) {
-                int start_stage = m_dispatch_reg->latency - m_dispatch_reg->initiation_interval;
-                move_warp(m_pipeline_reg[start_stage],m_dispatch_reg);
-            }
-        }
-        occupied >>=1;
-    }
+    virtual void cycle();
     virtual void issue( register_set& source_reg );
     virtual unsigned get_active_lanes_in_pipeline()
     {
@@ -1265,7 +1270,7 @@ struct shader_core_config : public core_config
     unsigned n_regfile_gating_group;
     unsigned max_warps_per_shader; 
     unsigned max_cta_per_core; //Limit on number of concurrent CTAs in shader core
-
+    unsigned max_barriers_per_cta;
     char * gpgpu_scheduler_string;
 
     char* pipeline_widths_string;
@@ -1572,6 +1577,7 @@ public:
     void cache_flush();
     void accept_fetch_response( mem_fetch *mf );
     void accept_ldst_unit_response( class mem_fetch * mf );
+    void broadcast_barrier_reduction(unsigned cta_id, unsigned bar_id,warp_set_t warps);
     void set_kernel( kernel_info_t *k ) 
     {
         assert(k);
@@ -1728,8 +1734,9 @@ public:
 	 void incfumemactivelanes_stat(unsigned active_count) {m_stats->m_active_fu_mem_lanes[m_sid]=m_stats->m_active_fu_mem_lanes[m_sid]+active_count;}
 
 	 void inc_simt_to_mem(unsigned n_flits){ m_stats->n_simt_to_mem[m_sid] += n_flits; }
+	 bool check_if_non_released_reduction_barrier(warp_inst_t &inst);
 
-private:
+	private:
 	 unsigned inactive_lanes_accesses_sfu(unsigned active_count,double latency){
       return  ( ((32-active_count)>>1)*latency) + ( ((32-active_count)>>3)*latency) + ( ((32-active_count)>>3)*latency);
 	 }