From fd3c208f3c51c5520dcd7150d8a521a61577a9d5 Mon Sep 17 00:00:00 2001 From: Ahmed El-Shafiey Date: Wed, 6 Aug 2014 16:57:40 -0800 Subject: Support for named bariers + bar.red + bar.arrive instructions [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 18452] --- src/cuda-sim/ptx.l | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/ptx.l') 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; -- cgit v1.3