diff options
| author | Ahmed El-Shafiey <[email protected]> | 2014-08-06 16:57:40 -0800 |
|---|---|---|
| committer | Andrew Boktor <[email protected]> | 2014-08-14 13:50:59 -0700 |
| commit | fd3c208f3c51c5520dcd7150d8a521a61577a9d5 (patch) | |
| tree | d17d3c9f47c67e1f56c9e567f51fe7f56dd599c3 /src/cuda-sim/ptx.l | |
| parent | 94a6b6c513c0afaf86770d2af0e41c8d7d0da3f5 (diff) | |
Support for named bariers + bar.red + bar.arrive instructions
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 18452]
Diffstat (limited to 'src/cuda-sim/ptx.l')
| -rw-r--r-- | src/cuda-sim/ptx.l | 12 |
1 files changed, 11 insertions, 1 deletions
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; |