- add placeholders for new (e.g., PTX 2.x) directives and instructions

- update regression script to use compute sdk location for cuda (both
  distributed in one download now)

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

1 files changed, 45 insertions, 10 deletions

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;