[OpenACC,0/8] Multi-dimensional dynamic array support for OpenACC data clauses

Hi Jakub, this patch set is supposed to be OpenACC functionality, but most of it touches shared code,
so I have CCed both you and Thomas.

This patch adds capability to handle C/C++ non-contiguous, dynamically
allocated multi-dimensional arrays in OpenACC data clauses:

int *a[100], **b;

#pragma acc parallel copyin (a[0:n][0:m], b[1:x][5:y]) // re-constructs array slices on GPU and copies data in

We currently only allow arrays (e.g. []) at the "outermost" dimension, for example:

// These are all okay
int **p;
int ***m;
int ****n;
int *x[100];

int (*y)[100];   // not allowed

Some of this was due to limiting the scope of implementation, but may actually be extended to support with reasonable effort.

I have added descriptions of each part of the implementation in the respective
patch mails. The test results are all okay, no regressions of any sort.
Asking for permission to apply to trunk.

Thanks,
Chung-Lin

2018-10-16  Chung-Lin Tang  <cltang@codesourcery.com>

	include/
	* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define.
	(enum gomp_map_kind): Add GOMP_MAP_DYNAMIC_ARRAY,
	GOMP_MAP_DYNAMIC_ARRAY_TO, GOMP_MAP_DYNAMIC_ARRAY_FROM,
	GOMP_MAP_DYNAMIC_ARRAY_TOFROM, GOMP_MAP_DYNAMIC_ARRAY_FORCE_TO,
	GOMP_MAP_DYNAMIC_ARRAY_FORCE_FROM, GOMP_MAP_DYNAMIC_ARRAY_FORCE_TOFROM,
	GOMP_MAP_DYNAMIC_ARRAY_ALLOC, GOMP_MAP_DYNAMIC_ARRAY_FORCE_ALLOC,
	GOMP_MAP_DYNAMIC_ARRAY_FORCE_PRESENT.
	(GOMP_MAP_DYNAMIC_ARRAY_P): Define.

	gcc/c/
	* c-typeck.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
	parameter, adjust recursive call site, add cases for allowing
	pointer based multi-dimensional arrays for OpenACC.
	(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
	handle non-contiguous case to create dynamic array map.

	gcc/cp/
	* semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
	parameter, adjust recursive call site, add cases for allowing
	pointer based multi-dimensional arrays for OpenACC.
	(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
	handle non-contiguous case to create dynamic array map.

	gcc/
	* gimplify.c (gimplify_scan_omp_clauses): For dynamic array map kinds,
	make sure bias in each dimension are put into firstprivate variables.
	* tree-pretty-print.c (dump_omp_clauses): Add cases for printing
	GOMP_MAP_DYNAMIC_ARRAY map kinds.
	* omp-low.c (struct omp_context):
	Add 'hash_map<tree_operand_hash, tree> *dynamic_arrays' field, also
	added include of "tree-hash-traits.h".
	(append_field_to_record_type): New function.
	(create_dynamic_array_descr_type): Likewise.
	(create_dynamic_array_descr_init_code): Likewise.
	(new_omp_context): Add initialize of dynamic_arrays field.
	(delete_omp_context): Add delete of dynamic_arrays field.
	(scan_sharing_clauses): For dynamic array map kinds, check for
	supported dimension structure, and install dynamic array variable into
	current omp_context.
	(lower_omp_target): Add handling for dynamic array map kinds.
	(dynamic_array_lookup): New function.
	(dynamic_array_reference_start): Likewise.
	(scan_for_op): Likewise.
	(scan_for_reference): Likewise.
	(da_create_bias): Likewise.
	(da_dimension_peel): Likewise.
	(lower_omp_1): Add case to look for start of dynamic array reference,
	and handle bias adjustments for the code sequence.

	libgomp/
	PR other/76739
	* target.c (struct da_dim): New struct declaration.
	(struct da_descr_type): Likewise.
	(struct da_info): Likewise.
	(gomp_dynamic_array_count_rows): New function.
	(gomp_dynamic_array_compute_info): Likewise.
	(gomp_dynamic_array_fill_rows_1): Likewise.
	(gomp_dynamic_array_fill_rows): Likewise.
	(gomp_dynamic_array_create_ptrblock): Likewise.
	(gomp_map_vars): Add code to handle dynamic array map kinds.
	* testsuite/libgomp.oacc-c-c++-common/da-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/da-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/da-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/da-4.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/da-utils.h: New test.