2015-08-25 Nathan Sidwell inlude/ * gomp-constants.h (GOMP_VERSION_NVIDIA_PTX): Increment. (GOMP_DIM_GANG, GOMP_DIM_WORKER, GOMP_DIM_VECTOR, GOMP_DIM_MAX, GOMP_DIM_MASK): New. (GOMP_LAUNCH_END, GOMP_LAUNCH_DIM, GOMP_LAUNCH_ASYNC, GOMP_LAUNCH_WAIT): New. (GOMP_LAUNCH_CODE_SHIFT, GOMP_LAUNCH_DEVICE_SHIFT, GOMP_LAUNCH_OP_SHIFT): New. (GOMP_LAUNCH_PACK, GOMP_LAUNCH_CODE, GOMP_LAUNCH_DEVICE, GOMP_LAUNCH_OP): New. (GOMP_LAUNCH_OP_MAX): New. libgomp/ * libgomp.h (acc_dispatch_t): Replace separate geometry args with array. * libgomp.map (GOACC_parallel_keyed): New. * oacc-parallel.c (goacc_wait): Take pointer to va_list. Adjust all callers. (GOACC_parallel_keyed): New interface. Lose geometry arguments and take keyed varargs list. Adjust call to exec_func. (GOACC_parallel): Forward to GACC_parallel_keyed. * libgomp_g.h (GOACC_parallel): Remove. (GOACC_parallel_keyed): Declare. * plugin/plugin-nvptx.c (struct targ_fn_launch): New struct. (stuct targ_gn_descriptor): Replace name field with launch field. (nvptx_exec): Lose separate geometry args, take array. Process dynamic dimensions and adjust. (struct nvptx_tdata): Replace fn_names field with fn_descs. (GOMP_OFFLOAD_load_image): Adjust for change in function table data. (GOMP_OFFLOAD_openacc_parallel): Adjust for change in dimension passing. * oacc-host.c (host_openacc_exec): Adjust for change in dimension passing. gcc/ * config/nvptx/nvptx.c: Include omp-low.h and gomp-constants.h. (nvptx_record_offload_symbol): Record function execution geometry. * config/nvptx/mkoffload.c (process): Include launch geometry in function data. * omp-low.c (oacc_launch_pack): New. (replace_oacc_fn_attrib): New. (set_oacc_fn_attrib): New. (get_oacc_fn_attrib): New. (expand_omp_target): Create keyed varargs for GOACC_parallel call generation. * omp-low.h (get_oacc_fn_attrib): Declare. * builtin-types.def (DEF_FUNCTION_TyPE_VAR_6): New. (DEF_FUNCTION_TYPE_VAR_11): Delete. * tree.h (OMP_CLAUSE_EXPR): New. * omp-builtins.def (BUILT_IN_GOACC_PARALLEL): Change target fn name. gcc/lto/ * lto-lang.c (DEF_FUNCTION_TYPE_VAR_6): New. (DEF_FUNCTION_TYPE_VAR_11): Delete. gcc/c-family/ * c-common.c (DEF_FUNCTION_TYPE_VAR_6): New. (DEF_FUNCTION_TYPE_VAR_11): Delete. gcc/fortran/ * f95-lang.c (DEF_FUNCTION_TYPE_VAR_6): New. (DEF_FUNCTION_TYPE_VAR_11): Delete. * types.def (DEF_FUNCTION_TYPE_VAR_6): New. (DEF_FUNCTION_TYPE_VAR_11): Delete. Index: include/gomp-constants.h =================================================================== --- include/gomp-constants.h (revision 227137) +++ include/gomp-constants.h (working copy) @@ -115,11 +115,34 @@ enum gomp_map_kind /* Versions of libgomp and device-specific plugins. */ #define GOMP_VERSION 0 -#define GOMP_VERSION_NVIDIA_PTX 0 +#define GOMP_VERSION_NVIDIA_PTX 1 #define GOMP_VERSION_INTEL_MIC 0 #define GOMP_VERSION_PACK(LIB, DEV) (((LIB) << 16) | (DEV)) #define GOMP_VERSION_LIB(PACK) (((PACK) >> 16) & 0xffff) #define GOMP_VERSION_DEV(PACK) ((PACK) & 0xffff) +#define GOMP_DIM_GANG 0 +#define GOMP_DIM_WORKER 1 +#define GOMP_DIM_VECTOR 2 +#define GOMP_DIM_MAX 3 +#define GOMP_DIM_MASK(X) (1u << (X)) + +/* Varadic launch arguments. */ +#define GOMP_LAUNCH_END 0 /* End of args, no dev or op */ +#define GOMP_LAUNCH_DIM 1 /* Launch dimensions, op = mask */ +#define GOMP_LAUNCH_ASYNC 2 /* Async, op = cst val if not MAX */ +#define GOMP_LAUNCH_WAIT 3 /* Waits, op = num waits. */ +#define GOMP_LAUNCH_CODE_SHIFT 28 +#define GOMP_LAUNCH_DEVICE_SHIFT 16 +#define GOMP_LAUNCH_OP_SHIFT 0 +#define GOMP_LAUNCH_PACK(CODE,DEVICE,OP) \ + (((CODE) << GOMP_LAUNCH_CODE_SHIFT) \ + | ((DEVICE) << GOMP_LAUNCH_DEVICE_SHIFT) \ + | ((OP) << GOMP_LAUNCH_OP_SHIFT)) +#define GOMP_LAUNCH_CODE(X) (((X) >> GOMP_LAUNCH_CODE_SHIFT) & 0xf) +#define GOMP_LAUNCH_DEVICE(X) (((X) >> GOMP_LAUNCH_DEVICE_SHIFT) & 0xfff) +#define GOMP_LAUNCH_OP(X) (((X) >> GOMP_LAUNCH_OP_SHIFT) & 0xffff) +#define GOMP_LAUNCH_OP_MAX 0xffff + #endif Index: libgomp/libgomp.h =================================================================== --- libgomp/libgomp.h (revision 227137) +++ libgomp/libgomp.h (working copy) @@ -693,7 +693,7 @@ typedef struct acc_dispatch_t /* Execute. */ void (*exec_func) (void (*) (void *), size_t, void **, void **, size_t *, - unsigned short *, int, int, int, int, void *); + unsigned short *, int, unsigned *, void *); /* Async cleanup callback registration. */ void (*register_async_cleanup_func) (void *); Index: libgomp/oacc-parallel.c =================================================================== --- libgomp/oacc-parallel.c (revision 227137) +++ libgomp/oacc-parallel.c (working copy) @@ -49,14 +49,12 @@ find_pset (int pos, size_t mapnum, unsig return kind == GOMP_MAP_TO_PSET; } -static void goacc_wait (int async, int num_waits, va_list ap); +static void goacc_wait (int async, int num_waits, va_list *ap); void -GOACC_parallel (int device, void (*fn) (void *), - size_t mapnum, void **hostaddrs, size_t *sizes, - unsigned short *kinds, - int num_gangs, int num_workers, int vector_length, - int async, int num_waits, ...) +GOACC_parallel_keyed (int device, void (*fn) (void *), + size_t mapnum, void **hostaddrs, size_t *sizes, + unsigned short *kinds, ...) { bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK; va_list ap; @@ -68,22 +66,16 @@ GOACC_parallel (int device, void (*fn) ( struct splay_tree_key_s k; splay_tree_key tgt_fn_key; void (*tgt_fn); - - if (num_gangs != 1) - gomp_fatal ("num_gangs (%d) different from one is not yet supported", - num_gangs); - if (num_workers != 1) - gomp_fatal ("num_workers (%d) different from one is not yet supported", - num_workers); + int async = GOMP_ASYNC_SYNC; + unsigned dims[GOMP_DIM_MAX]; + unsigned tag; #ifdef HAVE_INTTYPES_H - gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p, " - "async = %d\n", - __FUNCTION__, (uint64_t) mapnum, hostaddrs, sizes, kinds, async); + gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n", + __FUNCTION__, (uint64_t) mapnum, hostaddrs, sizes, kinds); #else - gomp_debug (0, "%s: mapnum=%lu, hostaddrs=%p, sizes=%p, kinds=%p, async=%d\n", - __FUNCTION__, (unsigned long) mapnum, hostaddrs, sizes, kinds, - async); + gomp_debug (0, "%s: mapnum=%lu, hostaddrs=%p, sizes=%p, kinds=%p\n", + __FUNCTION__, (unsigned long) mapnum, hostaddrs, sizes, kinds); #endif goacc_lazy_initialize (); @@ -105,12 +97,45 @@ GOACC_parallel (int device, void (*fn) ( return; } - if (num_waits) + va_start (ap, kinds); + /* TODO: This will need amending when device_type is implemented. */ + while (GOMP_LAUNCH_PACK (GOMP_LAUNCH_END, 0, 0) + != (tag = va_arg (ap, unsigned))) { - va_start (ap, num_waits); - goacc_wait (async, num_waits, ap); - va_end (ap); + assert (!GOMP_LAUNCH_DEVICE (tag)); + switch (GOMP_LAUNCH_CODE (tag)) + { + case GOMP_LAUNCH_DIM: + { + unsigned mask = GOMP_LAUNCH_OP (tag); + + for (i = 0; i != GOMP_DIM_MAX; i++) + if (mask & GOMP_DIM_MASK (i)) + dims[i] = va_arg (ap, unsigned); + } + break; + + case GOMP_LAUNCH_ASYNC: + { + /* Small constant values are encoded in the operand. */ + async = GOMP_LAUNCH_OP (tag); + + if (async == GOMP_LAUNCH_OP_MAX) + async = va_arg (ap, unsigned); + break; + } + + case GOMP_LAUNCH_WAIT: + { + unsigned num_waits = GOMP_LAUNCH_OP (tag); + + if (num_waits) + goacc_wait (async, num_waits, &ap); + break; + } + } } + va_end (ap); acc_dev->openacc.async_set_async_func (async); @@ -138,9 +163,8 @@ GOACC_parallel (int device, void (*fn) ( devaddrs[i] = (void *) (tgt->list[i]->tgt->tgt_start + tgt->list[i]->tgt_offset); - acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs, sizes, kinds, - num_gangs, num_workers, vector_length, async, - tgt); + acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs, sizes, + kinds, async, dims, tgt); /* If running synchronously, unmap immediately. */ if (async < acc_async_noval) @@ -154,6 +178,38 @@ GOACC_parallel (int device, void (*fn) ( acc_dev->openacc.async_set_async_func (acc_async_sync); } +/* Legacy entry point. */ +void +GOACC_parallel (int device, void (*fn) (void *), + size_t mapnum, void **hostaddrs, size_t *sizes, + unsigned short *kinds, + int num_gangs, int num_workers, int vector_length, + int async, int num_waits, ...) +{ + int waits[9]; + unsigned ix; + va_list ap; + + if (num_waits > 8) + gomp_fatal ("Too many waits for legacy interface"); + + va_start (ap, num_waits); + for (ix = 0; ix != num_waits; ix++) + waits[ix] = va_arg (ap, int); + waits[ix] = GOMP_LAUNCH_PACK (GOMP_LAUNCH_END, 0, 0); + va_end (ap); + + GOACC_parallel_keyed (device, fn, mapnum, hostaddrs, sizes, kinds, + GOMP_LAUNCH_PACK (GOMP_LAUNCH_DIM, 0, + GOMP_DIM_MASK (GOMP_DIM_MAX) - 1), + num_gangs, num_workers, vector_length, + GOMP_LAUNCH_PACK (GOMP_LAUNCH_ASYNC, 0, + GOMP_LAUNCH_OP_MAX), async, + GOMP_LAUNCH_PACK (GOMP_LAUNCH_WAIT, 0, num_waits), + async, waits[0], waits[1], waits[2], waits[3], + waits[4], waits[5], waits[6], waits[7], waits[8]); +} + void GOACC_data_start (int device, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned short *kinds) @@ -230,7 +286,7 @@ GOACC_enter_exit_data (int device, size_ va_list ap; va_start (ap, num_waits); - goacc_wait (async, num_waits, ap); + goacc_wait (async, num_waits, &ap); va_end (ap); } @@ -344,15 +400,15 @@ GOACC_enter_exit_data (int device, size_ } static void -goacc_wait (int async, int num_waits, va_list ap) +goacc_wait (int async, int num_waits, va_list *ap) { struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; while (num_waits--) { - int qid = va_arg (ap, int); - + int qid = va_arg (*ap, int); + if (acc_async_test (qid)) continue; @@ -389,7 +445,7 @@ GOACC_update (int device, size_t mapnum, va_list ap; va_start (ap, num_waits); - goacc_wait (async, num_waits, ap); + goacc_wait (async, num_waits, &ap); va_end (ap); } @@ -430,7 +486,7 @@ GOACC_wait (int async, int num_waits, .. va_list ap; va_start (ap, num_waits); - goacc_wait (async, num_waits, ap); + goacc_wait (async, num_waits, &ap); va_end (ap); } else if (async == acc_async_sync) Index: libgomp/libgomp_g.h =================================================================== --- libgomp/libgomp_g.h (revision 227137) +++ libgomp/libgomp_g.h (working copy) @@ -222,9 +222,8 @@ extern void GOACC_data_start (int, size_ extern void GOACC_data_end (void); extern void GOACC_enter_exit_data (int, size_t, void **, size_t *, unsigned short *, int, int, ...); -extern void GOACC_parallel (int, void (*) (void *), size_t, - void **, size_t *, unsigned short *, int, int, int, - int, int, ...); +extern void GOACC_parallel_keyed (int, void (*) (void *), size_t, + void **, size_t *, unsigned short *, ...); extern void GOACC_update (int, size_t, void **, size_t *, unsigned short *, int, int, ...); extern void GOACC_wait (int, int, ...); Index: libgomp/libgomp.map =================================================================== --- libgomp/libgomp.map (revision 227137) +++ libgomp/libgomp.map (working copy) @@ -332,6 +332,11 @@ GOACC_2.0 { GOACC_get_num_threads; }; +GOACC_2.0,1 { + global: + GOACC_parallel_keyed; +} GOACC_2.0; + GOMP_PLUGIN_1.0 { global: GOMP_PLUGIN_malloc; Index: libgomp/plugin/plugin-nvptx.c =================================================================== --- libgomp/plugin/plugin-nvptx.c (revision 227137) +++ libgomp/plugin/plugin-nvptx.c (working copy) @@ -282,12 +282,20 @@ map_push (struct ptx_stream *s, int asyn return; } +/* Target data function launch information. */ + +struct targ_fn_launch +{ + const char *fn; + unsigned short dim[GOMP_DIM_MAX]; +}; + /* Descriptor of a loaded function. */ struct targ_fn_descriptor { CUfunction fn; - const char *name; + const struct targ_fn_launch *launch; }; /* A loaded PTX image. */ @@ -929,8 +937,8 @@ event_add (enum ptx_event_type type, CUe void nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs, - size_t *sizes, unsigned short *kinds, int num_gangs, int num_workers, - int vector_length, int async, void *targ_mem_desc) + size_t *sizes, unsigned short *kinds, int async, unsigned *dims, + void *targ_mem_desc) { struct targ_fn_descriptor *targ_fn = (struct targ_fn_descriptor *) fn; CUfunction function; @@ -939,7 +947,6 @@ nvptx_exec (void (*fn), size_t mapnum, v struct ptx_stream *dev_str; void *kargs[1]; void *hp, *dp; - unsigned int nthreads_in_block; struct nvptx_thread *nvthd = nvptx_thread (); const char *maybe_abort_msg = "(perhaps abort was called)"; @@ -948,6 +955,20 @@ nvptx_exec (void (*fn), size_t mapnum, v dev_str = select_stream_for_async (async, pthread_self (), false, NULL); assert (dev_str == nvthd->current_stream); + /* Initialize the launch dimensions. Typically this is constant, + provided by the device compiler, but we must permit runtime + values. */ + for (i = 0; i != GOMP_DIM_MAX; i++) + if (targ_fn->launch->dim[i]) + dims[i] = targ_fn->launch->dim[i]; + + if (dims[GOMP_DIM_GANG] != 1) + GOMP_PLUGIN_fatal ("non-unity num_gangs (%d) not supported", + dims[GOMP_DIM_GANG]); + if (dims[GOMP_DIM_WORKER] != 1) + GOMP_PLUGIN_fatal ("non-unity num_workers (%d) not supported", + dims[GOMP_DIM_WORKER]); + /* This reserves a chunk of a pre-allocated page of memory mapped on both the host and the device. HP is a host pointer to the new chunk, and DP is the corresponding device pointer. */ @@ -965,35 +986,21 @@ nvptx_exec (void (*fn), size_t mapnum, v if (r != CUDA_SUCCESS) GOMP_PLUGIN_fatal ("cuMemcpy failed: %s", cuda_error (r)); - GOMP_PLUGIN_debug (0, " %s: kernel %s: launch\n", __FUNCTION__, targ_fn->name); + GOMP_PLUGIN_debug (0, " %s: kernel %s: launch" + " gangs=%u, workers=%u, vectors=%u\n", + __FUNCTION__, targ_fn->launch->fn, + dims[GOMP_DIM_GANG], dims[GOMP_DIM_WORKER], dims[GOMP_DIM_VECTOR]); // OpenACC CUDA // - // num_gangs blocks - // num_workers warps (where a warp is equivalent to 32 threads) - // vector length threads - // - - /* The openacc vector_length clause 'determines the vector length to use for - vector or SIMD operations'. The question is how to map this to CUDA. - - In CUDA, the warp size is the vector length of a CUDA device. However, the - CUDA interface abstracts away from that, and only shows us warp size - indirectly in maximum number of threads per block, which is a product of - warp size and the number of hyperthreads of a multiprocessor. - - We choose to map openacc vector_length directly onto the number of threads - in a block, in the x dimension. This is reflected in gcc code generation - that uses ThreadIdx.x to access vector elements. - - Attempting to use an openacc vector_length of more than the maximum number - of threads per block will result in a cuda error. */ - nthreads_in_block = vector_length; + // num_gangs nctaid.x + // num_workers ntid.y + // vector length ntid.x kargs[0] = &dp; r = cuLaunchKernel (function, - num_gangs, 1, 1, - nthreads_in_block, 1, 1, + dims[GOMP_DIM_GANG], 1, 1, + dims[GOMP_DIM_VECTOR], dims[GOMP_DIM_WORKER], 1, 0, dev_str->stream, kargs, 0); if (r != CUDA_SUCCESS) GOMP_PLUGIN_fatal ("cuLaunchKernel error: %s", cuda_error (r)); @@ -1039,7 +1046,7 @@ nvptx_exec (void (*fn), size_t mapnum, v #endif GOMP_PLUGIN_debug (0, " %s: kernel %s: finished\n", __FUNCTION__, - targ_fn->name); + targ_fn->launch->fn); #ifndef DISABLE_ASYNC if (async < acc_async_noval) @@ -1567,7 +1574,7 @@ typedef struct nvptx_tdata const char *const *var_names; size_t var_num; - const char *const *fn_names; + const struct targ_fn_launch *fn_descs; size_t fn_num; } nvptx_tdata_t; @@ -1588,7 +1595,8 @@ GOMP_OFFLOAD_load_image (int ord, unsign struct addr_pair **target_table) { CUmodule module; - const char *const *fn_names, *const *var_names; + const char *const *var_names; + const struct targ_fn_launch *fn_descs; unsigned int fn_entries, var_entries, i, j; CUresult r; struct targ_fn_descriptor *targ_fns; @@ -1617,7 +1625,7 @@ GOMP_OFFLOAD_load_image (int ord, unsign var_entries = img_header->var_num; var_names = img_header->var_names; fn_entries = img_header->fn_num; - fn_names = img_header->fn_names; + fn_descs = img_header->fn_descs; targ_tbl = GOMP_PLUGIN_malloc (sizeof (struct addr_pair) * (fn_entries + var_entries)); @@ -1640,12 +1648,12 @@ GOMP_OFFLOAD_load_image (int ord, unsign { CUfunction function; - r = cuModuleGetFunction (&function, module, fn_names[i]); + r = cuModuleGetFunction (&function, module, fn_descs[i].fn); if (r != CUDA_SUCCESS) GOMP_PLUGIN_fatal ("cuModuleGetFunction error: %s", cuda_error (r)); targ_fns->fn = function; - targ_fns->name = (const char *) fn_names[i]; + targ_fns->launch = &fn_descs[i]; targ_tbl->start = (uintptr_t) targ_fns; targ_tbl->end = targ_tbl->start + 1; @@ -1724,13 +1732,12 @@ void (*device_run) (int n, void *fn_ptr, void GOMP_OFFLOAD_openacc_parallel (void (*fn) (void *), size_t mapnum, - void **hostaddrs, void **devaddrs, size_t *sizes, - unsigned short *kinds, int num_gangs, - int num_workers, int vector_length, int async, - void *targ_mem_desc) + void **hostaddrs, void **devaddrs, + size_t *sizes, unsigned short *kinds, + int async, unsigned *dims, void *targ_mem_desc) { - nvptx_exec (fn, mapnum, hostaddrs, devaddrs, sizes, kinds, num_gangs, - num_workers, vector_length, async, targ_mem_desc); + nvptx_exec (fn, mapnum, hostaddrs, devaddrs, sizes, kinds, + async, dims, targ_mem_desc); } void Index: libgomp/oacc-host.c =================================================================== --- libgomp/oacc-host.c (revision 227137) +++ libgomp/oacc-host.c (working copy) @@ -137,10 +137,8 @@ host_openacc_exec (void (*fn) (void *), void **devaddrs __attribute__ ((unused)), size_t *sizes __attribute__ ((unused)), unsigned short *kinds __attribute__ ((unused)), - int num_gangs __attribute__ ((unused)), - int num_workers __attribute__ ((unused)), - int vector_length __attribute__ ((unused)), int async __attribute__ ((unused)), + unsigned *dims __attribute ((unused)), void *targ_mem_desc __attribute__ ((unused))) { fn (hostaddrs); Index: gcc/config/nvptx/nvptx.c =================================================================== --- gcc/config/nvptx/nvptx.c (revision 227137) +++ gcc/config/nvptx/nvptx.c (working copy) @@ -56,6 +56,8 @@ #include "cfgrtl.h" #include "stor-layout.h" #include "builtins.h" +#include "omp-low.h" +#include "gomp-constants.h" /* This file should be included last. */ #include "target-def.h" @@ -2064,9 +2066,51 @@ nvptx_vector_alignment (const_tree type) static void nvptx_record_offload_symbol (tree decl) { - fprintf (asm_out_file, "//:%s_MAP %s\n", - TREE_CODE (decl) == VAR_DECL ? "VAR" : "FUNC", - IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (decl))); + switch (TREE_CODE (decl)) + { + case VAR_DECL: + fprintf (asm_out_file, "//:VAR_MAP \"%s\"\n", + IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (decl))); + break; + + case FUNCTION_DECL: + { + tree attr = get_oacc_fn_attrib (decl); + tree dims = NULL_TREE; + unsigned ix; + + if (attr) + dims = TREE_VALUE (attr); + fprintf (asm_out_file, "//:FUNC_MAP \"%s\"", + IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (decl))); + + for (ix = 0; ix != GOMP_DIM_MAX; ix++) + { + int size = 1; + + /* TODO: This check can go away once the dimension default + machinery is merged to trunk. */ + if (dims) + { + tree dim = TREE_VALUE (dims); + + if (dim) + size = TREE_INT_CST_LOW (dim); + + gcc_assert (!TREE_PURPOSE (dims)); + dims = TREE_CHAIN (dims); + } + + fprintf (asm_out_file, ", %#x", size); + } + + fprintf (asm_out_file, "\n"); + } + break; + + default: + gcc_unreachable (); + } } /* Implement TARGET_ASM_FILE_START. Write the kinds of things ptxas expects Index: gcc/config/nvptx/mkoffload.c =================================================================== --- gcc/config/nvptx/mkoffload.c (revision 227137) +++ gcc/config/nvptx/mkoffload.c (working copy) @@ -842,6 +842,8 @@ process (FILE *in, FILE *out) { const char *input = read_file (in); Token *tok = tokenize (input); + const char *comma; + id_map const *id; do tok = parse_file (tok); @@ -853,21 +855,25 @@ process (FILE *in, FILE *out) write_stmts (out, rev_stmts (fns)); fprintf (out, ";\n\n"); - fprintf (out, "static const char *const var_mappings[] = {\n"); - for (id_map *id = var_ids; id; id = id->next) - fprintf (out, "\t\"%s\"%s\n", id->ptx_name, id->next ? "," : ""); - fprintf (out, "};\n\n"); - fprintf (out, "static const char *const func_mappings[] = {\n"); - for (id_map *id = func_ids; id; id = id->next) - fprintf (out, "\t\"%s\"%s\n", id->ptx_name, id->next ? "," : ""); - fprintf (out, "};\n\n"); + fprintf (out, "static const char *const var_mappings[] = {"); + for (comma = "", id = var_ids; id; comma = ",", id = id->next) + fprintf (out, "%s\n\t%s", comma, id->ptx_name); + fprintf (out, "\n};\n\n"); + + fprintf (out, "static const struct nvptx_fn {\n" + " const char *name;\n" + " unsigned short dim[3];\n" + "} func_mappings[] = {\n"); + for (comma = "", id = func_ids; id; comma = ",", id = id->next) + fprintf (out, "%s\n\t{%s}", comma, id->ptx_name); + fprintf (out, "\n};\n\n"); fprintf (out, "static const struct nvptx_tdata {\n" " const char *ptx_src;\n" " const char *const *var_names;\n" " __SIZE_TYPE__ var_num;\n" - " const char *const *fn_names;\n" + " const struct nvptx_fn *fn_names;\n" " __SIZE_TYPE__ fn_num;\n" "} target_data = {\n" " ptx_code,\n" Index: gcc/omp-low.c =================================================================== --- gcc/omp-low.c (revision 227137) +++ gcc/omp-low.c (working copy) @@ -8795,6 +8794,102 @@ expand_omp_atomic (struct omp_region *re } +/* Encode an oacc launc argument. This matches the GOMP_LAUNCH_PACK + macro on gomp-constants.h. We do not check for overflow. */ + +static tree +oacc_launch_pack (unsigned code, tree device, unsigned op) +{ + tree res; + + res = build_int_cst (unsigned_type_node, GOMP_LAUNCH_PACK (code, 0, op)); + if (device) + { + device = fold_build2 (LSHIFT_EXPR, unsigned_type_node, + device, build_int_cst (unsigned_type_node, + GOMP_LAUNCH_DEVICE_SHIFT)); + res = fold_build2 (BIT_IOR_EXPR, unsigned_type_node, res, device); + } + return res; +} + +/* Look for compute grid dimension clauses and convert to an attribute + attached to FN. This permits the target-side code to (a) massage + the dimensions, (b) emit that data and (c) optimize. Non-constant + dimensions are pushed onto ARGS. + + The attribute value is a TREE_LIST. A set of dimensions is + represented as a list of INTEGER_CST. Those that are runtime + expres are represented as an INTEGER_CST of zero. + + TOOO. Normally the attribute will just contain a single such list. If + however it contains a list of lists, this will represent the use of + device_type. Each member of the outer list is an assoc list of + dimensions, keyed by the device type. The first entry will be the + default. Well, that's the plan. */ + +#define OACC_FN_ATTRIB "oacc function" + +/* Replace any existing oacc fn attribute with updated dimensions. */ + +void +replace_oacc_fn_attrib (tree fn, tree dims) +{ + tree ident = get_identifier (OACC_FN_ATTRIB); + tree attribs = DECL_ATTRIBUTES (fn); + + /* If we happen to be present as the first attrib, drop it. */ + if (attribs && TREE_PURPOSE (attribs) == ident) + attribs = TREE_CHAIN (attribs); + DECL_ATTRIBUTES (fn) = tree_cons (ident, dims, attribs); +} + +static void +set_oacc_fn_attrib (tree fn, tree clauses, vec *args) +{ + /* Must match GOMP_DIM ordering. */ + static const omp_clause_code ids[] = + {OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH}; + unsigned ix; + tree dims[GOMP_DIM_MAX]; + tree attr = NULL_TREE; + unsigned non_const = 0; + + for (ix = GOMP_DIM_MAX; ix--;) + { + tree clause = find_omp_clause (clauses, ids[ix]); + tree dim = NULL_TREE; + + if (clause) + dim = OMP_CLAUSE_EXPR (clause, ids[ix]); + dims[ix] = dim; + if (dim && TREE_CODE (dim) != INTEGER_CST) + { + dim = integer_zero_node; + non_const |= GOMP_DIM_MASK (ix); + } + attr = tree_cons (NULL_TREE, dim, attr); + } + + replace_oacc_fn_attrib (fn, attr); + + if (non_const) + { + /* Push a dynamic argument set. */ + args->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM, + NULL_TREE, non_const)); + for (unsigned ix = 0; ix != GOMP_DIM_MAX; ix++) + if (non_const & GOMP_DIM_MASK (ix)) + args->safe_push (dims[ix]); + } +} + +tree +get_oacc_fn_attrib (tree fn) +{ + return lookup_attribute (OACC_FN_ATTRIB, DECL_ATTRIBUTES (fn)); +} + /* Expand the GIMPLE_OMP_TARGET starting at REGION. */ static void @@ -9150,6 +9245,7 @@ expand_omp_target (struct omp_region *re } gimple g; + bool tagging = false; /* The maximum number used by any start_ix, without varargs. */ auto_vec args; args.quick_push (device); @@ -9185,88 +9281,86 @@ expand_omp_target (struct omp_region *re break; case BUILT_IN_GOACC_PARALLEL: { - tree t_num_gangs, t_num_workers, t_vector_length; - - /* Default values for num_gangs, num_workers, and vector_length. */ - t_num_gangs = t_num_workers = t_vector_length - = fold_convert_loc (gimple_location (entry_stmt), - integer_type_node, integer_one_node); - /* ..., but if present, use the value specified by the respective - clause, making sure that are of the correct type. */ - c = find_omp_clause (clauses, OMP_CLAUSE_NUM_GANGS); - if (c) - t_num_gangs = fold_convert_loc (OMP_CLAUSE_LOCATION (c), - integer_type_node, - OMP_CLAUSE_NUM_GANGS_EXPR (c)); - c = find_omp_clause (clauses, OMP_CLAUSE_NUM_WORKERS); - if (c) - t_num_workers = fold_convert_loc (OMP_CLAUSE_LOCATION (c), - integer_type_node, - OMP_CLAUSE_NUM_WORKERS_EXPR (c)); - c = find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH); - if (c) - t_vector_length = fold_convert_loc (OMP_CLAUSE_LOCATION (c), - integer_type_node, - OMP_CLAUSE_VECTOR_LENGTH_EXPR (c)); - args.quick_push (t_num_gangs); - args.quick_push (t_num_workers); - args.quick_push (t_vector_length); + set_oacc_fn_attrib (child_fn, clauses, &args); + tagging = true; } /* FALLTHRU */ case BUILT_IN_GOACC_ENTER_EXIT_DATA: case BUILT_IN_GOACC_UPDATE: { - tree t_async; - int t_wait_idx; + tree t_async = NULL_TREE; - /* Default values for t_async. */ - t_async = fold_convert_loc (gimple_location (entry_stmt), - integer_type_node, - build_int_cst (integer_type_node, - GOMP_ASYNC_SYNC)); - /* ..., but if present, use the value specified by the respective + /* If present, use the value specified by the respective clause, making sure that is of the correct type. */ c = find_omp_clause (clauses, OMP_CLAUSE_ASYNC); if (c) t_async = fold_convert_loc (OMP_CLAUSE_LOCATION (c), integer_type_node, OMP_CLAUSE_ASYNC_EXPR (c)); - - args.quick_push (t_async); - /* Save the index, and... */ - t_wait_idx = args.length (); - /* ... push a default value. */ - args.quick_push (fold_convert_loc (gimple_location (entry_stmt), - integer_type_node, - integer_zero_node)); - c = find_omp_clause (clauses, OMP_CLAUSE_WAIT); - if (c) + else if (!tagging) + /* Default values for t_async. */ + t_async = fold_convert_loc (gimple_location (entry_stmt), + integer_type_node, + build_int_cst (integer_type_node, + GOMP_ASYNC_SYNC)); + if (tagging && t_async) { - int n = 0; + unsigned HOST_WIDE_INT i_async; - for (; c; c = OMP_CLAUSE_CHAIN (c)) + if (TREE_CODE (t_async) == INTEGER_CST) { - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WAIT) - { - args.safe_push (fold_convert_loc (OMP_CLAUSE_LOCATION (c), - integer_type_node, - OMP_CLAUSE_WAIT_EXPR (c))); - n++; - } + /* See if we can pack the async arg in to the tag's + operand. */ + i_async = TREE_INT_CST_LOW (t_async); + + if (i_async < GOMP_LAUNCH_OP_MAX) + t_async = NULL_TREE; } + if (t_async) + i_async = GOMP_LAUNCH_OP_MAX; + args.safe_push (oacc_launch_pack + (GOMP_LAUNCH_ASYNC, NULL_TREE, i_async)); + } + if (t_async) + args.safe_push (t_async); - /* Now that we know the number, replace the default value. */ - args.ordered_remove (t_wait_idx); - args.quick_insert (t_wait_idx, - fold_convert_loc (gimple_location (entry_stmt), - integer_type_node, - build_int_cst (integer_type_node, n))); + /* Save the argument index, and ... */ + unsigned t_wait_idx = args.length (); + unsigned num_waits = 0; + c = find_omp_clause (clauses, OMP_CLAUSE_WAIT); + if (!tagging || c) + /* ... push a placeholder. */ + args.safe_push (integer_zero_node); + + for (; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WAIT) + { + args.safe_push (fold_convert_loc (OMP_CLAUSE_LOCATION (c), + integer_type_node, + OMP_CLAUSE_WAIT_EXPR (c))); + num_waits++; + } + + if (!tagging || num_waits) + { + tree len; + + /* Now that we know the number, update the placeholder. */ + if (tagging) + len = oacc_launch_pack (GOMP_LAUNCH_WAIT, NULL_TREE, num_waits); + else + len = build_int_cst (integer_type_node, num_waits); + len = fold_convert_loc (gimple_location (entry_stmt), + unsigned_type_node, len); + args[t_wait_idx] = len; } } break; default: gcc_unreachable (); } + if (tagging) + args.safe_push (oacc_launch_pack (GOMP_LAUNCH_END, NULL_TREE, 0)); g = gimple_build_call_vec (builtin_decl_explicit (start_ix), args); gimple_set_location (g, gimple_location (entry_stmt)); Index: gcc/omp-low.h =================================================================== --- gcc/omp-low.h (revision 227137) +++ gcc/omp-low.h (working copy) @@ -28,6 +28,7 @@ extern void free_omp_regions (void); extern tree omp_reduction_init (tree, tree); extern bool make_gimple_omp_edges (basic_block, struct omp_region **, int *); extern void omp_finish_file (void); +extern tree get_oacc_fn_attrib (tree); extern GTY(()) vec *offload_funcs; extern GTY(()) vec *offload_vars; Index: gcc/builtin-types.def =================================================================== --- gcc/builtin-types.def (revision 227137) +++ gcc/builtin-types.def (working copy) @@ -590,15 +590,14 @@ DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_STRIN DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_INT_INT_INT_INT_INT_VAR, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT) +DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR, + BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, + BT_PTR, BT_PTR, BT_PTR) + DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR, BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT) -DEF_FUNCTION_TYPE_VAR_11 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR, - BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, - BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT, BT_INT, - BT_INT, BT_INT) - DEF_POINTER_TYPE (BT_PTR_FN_VOID_VAR, BT_FN_VOID_VAR) DEF_FUNCTION_TYPE_3 (BT_FN_PTR_PTR_FN_VOID_VAR_PTR_SIZE, BT_PTR, BT_PTR_FN_VOID_VAR, BT_PTR, BT_SIZE) Index: gcc/tree.h =================================================================== --- gcc/tree.h (revision 227137) +++ gcc/tree.h (working copy) @@ -1,3 +1,4 @@ + /* Definitions for the ubiquitous 'tree' type for GNU compilers. Copyright (C) 1989-2015 Free Software Foundation, Inc. @@ -1369,6 +1370,8 @@ extern void protected_set_expr_location OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_SCHEDULE), 0) /* OpenACC clause expressions */ +#define OMP_CLAUSE_EXPR(NODE, CLAUSE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, CLAUSE), 0) #define OMP_CLAUSE_GANG_EXPR(NODE) \ OMP_CLAUSE_OPERAND ( \ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GANG), 0) Index: gcc/omp-builtins.def =================================================================== --- gcc/omp-builtins.def (revision 227137) +++ gcc/omp-builtins.def (working copy) @@ -38,8 +38,8 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_E DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data", BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR, ATTR_NOTHROW_LIST) -DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel", - BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR, +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel_keyed", + BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR, ATTR_NOTHROW_LIST) DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update", BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR, Index: gcc/lto/lto-lang.c =================================================================== --- gcc/lto/lto-lang.c (revision 227137) +++ gcc/lto/lto-lang.c (working copy) @@ -160,10 +160,10 @@ enum lto_builtin_type #define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME, #define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG6) \ NAME, +#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ + ARG6) NAME, #define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ ARG6, ARG7) NAME, -#define DEF_FUNCTION_TYPE_VAR_11(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ - ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) NAME, #define DEF_POINTER_TYPE(NAME, TYPE) NAME, #include "builtin-types.def" #undef DEF_PRIMITIVE_TYPE @@ -182,8 +182,8 @@ enum lto_builtin_type #undef DEF_FUNCTION_TYPE_VAR_3 #undef DEF_FUNCTION_TYPE_VAR_4 #undef DEF_FUNCTION_TYPE_VAR_5 +#undef DEF_FUNCTION_TYPE_VAR_6 #undef DEF_FUNCTION_TYPE_VAR_7 -#undef DEF_FUNCTION_TYPE_VAR_11 #undef DEF_POINTER_TYPE BT_LAST }; @@ -668,13 +668,12 @@ lto_define_builtins (tree va_list_ref_ty def_fn_type (ENUM, RETURN, 1, 4, ARG1, ARG2, ARG3, ARG4); #define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \ def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5); +#define DEF_FUNCTION_TYPE_VAR_6(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ + ARG6) \ + def_fn_type (ENUM, RETURN, 1, 6, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6); #define DEF_FUNCTION_TYPE_VAR_7(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ ARG6, ARG7) \ def_fn_type (ENUM, RETURN, 1, 7, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, ARG7); -#define DEF_FUNCTION_TYPE_VAR_11(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ - ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) \ - def_fn_type (ENUM, RETURN, 1, 11, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \ - ARG7, ARG8, ARG9, ARG10, ARG11); #define DEF_POINTER_TYPE(ENUM, TYPE) \ builtin_types[(int) ENUM] = build_pointer_type (builtin_types[(int) TYPE]); @@ -696,8 +695,8 @@ lto_define_builtins (tree va_list_ref_ty #undef DEF_FUNCTION_TYPE_VAR_3 #undef DEF_FUNCTION_TYPE_VAR_4 #undef DEF_FUNCTION_TYPE_VAR_5 +#undef DEF_FUNCTION_TYPE_VAR_6 #undef DEF_FUNCTION_TYPE_VAR_7 -#undef DEF_FUNCTION_TYPE_VAR_11 #undef DEF_POINTER_TYPE builtin_types[(int) BT_LAST] = NULL_TREE; Index: gcc/c-family/c-common.c =================================================================== --- gcc/c-family/c-common.c (revision 227137) +++ gcc/c-family/c-common.c (working copy) @@ -5545,10 +5545,10 @@ enum c_builtin_type #define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME, #define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \ NAME, +#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ + ARG6) NAME, #define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ ARG6, ARG7) NAME, -#define DEF_FUNCTION_TYPE_VAR_11(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ - ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) NAME, #define DEF_POINTER_TYPE(NAME, TYPE) NAME, #include "builtin-types.def" #undef DEF_PRIMITIVE_TYPE @@ -5567,8 +5567,8 @@ enum c_builtin_type #undef DEF_FUNCTION_TYPE_VAR_3 #undef DEF_FUNCTION_TYPE_VAR_4 #undef DEF_FUNCTION_TYPE_VAR_5 +#undef DEF_FUNCTION_TYPE_VAR_6 #undef DEF_FUNCTION_TYPE_VAR_7 -#undef DEF_FUNCTION_TYPE_VAR_11 #undef DEF_POINTER_TYPE BT_LAST }; @@ -5661,13 +5661,12 @@ c_define_builtins (tree va_list_ref_type def_fn_type (ENUM, RETURN, 1, 4, ARG1, ARG2, ARG3, ARG4); #define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \ def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5); +#define DEF_FUNCTION_TYPE_VAR_6(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ + ARG6) \ + def_fn_type (ENUM, RETURN, 1, 6, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6); #define DEF_FUNCTION_TYPE_VAR_7(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ ARG6, ARG7) \ def_fn_type (ENUM, RETURN, 1, 7, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, ARG7); -#define DEF_FUNCTION_TYPE_VAR_11(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ - ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) \ - def_fn_type (ENUM, RETURN, 1, 11, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \ - ARG7, ARG8, ARG9, ARG10, ARG11); #define DEF_POINTER_TYPE(ENUM, TYPE) \ builtin_types[(int) ENUM] = build_pointer_type (builtin_types[(int) TYPE]); @@ -5689,8 +5688,8 @@ c_define_builtins (tree va_list_ref_type #undef DEF_FUNCTION_TYPE_VAR_3 #undef DEF_FUNCTION_TYPE_VAR_4 #undef DEF_FUNCTION_TYPE_VAR_5 +#undef DEF_FUNCTION_TYPE_VAR_6 #undef DEF_FUNCTION_TYPE_VAR_7 -#undef DEF_FUNCTION_TYPE_VAR_11 #undef DEF_POINTER_TYPE builtin_types[(int) BT_LAST] = NULL_TREE; Index: gcc/fortran/f95-lang.c =================================================================== --- gcc/fortran/f95-lang.c (revision 227137) +++ gcc/fortran/f95-lang.c (working copy) @@ -635,10 +635,10 @@ gfc_init_builtin_functions (void) ARG6, ARG7, ARG8) NAME, #define DEF_FUNCTION_TYPE_VAR_0(NAME, RETURN) NAME, #define DEF_FUNCTION_TYPE_VAR_2(NAME, RETURN, ARG1, ARG2) NAME, +#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ + ARG6) NAME, #define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ ARG6, ARG7) NAME, -#define DEF_FUNCTION_TYPE_VAR_11(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ - ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) NAME, #define DEF_POINTER_TYPE(NAME, TYPE) NAME, #include "types.def" #undef DEF_PRIMITIVE_TYPE @@ -653,8 +653,8 @@ gfc_init_builtin_functions (void) #undef DEF_FUNCTION_TYPE_8 #undef DEF_FUNCTION_TYPE_VAR_0 #undef DEF_FUNCTION_TYPE_VAR_2 +#undef DEF_FUNCTION_TYPE_VAR_6 #undef DEF_FUNCTION_TYPE_VAR_7 -#undef DEF_FUNCTION_TYPE_VAR_11 #undef DEF_POINTER_TYPE BT_LAST }; @@ -1096,8 +1096,8 @@ gfc_init_builtin_functions (void) builtin_types[(int) ARG1], \ builtin_types[(int) ARG2], \ NULL_TREE); -#define DEF_FUNCTION_TYPE_VAR_7(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ - ARG6, ARG7) \ +#define DEF_FUNCTION_TYPE_VAR_6(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ + ARG6) \ builtin_types[(int) ENUM] \ = build_varargs_function_type_list (builtin_types[(int) RETURN], \ builtin_types[(int) ARG1], \ @@ -1106,10 +1106,9 @@ gfc_init_builtin_functions (void) builtin_types[(int) ARG4], \ builtin_types[(int) ARG5], \ builtin_types[(int) ARG6], \ - builtin_types[(int) ARG7], \ NULL_TREE); -#define DEF_FUNCTION_TYPE_VAR_11(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ - ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) \ +#define DEF_FUNCTION_TYPE_VAR_7(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \ + ARG6, ARG7) \ builtin_types[(int) ENUM] \ = build_varargs_function_type_list (builtin_types[(int) RETURN], \ builtin_types[(int) ARG1], \ @@ -1119,10 +1118,6 @@ gfc_init_builtin_functions (void) builtin_types[(int) ARG5], \ builtin_types[(int) ARG6], \ builtin_types[(int) ARG7], \ - builtin_types[(int) ARG8], \ - builtin_types[(int) ARG9], \ - builtin_types[(int) ARG10], \ - builtin_types[(int) ARG11], \ NULL_TREE); #define DEF_POINTER_TYPE(ENUM, TYPE) \ builtin_types[(int) ENUM] \ @@ -1140,8 +1135,8 @@ gfc_init_builtin_functions (void) #undef DEF_FUNCTION_TYPE_8 #undef DEF_FUNCTION_TYPE_VAR_0 #undef DEF_FUNCTION_TYPE_VAR_2 +#undef DEF_FUNCTION_TYPE_VAR_6 #undef DEF_FUNCTION_TYPE_VAR_7 -#undef DEF_FUNCTION_TYPE_VAR_11 #undef DEF_POINTER_TYPE builtin_types[(int) BT_LAST] = NULL_TREE; Index: gcc/fortran/types.def =================================================================== --- gcc/fortran/types.def (revision 227137) +++ gcc/fortran/types.def (working copy) @@ -219,7 +219,6 @@ DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT) -DEF_FUNCTION_TYPE_VAR_11 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR, +DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR, BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, - BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT, BT_INT, - BT_INT, BT_INT) + BT_PTR, BT_PTR, BT_PTR)