summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CHANGES2
-rw-r--r--src/abstract_hardware_model.cc15
-rw-r--r--src/abstract_hardware_model.h58
-rw-r--r--src/cuda-sim/cuda-sim.cc44
-rw-r--r--src/cuda-sim/instructions.cc118
-rw-r--r--src/cuda-sim/opcodes.def2
-rw-r--r--src/cuda-sim/ptx.l12
-rw-r--r--src/cuda-sim/ptx.y12
-rw-r--r--src/cuda-sim/ptx_ir.cc10
-rw-r--r--src/cuda-sim/ptx_ir.h6
-rw-r--r--src/cuda-sim/ptx_sim.h4
-rw-r--r--src/gpgpu-sim/gpu-sim.cc3
-rw-r--r--src/gpgpu-sim/shader.cc150
-rw-r--r--src/gpgpu-sim/shader.h59
14 files changed, 427 insertions, 68 deletions
diff --git a/CHANGES b/CHANGES
index c994737..7352d28 100644
--- a/CHANGES
+++ b/CHANGES
@@ -3,6 +3,8 @@ Version 3.2.3+edits (development branch) versus 3.2.3
- Added a parameter to the cache configuration to configure the set index
function. Added a hash set index function to the Fermi L1 data cache for
the two default cache sizes, 16KB/48KB with 32/64 sets.
+- Added support for named barriers.
+- Added support for bar.arrive and bar.red instructions.
- Bug fixes:
- Fixed bug #81, fix ordering of pushing branch entries to the stack
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);
}
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:26:06 +0000