summaryrefslogtreecommitdiff
path: root/src/cuda-sim
diff options
context:
space:
mode:
Diffstat (limited to 'src/cuda-sim')
-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
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;
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 00:36:40 +0000