From 9ba66bf5b9c69e0e2bcd1b2ab88160bf9b2aa417 Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Thu, 31 Oct 2019 11:09:43 +0100 Subject: [PATCH] configure.ac: Compute and substitute omp_device_properties and omp_device_property_deps. * configure.ac: Compute and substitute omp_device_properties and omp_device_property_deps. * Makefile.in (generated_files): Add omp-device-properties.h. (omp-general.o): Depend on omp-device-properties.h. (omp_device_properties): New make variable. (omp-device-properties.h, s-omp-device-properties-h, install-omp-device-properties): New goals. (install): Depend on install-omp-device-properties for accelerators. * target.def (TARGET_OMP_DEVICE_KIND_ARCH_ISA): New target hook. * target.h (enum omp_device_kind_arch_isa): New enum. * doc/tm.texi.in: Add placeholder for TARGET_OMP_DEVICE_KIND_ARCH_ISA documentation. * omp-general.c: Include omp-device-properties.h. (omp_max_simt_vf): Expect OFFLOAD_TARGET_NAMES to be separated by colon instead of comma. (omp_offload_device_kind_arch_isa, omp_maybe_offloaded): New functions. (omp_context_selector_matches): Implement device set arch/isa selectors, improve device set kind selector handling. * config/i386/i386-options.h (ix86_omp_device_kind_arch_isa): Declare. * config/i386/i386.c (TARGET_SIMD_CLONE_ADJUST, TARGET_SIMD_CLONE_USABLE): Formatting fix. (TARGET_OMP_DEVICE_KIND_ARCH_ISA): Redefine to ix86_omp_device_kind_arch_isa. * config/i386/i386-options.c (struct ix86_target_opts): Move type definition from ix86_target_string to file scope. (isa2_opts, isa_opts): Moved arrays from ix86_target_string function to file scope. (ix86_omp_device_kind_arch_isa): New function. (ix86_target_string): Moved struct ix86_target_opts, isa2_opts and isa_opts definitions to file scope. * config/i386/t-intelmic (omp-device-properties): New goal. * config/nvptx/t-nvptx (omp-device-properties): Likewise. * config/nvptx/nvptx.c (nvptx_omp_device_kind_arch_isa): New function. (TARGET_OMP_DEVICE_KIND_ARCH_ISA): Redefine to nvptx_omp_device_kind_arch_isa. * configure: Regenerate. * doc/tm.texi: Regenerate. testsuite/ * c-c++-common/gomp/declare-variant-9.c: New test. * c-c++-common/gomp/declare-variant-10.c: New test. From-SVN: r277662 --- gcc/ChangeLog | 39 +++ gcc/Makefile.in | 35 ++- gcc/config/i386/i386-options.c | 256 +++++++++++------- gcc/config/i386/i386-options.h | 3 + gcc/config/i386/i386.c | 9 +- gcc/config/i386/t-intelmic | 7 + gcc/config/nvptx/nvptx.c | 29 ++ gcc/config/nvptx/t-nvptx | 5 + gcc/configure | 17 +- gcc/configure.ac | 11 + gcc/doc/tm.texi | 7 + gcc/doc/tm.texi.in | 2 + gcc/omp-general.c | 214 ++++++++++++--- gcc/target.def | 15 + gcc/target.h | 7 + gcc/testsuite/ChangeLog | 5 + .../c-c++-common/gomp/declare-variant-10.c | 77 ++++++ .../c-c++-common/gomp/declare-variant-9.c | 63 +++++ 18 files changed, 665 insertions(+), 136 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-variant-10.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-variant-9.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 8265f4ce365..c068ac865c3 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,5 +1,44 @@ 2019-10-31 Jakub Jelinek + * configure.ac: Compute and substitute omp_device_properties and + omp_device_property_deps. + * Makefile.in (generated_files): Add omp-device-properties.h. + (omp-general.o): Depend on omp-device-properties.h. + (omp_device_properties): New make variable. + (omp-device-properties.h, s-omp-device-properties-h, + install-omp-device-properties): New goals. + (install): Depend on install-omp-device-properties for accelerators. + * target.def (TARGET_OMP_DEVICE_KIND_ARCH_ISA): New target hook. + * target.h (enum omp_device_kind_arch_isa): New enum. + * doc/tm.texi.in: Add placeholder for TARGET_OMP_DEVICE_KIND_ARCH_ISA + documentation. + * omp-general.c: Include omp-device-properties.h. + (omp_max_simt_vf): Expect OFFLOAD_TARGET_NAMES to be separated by + colon instead of comma. + (omp_offload_device_kind_arch_isa, omp_maybe_offloaded): New + functions. + (omp_context_selector_matches): Implement device set arch/isa + selectors, improve device set kind selector handling. + * config/i386/i386-options.h (ix86_omp_device_kind_arch_isa): Declare. + * config/i386/i386.c (TARGET_SIMD_CLONE_ADJUST, + TARGET_SIMD_CLONE_USABLE): Formatting fix. + (TARGET_OMP_DEVICE_KIND_ARCH_ISA): Redefine to + ix86_omp_device_kind_arch_isa. + * config/i386/i386-options.c (struct ix86_target_opts): Move type + definition from ix86_target_string to file scope. + (isa2_opts, isa_opts): Moved arrays from ix86_target_string function + to file scope. + (ix86_omp_device_kind_arch_isa): New function. + (ix86_target_string): Moved struct ix86_target_opts, isa2_opts and + isa_opts definitions to file scope. + * config/i386/t-intelmic (omp-device-properties): New goal. + * config/nvptx/t-nvptx (omp-device-properties): Likewise. + * config/nvptx/nvptx.c (nvptx_omp_device_kind_arch_isa): New function. + (TARGET_OMP_DEVICE_KIND_ARCH_ISA): Redefine to + nvptx_omp_device_kind_arch_isa. + * configure: Regenerate. + * doc/tm.texi: Regenerate. + PR middle-end/92231 * tree.h (fndecl_built_in_p): Use fndecl_built_in_p instead of DECL_BUILT_IN in comment. Remove redundant ()s around return diff --git a/gcc/Makefile.in b/gcc/Makefile.in index c82858fa93e..551674b0d39 100644 --- a/gcc/Makefile.in +++ b/gcc/Makefile.in @@ -2645,7 +2645,7 @@ generated_files = config.h tm.h $(TM_P_H) $(TM_D_H) $(TM_H) multilib.h \ common/common-target-hooks-def.h pass-instances.def \ c-family/c-target-hooks-def.h d/d-target-hooks-def.h \ params.list params.options case-cfn-macros.h \ - cfn-operators.pd + cfn-operators.pd omp-device-properties.h # # How to compile object files to run on the build machine. @@ -2854,6 +2854,30 @@ $(genprog:%=build/gen%$(build_exeext)): build/gen%$(build_exeext): build/gen%.o +$(LINKER_FOR_BUILD) $(BUILD_LINKERFLAGS) $(BUILD_LDFLAGS) -o $@ \ $(filter-out $(BUILD_LIBDEPS), $^) $(BUILD_LIBS) +omp-general.o: omp-device-properties.h + +omp_device_properties = @omp_device_properties@ +omp-device-properties.h: s-omp-device-properties-h ; @true +s-omp-device-properties-h: @omp_device_property_deps@ + -rm -f tmp-omp-device-properties.h; \ + for kind in kind arch isa; do \ + echo 'const char omp_offload_device_'$${kind}'[] = ' \ + >> tmp-omp-device-properties.h; \ + for prop in none $(omp_device_properties); do \ + [ "$$prop" = "none" ] && continue; \ + tgt=`echo "$$prop" | sed 's/=.*$$//'`; \ + props=`echo "$$prop" | sed 's/.*=//'`; \ + echo "\"$$tgt\\0\"" >> tmp-omp-device-properties.h; \ + sed -n 's/^'$${kind}': //p' $${props} \ + | sed 's/[[:blank:]]/ /g;s/ */ /g;s/^ //;s/ $$//;s/ /\\0/g;s/^/"/;s/$$/\\0\\0"/' \ + >> tmp-omp-device-properties.h; \ + done; \ + echo '"";' >> tmp-omp-device-properties.h; \ + done; \ + $(SHELL) $(srcdir)/../move-if-change tmp-omp-device-properties.h \ + omp-device-properties.h + $(STAMP) s-omp-device-properties-h + # Generated source files for gengtype. Prepend inclusion of # config.h/bconfig.h because AIX requires _LARGE_FILES to be defined before # any system header is included. @@ -3452,6 +3476,10 @@ ifeq ($(enable_plugin),yes) install: install-plugin endif +ifeq ($(enable_as_accelerator),yes) +install: install-omp-device-properties +endif + install-strip: override INSTALL_PROGRAM = $(INSTALL_STRIP_PROGRAM) ifneq ($(STRIP),) install-strip: STRIPPROG = $(STRIP) @@ -3637,6 +3665,11 @@ install-driver: installdirs xgcc$(exeext) fi; \ fi +# Install omp-device-properties file for accelerator compilers. +install-omp-device-properties: omp-device-properties installdirs + $(INSTALL_DATA) omp-device-properties \ + $(DESTDIR)$(libsubdir)/omp-device-properties + # Install the info files. # $(INSTALL_DATA) might be a relative pathname, so we can't cd into srcdir # to do the install. diff --git a/gcc/config/i386/i386-options.c b/gcc/config/i386/i386-options.c index ed286bffaaa..dfc8ae23ba0 100644 --- a/gcc/config/i386/i386-options.c +++ b/gcc/config/i386/i386-options.c @@ -178,6 +178,167 @@ static unsigned HOST_WIDE_INT initial_ix86_tune_features[X86_TUNE_LAST] = { /* Feature tests against the various architecture variations. */ unsigned char ix86_arch_features[X86_ARCH_LAST]; +struct ix86_target_opts +{ + const char *option; /* option string */ + HOST_WIDE_INT mask; /* isa mask options */ +}; + +/* This table is ordered so that options like -msse4.2 that imply other + ISAs come first. Target string will be displayed in the same order. */ +static struct ix86_target_opts isa2_opts[] = +{ + { "-mcx16", OPTION_MASK_ISA_CX16 }, + { "-mvaes", OPTION_MASK_ISA_VAES }, + { "-mrdpid", OPTION_MASK_ISA_RDPID }, + { "-mpconfig", OPTION_MASK_ISA_PCONFIG }, + { "-mwbnoinvd", OPTION_MASK_ISA_WBNOINVD }, + { "-mavx512vp2intersect", OPTION_MASK_ISA_AVX512VP2INTERSECT }, + { "-msgx", OPTION_MASK_ISA_SGX }, + { "-mavx5124vnniw", OPTION_MASK_ISA_AVX5124VNNIW }, + { "-mavx5124fmaps", OPTION_MASK_ISA_AVX5124FMAPS }, + { "-mhle", OPTION_MASK_ISA_HLE }, + { "-mmovbe", OPTION_MASK_ISA_MOVBE }, + { "-mclzero", OPTION_MASK_ISA_CLZERO }, + { "-mmwaitx", OPTION_MASK_ISA_MWAITX }, + { "-mmovdir64b", OPTION_MASK_ISA_MOVDIR64B }, + { "-mwaitpkg", OPTION_MASK_ISA_WAITPKG }, + { "-mcldemote", OPTION_MASK_ISA_CLDEMOTE }, + { "-mptwrite", OPTION_MASK_ISA_PTWRITE }, + { "-mavx512bf16", OPTION_MASK_ISA_AVX512BF16 }, + { "-menqcmd", OPTION_MASK_ISA_ENQCMD } +}; +static struct ix86_target_opts isa_opts[] = +{ + { "-mavx512vpopcntdq", OPTION_MASK_ISA_AVX512VPOPCNTDQ }, + { "-mavx512bitalg", OPTION_MASK_ISA_AVX512BITALG }, + { "-mvpclmulqdq", OPTION_MASK_ISA_VPCLMULQDQ }, + { "-mgfni", OPTION_MASK_ISA_GFNI }, + { "-mavx512vnni", OPTION_MASK_ISA_AVX512VNNI }, + { "-mavx512vbmi2", OPTION_MASK_ISA_AVX512VBMI2 }, + { "-mavx512vbmi", OPTION_MASK_ISA_AVX512VBMI }, + { "-mavx512ifma", OPTION_MASK_ISA_AVX512IFMA }, + { "-mavx512vl", OPTION_MASK_ISA_AVX512VL }, + { "-mavx512bw", OPTION_MASK_ISA_AVX512BW }, + { "-mavx512dq", OPTION_MASK_ISA_AVX512DQ }, + { "-mavx512er", OPTION_MASK_ISA_AVX512ER }, + { "-mavx512pf", OPTION_MASK_ISA_AVX512PF }, + { "-mavx512cd", OPTION_MASK_ISA_AVX512CD }, + { "-mavx512f", OPTION_MASK_ISA_AVX512F }, + { "-mavx2", OPTION_MASK_ISA_AVX2 }, + { "-mfma", OPTION_MASK_ISA_FMA }, + { "-mxop", OPTION_MASK_ISA_XOP }, + { "-mfma4", OPTION_MASK_ISA_FMA4 }, + { "-mf16c", OPTION_MASK_ISA_F16C }, + { "-mavx", OPTION_MASK_ISA_AVX }, +/*{ "-msse4" OPTION_MASK_ISA_SSE4 }, */ + { "-msse4.2", OPTION_MASK_ISA_SSE4_2 }, + { "-msse4.1", OPTION_MASK_ISA_SSE4_1 }, + { "-msse4a", OPTION_MASK_ISA_SSE4A }, + { "-mssse3", OPTION_MASK_ISA_SSSE3 }, + { "-msse3", OPTION_MASK_ISA_SSE3 }, + { "-maes", OPTION_MASK_ISA_AES }, + { "-msha", OPTION_MASK_ISA_SHA }, + { "-mpclmul", OPTION_MASK_ISA_PCLMUL }, + { "-msse2", OPTION_MASK_ISA_SSE2 }, + { "-msse", OPTION_MASK_ISA_SSE }, + { "-m3dnowa", OPTION_MASK_ISA_3DNOW_A }, + { "-m3dnow", OPTION_MASK_ISA_3DNOW }, + { "-mmmx", OPTION_MASK_ISA_MMX }, + { "-mrtm", OPTION_MASK_ISA_RTM }, + { "-mprfchw", OPTION_MASK_ISA_PRFCHW }, + { "-mrdseed", OPTION_MASK_ISA_RDSEED }, + { "-madx", OPTION_MASK_ISA_ADX }, + { "-mprefetchwt1", OPTION_MASK_ISA_PREFETCHWT1 }, + { "-mclflushopt", OPTION_MASK_ISA_CLFLUSHOPT }, + { "-mxsaves", OPTION_MASK_ISA_XSAVES }, + { "-mxsavec", OPTION_MASK_ISA_XSAVEC }, + { "-mxsaveopt", OPTION_MASK_ISA_XSAVEOPT }, + { "-mxsave", OPTION_MASK_ISA_XSAVE }, + { "-mabm", OPTION_MASK_ISA_ABM }, + { "-mbmi", OPTION_MASK_ISA_BMI }, + { "-mbmi2", OPTION_MASK_ISA_BMI2 }, + { "-mlzcnt", OPTION_MASK_ISA_LZCNT }, + { "-mtbm", OPTION_MASK_ISA_TBM }, + { "-mpopcnt", OPTION_MASK_ISA_POPCNT }, + { "-msahf", OPTION_MASK_ISA_SAHF }, + { "-mcrc32", OPTION_MASK_ISA_CRC32 }, + { "-mfsgsbase", OPTION_MASK_ISA_FSGSBASE }, + { "-mrdrnd", OPTION_MASK_ISA_RDRND }, + { "-mpku", OPTION_MASK_ISA_PKU }, + { "-mlwp", OPTION_MASK_ISA_LWP }, + { "-mfxsr", OPTION_MASK_ISA_FXSR }, + { "-mclwb", OPTION_MASK_ISA_CLWB }, + { "-mshstk", OPTION_MASK_ISA_SHSTK }, + { "-mmovdiri", OPTION_MASK_ISA_MOVDIRI } +}; + +/* Return 1 if TRAIT NAME is present in the OpenMP context's + device trait set, return 0 if not present in any OpenMP context in the + whole translation unit, or -1 if not present in the current OpenMP context + but might be present in another OpenMP context in the same TU. */ + +int +ix86_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait, + const char *name) +{ + switch (trait) + { + case omp_device_kind: + return strcmp (name, "cpu") == 0; + case omp_device_arch: + if (strcmp (name, "x86") == 0) + return 1; + if (TARGET_64BIT) + { + if (TARGET_X32) + return strcmp (name, "x32") == 0; + else + return strcmp (name, "x86_64") == 0; + } + if (strcmp (name, "ia32") == 0 || strcmp (name, "i386") == 0) + return 1; + if (strcmp (name, "i486") == 0) + return ix86_arch != PROCESSOR_I386 ? 1 : -1; + if (strcmp (name, "i586") == 0) + return (ix86_arch != PROCESSOR_I386 + && ix86_arch != PROCESSOR_I486) ? 1 : -1; + if (strcmp (name, "i686") == 0) + return (ix86_arch != PROCESSOR_I386 + && ix86_arch != PROCESSOR_I486 + && ix86_arch != PROCESSOR_LAKEMONT + && ix86_arch != PROCESSOR_PENTIUM) ? 1 : -1; + return 0; + case omp_device_isa: + for (int i = 0; i < 2; i++) + { + struct ix86_target_opts *opts = i ? isa2_opts : isa_opts; + size_t nopts = i ? ARRAY_SIZE (isa2_opts) : ARRAY_SIZE (isa_opts); + HOST_WIDE_INT mask = i ? ix86_isa_flags2 : ix86_isa_flags; + for (size_t n = 0; n < nopts; n++) + { + const char *option = opts[n].option + 2; + /* -msse4.2 and -msse4.1 options contain dot, which is not valid + in identifiers. Use underscore instead, and handle sse4 + as an alias to sse4_2. */ + if (opts[n].mask == OPTION_MASK_ISA_SSE4_2) + { + option = "sse4_2"; + if (strcmp (name, "sse4") == 0) + return (mask & opts[n].mask) != 0 ? 1 : -1; + } + else if (opts[n].mask == OPTION_MASK_ISA_SSE4_1) + option = "sse4_1"; + if (strcmp (name, option) == 0) + return (mask & opts[n].mask) != 0 ? 1 : -1; + } + } + return 0; + default: + gcc_unreachable (); + } +} + /* Return a string that documents the current -m options. The caller is responsible for freeing the string. */ @@ -187,101 +348,6 @@ ix86_target_string (HOST_WIDE_INT isa, HOST_WIDE_INT isa2, const char *arch, const char *tune, enum fpmath_unit fpmath, bool add_nl_p, bool add_abi_p) { - struct ix86_target_opts - { - const char *option; /* option string */ - HOST_WIDE_INT mask; /* isa mask options */ - }; - - /* This table is ordered so that options like -msse4.2 that imply other - ISAs come first. Target string will be displayed in the same order. */ - static struct ix86_target_opts isa2_opts[] = - { - { "-mcx16", OPTION_MASK_ISA_CX16 }, - { "-mvaes", OPTION_MASK_ISA_VAES }, - { "-mrdpid", OPTION_MASK_ISA_RDPID }, - { "-mpconfig", OPTION_MASK_ISA_PCONFIG }, - { "-mwbnoinvd", OPTION_MASK_ISA_WBNOINVD }, - { "-mavx512vp2intersect", OPTION_MASK_ISA_AVX512VP2INTERSECT }, - { "-msgx", OPTION_MASK_ISA_SGX }, - { "-mavx5124vnniw", OPTION_MASK_ISA_AVX5124VNNIW }, - { "-mavx5124fmaps", OPTION_MASK_ISA_AVX5124FMAPS }, - { "-mhle", OPTION_MASK_ISA_HLE }, - { "-mmovbe", OPTION_MASK_ISA_MOVBE }, - { "-mclzero", OPTION_MASK_ISA_CLZERO }, - { "-mmwaitx", OPTION_MASK_ISA_MWAITX }, - { "-mmovdir64b", OPTION_MASK_ISA_MOVDIR64B }, - { "-mwaitpkg", OPTION_MASK_ISA_WAITPKG }, - { "-mcldemote", OPTION_MASK_ISA_CLDEMOTE }, - { "-mptwrite", OPTION_MASK_ISA_PTWRITE }, - { "-mavx512bf16", OPTION_MASK_ISA_AVX512BF16 }, - { "-menqcmd", OPTION_MASK_ISA_ENQCMD } - }; - static struct ix86_target_opts isa_opts[] = - { - { "-mavx512vpopcntdq", OPTION_MASK_ISA_AVX512VPOPCNTDQ }, - { "-mavx512bitalg", OPTION_MASK_ISA_AVX512BITALG }, - { "-mvpclmulqdq", OPTION_MASK_ISA_VPCLMULQDQ }, - { "-mgfni", OPTION_MASK_ISA_GFNI }, - { "-mavx512vnni", OPTION_MASK_ISA_AVX512VNNI }, - { "-mavx512vbmi2", OPTION_MASK_ISA_AVX512VBMI2 }, - { "-mavx512vbmi", OPTION_MASK_ISA_AVX512VBMI }, - { "-mavx512ifma", OPTION_MASK_ISA_AVX512IFMA }, - { "-mavx512vl", OPTION_MASK_ISA_AVX512VL }, - { "-mavx512bw", OPTION_MASK_ISA_AVX512BW }, - { "-mavx512dq", OPTION_MASK_ISA_AVX512DQ }, - { "-mavx512er", OPTION_MASK_ISA_AVX512ER }, - { "-mavx512pf", OPTION_MASK_ISA_AVX512PF }, - { "-mavx512cd", OPTION_MASK_ISA_AVX512CD }, - { "-mavx512f", OPTION_MASK_ISA_AVX512F }, - { "-mavx2", OPTION_MASK_ISA_AVX2 }, - { "-mfma", OPTION_MASK_ISA_FMA }, - { "-mxop", OPTION_MASK_ISA_XOP }, - { "-mfma4", OPTION_MASK_ISA_FMA4 }, - { "-mf16c", OPTION_MASK_ISA_F16C }, - { "-mavx", OPTION_MASK_ISA_AVX }, -/* { "-msse4" OPTION_MASK_ISA_SSE4 }, */ - { "-msse4.2", OPTION_MASK_ISA_SSE4_2 }, - { "-msse4.1", OPTION_MASK_ISA_SSE4_1 }, - { "-msse4a", OPTION_MASK_ISA_SSE4A }, - { "-mssse3", OPTION_MASK_ISA_SSSE3 }, - { "-msse3", OPTION_MASK_ISA_SSE3 }, - { "-maes", OPTION_MASK_ISA_AES }, - { "-msha", OPTION_MASK_ISA_SHA }, - { "-mpclmul", OPTION_MASK_ISA_PCLMUL }, - { "-msse2", OPTION_MASK_ISA_SSE2 }, - { "-msse", OPTION_MASK_ISA_SSE }, - { "-m3dnowa", OPTION_MASK_ISA_3DNOW_A }, - { "-m3dnow", OPTION_MASK_ISA_3DNOW }, - { "-mmmx", OPTION_MASK_ISA_MMX }, - { "-mrtm", OPTION_MASK_ISA_RTM }, - { "-mprfchw", OPTION_MASK_ISA_PRFCHW }, - { "-mrdseed", OPTION_MASK_ISA_RDSEED }, - { "-madx", OPTION_MASK_ISA_ADX }, - { "-mprefetchwt1", OPTION_MASK_ISA_PREFETCHWT1 }, - { "-mclflushopt", OPTION_MASK_ISA_CLFLUSHOPT }, - { "-mxsaves", OPTION_MASK_ISA_XSAVES }, - { "-mxsavec", OPTION_MASK_ISA_XSAVEC }, - { "-mxsaveopt", OPTION_MASK_ISA_XSAVEOPT }, - { "-mxsave", OPTION_MASK_ISA_XSAVE }, - { "-mabm", OPTION_MASK_ISA_ABM }, - { "-mbmi", OPTION_MASK_ISA_BMI }, - { "-mbmi2", OPTION_MASK_ISA_BMI2 }, - { "-mlzcnt", OPTION_MASK_ISA_LZCNT }, - { "-mtbm", OPTION_MASK_ISA_TBM }, - { "-mpopcnt", OPTION_MASK_ISA_POPCNT }, - { "-msahf", OPTION_MASK_ISA_SAHF }, - { "-mcrc32", OPTION_MASK_ISA_CRC32 }, - { "-mfsgsbase", OPTION_MASK_ISA_FSGSBASE }, - { "-mrdrnd", OPTION_MASK_ISA_RDRND }, - { "-mpku", OPTION_MASK_ISA_PKU }, - { "-mlwp", OPTION_MASK_ISA_LWP }, - { "-mfxsr", OPTION_MASK_ISA_FXSR }, - { "-mclwb", OPTION_MASK_ISA_CLWB }, - { "-mshstk", OPTION_MASK_ISA_SHSTK }, - { "-mmovdiri", OPTION_MASK_ISA_MOVDIRI } - }; - /* Flag options. */ static struct ix86_target_opts flag_opts[] = { diff --git a/gcc/config/i386/i386-options.h b/gcc/config/i386/i386-options.h index 966eb766ad9..1d46aa1d1ce 100644 --- a/gcc/config/i386/i386-options.h +++ b/gcc/config/i386/i386-options.h @@ -19,6 +19,9 @@ along with GCC; see the file COPYING3. If not see #ifndef GCC_I386_OPTIONS_H #define GCC_I386_OPTIONS_H +extern int ix86_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait, + const char *name); + char *ix86_target_string (HOST_WIDE_INT isa, HOST_WIDE_INT isa2, int flags, int flags2, const char *arch, const char *tune, diff --git a/gcc/config/i386/i386.c b/gcc/config/i386/i386.c index b9122a4ded9..03a7082d2fc 100644 --- a/gcc/config/i386/i386.c +++ b/gcc/config/i386/i386.c @@ -23034,12 +23034,13 @@ ix86_run_selftests (void) ix86_simd_clone_compute_vecsize_and_simdlen #undef TARGET_SIMD_CLONE_ADJUST -#define TARGET_SIMD_CLONE_ADJUST \ - ix86_simd_clone_adjust +#define TARGET_SIMD_CLONE_ADJUST ix86_simd_clone_adjust #undef TARGET_SIMD_CLONE_USABLE -#define TARGET_SIMD_CLONE_USABLE \ - ix86_simd_clone_usable +#define TARGET_SIMD_CLONE_USABLE ix86_simd_clone_usable + +#undef TARGET_OMP_DEVICE_KIND_ARCH_ISA +#define TARGET_OMP_DEVICE_KIND_ARCH_ISA ix86_omp_device_kind_arch_isa #undef TARGET_FLOAT_EXCEPTIONS_ROUNDING_SUPPORTED_P #define TARGET_FLOAT_EXCEPTIONS_ROUNDING_SUPPORTED_P \ diff --git a/gcc/config/i386/t-intelmic b/gcc/config/i386/t-intelmic index 9de4b76e310..58f166803bf 100644 --- a/gcc/config/i386/t-intelmic +++ b/gcc/config/i386/t-intelmic @@ -8,3 +8,10 @@ ALL_HOST_OBJS += mkoffload.o mkoffload$(exeext): mkoffload.o collect-utils.o libcommon-target.a $(LIBIBERTY) $(LIBDEPS) $(LINKER) $(ALL_LINKERFLAGS) $(LDFLAGS) -o $@ \ mkoffload.o collect-utils.o libcommon-target.a $(LIBIBERTY) $(LIBS) + +omp-device-properties: $(srcdir)/config/i386/i386-options.c + echo kind: cpu > omp-device-properties + echo arch: x86 x86_64 i386 i486 i586 i686 ia32 >> omp-device-properties + echo isa: sse4 `sed -n '/^static struct ix86_target_opts isa2\?_opts\[\] =/,/^};/p' \ + $(srcdir)/config/i386/i386-options.c | \ + sed -n 's/",.*$$//;s/\./_/;s/^ { "-m//p'` >> omp-device-properties diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 90171a95784..0d6e8840852 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -5474,6 +5474,32 @@ nvptx_simt_vf () return PTX_WARP_SIZE; } +/* Return 1 if TRAIT NAME is present in the OpenMP context's + device trait set, return 0 if not present in any OpenMP context in the + whole translation unit, or -1 if not present in the current OpenMP context + but might be present in another OpenMP context in the same TU. */ + +int +nvptx_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait, + const char *name) +{ + switch (trait) + { + case omp_device_kind: + return strcmp (name, "gpu") == 0; + case omp_device_arch: + return strcmp (name, "nvptx") == 0; + case omp_device_isa: + if (strcmp (name, "sm_30") == 0) + return !TARGET_SM35; + if (strcmp (name, "sm_35") == 0) + return TARGET_SM35; + return 0; + default: + gcc_unreachable (); + } +} + static bool nvptx_welformed_vector_length_p (int l) { @@ -6539,6 +6565,9 @@ nvptx_set_current_function (tree fndecl) #undef TARGET_SIMT_VF #define TARGET_SIMT_VF nvptx_simt_vf +#undef TARGET_OMP_DEVICE_KIND_ARCH_ISA +#define TARGET_OMP_DEVICE_KIND_ARCH_ISA nvptx_omp_device_kind_arch_isa + #undef TARGET_GOACC_VALIDATE_DIMS #define TARGET_GOACC_VALIDATE_DIMS nvptx_goacc_validate_dims diff --git a/gcc/config/nvptx/t-nvptx b/gcc/config/nvptx/t-nvptx index 6c1010ddd66..d692594a377 100644 --- a/gcc/config/nvptx/t-nvptx +++ b/gcc/config/nvptx/t-nvptx @@ -10,3 +10,8 @@ mkoffload$(exeext): mkoffload.o collect-utils.o libcommon-target.a $(LIBIBERTY) mkoffload.o collect-utils.o libcommon-target.a $(LIBIBERTY) $(LIBS) MULTILIB_OPTIONS = mgomp + +omp-device-properties: $(srcdir)/config/nvptx/nvptx.c + echo kind: gpu > omp-device-properties + echo arch: nvptx >> omp-device-properties + echo isa: sm_30 sm_35 >> omp-device-properties diff --git a/gcc/configure b/gcc/configure index 9de9ef85f24..13935222228 100755 --- a/gcc/configure +++ b/gcc/configure @@ -811,6 +811,8 @@ LN LN_S AWK SET_MAKE +omp_device_property_deps +omp_device_properties accel_dir_suffix real_target_noncanonical enable_as_accelerator @@ -7879,12 +7881,20 @@ fi for tgt in `echo $enable_offload_targets | sed 's/,/ /g'`; do + tgt_dir=`echo $tgt | sed -n 's/.*=//p'` tgt=`echo $tgt | sed 's/=.*//'` if echo "$tgt" | grep "^hsa" > /dev/null ; then enable_hsa=1 else enable_offloading=1 + if test -n "$tgt_dir"; then + omp_device_property="${tgt_dir}/lib/gcc/\$(real_target_noncanonical)/\$(version)/accel/${tgt}/omp-device-properties" + else + omp_device_property="\$(libsubdir)/accel/${tgt}/omp-device-properties" + fi + omp_device_properties="${omp_device_properties} ${tgt}=${omp_device_property}" + omp_device_property_deps="${omp_device_property_deps} ${omp_device_property}" fi if test x"$offload_targets" = x; then @@ -7894,6 +7904,9 @@ for tgt in `echo $enable_offload_targets | sed 's/,/ /g'`; do fi done + + + cat >>confdefs.h <<_ACEOF #define OFFLOAD_TARGETS "$offload_targets" _ACEOF @@ -18851,7 +18864,7 @@ else lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2 lt_status=$lt_dlunknown cat > conftest.$ac_ext <<_LT_EOF -#line 18854 "configure" +#line 18867 "configure" #include "confdefs.h" #if HAVE_DLFCN_H @@ -18957,7 +18970,7 @@ else lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2 lt_status=$lt_dlunknown cat > conftest.$ac_ext <<_LT_EOF -#line 18960 "configure" +#line 18973 "configure" #include "confdefs.h" #if HAVE_DLFCN_H diff --git a/gcc/configure.ac b/gcc/configure.ac index 62f4b2651cc..695e994ec9f 100644 --- a/gcc/configure.ac +++ b/gcc/configure.ac @@ -1026,12 +1026,20 @@ AC_SUBST(real_target_noncanonical) AC_SUBST(accel_dir_suffix) for tgt in `echo $enable_offload_targets | sed 's/,/ /g'`; do + tgt_dir=`echo $tgt | sed -n 's/.*=//p'` tgt=`echo $tgt | sed 's/=.*//'` if echo "$tgt" | grep "^hsa" > /dev/null ; then enable_hsa=1 else enable_offloading=1 + if test -n "$tgt_dir"; then + omp_device_property="${tgt_dir}/lib/gcc/\$(real_target_noncanonical)/\$(version)/accel/${tgt}/omp-device-properties" + else + omp_device_property="\$(libsubdir)/accel/${tgt}/omp-device-properties" + fi + omp_device_properties="${omp_device_properties} ${tgt}=${omp_device_property}" + omp_device_property_deps="${omp_device_property_deps} ${omp_device_property}" fi if test x"$offload_targets" = x; then @@ -1040,6 +1048,9 @@ for tgt in `echo $enable_offload_targets | sed 's/,/ /g'`; do offload_targets="$offload_targets,$tgt" fi done +AC_SUBST(omp_device_properties) +AC_SUBST(omp_device_property_deps) + AC_DEFINE_UNQUOTED(OFFLOAD_TARGETS, "$offload_targets", [Define to offload targets, separated by commas.]) if test x"$enable_offloading" != x; then diff --git a/gcc/doc/tm.texi b/gcc/doc/tm.texi index 915e9612208..cd9aed9874f 100644 --- a/gcc/doc/tm.texi +++ b/gcc/doc/tm.texi @@ -6103,6 +6103,13 @@ to use it. Return number of threads in SIMT thread group on the target. @end deftypefn +@deftypefn {Target Hook} int TARGET_OMP_DEVICE_KIND_ARCH_ISA (enum omp_device_kind_arch_isa @var{trait}, const char *@var{name}) +Return 1 if @var{trait} @var{name} is present in the OpenMP context's +device trait set, return 0 if not present in any OpenMP context in the +whole translation unit, or -1 if not present in the current OpenMP context +but might be present in another OpenMP context in the same TU. +@end deftypefn + @deftypefn {Target Hook} bool TARGET_GOACC_VALIDATE_DIMS (tree @var{decl}, int *@var{dims}, int @var{fn_level}, unsigned @var{used}) This hook should check the launch dimensions provided for an OpenACC compute region, or routine. Defaulted values are represented as -1 diff --git a/gcc/doc/tm.texi.in b/gcc/doc/tm.texi.in index ac0f0494992..2739e9ceec5 100644 --- a/gcc/doc/tm.texi.in +++ b/gcc/doc/tm.texi.in @@ -4205,6 +4205,8 @@ address; but often a machine-dependent strategy can generate better code. @hook TARGET_SIMT_VF +@hook TARGET_OMP_DEVICE_KIND_ARCH_ISA + @hook TARGET_GOACC_VALIDATE_DIMS @hook TARGET_GOACC_DIM_LIMIT diff --git a/gcc/omp-general.c b/gcc/omp-general.c index 9397b1951f8..82f47467499 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -40,6 +40,7 @@ along with GCC; see the file COPYING3. If not see #include "symbol-summary.h" #include "hsa-common.h" #include "tree-pass.h" +#include "omp-device-properties.h" enum omp_requires omp_requires_mask; @@ -537,7 +538,7 @@ omp_max_simt_vf (void) { if (!strncmp (c, "nvptx", strlen ("nvptx"))) return 32; - else if ((c = strchr (c, ','))) + else if ((c = strchr (c, ':'))) c++; } return 0; @@ -571,6 +572,79 @@ omp_constructor_traits_to_codes (tree ctx, enum tree_code *constructs) return nconstructs; } +/* Return true if PROP is possibly present in one of the offloading target's + OpenMP contexts. The format of PROPS string is always offloading target's + name terminated by '\0', followed by properties for that offloading + target separated by '\0' and terminated by another '\0'. The strings + are created from omp-device-properties installed files of all configured + offloading targets. */ + +static bool +omp_offload_device_kind_arch_isa (const char *props, const char *prop) +{ + const char *names = getenv ("OFFLOAD_TARGET_NAMES"); + if (names == NULL || *names == '\0') + return false; + while (*props != '\0') + { + size_t name_len = strlen (props); + bool matches = false; + for (const char *c = names; c; ) + { + if (strncmp (props, c, name_len) == 0 + && (c[name_len] == '\0' + || c[name_len] == ':' + || c[name_len] == '=')) + { + matches = true; + break; + } + else if ((c = strchr (c, ':'))) + c++; + } + props = props + name_len + 1; + while (*props != '\0') + { + if (matches && strcmp (props, prop) == 0) + return true; + props = strchr (props, '\0') + 1; + } + props++; + } + return false; +} + +/* Return true if the current code location is or might be offloaded. + Return true in declare target functions, or when nested in a target + region or when unsure, return false otherwise. */ + +static bool +omp_maybe_offloaded (void) +{ + if (!hsa_gen_requested_p ()) + { + if (!ENABLE_OFFLOADING) + return false; + const char *names = getenv ("OFFLOAD_TARGET_NAMES"); + if (names == NULL || *names == '\0') + return false; + } + if (symtab->state == PARSING) + /* Maybe. */ + return true; + if (current_function_decl + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (current_function_decl))) + return true; + if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0) + { + enum tree_code construct = OMP_TARGET; + if (omp_construct_selector_matches (&construct, 1)) + return true; + } + return false; +} + /* Return 1 if context selector matches the current OpenMP context, 0 if it does not and -1 if it is unknown and need to be determined later. Some properties can be checked right away during parsing (this routine), @@ -667,8 +741,45 @@ omp_context_selector_matches (tree ctx) return 0; } if (set == 'd' && !strcmp (sel, "arch")) - /* For now, need a target hook. */ - ret = -1; + for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3)) + { + const char *arch = IDENTIFIER_POINTER (TREE_PURPOSE (t3)); + int r = 0; + if (targetm.omp.device_kind_arch_isa != NULL) + r = targetm.omp.device_kind_arch_isa (omp_device_arch, + arch); + if (r == 0 || (r == -1 && symtab->state != PARSING)) + { + /* If we are or might be in a target region or + declare target function, need to take into account + also offloading values. */ + if (!omp_maybe_offloaded ()) + return 0; + if (strcmp (arch, "hsa") == 0 + && hsa_gen_requested_p ()) + { + ret = -1; + continue; + } + if (ENABLE_OFFLOADING) + { + const char *arches = omp_offload_device_arch; + if (omp_offload_device_kind_arch_isa (arches, + arch)) + { + ret = -1; + continue; + } + } + return 0; + } + else if (r == -1) + ret = -1; + /* If arch matches on the host, it still might not match + in the offloading region. */ + else if (omp_maybe_offloaded ()) + ret = -1; + } break; case 'u': if (set == 'i' && !strcmp (sel, "unified_address")) @@ -729,57 +840,92 @@ omp_context_selector_matches (tree ctx) const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3)); if (!strcmp (prop, "any")) continue; - if (!strcmp (prop, "fpga")) - return 0; /* Right now GCC doesn't support any fpgas. */ if (!strcmp (prop, "host")) { - if (ENABLE_OFFLOADING || hsa_gen_requested_p ()) + if (omp_maybe_offloaded ()) ret = -1; continue; } if (!strcmp (prop, "nohost")) { - if (ENABLE_OFFLOADING || hsa_gen_requested_p ()) + if (omp_maybe_offloaded ()) ret = -1; else return 0; continue; } - if (!strcmp (prop, "cpu") || !strcmp (prop, "gpu")) + int r = 0; + if (targetm.omp.device_kind_arch_isa != NULL) + r = targetm.omp.device_kind_arch_isa (omp_device_kind, + prop); + else + r = strcmp (prop, "cpu") == 0; + if (r == 0 || (r == -1 && symtab->state != PARSING)) { - bool maybe_gpu = false; - if (hsa_gen_requested_p ()) - maybe_gpu = true; - else if (ENABLE_OFFLOADING) - for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); - c; ) - { - if (!strncmp (c, "nvptx", strlen ("nvptx")) - || !strncmp (c, "amdgcn", strlen ("amdgcn"))) - { - maybe_gpu = true; - break; - } - else if ((c = strchr (c, ','))) - c++; - } - if (!maybe_gpu) + /* If we are or might be in a target region or + declare target function, need to take into account + also offloading values. */ + if (!omp_maybe_offloaded ()) + return 0; + if (strcmp (prop, "gpu") == 0 + && hsa_gen_requested_p ()) { - if (prop[0] == 'g') - return 0; + ret = -1; + continue; } - else - ret = -1; - continue; + if (ENABLE_OFFLOADING) + { + const char *kinds = omp_offload_device_kind; + if (omp_offload_device_kind_arch_isa (kinds, prop)) + { + ret = -1; + continue; + } + } + return 0; } - /* Any other kind doesn't match. */ - return 0; + else if (r == -1) + ret = -1; + /* If kind matches on the host, it still might not match + in the offloading region. */ + else if (omp_maybe_offloaded ()) + ret = -1; } break; case 'i': if (set == 'd' && !strcmp (sel, "isa")) - /* For now, need a target hook. */ - ret = -1; + for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3)) + { + const char *isa = IDENTIFIER_POINTER (TREE_PURPOSE (t3)); + int r = 0; + if (targetm.omp.device_kind_arch_isa != NULL) + r = targetm.omp.device_kind_arch_isa (omp_device_isa, + isa); + if (r == 0 || (r == -1 && symtab->state != PARSING)) + { + /* If we are or might be in a target region or + declare target function, need to take into account + also offloading values. */ + if (!omp_maybe_offloaded ()) + return 0; + if (ENABLE_OFFLOADING) + { + const char *isas = omp_offload_device_isa; + if (omp_offload_device_kind_arch_isa (isas, isa)) + { + ret = -1; + continue; + } + } + return 0; + } + else if (r == -1) + ret = -1; + /* If isa matches on the host, it still might not match + in the offloading region. */ + else if (omp_maybe_offloaded ()) + ret = -1; + } break; case 'c': if (set == 'u' && !strcmp (sel, "condition")) diff --git a/gcc/target.def b/gcc/target.def index 1f011edf88b..8e83c2c7a71 100644 --- a/gcc/target.def +++ b/gcc/target.def @@ -1669,6 +1669,21 @@ int, (void), NULL) HOOK_VECTOR_END (simt) +/* Functions relating to OpenMP. */ +#undef HOOK_PREFIX +#define HOOK_PREFIX "TARGET_OMP_" +HOOK_VECTOR (TARGET_OMP, omp) + +DEFHOOK +(device_kind_arch_isa, +"Return 1 if @var{trait} @var{name} is present in the OpenMP context's\n\ +device trait set, return 0 if not present in any OpenMP context in the\n\ +whole translation unit, or -1 if not present in the current OpenMP context\n\ +but might be present in another OpenMP context in the same TU.", +int, (enum omp_device_kind_arch_isa trait, const char *name), NULL) + +HOOK_VECTOR_END (omp) + /* Functions relating to openacc. */ #undef HOOK_PREFIX #define HOOK_PREFIX "TARGET_GOACC_" diff --git a/gcc/target.h b/gcc/target.h index 9f80658d8c8..843c3d7887f 100644 --- a/gcc/target.h +++ b/gcc/target.h @@ -211,6 +211,13 @@ typedef vec vector_sizes; automatically freed. */ typedef auto_vec auto_vector_sizes; +/* First argument of targetm.omp.device_kind_arch_isa. */ +enum omp_device_kind_arch_isa { + omp_device_kind, + omp_device_arch, + omp_device_isa +}; + /* The target structure. This holds all the backend hooks. */ #define DEFHOOKPOD(NAME, DOC, TYPE, INIT) TYPE NAME; #define DEFHOOK(NAME, DOC, TYPE, PARAMS, INIT) TYPE (* NAME) PARAMS; diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 6b9451b8ede..8e6588a8b02 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,8 @@ +2019-10-31 Jakub Jelinek + + * c-c++-common/gomp/declare-variant-9.c: New test. + * c-c++-common/gomp/declare-variant-10.c: New test. + 2019-10-31 Tobias Burnus PR fortran/92277 diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-10.c b/gcc/testsuite/c-c++-common/gomp/declare-variant-10.c new file mode 100644 index 00000000000..ad886cfb887 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-10.c @@ -0,0 +1,77 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-foffload=disable -fdump-tree-gimple" } */ +/* { dg-additional-options "-mavx512bw" { target { i?86-*-* x86_64-*-* } } } */ + +#undef i386 +void f01 (void); +#pragma omp declare variant (f01) match (device={isa(avx512f,avx512bw)}) +void f02 (void); +void f03 (void); +#pragma omp declare variant (f03) match (device={kind(any),arch(x86_64),isa(avx512f,avx512bw)}) +void f04 (void); +void f05 (void); +#pragma omp declare variant (f05) match (device={kind(gpu)}) +void f06 (void); +void f07 (void); +#pragma omp declare variant (f07) match (device={kind(cpu)}) +void f08 (void); +void f09 (void); +#pragma omp declare variant (f09) match (device={isa(sm_35)}) +void f10 (void); +void f11 (void); +#pragma omp declare variant (f11) match (device={arch(nvptx)}) +void f12 (void); +void f13 (void); +#pragma omp declare variant (f13) match (device={arch(i386),isa(sse4)}) +void f14 (void); +void f15 (void); +#pragma omp declare variant (f15) match (device={isa(sse4,ssse3),arch(i386)}) +void f16 (void); +void f17 (void); +#pragma omp declare variant (f17) match (device={kind(any,fpga)}) +void f18 (void); + +#pragma omp declare target +void +test1 (void) +{ + int i; + f02 (); /* { dg-final { scan-tree-dump-times "f01 \\\(\\\);" 1 "gimple" { target i?86-*-* x86_64-*-* } } } */ + /* { dg-final { scan-tree-dump-times "f02 \\\(\\\);" 1 "gimple" { target { ! { i?86-*-* x86_64-*-* } } } } } */ + f14 (); /* { dg-final { scan-tree-dump-times "f13 \\\(\\\);" 1 "gimple" { target ia32 } } } */ + /* { dg-final { scan-tree-dump-times "f14 \\\(\\\);" 1 "gimple" { target { ! ia32 } } } } */ + f18 (); /* { dg-final { scan-tree-dump-times "f18 \\\(\\\);" 1 "gimple" } } */ +} +#pragma omp end declare target + +#if defined(__i386__) || defined(__x86_64__) +__attribute__((target ("avx512f,avx512bw"))) +#endif +void +test2 (void) +{ + #pragma omp target + f04 (); /* { dg-final { scan-tree-dump-times "f03 \\\(\\\);" 1 "gimple" { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */ + /* { dg-final { scan-tree-dump-times "f04 \\\(\\\);" 1 "gimple" { target { { ! lp64 } || { ! { i?86-*-* x86_64-*-* } } } } } } */ + #pragma omp target + f16 (); /* { dg-final { scan-tree-dump-times "f15 \\\(\\\);" 1 "gimple" { target ia32 } } } */ + /* { dg-final { scan-tree-dump-times "f16 \\\(\\\);" 1 "gimple" { target { ! ia32 } } } } */ +} + +void +test3 (void) +{ + f06 (); /* { dg-final { scan-tree-dump-times "f06 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */ + f08 (); /* { dg-final { scan-tree-dump-times "f07 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */ +} +#pragma omp declare target to (test3) + +void +test4 (void) +{ + #pragma omp target + f10 (); /* { dg-final { scan-tree-dump-times "f10 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */ + #pragma omp target + f12 (); /* { dg-final { scan-tree-dump-times "f12 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* } } } } } */ + /* { dg-final { scan-tree-dump-times "f11 \\\(\\\);" 1 "gimple" { target { nvptx*-*-* } } } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-9.c b/gcc/testsuite/c-c++-common/gomp/declare-variant-9.c new file mode 100644 index 00000000000..21ddc56681b --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-9.c @@ -0,0 +1,63 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ +/* { dg-additional-options "-mno-sse3" { target { i?86-*-* x86_64-*-* } } } */ + +#undef i386 +void f01 (void); +#pragma omp declare variant (f01) match (device={isa(avx512f,avx512bw)}) +void f02 (void); +void f03 (void); +#pragma omp declare variant (f03) match (device={kind(any),arch(x86_64),isa(avx512f,avx512bw)}) +void f04 (void); +void f05 (void); +#pragma omp declare variant (f05) match (device={kind(gpu)}) +void f06 (void); +void f07 (void); +#pragma omp declare variant (f07) match (device={kind(cpu)}) +void f08 (void); +void f09 (void); +#pragma omp declare variant (f09) match (device={isa(sm_35)}) +void f10 (void); +void f11 (void); +#pragma omp declare variant (f11) match (device={arch(nvptx)}) +void f12 (void); +void f13 (void); +#pragma omp declare variant (f13) match (device={arch(i386),isa(sse4)}) +void f14 (void); +void f15 (void); +#pragma omp declare variant (f15) match (device={isa(sse4,ssse3),arch(i386)}) +void f16 (void); +void f17 (void); +#pragma omp declare variant (f17) match (device={kind(any,fpga)}) +void f18 (void); + +void +test1 (void) +{ + int i; + f02 (); /* { dg-final { scan-tree-dump-times "f02 \\\(\\\);" 1 "gimple" } } */ + f14 (); /* { dg-final { scan-tree-dump-times "f14 \\\(\\\);" 1 "gimple" } } */ + f18 (); /* { dg-final { scan-tree-dump-times "f18 \\\(\\\);" 1 "gimple" } } */ +} + +#if defined(__i386__) || defined(__x86_64__) +__attribute__((target ("avx512f,avx512bw"))) +#endif +void +test2 (void) +{ + f04 (); /* { dg-final { scan-tree-dump-times "f03 \\\(\\\);" 1 "gimple" { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */ + /* { dg-final { scan-tree-dump-times "f04 \\\(\\\);" 1 "gimple" { target { { ! lp64 } || { ! { i?86-*-* x86_64-*-* } } } } } } */ + f16 (); /* { dg-final { scan-tree-dump-times "f15 \\\(\\\);" 1 "gimple" { target ia32 } } } */ + /* { dg-final { scan-tree-dump-times "f16 \\\(\\\);" 1 "gimple" { target { ! ia32 } } } } */ +} + +void +test3 (void) +{ + f06 (); /* { dg-final { scan-tree-dump-times "f06 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */ + f08 (); /* { dg-final { scan-tree-dump-times "f07 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */ + f10 (); /* { dg-final { scan-tree-dump-times "f10 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */ + f12 (); /* { dg-final { scan-tree-dump-times "f12 \\\(\\\);" 1 "gimple" { target { ! { nvptx*-*-* } } } } } */ + /* { dg-final { scan-tree-dump-times "f11 \\\(\\\);" 1 "gimple" { target { nvptx*-*-* } } } } */ +} -- 2.30.2