diff mbox

[6/6,GOMP4] OpenACC 1.0+ support in fortran front-end

Message ID 52E159F9.4060209@samsung.com
State New
Headers show

Commit Message

Ilmir Usmanov Jan. 23, 2014, 6:05 p.m. UTC

Comments

Ilmir Usmanov Jan. 23, 2014, 6:09 p.m. UTC | #1 
Finally, ChangeLog entry.
23-01-2014 Ilmir Usmanov <i.usmanov@samsung.com>
	Add OpenACC 1.0 support to fortran FE and GENERIC, except loop directive
	and subarrays.

	gcc/fortran/
	* decl.c (gfc_match_end): Match end of OpenACC region.
	* dump-parse-tree.c 
	(show_oacc_node): New function to dump OpenACC executable statements.
	(show_code_node): Call it.
	* gfortran.h 
	(ST_OACC_PARALLEL_LOOP, ST_OACC_END_PARALLEL_LOOP, ST_OACC_PARALLEL, 
	ST_OACC_END_PARALLEL, ST_OACC_KERNELS, ST_OACC_END_KERNELS, 
	ST_OACC_DATA, ST_OACC_END_DATA, ST_OACC_HOST_DATA, 
	ST_OACC_END_HOST_DATA, ST_OACC_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, 
	ST_OACC_WAIT, ST_OACC_CACHE, ST_OACC_KERNELS_LOOP, 
	ST_OACC_END_KERNELS_LOOP, ST_OACC_ENTER_DATA, 
	ST_OACC_EXIT_DATA): New statements.
	(gfc_exprlist): New structure to hold list of expressions.
	(OACC_LIST_PRIVATE, OACC_LIST_REDUCTION_FIRST, OACC_LIST_REDUCTION_LAST, 
	OACC_LIST_COPY, OACC_LIST_FIRST, OACC_LIST_DATA_CLAUSE_FIRST, 
	OACC_LIST_COPYIN, OACC_LIST_COPYOUT, OACC_LIST_CREATE, OACC_LIST_DELETE,
	OACC_LIST_PRESENT, OACC_LIST_PRESENT_OR_COPY, 
	OACC_LIST_PRESENT_OR_COPYIN, OACC_LIST_PRESENT_OR_COPYOUT, 
	OACC_LIST_PRESENT_OR_CREATE, OACC_LIST_DEVICEPTR, 
	OACC_LIST_DATA_CLAUSE_LAST, OACC_LIST_USE_DEVICE,
	OACC_LIST_DEVICE_RESIDENT, OACC_LIST_HOST, OACC_LIST_DEVICE, 
	OACC_LIST_CACHE, OACC_LIST_NUM): New types of list, allowed in clauses.
	(gfc_omp_clauses): Add OpenACC clauses.
	(gfc_oacc_clauses): Pseudo structure.
	(gfc_get_oacc_clauses): New function.
	(gfc_namespace): Add OpenACC declare directive clauses.
	(EXEC_OACC_KERNELS_LOOP, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_PARALLEL,
	EXEC_OACC_KERNELS, EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
	EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE, EXEC_OACC_ENTER_DATA,
	EXEC_OACC_EXIT_DATA): New executable statements.
	(gfc_free_exprlist): New function declaration.
	(gfc_resolve_oacc_directive): Likewise.
	(gfc_resolve_oacc_parallel_loop_blocks): Likewise.
	(gfc_resolve_oacc_blocks): Likewise.
	* match.c (match_exit_cycle): Add support of OpenACC regions and loops.
	* match.h (gfc_match_oacc_cache): New function declaration.
	(gfc_match_oacc_wait, gfc_match_oacc_update): Likewise.
	(gfc_match_oacc_declare, gfc_match_oacc_loop): Likewise.
	(gfc_match_oacc_host_data, gfc_match_oacc_data): Likewise.
	(gfc_match_oacc_kernels, gfc_match_oacc_kernels_loop): Likewise.
	(gfc_match_oacc_parallel, gfc_match_oacc_parallel_loop): Likewise.
	(gfc_match_oacc_enter_data, gfc_match_oacc_exit_data): Likewise.
	* parse.c (decode_oacc_directive): New function.
	(verify_token_free, verify_token_fixed): New helper functions.
	(next_free, next_fixed): Decode !$ACC sentinel.
	(case_executable): Add ST_OACC_UPDATE, ST_OACC_WAIT, ST_OACC_CACHE, 
	ST_OACC_ENTER_DATA and ST_OACC_EXIT_DATA directives.
	(case_exec_markers): Add ST_OACC_PARALLEL_LOOP, ST_OACC_PARALLEL, 
	ST_OACC_KERNELS, ST_OACC_DATA, ST_OACC_HOST_DATA, ST_OACC_LOOP and 
	ST_OACC_KERNELS_LOOP directives.
	(push_state): Initialize OpenACC declare clauses.
	(gfc_ascii_statement): Dump names of OpenACC directives.
	(verify_st_order): Verify OpenACC declare directive as declarative.
	(parse_spec): Push clauses to state stack when declare directive is 
	parsed.
	(parse_oacc_structured_block, parse_oacc_loop): New functions.
	(parse_executable): Call them.
	(parse_progunit): Move declare clauses from state stack to namespace.
	* parse.h 
	(COMP_OACC_STRUCTURED_BLOCK): New structured block to represent OpenACC 
	region.
	(gfc_state_data): Add declare directive's clauses.
	* resolve.c (gfc_resolve_blocks): Resolve OpenACC directives.
	(resolve_code): Likewise.
	* scanner.c (openacc_flag, openacc_locus): New static variables.
	(skip_oacc_attribute, skip_omp_attribute): New helper functions.
	(skip_free_comments, skip_fixed_comments): Don't skip !$ACC sentinel.
	(gfc_next_char_literal): Support OpenACC directives.
	* st.c (gfc_free_statement): Free also OpenACC directives.
	* openmp.c (gfc_free_omp_clauses): Remove also OpenACC clauses.
	(gfc_free_exprlist): New function to clear expression list.
	(match_oacc_exprlist): New function to match expression list.
	(match_oacc_clause_gang): New function to match OpenACC 2.0 gang clauses.
	(OACC_CLAUSE_IF, OACC_CLAUSE_ASYNC, OACC_CLAUSE_NUM_GANGS, 
	OACC_CLAUSE_NUM_WORKERS, OACC_CLAUSE_VECTOR_LENGTH, 
	OACC_CLAUSE_REDUCTION, OACC_CLAUSE_COPY, OACC_CLAUSE_COPYIN, 
	OACC_CLAUSE_COPYOUT, OACC_CLAUSE_CREATE, OACC_CLAUSE_PRESENT, 
	OACC_CLAUSE_PRESENT_OR_COPY, OACC_CLAUSE_PRESENT_OR_COPYIN, 
	OACC_CLAUSE_PRESENT_OR_COPYOUT, OACC_CLAUSE_PRESENT_OR_CREATE, 
	OACC_CLAUSE_DEVICEPTR, OACC_CLAUSE_PRIVATE, OACC_CLAUSE_FIRSTPRIVATE, 
	OACC_CLAUSE_COLLAPSE, OACC_CLAUSE_GANG, OACC_CLAUSE_WORKER, 
	OACC_CLAUSE_VECTOR, OACC_CLAUSE_SEQ, OACC_CLAUSE_INDEPENDENT, 
	OACC_CLAUSE_USE_DEVICE, OACC_CLAUSE_HOST, OACC_CLAUSE_DEVICE_RESIDENT, 
	OACC_CLAUSE_DEVICE, OACC_CLAUSE_DEFAULT, OACC_CLAUSE_WAIT, 
	OACC_CLAUSE_DELETE, OACC_CLAUSE_AUTO, OACC_CLAUSE_TILE): New clauses.
	(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES, 
	OACC_LOOP_CLAUSES, OACC_PARALLEL_LOOP_CLAUSES, 
	OACC_KERNELS_LOOP_CLAUSES, OACC_HOST_DATA_CLAUSES, OACC_DECLARE_CLAUSES,
	OACC_UPDATE_CLAUSES, OACC_ENTER_DATA_CLAUSES, 
	OACC_EXIT_DATA_CLAUSES): New defines.
	(gfc_match_oacc_parallel_loop): New matcher function.
	(gfc_match_oacc_parallel, gfc_match_oacc_kernels_loop): Likewise.
	(gfc_match_oacc_kernels, gfc_match_oacc_data): Likewise.
	(gfc_match_oacc_host_data, gfc_match_oacc_loop): Likewise.
	(gfc_match_oacc_declare, gfc_match_oacc_update): Likewise.
	(gfc_match_oacc_enter_data, gfc_match_oacc_exit_data): Likewise.
	(gfc_match_oacc_wait, gfc_match_oacc_cache): Likewise.
	(gfc_match_omp_clauses): Match also OpenACC clauses. Add parameter.
	(gfc_match_omp_parallel): Update. 
	(gfc_match_omp_task, gfc_match_omp_do): Likewise.
	(gfc_match_omp_parallel_do, gfc_match_omp_parallel_sections): Likewise.
	(gfc_match_omp_parallel_workshare, gfc_match_omp_sections): Likewise.
	(gfc_match_omp_single, gfc_match_omp_end_single): Likewise.
	(resolve_oacc_scalar_int_expr): New function.
	(resolve_oacc_positive_int_expr): Likewise.
	(resolve_omp_clauses): Resolve also OpenACC clauses.
	(oacc_is_parallel): New helper function.
	(oacc_is_kernels, oacc_is_loop): Likewise.
	(resolve_oacc_nested_loops): New resolve function.
	(resolve_oacc_params_in_parallel, resolve_oacc_loop_blocks): Likewise.
	(gfc_resolve_oacc_blocks, resolve_oacc_loop): Likewise.
	(resolve_oacc_cache, resolve_oacc_wait): Likewise.
	(gfc_resolve_oacc_directive): Likewise.
	* trans-decl.c
	(gfc_generate_function_code): Insert OACC_DECLARE GENERIC node.
	* trans-openmp.c (gfc_trans_oacc_add_clause): New helper function.
	(gfc_trans_oacc_variable, gfc_convert_oacc_expr_to_tree): Likewise.
	(gfc_trans_oacc_variable_list, gfc_trans_oacc_code): Likewise.
	(gfc_trans_omp_array_reduction): Support also OpenACC. Add parameter.
	(gfc_trans_omp_reduction_list): Update.
	(gfc_trans_oacc_reduction_list): New transform function.
	(gfc_trans_oacc_construct): Likewise.
	(gfc_trans_oacc_executable_directive): Likewise.
	(gfc_trans_oacc_combined_directive, gfc_trans_oacc_declare): Likewise.
	(gfc_trans_oacc_directive): Use them.
	(gfc_trans_oacc_loop): Stub.
	(gfc_trans_oacc_clauses): Transform OpenACC clauses.
	(gfc_trans_omp_clauses): Use it.
	* trans-stmt.h  (gfc_trans_oacc_directive): New function prototype.
	(gfc_trans_oacc_declare): Likewise.
	* trans.c (trans_code): Transform also OpenACC directives.
	gcc/
	* gimplify.c (is_gimple_stmt): Stub OpenACC directives and clauses.
	(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Likewise.
	(gimplify_expr): Likewise.
	* omp-low.c (scan_sharing_clauses): Likewise.
	* tree-core.h 
	(OACC_CLAUSE_IF, OACC_CLAUSE_ASYNC, OACC_CLAUSE_NUM_GANGS, 
	OACC_CLAUSE_NUM_WORKERS, OACC_CLAUSE_VECTOR_LENGTH, 
	OACC_CLAUSE_REDUCTION, OACC_CLAUSE_COPY, OACC_CLAUSE_COPYIN, 
	OACC_CLAUSE_COPYOUT, OACC_CLAUSE_CREATE, OACC_CLAUSE_PRESENT, 
	OACC_CLAUSE_PRESENT_OR_COPY, OACC_CLAUSE_PRESENT_OR_COPYIN, 
	OACC_CLAUSE_PRESENT_OR_COPYOUT, OACC_CLAUSE_PRESENT_OR_CREATE, 
	OACC_CLAUSE_DEVICEPTR, OACC_CLAUSE_PRIVATE, OACC_CLAUSE_FIRSTPRIVATE, 
	OACC_CLAUSE_COLLAPSE, OACC_CLAUSE_GANG, OACC_CLAUSE_WORKER, 
	OACC_CLAUSE_VECTOR, OACC_CLAUSE_SEQ, OACC_CLAUSE_INDEPENDENT, 
	OACC_CLAUSE_USE_DEVICE, OACC_CLAUSE_HOST, OACC_CLAUSE_DEVICE_RESIDENT, 
	OACC_CLAUSE_DEVICE, OACC_CLAUSE_DEFAULT, OACC_CLAUSE_WAIT, 
	OACC_CLAUSE_DELETE): New clauses.
	* tree-pretty-print.c (dump_oacc_body): Print OpenACC region.
	(dump_oacc_clause_remap): New helper function.
	(dump_oacc_clause): Print OpenACC clause.
	(dump_oacc_clauses): Print OpenACC clauses.
	(dump_generic_node): Print OpenACC directives and its clauses.
	* tree-pretty-print.h: (dump_oacc_body): New function prototype.
	(dump_oacc_clause_remap, dump_oacc_clause, dump_oacc_clauses): Likewise.
	* tree.c (tree_node_kind_names): New name.
	(omp_clause_num_ops): Add OpenACC clauses.
	(omp_clause_code_name): Likewise.
	(walk_tree_1): Likewise.
	* tree.def (OACC_KERNELS): New tree node.
	(OACC_DATA, OACC_HOST_DATA, OACC_DECLARE, OACC_UPDATE): Likewise.
	(OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_WAIT, OACC_CACHE): Likewise.
	* tree.h (OACC_BODY): New macros.
	(OACC_CLAUSE_CHAIN, OACC_CLAUSE, OACC_CLAUSE_DECL, OACC_CLAUSE_CODE, 
	OACC_CLAUSE_OPERAND, OACC_CLAUSE_SET_CODE, OACC_CLAUSE_LOCATION, 
	OACC_PARALLEL_BODY, OACC_PARALLEL_CLAUSES, OACC_KERNELS_BODY, 
	OACC_KERNELS_CLAUSES, OACC_CLAUSE_NUM_GANGS_EXPR, 
	OACC_CLAUSE_NUM_WORKERS_EXPR, OACC_CLAUSE_VECTOR_LENGTH_EXPR, 
	OACC_CLAUSE_VECTOR_EXPR, OACC_CLAUSE_WORKER_EXPR, OACC_CLAUSE_GANG_EXPR, 
	OACC_CLAUSE_COLLAPSE_EXPR, OACC_CLAUSE_IF_EXPR, OACC_CLAUSE_ASYNC_EXPR, 
	OACC_WAIT_EXPR, OACC_DATA_BODY, OACC_DATA_CLAUSES, OACC_DECLARE_CLAUSES, 
	OACC_UPDATE_CLAUSES, OACC_WAIT_CLAUSES, OACC_CACHE_CLAUSES, 
	OACC_HOST_DATA_BODY, OACC_HOST_DATA_CLAUSES, OACC_CLAUSE_REDUCTION_CODE, 
	OACC_CLAUSE_REDUCTION_INIT, OACC_CLAUSE_REDUCTION_MERGE, 
	OACC_CLAUSE_REDUCTION_PLACEHOLDER): Likewise.
	(build_oacc_clause): New helper function.
	gcc/testsuite/gfortran.dg/goacc/
	* goacc.exp: New test directory.
	* branch.f95: New test.
	* continuation-free-form.f95: Likewise.
	* data-clauses.f95: Likewise.
	* data-tree.f95: Likewise.
	* declare-1.f95: Likewise.
	* declare.f95: Likewise.
	* directive-names.f95: Likewise.
	* enter-exit-data.f95: Likewise.
	* host_data-tree.f95: Likewise.
	* if.f95: Likewise.
	* kernels-tree.f95: Likewise.
	* list.f95: Likewise.
	* parallel-kernels-clauses.f95: Likewise.
	* parallel-kernels-regions.f95: Likewise.
	* parallel-tree.f95: Likewise.
	* pure-elemental-procedures.f95: Likewise.
	* reduction.f95: Likewise.
	* sentinel-free-form.f95: Likewise.
	* several-directives.f95: Likewise.
	* sie.f95: Likewise.
	gcc/doc/
	* generic.texi: Document OACC_KERNELS, OACC_DATA, OACC_HOST_DATA, 
	OACC_DECLARE, OACC_UPDATE, OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_WAIT, 
	OACC_CACHE.
diff mbox

Patch

From 668a50443d0a70a633707ec49759d9e8b6d00b1e Mon Sep 17 00:00:00 2001
From: Ilmir Usmanov <i.usmanov@samsung.com>
Date: Thu, 23 Jan 2014 21:26:50 +0400
Subject: [PATCH 6/6] OpenACC GENERIC docs

---
 gcc/doc/generic.texi | 40 ++++++++++++++++++++++++++++++++++++++++
 1 file changed, 40 insertions(+)

diff --git a/gcc/doc/generic.texi b/gcc/doc/generic.texi
index d51033e..dd88489 100644
--- a/gcc/doc/generic.texi
+++ b/gcc/doc/generic.texi
@@ -2052,6 +2052,14 @@  edge.  Rethrowing the exception is represented using @code{RESX_EXPR}.
 @node OpenMP
 @subsection OpenMP
 @tindex OACC_PARALLEL
+@tindex OACC_KERNELS
+@tindex OACC_DATA
+@tindex OACC_HOST_DATA
+@tindex OACC_UPDATE
+@tindex OACC_ENTER_DATA
+@tindex OACC_EXIT_DATA
+@tindex OACC_WAIT
+@tindex OACC_CACHE
 @tindex OMP_PARALLEL
 @tindex OMP_FOR
 @tindex OMP_SECTIONS
@@ -2073,6 +2081,38 @@  clauses used by the OpenMP API @w{@uref{http://www.openmp.org/}}.
 
 Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}.
 
+@item OACC_KERNELS
+
+Represents @code{#pragma acc kernels [clause1 @dots{} clauseN]}.
+
+@item OACC_DATA
+
+Represents @code{#pragma acc data [clause1 @dots{} clauseN]}.
+
+@item OACC_HOST_DATA
+
+Represents @code{#pragma acc host_data [clause1 @dots{} clauseN]}.
+
+@item OACC_UPDATE
+
+Represents @code{#pragma acc update [clause1 @dots{} clauseN]}.
+
+@item OACC_ENTER_DATA
+
+Represents @code{#pragma acc enter data [clause1 @dots{} clauseN]}.
+
+@item OACC_EXIT_DATA
+
+Represents @code{#pragma acc exit data [clause1 @dots{} clauseN]}.
+
+@item OACC_WAIT
+
+Represents @code{#pragma acc wait [(num @dots{})]}.
+
+@item OACC_CACHE
+
+Represents @code{#pragma acc cache (var @dots{})}.
+
 @item OMP_PARALLEL
 
 Represents @code{#pragma omp parallel [clause1 @dots{} clauseN]}. It
-- 
1.8.3.2