Hi! While working on something else, I came across the following. Cesar, can you please verify that this is really dead code in the Fortran front end, which currently is the only producer of OMP_CLAUSE_DEVICE_RESIDENT?
Also, I noticed that the Fortran front end never creates OMP_CLAUSE_MAP with GOMP_MAP_LINK or GOMP_MAP_DEVICE_RESIDENT -- how is OpenACC declare working at all? Cesar, if this is not a simple question to answer, please file a GCC PR, for somebody to have a look later. Jakub, if the following patch is OK, should I also clean it up on release branches as applicable, or just on trunk? commit 65f613d59aca51bc6460eaae7ea19871577a7b26 Author: Thomas Schwinge <tho...@codesourcery.com> Date: Tue May 31 15:59:24 2016 +0200 Remove the unused OMP_CLAUSE_DEVICE_RESIDENT gcc/ * tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_DEVICE_RESIDENT. Adjust all users. --- gcc/fortran/trans-openmp.c | 3 --- gcc/gimplify.c | 5 ----- gcc/omp-low.c | 2 -- gcc/tree-core.h | 3 --- gcc/tree-pretty-print.c | 3 --- gcc/tree.c | 3 --- 6 files changed, 19 deletions(-) diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c index c2d89eb..d3276f9 100644 --- gcc/fortran/trans-openmp.c +++ gcc/fortran/trans-openmp.c @@ -1773,9 +1773,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_LIST_USE_DEVICE: clause_code = OMP_CLAUSE_USE_DEVICE_PTR; goto add_clause; - case OMP_LIST_DEVICE_RESIDENT: - clause_code = OMP_CLAUSE_DEVICE_RESIDENT; - goto add_clause; add_clause: omp_clauses diff --git gcc/gimplify.c gcc/gimplify.c index 8316bb8..e53a2d3 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -7531,10 +7531,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } break; - case OMP_CLAUSE_DEVICE_RESIDENT: - remove = true; - break; - case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_ORDERED: case OMP_CLAUSE_UNTIED: @@ -8268,7 +8264,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE__CILK_FOR_COUNT_: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: - case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_NUM_GANGS: case OMP_CLAUSE_NUM_WORKERS: diff --git gcc/omp-low.c gcc/omp-low.c index a11f44b..77bdb18 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -2200,7 +2200,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, install_var_local (decl, ctx); break; - case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE__CACHE_: sorry ("Clause not supported yet"); break; @@ -2368,7 +2367,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE__GRIDDIM_: break; - case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE__CACHE_: sorry ("Clause not supported yet"); break; diff --git gcc/tree-core.h gcc/tree-core.h index b069928..db5b470 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -316,9 +316,6 @@ enum omp_clause_code { #pragma acc cache (variable-list). */ OMP_CLAUSE__CACHE_, - /* OpenACC clause: device_resident (variable_list). */ - OMP_CLAUSE_DEVICE_RESIDENT, - /* OpenACC clause: gang [(gang-argument-list)]. Where gang-argument-list: [gang-argument-list, ] gang-argument diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index 0e7fdd1..734ecda 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -407,9 +407,6 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) case OMP_CLAUSE__LOOPTEMP_: name = "_looptemp_"; goto print_remap; - case OMP_CLAUSE_DEVICE_RESIDENT: - name = "device_resident"; - goto print_remap; case OMP_CLAUSE_TO_DECLARE: name = "to"; goto print_remap; diff --git gcc/tree.c gcc/tree.c index f4b470b..7511d0a 100644 --- gcc/tree.c +++ gcc/tree.c @@ -281,7 +281,6 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_USE_DEVICE_PTR */ 1, /* OMP_CLAUSE_IS_DEVICE_PTR */ 2, /* OMP_CLAUSE__CACHE_ */ - 1, /* OMP_CLAUSE_DEVICE_RESIDENT */ 2, /* OMP_CLAUSE_GANG */ 1, /* OMP_CLAUSE_ASYNC */ 1, /* OMP_CLAUSE_WAIT */ @@ -353,7 +352,6 @@ const char * const omp_clause_code_name[] = "use_device_ptr", "is_device_ptr", "_cache_", - "device_resident", "gang", "async", "wait", @@ -11764,7 +11762,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 1)); /* FALLTHRU */ - case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: case OMP_CLAUSE_WORKER: Grüße Thomas
signature.asc
Description: PGP signature