diff options
| -rw-r--r-- | src/cuda-sim/instructions.cc | 41 | ||||
| -rw-r--r-- | src/cuda-sim/opcodes.def | 29 | ||||
| -rw-r--r-- | src/cuda-sim/ptx.l | 55 | ||||
| -rw-r--r-- | src/cuda-sim/ptx.y | 10 |
4 files changed, 119 insertions, 16 deletions
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 226e965..f3aa24a 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -670,6 +670,10 @@ void bar_sync_impl( const ptx_instruction *pI, ptx_thread_info *thread ) assert( b.u32 == 0 ); // not clear what should happen if this is not zero } +void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void bfi_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void bfind_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } + void bra_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { const operand_info &target = pI->dst(); @@ -679,6 +683,7 @@ void bra_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_npc(target_pc); } +void brev_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } void brkpt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } extern int gpgpu_simd_model; @@ -756,6 +761,8 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_npc(target_func); } +void clz_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } + void cnot_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t a, b, d; @@ -1319,6 +1326,8 @@ void cvt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_operand_value(dst,data); } +void cvta_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } + void div_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t data; @@ -1393,6 +1402,8 @@ void fma_impl( const ptx_instruction *pI, ptx_thread_info *thread ) mad_def(pI,thread); } +void isspacep_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } + void ld_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { const operand_info &dst = pI->dst(); @@ -1444,6 +1455,8 @@ void ld_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->m_last_memory_space = space; } +void ldu_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } + void lg2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t a, d; @@ -1648,10 +1661,7 @@ void max_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_operand_value(dst,d); } -void membar_impl( const ptx_instruction *pI, ptx_thread_info *thread ) -{ - /* TODO: add timing model support */ -} +void membar_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } void min_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { @@ -1955,6 +1965,12 @@ void or_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_operand_value(dst,data); } +void pmevent_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void popc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void prefetch_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void prefetchu_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void prmt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } + void rcp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t src1_data, src2_data, data; @@ -2632,6 +2648,12 @@ void sub_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_operand_value(dst,data); } +void subc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void suld_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void sured_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void sust_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void suq_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } + ptx_reg_t* ptx_tex_regs = NULL; union intfloat { @@ -2867,7 +2889,18 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->m_last_memory_space = TEX_DIRECTIVE; thread->set_vector_operand_values(dst,data1,data2,data3,data4,4); } + +void txq_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } void trap_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vabsdiff_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vadd_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vmad_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vmax_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vmin_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vset_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vshl_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vshr_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void vsub_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } extern unsigned g_warp_active_mask; diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 5a34cff..143bb93 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -69,6 +69,7 @@ Control 3 SFU 4 Mem(except Tex) 5 Tex 6 +Other 10 */ OP_DEF(ABS_OP,abs_impl,"abs",1,1) OP_DEF(ADD_OP,add_impl,"add",1,1) @@ -76,16 +77,23 @@ OP_DEF(ADDC_OP,addc_impl,"addc",1,1) OP_DEF(AND_OP,and_impl,"and",1,1) OP_DEF(ATOM_OP,atom_impl,"atom",0,3) OP_DEF(BAR_OP,bar_sync_impl,"bar.sync",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) OP_DEF(BRA_OP,bra_impl,"bra",0,3) +OP_DEF(BREV_OP,brev_impl,"brev",1,1) OP_DEF(BRKPT_OP,brkpt_impl,"brkpt",1,9) -OP_DEF(CALL_OP,call_impl,"call",1,3) /*call may actually return an value if the syntax is call (ret-param),fname,(param-list)*/ +OP_DEF(CALL_OP,call_impl,"call",1,3) +OP_DEF(CLZ_OP,clz_impl,"clz",1,1) OP_DEF(CNOT_OP,cnot_impl,"cnot",1,1) OP_DEF(COS_OP,cos_impl,"cos",1,4) OP_DEF(CVT_OP,cvt_impl,"cvt",1,1) +OP_DEF(CVTA_OP,cvta_impl,"cvta",1,1) OP_DEF(DIV_OP,div_impl,"div",1,1) OP_DEF(EX2_OP,ex2_impl,"ex2",1,4) OP_DEF(EXIT_OP,exit_impl,"exit",1,3) OP_DEF(FMA_OP,fma_impl,"fma",1,2) +OP_DEF(ISSPACEP_OP,isspacep_impl,"isspacep",1,1) OP_DEF(LD_OP,ld_impl,"ld",1,5) OP_DEF(LG2_OP,lg2_impl,"lg2",1,4) OP_DEF(MAD24_OP,mad24_impl,"mad24",1,2) @@ -99,6 +107,11 @@ OP_DEF(MUL_OP,mul_impl,"mul",1,1) OP_DEF(NEG_OP,neg_impl,"neg",1,1) OP_DEF(NOT_OP,not_impl,"not",1,1) OP_DEF(OR_OP,or_impl,"or",1,1) +OP_DEF(PMEVENT_OP,pmevent_impl,"pmevent",1,10) +OP_DEF(POPC_OP,popc_impl,"popc",1,1) +OP_DEF(PREFETCH_OP,prefetch_impl,"prefetch",1,5) +OP_DEF(PREFETCHU_OP,prefetchu_impl,"prefetchu",1,5) +OP_DEF(PRMT_OP,prmt_impl,"prmt",1,1) OP_DEF(RCP_OP,rcp_impl,"rcp",1,4) OP_DEF(RED_OP,red_impl,"red",1,7) OP_DEF(REM_OP,rem_impl,"rem",1,1) @@ -115,7 +128,21 @@ OP_DEF(SLCT_OP,slct_impl,"slct",1,1) OP_DEF(SQRT_OP,sqrt_impl,"sqrt",1,4) OP_DEF(ST_OP,st_impl,"st",0,5) OP_DEF(SUB_OP,sub_impl,"sub",1,1) +OP_DEF(SUBC_OP,subc_impl,"subc",1,1) +OP_DEF(SULD_OP,suld_impl,"suld",1,6) +OP_DEF(SURED_OP,sured_impl,"sured",1,6) +OP_DEF(SUST_OP,sust_impl,"sust",1,6) +OP_DEF(SUQ_OP,suq_impl,"suq",1,6) OP_DEF(TEX_OP,tex_impl,"tex",1,6) OP_DEF(TRAP_OP,trap_impl,"trap",1,3) +OP_DEF(VABSDIFF_OP,vabsdiff_impl,"vabsdiff",0,11) +OP_DEF(VADD_OP,vadd_impl,"vadd",0,11) +OP_DEF(VMAD_OP,vmad_impl,"vmad",0,11) +OP_DEF(VMAX_OP,vmax_impl,"vmax",0,11) +OP_DEF(VMIN_OP,vmin_impl,"vmin",0,11) +OP_DEF(VSET_OP,vset_impl,"vset",0,11) +OP_DEF(VSHL_OP,vshl_impl,"vshl",0,11) +OP_DEF(VSHR_OP,vshr_impl,"vshr",0,11) +OP_DEF(VSUB_OP,vsub_impl,"vsub",0,11) OP_DEF(VOTE_OP,vote_impl,"vote",0,3) OP_DEF(XOR_OP,xor_impl,"xor",1,1) diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 3ef1071..e4c9fef 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -81,23 +81,34 @@ int ptx_error( const char *s ); abs TC; ptx_lval.int_value = ABS_OP; return OPCODE; add TC; ptx_lval.int_value = ADD_OP; return OPCODE; +addc TC; ptx_lval.int_value = ADDC_OP; return OPCODE; and TC; ptx_lval.int_value = AND_OP; return OPCODE; atom TC; ptx_lval.int_value = ATOM_OP; return OPCODE; bar.sync 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; bra TC; ptx_lval.int_value = BRA_OP; return OPCODE; +brev TC; ptx_lval.int_value = BREV_OP; return OPCODE; +brkpt TC; ptx_lval.int_value = BRKPT_OP; return OPCODE; call TC; ptx_lval.int_value = CALL_OP; return OPCODE; +clz TC; ptx_lval.int_value = CLZ_OP; return OPCODE; cnot TC; ptx_lval.int_value = CNOT_OP; return OPCODE; cos TC; ptx_lval.int_value = COS_OP; return OPCODE; cvt TC; ptx_lval.int_value = CVT_OP; return OPCODE; +cvta TC; ptx_lval.int_value = CVTA_OP; return OPCODE; div TC; ptx_lval.int_value = DIV_OP; return OPCODE; ex2 TC; ptx_lval.int_value = EX2_OP; return OPCODE; exit TC; ptx_lval.int_value = EXIT_OP; return OPCODE; +fma TC; ptx_lval.int_value = FMA_OP; return OPCODE; +isspacep TC; ptx_lval.int_value = ISSPACEP_OP; return OPCODE; ld TC; ptx_lval.int_value = LD_OP; return OPCODE; ld.volatile TC; ptx_lval.int_value = LD_OP; return OPCODE; lg2 TC; ptx_lval.int_value = LG2_OP; return OPCODE; mad24 TC; ptx_lval.int_value = MAD24_OP; return OPCODE; mad TC; ptx_lval.int_value = MAD_OP; return OPCODE; max TC; ptx_lval.int_value = MAX_OP; return OPCODE; +membar TC; ptx_lval.int_value = MEMBAR_OP; return OPCODE; min TC; ptx_lval.int_value = MIN_OP; return OPCODE; mov TC; ptx_lval.int_value = MOV_OP; return OPCODE; mul24 TC; ptx_lval.int_value = MUL24_OP; return OPCODE; @@ -105,7 +116,13 @@ mul TC; ptx_lval.int_value = MUL_OP; return OPCODE; neg TC; ptx_lval.int_value = NEG_OP; return OPCODE; not TC; ptx_lval.int_value = NOT_OP; return OPCODE; or TC; ptx_lval.int_value = OR_OP; return OPCODE; +pmevent TC; ptx_lval.int_value = PMEVENT_OP; return OPCODE; +popc TC; ptx_lval.int_value = POPC_OP; return OPCODE; +prefetch TC; ptx_lval.int_value = PREFETCH_OP; return OPCODE; +prefetchu TC; ptx_lval.int_value = PREFETCHU_OP; return OPCODE; +prmt TC; ptx_lval.int_value = PRMT_OP; return OPCODE; rcp TC; ptx_lval.int_value = RCP_OP; return OPCODE; +red TC; ptx_lval.int_value = RED_OP; return OPCODE; rem TC; ptx_lval.int_value = REM_OP; return OPCODE; ret TC; ptx_lval.int_value = RET_OP; return OPCODE; rsqrt TC; ptx_lval.int_value = RSQRT_OP; return OPCODE; @@ -121,20 +138,33 @@ sqrt TC; ptx_lval.int_value = SQRT_OP; return OPCODE; st TC; ptx_lval.int_value = ST_OP; return OPCODE; st.volatile TC; ptx_lval.int_value = ST_OP; return OPCODE; sub TC; ptx_lval.int_value = SUB_OP; return OPCODE; +subc TC; ptx_lval.int_value = SUBC_OP; return OPCODE; +suld TC; ptx_lval.int_value = SULD_OP; return OPCODE; +sured TC; ptx_lval.int_value = SURED_OP; return OPCODE; +surst TC; ptx_lval.int_value = SUST_OP; return OPCODE; +suq TC; ptx_lval.int_value = SUQ_OP; return OPCODE; tex TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = TEX_OP; return OPCODE; +txq TC; ptx_lval.int_value = TEX_OP; return OPCODE; trap TC; ptx_lval.int_value = TRAP_OP; return OPCODE; +vabsdiff TC; ptx_lval.int_value = VABSDIFF_OP; return OPCODE; +vadd TC; ptx_lval.int_value = VADD_OP; return OPCODE; +vmad TC; ptx_lval.int_value = VMAD_OP; return OPCODE; +vmax TC; ptx_lval.int_value = VMAX_OP; return OPCODE; +vmin TC; ptx_lval.int_value = VMIN_OP; return OPCODE; +vset TC; ptx_lval.int_value = VSET_OP; return OPCODE; +vshl TC; ptx_lval.int_value = VSHL_OP; return OPCODE; +vshr TC; ptx_lval.int_value = VSHR_OP; return OPCODE; +vsub TC; ptx_lval.int_value = VSUB_OP; return OPCODE; vote TC; ptx_lval.int_value = VOTE_OP; return OPCODE; xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; -membar TC; ptx_lval.int_value = MEMBAR_OP; return OPCODE; -red TC; ptx_lval.int_value = RED_OP; return OPCODE; -brkpt TC; ptx_lval.int_value = BRKPT_OP; return OPCODE; -addc TC; ptx_lval.int_value = ADDC_OP; return OPCODE; -fma TC; ptx_lval.int_value = FMA_OP; return OPCODE; <INITIAL,NOT_OPCODE>{ \.align TC; return ALIGN_DIRECTIVE; -\.byte TC; return BYTE_DIRECTIVE; +\.branchtargets TC; return BRANCHTARGETS_DIRECTIVE; +\.byte TC; return BYTE_DIRECTIVE; /* not in PTX 2.1 */ +\.callprototype TC; return CALLPROTOTYPE_DIRECTIVE; +\.calltargets TC; return CALLTARGETS_DIRECTIVE; \.const\[[0-9]+\] TC; return CONST_DIRECTIVE; \.const TC; return CONST_DIRECTIVE; \.entry TC; return ENTRY_DIRECTIVE; @@ -144,20 +174,25 @@ fma TC; ptx_lval.int_value = FMA_OP; return OPCODE; \.global TC; return GLOBAL_DIRECTIVE; \.local TC; return LOCAL_DIRECTIVE; \.loc TC; return LOC_DIRECTIVE; +\.maxnctapersm TC; return MAXNCTAPERSM_DIRECTIVE; +\.maxnreg TC; return MAXNNREG_DIRECTIVE; +\.maxntid TC; return MAXNTID_DIRECTIVE; +\.minnctapersm TC; return MINNCTAPERSM_DIRECTIVE; \.param TC; return PARAM_DIRECTIVE; +\.pragma TC; return PRAGMA_DIRECTIVE; \.reg TC; return REG_DIRECTIVE; +\.reqntid TC; return REQNTID_DIRECTIVE; \.section TC; return SECTION_DIRECTIVE; \.shared TC; return SHARED_DIRECTIVE; \.sreg TC; return SREG_DIRECTIVE; -\.struct TC; return STRUCT_DIRECTIVE; -\.surf TC; return SURF_DIRECTIVE; +\.struct TC; return STRUCT_DIRECTIVE; /* not in PTX 2.1 */ +\.surf TC; return SURF_DIRECTIVE; /* not in PTX 2.1 */ \.target TC; return TARGET_DIRECTIVE; \.tex TC; BEGIN(NOT_OPCODE); return TEX_DIRECTIVE; -\.union TC; return UNION_DIRECTIVE; +\.union TC; return UNION_DIRECTIVE; /* not in PTX 2.1 */ \.version TC; return VERSION_DIRECTIVE; \.visible TC; return VISIBLE_DIRECTIVE; -\.maxntid TC; return MAXNTID_DIRECTIVE; "%clock" TC; ptx_lval.int_value = CLOCK_ID; return SPECIAL_REGISTER; "%ctaid" TC; ptx_lval.int_value = CTA_ID; return SPECIAL_REGISTER; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 2e1cb0f..bb10ae1 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -72,7 +72,10 @@ %token <string_value> STRING %token <int_value> OPCODE %token ALIGN_DIRECTIVE +%token BRANCHTARGETS_DIRECTIVE %token BYTE_DIRECTIVE +%token CALLPROTOTYPE_DIRECTIVE +%token CALLTARGETS_DIRECTIVE %token CONST_DIRECTIVE %token ENTRY_DIRECTIVE %token EXTERN_DIRECTIVE @@ -81,8 +84,14 @@ %token GLOBAL_DIRECTIVE %token LOCAL_DIRECTIVE %token LOC_DIRECTIVE +%token MAXNCTAPERSM_DIRECTIVE +%token MAXNNREG_DIRECTIVE +%token MAXNTID_DIRECTIVE +%token MINNCTAPERSM_DIRECTIVE %token PARAM_DIRECTIVE +%token PRAGMA_DIRECTIVE %token REG_DIRECTIVE +%token REQNTID_DIRECTIVE %token SECTION_DIRECTIVE %token SHARED_DIRECTIVE %token SREG_DIRECTIVE @@ -93,7 +102,6 @@ %token UNION_DIRECTIVE %token VERSION_DIRECTIVE %token VISIBLE_DIRECTIVE -%token MAXNTID_DIRECTIVE %token <string_value> IDENTIFIER %token <int_value> INT_OPERAND %token <float_value> FLOAT_OPERAND |