diff options
| author | Martin Jambor <mjambor@suse.cz> | 2020-08-03 18:13:00 +0200 |
|---|---|---|
| committer | Martin Jambor <mjambor@suse.cz> | 2020-08-03 18:13:00 +0200 |
| commit | c56684fd61223abd45854270cd1e83ab2f07148c (patch) | |
| tree | f1ebfbe4dfd00b87f9522b84f78f02ee512e83a8 /libgomp/testsuite | |
| parent | 9623f61b142174b87760c81f78928dd14af7cbc6 (diff) | |
| download | gcc-c56684fd61223abd45854270cd1e83ab2f07148c.zip gcc-c56684fd61223abd45854270cd1e83ab2f07148c.tar.gz gcc-c56684fd61223abd45854270cd1e83ab2f07148c.tar.bz2 | |
Removal of HSA offloading from gcc and libgomp
This patch removes the generation of HSAIL from the compiler, the HSA
offloading plugin from libgomp and the associated testsuite tests and
infrastructure bits from the respective testsuites.
Apart from removal of the obvious files, I removed bits that I found
by searching for HSA related terms and by re-tracing my steps and
looking at the patches that introduced HSA in the first place. I did
not remove everything these patches brought in, for example:
- the mechanism to pass offload-target specific info from the application to
the offloading plugin - but the same mechanism is also used to
communicate number of teams and the thread limit to all offload targets.
- run_func hook in gomp_device_descr stays too, although now it is
not used. If some future offload target would like the ability to
refuse to offload some functions, it can use it. It is easy to
remove as a follow-up if it is considered clutter, though.
- configure options --with-hsa-runtime=PATH, -with-hsa-runtime-include=PATH
and --with-hsa-runtime-lib=PATH rmeain because GCN uses them too.
- Surprisingly, GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES (a constant
from gomp-constants.h) appears in the source of the amdgcn libgomp
plugin, although I tend to think that code path is not ever used
and this patch certainly removes it from the compiler.
Nevertheless, it seems it has potential value beyond HSAIL and so
I've kept it, it can of course always be easily removed in the
future of GCN folk abandon it too.
- I assume constants OFFLOAD_TARGET_TYPE_HSA and GOMP_DEVICE_HSA
need to stay indefinitely too just so that no future offload
target picks that number.
- I have kept dg-require-effective-target
offload_device_nonshared_as requirement of thests which have it.
It is quite probable I missed some small HSA artifacts but those
should be easy to remove later as we find them.
include/ChangeLog:
2020-07-24 Martin Jambor <mjambor@suse.cz>
* gomp-constants.h (GOMP_VERSION_HSA): Remove.
gcc/ChangeLog:
2020-07-24 Martin Jambor <mjambor@suse.cz>
* hsa-brig-format.h: Moved to brig/brigfrontend.
* hsa-brig.c: Removed.
* hsa-builtins.def: Likewise.
* hsa-common.c: Likewise.
* hsa-common.h: Likewise.
* hsa-dump.c: Likewise.
* hsa-gen.c: Likewise.
* hsa-regalloc.c: Likewise.
* ipa-hsa.c: Likewise.
* omp-grid.c: Likewise.
* omp-grid.h: Likewise.
* Makefile.in (BUILTINS_DEF): Remove hsa-builtins.def.
(OBJS): Remove hsa-common.o, hsa-gen.o, hsa-regalloc.o, hsa-brig.o,
hsa-dump.o, ipa-hsa.c and omp-grid.o.
(GTFILES): Removed hsa-common.c and omp-expand.c.
* builtins.def: Remove processing of hsa-builtins.def.
(DEF_HSA_BUILTIN): Remove.
* common.opt (flag_disable_hsa): Remove.
(-Whsa): Ignore.
* config.in (ENABLE_HSA): Removed.
* configure.ac: Removed handling configuration for hsa offloading.
(ENABLE_HSA): Removed.
* configure: Regenerated.
* doc/install.texi (--enable-offload-targets): Remove hsa from the
example.
(--with-hsa-runtime): Reword to reference any HSA run-time, not
specifically HSA offloading.
* doc/invoke.texi (Option Summary): Remove -Whsa.
(Warning Options): Likewise.
(Optimize Options): Remove hsa-gen-debug-stores.
* doc/passes.texi (Regular IPA passes): Remove section on IPA HSA
pass.
* gimple-low.c (lower_stmt): Remove GIMPLE_OMP_GRID_BODY case.
* gimple-pretty-print.c (dump_gimple_omp_for): Likewise.
(dump_gimple_omp_block): Likewise.
(pp_gimple_stmt_1): Likewise.
* gimple-walk.c (walk_gimple_stmt): Likewise.
* gimple.c (gimple_build_omp_grid_body): Removed function.
(gimple_copy): Remove GIMPLE_OMP_GRID_BODY case.
* gimple.def (GIMPLE_OMP_GRID_BODY): Removed.
* gimple.h (gf_mask): Removed GF_OMP_PARALLEL_GRID_PHONY,
OMP_FOR_KIND_GRID_LOOP, GF_OMP_FOR_GRID_PHONY,
GF_OMP_FOR_GRID_INTRA_GROUP, GF_OMP_FOR_GRID_GROUP_ITER and
GF_OMP_TEAMS_GRID_PHONY. Renumbered GF_OMP_FOR_KIND_SIMD and
GF_OMP_TEAMS_HOST.
(gimple_build_omp_grid_body): Removed declaration.
(gimple_has_substatements): Remove GIMPLE_OMP_GRID_BODY case.
(gimple_omp_for_grid_phony): Removed.
(gimple_omp_for_set_grid_phony): Likewise.
(gimple_omp_for_grid_intra_group): Likewise.
(gimple_omp_for_grid_intra_group): Likewise.
(gimple_omp_for_grid_group_iter): Likewise.
(gimple_omp_for_set_grid_group_iter): Likewise.
(gimple_omp_parallel_grid_phony): Likewise.
(gimple_omp_parallel_set_grid_phony): Likewise.
(gimple_omp_teams_grid_phony): Likewise.
(gimple_omp_teams_set_grid_phony): Likewise.
(CASE_GIMPLE_OMP): Remove GIMPLE_OMP_GRID_BODY case.
* lto-section-in.c (lto_section_name): Removed hsa.
* lto-streamer.h (lto_section_type): Removed LTO_section_ipa_hsa.
* lto-wrapper.c (compile_images_for_offload_targets): Remove special
handling of hsa.
* omp-expand.c: Do not include hsa-common.h and gt-omp-expand.h.
(parallel_needs_hsa_kernel_p): Removed.
(grid_launch_attributes_trees): Likewise.
(grid_launch_attributes_trees): Likewise.
(grid_create_kernel_launch_attr_types): Likewise.
(grid_insert_store_range_dim): Likewise.
(grid_get_kernel_launch_attributes): Likewise.
(get_target_arguments): Remove code passing HSA grid sizes.
(grid_expand_omp_for_loop): Remove.
(grid_arg_decl_map): Likewise.
(grid_remap_kernel_arg_accesses): Likewise.
(grid_expand_target_grid_body): Likewise.
(expand_omp): Remove call to grid_expand_target_grid_body.
(omp_make_gimple_edges): Remove GIMPLE_OMP_GRID_BODY case.
* omp-general.c: Do not include hsa-common.h.
(omp_maybe_offloaded): Do not check for HSA offloading.
(omp_context_selector_matches): Likewise.
* omp-low.c: Do not include hsa-common.h and omp-grid.h.
(build_outer_var_ref): Remove handling of GIMPLE_OMP_GRID_BODY.
(scan_sharing_clauses): Remove handling of OMP_CLAUSE__GRIDDIM_.
(scan_omp_parallel): Remove handling of the phoney variant.
(check_omp_nesting_restrictions): Remove handling of
GIMPLE_OMP_GRID_BODY and GF_OMP_FOR_KIND_GRID_LOOP.
(scan_omp_1_stmt): Remove handling of GIMPLE_OMP_GRID_BODY.
(lower_omp_for_lastprivate): Remove handling of gridified loops.
(lower_omp_for): Remove phony loop handling.
(lower_omp_taskreg): Remove phony construct handling.
(lower_omp_teams): Likewise.
(lower_omp_grid_body): Removed.
(lower_omp_1): Remove GIMPLE_OMP_GRID_BODY case.
(execute_lower_omp): Do not call omp_grid_gridify_all_targets.
* opts.c (common_handle_option): Do not handle hsa when processing
OPT_foffload_.
* params.opt (hsa-gen-debug-stores): Remove.
* passes.def: Remove pass_ipa_hsa and pass_gen_hsail.
* timevar.def: Remove TV_IPA_HSA.
* toplev.c: Do not include hsa-common.h.
(compile_file): Do not call hsa_output_brig.
* tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE__GRIDDIM_.
(tree_omp_clause): Remove union field dimension.
* tree-nested.c (convert_nonlocal_omp_clauses): Remove the
OMP_CLAUSE__GRIDDIM_ case.
(convert_local_omp_clauses): Likewise.
* tree-pass.h (make_pass_gen_hsail): Remove declaration.
(make_pass_ipa_hsa): Likewise.
* tree-pretty-print.c (dump_omp_clause): Remove GIMPLE_OMP_GRID_BODY
case.
* tree.c (omp_clause_num_ops): Remove the element corresponding to
OMP_CLAUSE__GRIDDIM_.
(omp_clause_code_name): Likewise.
(walk_tree_1): Remove GIMPLE_OMP_GRID_BODY case.
* tree.h (OMP_CLAUSE__GRIDDIM__DIMENSION): Remove.
(OMP_CLAUSE__GRIDDIM__SIZE): Likewise.
(OMP_CLAUSE__GRIDDIM__GROUP): Likewise.
gcc/fortran/ChangeLog:
2020-07-24 Martin Jambor <mjambor@suse.cz>
* f95-lang.c (gfc_init_builtin_functions): Remove processing of
hsa-builtins.def.
gcc/brig/ChangeLog:
2020-07-24 Martin Jambor <mjambor@suse.cz>
* brigfrontend/brig-util.h (hsa_type_packed_p): Declared.
* brigfrontend/brig-util.cc (hsa_type_packed_p): Moved here from
removed gcc/hsa-common.c.
libgomp/ChangeLog:
2020-07-24 Martin Jambor <mjambor@suse.cz>
* plugin/Makefrag.am: Remove configuration of HSA plugin.
* aclocal.m4: Regenerated.
* Makefile.in: Regenerated.
* config.h.in: Regenerated.
* configure: Regenerated.
* plugin/configfrag.ac: Likewise.
* plugin/hsa_ext_finalize.h: Removed.
* plugin/plugin-hsa.c: Likewise.
* testsuite/Makefile.in: Regenerated.
* testsuite/lib/libgomp.exp
(offload_target_to_openacc_device_type): Remove hsa case.
(check_effective_target_hsa_offloading_selected_nocache): Removed
(check_effective_target_hsa_offloading_selected): Likewise.
(libgomp_init): Do not add -Wno-hsa to additional_flags.
* testsuite/libgomp.hsa.c/alloca-1.c: Removed test.
* testsuite/libgomp.hsa.c/bitfield-1.c: Likewise.
* testsuite/libgomp.hsa.c/bits-insns.c: Likewise.
* testsuite/libgomp.hsa.c/builtins-1.c: Likewise.
* testsuite/libgomp.hsa.c/c.exp: Likewise.
* testsuite/libgomp.hsa.c/complex-1.c: Likewise.
* testsuite/libgomp.hsa.c/complex-align-2.c: Likewise.
* testsuite/libgomp.hsa.c/formal-actual-args-1.c: Likewise.
* testsuite/libgomp.hsa.c/function-call-1.c: Likewise.
* testsuite/libgomp.hsa.c/get-level-1.c: Likewise.
* testsuite/libgomp.hsa.c/gridify-1.c: Likewise.
* testsuite/libgomp.hsa.c/gridify-2.c: Likewise.
* testsuite/libgomp.hsa.c/gridify-3.c: Likewise.
* testsuite/libgomp.hsa.c/gridify-4.c: Likewise.
* testsuite/libgomp.hsa.c/memory-operations-1.c: Likewise.
* testsuite/libgomp.hsa.c/pr69568.c: Likewise.
* testsuite/libgomp.hsa.c/pr82416.c: Likewise.
* testsuite/libgomp.hsa.c/rotate-1.c: Likewise.
* testsuite/libgomp.hsa.c/staticvar.c: Likewise.
* testsuite/libgomp.hsa.c/switch-1.c: Likewise.
* testsuite/libgomp.hsa.c/switch-branch-1.c: Likewise.
* testsuite/libgomp.hsa.c/switch-sbr-2.c: Likewise.
* testsuite/libgomp.hsa.c/tiling-1.c: Likewise.
* testsuite/libgomp.hsa.c/tiling-2.c: Likewise.
gcc/testsuite/ChangeLog:
2020-07-24 Martin Jambor <mjambor@suse.cz>
* lib/target-supports.exp (check_effective_target_offload_hsa):
Removed.
* c-c++-common/gomp/gridify-1.c: Removed test.
* c-c++-common/gomp/gridify-2.c: Likewise.
* c-c++-common/gomp/gridify-3.c: Likewise.
* c-c++-common/gomp/hsa-indirect-call-1.c: Likewise.
* gfortran.dg/gomp/gridify-1.f90: Likewise.
* gcc.dg/gomp/gomp.exp: Do not pass -Wno-hsa to tests.
* g++.dg/gomp/gomp.exp: Likewise.
* gfortran.dg/gomp/gomp.exp: Likewise.
Diffstat (limited to 'libgomp/testsuite')
26 files changed, 4 insertions, 1880 deletions
diff --git a/libgomp/testsuite/Makefile.in b/libgomp/testsuite/Makefile.in index 58ba1f3..26e925b 100644 --- a/libgomp/testsuite/Makefile.in +++ b/libgomp/testsuite/Makefile.in @@ -1,7 +1,7 @@ -# Makefile.in generated by automake 1.15.1 from Makefile.am. +# Makefile.in generated by automake 1.16.1 from Makefile.am. # @configure_input@ -# Copyright (C) 1994-2017 Free Software Foundation, Inc. +# Copyright (C) 1994-2018 Free Software Foundation, Inc. # This Makefile.in is free software; the Free Software Foundation # gives unlimited permission to copy and/or distribute it, @@ -215,10 +215,6 @@ PLUGIN_GCN = @PLUGIN_GCN@ PLUGIN_GCN_CPPFLAGS = @PLUGIN_GCN_CPPFLAGS@ PLUGIN_GCN_LDFLAGS = @PLUGIN_GCN_LDFLAGS@ PLUGIN_GCN_LIBS = @PLUGIN_GCN_LIBS@ -PLUGIN_HSA = @PLUGIN_HSA@ -PLUGIN_HSA_CPPFLAGS = @PLUGIN_HSA_CPPFLAGS@ -PLUGIN_HSA_LDFLAGS = @PLUGIN_HSA_LDFLAGS@ -PLUGIN_HSA_LIBS = @PLUGIN_HSA_LIBS@ PLUGIN_NVPTX = @PLUGIN_NVPTX@ PLUGIN_NVPTX_CPPFLAGS = @PLUGIN_NVPTX_CPPFLAGS@ PLUGIN_NVPTX_LDFLAGS = @PLUGIN_NVPTX_LDFLAGS@ @@ -335,8 +331,8 @@ Makefile: $(srcdir)/Makefile.in $(top_builddir)/config.status *config.status*) \ cd $(top_builddir) && $(MAKE) $(AM_MAKEFLAGS) am--refresh;; \ *) \ - echo ' cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@ $(am__depfiles_maybe)'; \ - cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@ $(am__depfiles_maybe);; \ + echo ' cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@ $(am__maybe_remake_depfiles)'; \ + cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@ $(am__maybe_remake_depfiles);; \ esac; $(top_builddir)/config.status: $(top_srcdir)/configure $(CONFIG_STATUS_DEPENDENCIES) diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index 8ccb78f4..5d86e2a 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -233,9 +233,6 @@ proc libgomp_init { args } { # Disable caret lappend ALWAYS_CFLAGS "additional_flags=-fno-diagnostics-show-caret" - # Disable HSA warnings by default. - lappend ALWAYS_CFLAGS "additional_flags=-Wno-hsa" - # Disable color diagnostics lappend ALWAYS_CFLAGS "additional_flags=-fdiagnostics-color=never" @@ -325,9 +322,6 @@ proc offload_target_to_openacc_device_type { offload_target } { disable { return "host" } - hsa* { - return "" - } *-intelmic* { return "" } @@ -430,60 +424,6 @@ proc check_effective_target_openacc_host_selected { } { return [string match "host" $openacc_device_type] } -# Return 1 if the selected OMP device is actually a HSA device - -proc check_effective_target_hsa_offloading_selected_nocache {} { - global tool - - set src { - int main () { - int v = 1; - #pragma omp target map(from:v) - v = 0; - return v; - } - } - - set result [check_compile hsa_offloading_src executable $src] - set lines [lindex $result 0] - set exe [lindex $result 1] - - set ok 0 - if { [string match "" $lines] } { - # No error messages, let us switch on HSA debugging output and run it - set prev_HSA_DEBUG [getenv HSA_DEBUG] - setenv HSA_DEBUG "1" - set result [remote_load target "./$exe"] - if { [string match "" $prev_HSA_DEBUG] } { - unsetenv HSA_DEBUG - } else { - setenv HSA_DEBUG $prev_HSA_DEBUG - } - set status [lindex $result 0] - if { $status != "pass" } { - remote_file build delete $exe - verbose "HSA availability test failed" - return 0 - } - set output [lindex $result 1] - if { [string match "*HSA debug: Going to dispatch kernel*" $output] } { - verbose "HSA availability detected" - set ok 1 - } - } - remote_file build delete $exe - return $ok -} - -# Return 1 if the selected OMP device is actually a HSA device and -# cache the result - -proc check_effective_target_hsa_offloading_selected {} { - return [check_cached_effective_target hsa_offloading_selected { - check_effective_target_hsa_offloading_selected_nocache - }] -} - # Return 1 if at least one AMD GPU is accessible. proc check_effective_target_openacc_radeon_accel_present { } { diff --git a/libgomp/testsuite/libgomp.hsa.c/alloca-1.c b/libgomp/testsuite/libgomp.hsa.c/alloca-1.c deleted file mode 100644 index 48dca94..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/alloca-1.c +++ /dev/null @@ -1,25 +0,0 @@ -#define size 10 -int i, j, k; - -int -main () -{ - char *s = __builtin_malloc (size + 1); - -#pragma omp target teams - { -#pragma omp distribute parallel for default(none) private(i) shared(s) - for (i = 0; i < size; ++i) - { - char *buffer = __builtin_alloca (10); - buffer[5] = 97 + i; - s[i] = buffer[5]; - } - } - - for (i = 0; i < size; ++i) - if (s[i] != 97 + i) - __builtin_abort (); - - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/bitfield-1.c b/libgomp/testsuite/libgomp.hsa.c/bitfield-1.c deleted file mode 100644 index 4dbf348..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/bitfield-1.c +++ /dev/null @@ -1,160 +0,0 @@ -#include <assert.h> - -#define ASSIGN_SX(N) \ - s##N.a1 = 1; \ - s##N.a2 = 2; \ - s##N.a3 = 3; \ - s##N.a4 = 4; \ - s##N.a5 = 5; \ - s##N.a6 = 6; \ - s##N.a7 = 7; \ - s##N.a8 = 8; \ - s##N.a9 = 9; \ - s##N.a10 = 10; - -#define ASSERT_SX(N) \ - assert (s##N.a1 == 1); \ - assert (s##N.a2 == 2); \ - assert (s##N.a3 == 3); \ - assert (s##N.a4 == 4); \ - assert (s##N.a5 == 5); \ - assert (s##N.a6 == 6); \ - assert (s##N.a7 == 7); \ - assert (s##N.a8 == 8); \ - assert (s##N.a9 == 9); \ - assert (s##N.a10 == 10); - -struct S1 -{ - unsigned a : 10; - unsigned b : 20; -}; - -struct S2 -{ - unsigned a1 : 10; - unsigned a2 : 10; - unsigned a3 : 10; - unsigned a4 : 10; - unsigned a5 : 10; - unsigned a6 : 10; - unsigned a7 : 10; - unsigned a8 : 10; - unsigned a9 : 10; - unsigned a10 : 10; -}; - -struct S3 -{ - unsigned a1 : 10; - unsigned a2 : 9; - unsigned a3 : 8; - unsigned a4 : 7; - unsigned a5 : 6; - unsigned a6 : 5; - unsigned a7 : 6; - unsigned a8 : 7; - unsigned a9 : 8; - unsigned a10 : 9; -}; - -struct S4 -{ - unsigned a1 : 10; - int a2 : 9; - unsigned a3 : 8; - int a4 : 7; - unsigned a5 : 6; - int a6 : 5; - unsigned a7 : 6; - int a8 : 7; - unsigned a9 : 8; - int a10 : 9; -}; - -struct S5 -{ - unsigned a1 : 31; - int a2 : 9; - unsigned a3 : 17; - int a4 : 7; - unsigned a5 : 6; - int a6 : 5; - unsigned long a7 : 55; - int a8 : 7; - unsigned a9 : 8; - int a10 : 9; -}; - -int -main () -{ - struct S1 s1; - -#pragma omp target map(to: s1) - { - s1.a = 2; - s1.b = 3; - } - - assert (s1.a == 2); - assert (s1.b == 3); - - struct S2 s2; - -#pragma omp target map(to: s2) - { - ASSIGN_SX (2) - } - - ASSERT_SX (2) - - struct S3 s3; - -#pragma omp target map(to: s3) - { - ASSIGN_SX (3) - } - - ASSERT_SX (3) - - struct S4 s4; - -#pragma omp target map(to: s4) - { - ASSIGN_SX (4) - } - - ASSERT_SX (4) - - struct S4 s5; - - s5.a1 = 0; - s5.a2 = 1; - s5.a3 = 2; - s5.a4 = 3; - s5.a5 = 4; - s5.a6 = 5; - s5.a7 = 6; - s5.a8 = 7; - s5.a9 = 8; - s5.a10 = 9; - -#pragma omp target map(to: s5) - { - s5.a1++; - s5.a2++; - s5.a3++; - s5.a4++; - s5.a5++; - s5.a6++; - s5.a7++; - s5.a8++; - s5.a9++; - s5.a10++; - } - - ASSERT_SX (5) - - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/bits-insns.c b/libgomp/testsuite/libgomp.hsa.c/bits-insns.c deleted file mode 100644 index 21cac72..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/bits-insns.c +++ /dev/null @@ -1,73 +0,0 @@ -#include <math.h> - -#define N 12 - -int main() -{ - unsigned int arguments[N] = {0u, 1u, 2u, 3u, 111u, 333u, 444u, 0x80000000u, 0x0000ffffu, 0xf0000000u, 0xff000000u, 0xffffffffu}; - int clrsb[N] = {}; - int clz[N] = {}; - int ctz[N] = {}; - int ffs[N] = {}; - int parity[N] = {}; - int popcount[N] = {}; - - int ref_clrsb[N] = {}; - int ref_clz[N] = {}; - int ref_ctz[N] = {}; - int ref_ffs[N] = {}; - int ref_parity[N] = {}; - int ref_popcount[N] = {}; - - for (unsigned i = 0; i < N; i++) - { - ref_clrsb[i] = __builtin_clrsb (arguments[i]); - ref_clz[i] = __builtin_clz (arguments[i]); - ref_ctz[i] = __builtin_ctz (arguments[i]); - ref_ffs[i] = __builtin_ffs (arguments[i]); - ref_parity[i] = __builtin_parity (arguments[i]); - ref_popcount[i] = __builtin_popcount (arguments[i]); - } - - #pragma omp target map(from:clz, ctz, ffs, parity, popcount) - { - for (unsigned i = 0; i < N; i++) - { - clrsb[i] = __builtin_clrsb (arguments[i]); - clz[i] = __builtin_clz (arguments[i]); - ctz[i] = __builtin_ctz (arguments[i]); - ffs[i] = __builtin_ffs (arguments[i]); - parity[i] = __builtin_parity (arguments[i]); - popcount[i] = __builtin_popcount (arguments[i]); - } - } - - for (unsigned i = 0; i < N; i++) - if (ref_clrsb[i] != clrsb[i]) - __builtin_abort (); - - /* CLZ of zero is undefined for zero. */ - for (unsigned i = 1; i < N; i++) - if (ref_clz[i] != clz[i]) - __builtin_abort (); - - /* Likewise for ctz */ - for (unsigned i = 1; i < N; i++) - if (ref_ctz[i] != ctz[i]) - __builtin_abort (); - - for (unsigned i = 0; i < N; i++) - if (ref_ffs[i] != ffs[i]) - __builtin_abort (); - - for (unsigned i = 0; i < N; i++) - if (ref_parity[i] != parity[i]) - __builtin_abort (); - - for (unsigned i = 0; i < N; i++) - if (ref_popcount[i] != popcount[i]) - __builtin_abort (); - - return 0; -} - diff --git a/libgomp/testsuite/libgomp.hsa.c/builtins-1.c b/libgomp/testsuite/libgomp.hsa.c/builtins-1.c deleted file mode 100644 index e603c21..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/builtins-1.c +++ /dev/null @@ -1,97 +0,0 @@ -/* { dg-additional-options "-ffast-math" } */ - -#include <assert.h> -#include <math.h> - -#define N 10 -#define N2 14 - -#define c1 1.2345f -#define c2 1.2345 - -#define DELTA 0.001 - -#define TEST_BIT_BUILTINS(T, S, S2) \ - { \ - T arguments[N2] \ - = {0##S, 1##S, 2##S, 3##S, \ - 111##S, 333##S, 444##S, 0x80000000##S, \ - 0x0000ffff##S, 0xf0000000##S, 0xff000000##S, 0xffffffff##S}; \ - int clrsb[N2] = {}; \ - int clz[N2] = {}; \ - int ctz[N2] = {}; \ - int ffs[N2] = {}; \ - int parity[N2] = {}; \ - int popcount[N2] = {}; \ - \ - _Pragma ("omp target map(to:clz[:N2], ctz[:N2], ffs[:N2], parity[:N2], popcount[:N2])") \ - { \ - for (unsigned i = 0; i < N2; i++) \ - { \ - clrsb[i] = __builtin_clrsb##S2 (arguments[i]); \ - clz[i] = __builtin_clz##S2 (arguments[i]); \ - ctz[i] = __builtin_ctz##S2 (arguments[i]); \ - ffs[i] = __builtin_ffs##S2 (arguments[i]); \ - parity[i] = __builtin_parity##S2 (arguments[i]); \ - popcount[i] = __builtin_popcount##S2 (arguments[i]); \ - } \ - } \ - \ - for (unsigned i = 0; i < N2; i++) \ - { \ - assert (clrsb[i] == __builtin_clrsb##S2 (arguments[i])); \ - if (arguments[0] != 0) \ - { \ - assert (clz[i] == __builtin_clz##S2 (arguments[i])); \ - assert (ctz[i] == __builtin_ctz##S2 (arguments[i])); \ - } \ - assert (ffs[i] == __builtin_ffs##S2 (arguments[i])); \ - assert (parity[i] == __builtin_parity##S2 (arguments[i])); \ - assert (popcount[i] == __builtin_popcount##S2 (arguments[i])); \ - } \ - } - -#define ASSERT(v1, v2) assert (fabs (v1 - v2) < DELTA) - -int -main () -{ - float f[N] = {}; - float d[N] = {}; - -/* 1) test direct mapping to HSA insns. */ - -#pragma omp target map(to: f[ : N], d[ : N]) - { - f[0] = sinf (c1); - f[1] = cosf (c1); - f[2] = exp2f (c1); - f[3] = log2f (c1); - f[4] = truncf (c1); - f[5] = sqrtf (c1); - - d[0] = trunc (c2); - d[1] = sqrt (c2); - } - - ASSERT (f[0], sinf (c1)); - ASSERT (f[1], cosf (c1)); - ASSERT (f[2], exp2f (c1)); - ASSERT (f[3], log2f (c1)); - ASSERT (f[4], truncf (c1)); - ASSERT (f[5], sqrtf (c1)); - - ASSERT (d[0], trunc (c2)); - ASSERT (d[1], sqrt (c2)); - - /* 2) test bit builtins for unsigned int. */ - TEST_BIT_BUILTINS (int, , ); - - /* 3) test bit builtins for unsigned long int. */ - TEST_BIT_BUILTINS (long, l, l); - - /* 4) test bit builtins for unsigned long long int. */ - TEST_BIT_BUILTINS (long long, ll, ll); - - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/c.exp b/libgomp/testsuite/libgomp.hsa.c/c.exp deleted file mode 100644 index 4614192..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/c.exp +++ /dev/null @@ -1,42 +0,0 @@ -if [info exists lang_library_path] then { - unset lang_library_path - unset lang_link_flags -} -if [info exists lang_test_file] then { - unset lang_test_file -} -if [info exists lang_include_flags] then { - unset lang_include_flags -} - -load_lib libgomp-dg.exp -load_gcc_lib gcc-dg.exp - -# Initialize dg. -dg-init - -# Turn on OpenMP. -lappend ALWAYS_CFLAGS "additional_flags=-fopenmp" - -set ld_library_path $always_ld_library_path -append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] -set_ld_library_path_env_vars - -global DEFAULT_CFLAGS -if [info exists DEFAULT_CFLAGS] then { - set CFLAGS_list [list "-O0" $DEFAULT_CFLAGS] -} else { - set CFLAGS_list [list "-O0" "-O2"] -} - -if [check_effective_target_hsa_offloading_selected] { - foreach USE_CFLAGS $CFLAGS_list { - # Gather a list of all tests. - set tests [lsort [find $srcdir/$subdir *.c]] - # Main loop. - dg-runtest $tests "" [concat $USE_CFLAGS "-Whsa"] - } -} - -# All done. -dg-finish diff --git a/libgomp/testsuite/libgomp.hsa.c/complex-1.c b/libgomp/testsuite/libgomp.hsa.c/complex-1.c deleted file mode 100644 index 438c64a..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/complex-1.c +++ /dev/null @@ -1,65 +0,0 @@ -#include <assert.h> -#include <complex.h> -#include <math.h> - -#define uchar unsigned char -#define C 123 - -#define TEST(type) \ - type foo_##type (void) \ - { \ - _Complex type a = C + 45I; \ - return __real__ a; \ - } - -#pragma omp declare target -TEST (char) -TEST (uchar) -TEST (short) -TEST (int) - -float -bar (float a, float b) -{ - _Complex float c = a + b * I; - - c += 11.f + 12.f * I; - - _Complex float d = 2.f + 4.44f * I; - - return __real__(crealf (c + d) + cimag (d) * I); -} - -#pragma omp end declare target - -int -main (void) -{ - int v = 0; - float v2 = 0.0f; - -#pragma omp target map(to: v) - v = foo_char (); - - assert (v == C); - -#pragma omp target map(to: v) - v = foo_uchar (); - - assert (v == C); - -#pragma omp target map(to: v) - v = foo_short (); - - assert (v == C); - -#pragma omp target map(to: v) - v = foo_int (); - - assert (v == C); - -#pragma omp target map(to: v2) - v2 = bar (1.12f, 4.44f); - - assert (fabs (v2 - 14.12) < 0.0001f); -} diff --git a/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c b/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c deleted file mode 100644 index b2d7acf..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c +++ /dev/null @@ -1,27 +0,0 @@ -#pragma omp declare target - _Complex int *g; -#pragma omp end declare target - - - -_Complex float f(void); - -int -main () -{ - _Complex int y; -#pragma omp target map(from:y) - { - _Complex int x; - g = &x; - __imag__ x = 1; - __real__ x = 2; - y = x; - } - - if ((__imag__ y != 1) - || (__real__ y != 2)) - __builtin_abort (); - return 0; -} - diff --git a/libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c b/libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c deleted file mode 100644 index 058a036..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/formal-actual-args-1.c +++ /dev/null @@ -1,83 +0,0 @@ -#include <assert.h> - -struct Cube -{ - int x; - int y; - int z; -}; - -#pragma omp declare target -int -foo (short a) -{ - switch (a) - { - case 1: - return 11; - break; - case 33: - return 333; - break; - case 55: - return 55; - break; - default: - return -1; - } -} - -int -bar (int a) -{ - int *ptr = &a; - - *ptr = 100; - return a + *ptr; -} - -struct Cube -baz (struct Cube c) -{ - c.x = 11; - return c; -} - -#pragma omp end declare target - -#define s 100 - -int -main (int argc) -{ - /* Test 1: argument types: char to short. */ - - int array[s]; -#pragma omp target map(tofrom : array[ : s]) - { - for (char i = 0; i < s; i++) - array[i] = foo (i); - } - - for (int i = 0; i < s; i++) - assert (array[i] == foo (i)); - - /* Test 2: argument address is taken. */ - int v = 2; - -#pragma omp target map(tofrom : v) - v = bar (v); - - assert (v == 200); - - /* Test 3: passing a structure as a function argument. */ - struct Cube r; - struct Cube c = {.x = 1, .y = 2, .z = 3}; - -#pragma omp target map(to : r) map(from : c) - r = baz (c); - - assert (r.x == 11); - assert (r.y == c.y); - assert (r.z == c.z); -} diff --git a/libgomp/testsuite/libgomp.hsa.c/function-call-1.c b/libgomp/testsuite/libgomp.hsa.c/function-call-1.c deleted file mode 100644 index 7f15dff..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/function-call-1.c +++ /dev/null @@ -1,50 +0,0 @@ -#define size 8 - -#pragma omp declare target -int -identity (int x) -{ - return x; -} - -int -expx (int x, int n) -{ - for (int i = 0; i < n - 1; i++) - x *= x; - - return x; -} - -float -init (int x, int y) -{ - int x1 = identity (identity (identity (identity (x)))); - int y1 = identity (identity (identity (identity (y)))); - - int x2 = expx (x1, 2); - int y2 = expx (y1, 2); - - return (x2 + y2); -} -#pragma omp end declare target - -int -main () -{ - int i, j; - int a[size][size]; - -#pragma omp target teams map(to:a[:size][:size]) -#pragma omp distribute parallel for default(none) private(i, j) shared(a) - for (i = 0; i < size; ++i) - for (j = 0; j < size; ++j) - a[i][j] = init (i, j); - - for (i = 0; i < size; ++i) - for (j = 0; j < size; ++j) - if (i * i + j * j != a[i][j]) - __builtin_abort (); - - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/get-level-1.c b/libgomp/testsuite/libgomp.hsa.c/get-level-1.c deleted file mode 100644 index 81c9df0..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/get-level-1.c +++ /dev/null @@ -1,26 +0,0 @@ -#include <omp.h> - -int -main () -{ - int i; - int level = -1; - -#pragma omp target map(tofrom : level) - { - level = omp_get_level (); - } - - if (level != 0) - __builtin_abort (); - -#pragma omp target teams map(tofrom : level) -#pragma omp distribute parallel for default(none) private(i) shared(level) - for (i = 0; i < 1; ++i) - level += omp_get_level (); - - if (level != 1) - __builtin_abort (); - - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/gridify-1.c b/libgomp/testsuite/libgomp.hsa.c/gridify-1.c deleted file mode 100644 index b670b9b..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/gridify-1.c +++ /dev/null @@ -1,26 +0,0 @@ -void __attribute__((noinline, noclone)) -foo (int n, int *a, int workgroup_size) -{ - int i; -#pragma omp target -#pragma omp teams thread_limit(workgroup_size) -#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) - for (i = 0; i < n; i++) - a[i]++; -} - -int main (int argc, char **argv) -{ - int n = 32; - int *a = __builtin_malloc (sizeof (int) * n); - int i; - - __builtin_memset (a, 0, sizeof (int) * n); - foo (n, a, 32); - for (i = 0; i < n; i ++) - { - if (a[i] != 1) - __builtin_abort (); - } - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/gridify-2.c b/libgomp/testsuite/libgomp.hsa.c/gridify-2.c deleted file mode 100644 index 3692eb0..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/gridify-2.c +++ /dev/null @@ -1,26 +0,0 @@ -void __attribute__((noinline, noclone)) -foo (int j, int n, int *a) -{ - int i; -#pragma omp target -#pragma omp teams -#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) firstprivate(j) - for (i = j + 1; i < n; i++) - a[i] = i; -} - -int main (int argc, char **argv) -{ - int n = 32; - int *a = __builtin_malloc (sizeof (int) * n); - int i, j = 4; - - __builtin_memset (a, 0, sizeof (int) * n); - foo (j, n, a); - for (i = j + 1; i < n; i ++) - { - if (a[i] != i) - __builtin_abort (); - } - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/gridify-3.c b/libgomp/testsuite/libgomp.hsa.c/gridify-3.c deleted file mode 100644 index f881d81..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/gridify-3.c +++ /dev/null @@ -1,39 +0,0 @@ -#define THE_LOOP \ - for (i = j + 1; i < n; i += 3) \ - a[i] = i - -void __attribute__((noinline, noclone)) -foo (int j, int n, int *a) -{ - int i; -#pragma omp target -#pragma omp teams -#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) firstprivate(j) - THE_LOOP; -} - -void __attribute__((noinline, noclone)) -bar (int j, int n, int *a) -{ - int i; - THE_LOOP; -} - -int main (int argc, char **argv) -{ - int n = 32; - int *a = __builtin_malloc (sizeof (int) * n); - int *ref = __builtin_malloc (sizeof (int) * n); - int i, j = 4; - - __builtin_memset (a, 0, sizeof (int) * n); - __builtin_memset (ref, 0, sizeof (int) * n); - bar (j, n, ref); - foo (j, n, a); - for (i = 0; i < n; i ++) - { - if (a[i] != ref[i]) - __builtin_abort (); - } - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/gridify-4.c b/libgomp/testsuite/libgomp.hsa.c/gridify-4.c deleted file mode 100644 index c3fbdbf..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/gridify-4.c +++ /dev/null @@ -1,45 +0,0 @@ -#define THE_LOOP \ - for (i = j + 1; i < n; i += 3) \ - a[i] = i - -void __attribute__((noinline, noclone)) -foo (int j, int n, int *a) -{ -#pragma omp parallel - { - #pragma omp single - { - int i; -#pragma omp target -#pragma omp teams -#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) firstprivate(j) - THE_LOOP; - } - } -} - -void __attribute__((noinline, noclone)) -bar (int j, int n, int *a) -{ - int i; - THE_LOOP; -} - -int main (int argc, char **argv) -{ - int n = 32; - int *a = __builtin_malloc (sizeof (int) * n); - int *ref = __builtin_malloc (sizeof (int) * n); - int i, j = 4; - - __builtin_memset (a, 0, sizeof (int) * n); - __builtin_memset (ref, 0, sizeof (int) * n); - bar (j, n, ref); - foo (j, n, a); - for (i = 0; i < n; i ++) - { - if (a[i] != ref[i]) - __builtin_abort (); - } - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c b/libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c deleted file mode 100644 index a17be93..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/memory-operations-1.c +++ /dev/null @@ -1,92 +0,0 @@ -#include <assert.h> - -#define C 55 - -int i, j, k; - -static void -test_bzero (unsigned size) -{ - unsigned bsize = size * sizeof (int); - int *x = __builtin_malloc (bsize); - __builtin_memset (x, C, bsize); - -#pragma omp target map(tofrom: x[:size]) map(from: bsize) - { - __builtin_bzero (x, bsize); - } - - char *buffer = (char *) x; - for (unsigned i = 0; i < bsize; ++i) - assert (buffer[i] == 0); -} - -static void -test_memcpy (unsigned size) -{ - unsigned bsize = size * sizeof (int); - int *x = __builtin_malloc (bsize); - __builtin_memset (x, C, bsize); - int *y = __builtin_malloc (bsize); - -#pragma omp target map(tofrom: x[:size], y[:size]) map(from: bsize) - { - __builtin_memcpy (y, x, bsize); - } - - char *buffer = (char *) y; - for (unsigned i = 0; i < bsize; ++i) - assert (buffer[i] == C); -} - -static void -test_mempcpy (unsigned size) -{ - unsigned bsize = size * sizeof (int); - int *x = __builtin_malloc (bsize); - __builtin_memset (x, C, bsize); - int *y = __builtin_malloc (bsize); - int *ptr = 0; - -#pragma omp target map(tofrom :x[:size], y[:size], ptr) map(from: bsize) - { - ptr = __builtin_mempcpy (y, x, bsize); - } - - char *buffer = (char *) y; - for (unsigned i = 0; i < bsize; ++i) - assert (buffer[i] == C); - - assert (ptr == y + size); -} - -static void -test_memset (unsigned size) -{ - unsigned bsize = size * sizeof (int); - int *x = __builtin_malloc (bsize); - __builtin_bzero (x, bsize); - -#pragma omp target map(tofrom : x[:size]) map(from: bsize) - { - __builtin_memset (x, C, bsize); - } - - char *buffer = (char *) x; - for (unsigned i = 0; i < bsize; ++i) - assert (buffer[i] == C); -} - -int -main (void) -{ - unsigned tests[] = {1, 2, 3, 4, 5, 8, 15, 17, 23, 33, 0}; - - for (unsigned i = 0; tests[i]; i++) - { - test_bzero (tests[i]); - test_memset (tests[i]); - test_memcpy (tests[i]); - test_mempcpy (tests[i]); - } -} diff --git a/libgomp/testsuite/libgomp.hsa.c/pr69568.c b/libgomp/testsuite/libgomp.hsa.c/pr69568.c deleted file mode 100644 index 6262eee..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/pr69568.c +++ /dev/null @@ -1,41 +0,0 @@ -/* PR hsa/69568 */ - -typedef float float2 __attribute__ ((vector_size (8))); -float2 *output; - -void __attribute__((noinline, noclone)) -foo (int n, float2 *a, int workgroup_size) -{ - int i; -#pragma omp target map(from:a[:n]) firstprivate(n, workgroup_size) -#pragma omp teams thread_limit(workgroup_size) -#pragma omp distribute parallel for shared(a) firstprivate(n) private(i) - for (i = 0; i < n; i++) - { float2 v; - v[0] = i; - v[1] = 1+i; - a[i] = v; - } -} - -int main (int argc, char **argv) -{ - int n = 32; - float2 *a = __builtin_malloc (sizeof (float2) * n); - int i; - - __builtin_memset (a, 0, sizeof (float2) * n); - foo (n, a, 32); - for (i = 0; i < n; i++) - { - float2 v = a[i]; - if (__builtin_abs (v[0] - i) > 0.1 - || __builtin_abs (v[1] - i - 1) > 0.1) - { - __builtin_abort (); - return 1; - } - } - return 0; -} - diff --git a/libgomp/testsuite/libgomp.hsa.c/pr82416.c b/libgomp/testsuite/libgomp.hsa.c/pr82416.c deleted file mode 100644 index 40378ab..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/pr82416.c +++ /dev/null @@ -1,43 +0,0 @@ -char __attribute__ ((noipa)) -toup (char X) -{ - if (X >= 97 && X <= 122) - return X - 32; - else - return X; -} - -char -target_toup_1 (char X) -{ - char r; -#pragma omp target map(to:X) map(from:r) - { - if (X >= 97 && X <= 122) - r = X - 32; - else - r = X; - } - return r; -} - -char __attribute__ ((noipa)) -target_toup (char X) -{ - return target_toup_1 (X); -} - -int main (int argc, char **argv) -{ - char a = 'a'; - if (toup (a) != target_toup (a)) - __builtin_abort (); - a = 'Z'; - if (toup (a) != target_toup (a)) - __builtin_abort (); - a = 5; - if (toup (a) != target_toup (a)) - __builtin_abort (); - - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/rotate-1.c b/libgomp/testsuite/libgomp.hsa.c/rotate-1.c deleted file mode 100644 index 494388b..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/rotate-1.c +++ /dev/null @@ -1,39 +0,0 @@ -#include <assert.h> -#include <limits.h> - -#define T unsigned int -#define BITSIZE CHAR_BIT * sizeof (T) - -#define C1 123u - -#pragma omp declare target -T -rotate (T value, T shift) -{ - T r = (value << shift) | (value >> (BITSIZE - shift)); - return (r >> shift) | (r << (BITSIZE - shift)); -} -#pragma omp end declare target - -int -main (int argc) -{ - T v1, v2, v3, v4, v5; - -#pragma omp target map(to: v1, v2, v3, v4, v5) - { - v1 = rotate (C1, 10); - v2 = rotate (C1, 2); - v3 = rotate (C1, 5); - v4 = rotate (C1, 16); - v5 = rotate (C1, 32); - } - - assert (v1 == C1); - assert (v2 == C1); - assert (v3 == C1); - assert (v4 == C1); - assert (v5 == C1); - - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/staticvar.c b/libgomp/testsuite/libgomp.hsa.c/staticvar.c deleted file mode 100644 index 6d20c9a..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/staticvar.c +++ /dev/null @@ -1,23 +0,0 @@ -extern void abort (void); - -#pragma omp declare target -int -foo (void) -{ - static int s; - return ++s; -} -#pragma omp end declare target - -int -main () -{ - int r; - #pragma omp target map(from:r) - { - r = foo (); - } - if (r != 1) - abort (); - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/switch-1.c b/libgomp/testsuite/libgomp.hsa.c/switch-1.c deleted file mode 100644 index a180cf6..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/switch-1.c +++ /dev/null @@ -1,145 +0,0 @@ -#include <assert.h> - -#define s 100 - -#pragma omp declare target -int -switch1 (int a) -{ - switch (a) - { - case 1: - return 11; - case 33: - return 333; - case 55: - return 55; - default: - return -1; - } -} - -int -switch2 (int a) -{ - switch (a) - { - case 1 ... 11: - return 11; - break; - case 33: - return 333; - break; - case 55: - return 55; - break; - default: - return -1; - } -} - -int -switch3 (int a) -{ - switch (a) - { - case 1 ... 11: - return 11; - case 12 ... 22: - return 22; - case 23 ... 33: - return 33; - case 34 ... 44: - return 44; - default: - return 44; - } -} - -int -switch4 (int a, int b) -{ - switch (a) - { - case 1 ... 11: - return a; - case 12 ... 22: - return b; - case 23 ... 33: - return a; - case 34 ... 44: - return b; - default: - return 12345; - } -} - -int -switch5 (int a, int b) -{ - switch (a) - { - case 1 ... 2: - return 1; - case 3 ... 4: - return 2; - case 5 ... 6: - return 3; - case 7 ... 11: - return 4; - } - - return -1; -} -#pragma omp end declare target - -int -main (int argc) -{ - int array[s]; - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = 0; i < s; i++) - array[i] = switch1 (i); - } - - for (int i = 0; i < s; i++) - assert (array[i] == switch1 (i)); - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = 0; i < s; i++) - array[i] = switch2 (i); - } - - for (int i = 0; i < s; i++) - assert (array[i] == switch2 (i)); - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = 0; i < s; i++) - array[i] = switch3 (i); - } - - for (int i = 0; i < s; i++) - assert (array[i] == switch3 (i)); - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = 0; i < s; i++) - array[i] = switch4 (i, i + 1); - } - - for (int i = 0; i < s; i++) - assert (array[i] == switch4 (i, i + 1)); - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = 0; i < s; i++) - array[i] = switch5 (i, i + 1); - } - - for (int i = 0; i < s; i++) - assert (array[i] == switch5 (i, i + 1)); -} diff --git a/libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c b/libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c deleted file mode 100644 index 9af1d6d..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/switch-branch-1.c +++ /dev/null @@ -1,116 +0,0 @@ -#include <assert.h> - -#define s 100 - -#pragma omp declare target -int -switch1 (unsigned a) -{ - switch (a) - { - case 1 ... 11: - return 11; - case 12 ... 13: - return 22; - default: - return 44; - } -} - -int -switch2 (unsigned a) -{ - switch (a) - { - case 1 ... 5: - return 1; - case 9 ... 11: - return a + 3; - case 12 ... 13: - return a + 3; - default: - return 44; - } -} - -#define OFFSET 12 - -int -switch3 (unsigned a) -{ - switch (a) - { - case (OFFSET + 0): - return 1; - case (OFFSET + 1)...(OFFSET + 11): - return 11; - case (OFFSET + 12)...(OFFSET + 13): - return (OFFSET + 22); - default: - return (OFFSET + 44); - } -} - -int -switch4 (unsigned a) -{ - switch (a) - { - case -2: - return 1; - case -1: - return a + 3; - case 3: - return a + 3; - default: - return 44; - } -} -#pragma omp end declare target - -#define low -33 -#define high 55 - -int -main (int argc) -{ - int array[s]; - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = low; i < high; i++) - array[i - low] = switch1 (i); - } - - for (int i = low; i < high; i++) - assert (array[i - low] == switch1 (i)); - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = low; i < high; i++) - array[i - low] = switch2 (i); - } - - for (int i = low; i < high; i++) - assert (array[i - low] == switch2 (i)); - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = low; i < high; i++) - array[i - low] = switch3 (i); - } - - for (int i = low; i < high; i++) - assert (array[i - low] == switch3 (i)); - -#pragma omp target map(tofrom : array[:s]) - { - for (int i = low; i < high; i++) - array[i - low] = switch4 (i); - } - - for (int i = low; i < high; i++) - assert (array[i - low] == switch4 (i)); - - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c b/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c deleted file mode 100644 index 06990d1..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c +++ /dev/null @@ -1,59 +0,0 @@ -/* { dg-additional-options "-fno-tree-switch-conversion" } */ - -#pragma omp declare target -int -foo (unsigned a) -{ - switch (a) - { - case 1 ... 5: - return 1; - case 9 ... 11: - return a + 3; - case 12 ... 13: - return a + 3; - default: - return 44; - } -} -#pragma omp end declare target - -#define s 100 - -void __attribute__((noinline, noclone)) -verify(int *a) -{ - if (a[0] != 44) - __builtin_abort (); - - for (int i = 1; i <= 5; i++) - if (a[i] != 1) - __builtin_abort (); - - for (int i = 6; i <= 8; i++) - if (a[i] != 44) - __builtin_abort (); - - for (int i = 9; i <= 13; i++) - if (a[i] != i + 3) - __builtin_abort (); - - for (int i = 14; i < s; i++) - if (a[i] != 44) - __builtin_abort (); -} - -int main(int argc) -{ - int array[s]; -#pragma omp target - { - for (int i = 0; i < s; i++) - { - int v = foo (i); - array[i] = v; - } - } - verify (array); - return 0; -} diff --git a/libgomp/testsuite/libgomp.hsa.c/tiling-1.c b/libgomp/testsuite/libgomp.hsa.c/tiling-1.c deleted file mode 100644 index 9149adc..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/tiling-1.c +++ /dev/null @@ -1,212 +0,0 @@ -/* - - matmul.c : Matrix Multiplication with tiling for openmp4 example - -*/ - -#include <stdlib.h> -#include <math.h> - -#define BLOCK_SIZE 16 -/* - #define BLOCK_SIZE 32 -*/ -#define NSECPERSEC 1000000000L - -typedef struct { - int width; - int height; - int stride; - int hpad; - float* elements; -} Matrix; - -/* Correctly extract the number of nanoseconds from the two time structures */ -long int get_nanosecs( struct timespec start_time, struct timespec end_time) { - long int nanosecs; - if ((end_time.tv_nsec-start_time.tv_nsec)<0) nanosecs = - ((((long int) end_time.tv_sec- (long int) start_time.tv_sec )-1)*NSECPERSEC ) + - ( NSECPERSEC + (long int) end_time.tv_nsec - (long int) start_time.tv_nsec) ; - else nanosecs = - (((long int) end_time.tv_sec- (long int) start_time.tv_sec )*NSECPERSEC ) + - ( (long int) end_time.tv_nsec - (long int) start_time.tv_nsec ); - return nanosecs; -} - -void simple_sgemm_tt(const int M,const int N,const int K,const float alpha, const float* A,const int LDA, - const float* B,const int LDB, const float beta,float* C, const int LDC) ; -void simple_sgemm_tn(const int M,const int N,const int K,const float alpha, const float* A,const int LDA, - const float* B,const int LDB, const float beta,float* C, const int LDC) ; -void tiled_sgemm_tt(const int M,const int N,const int K,const float alpha, const float*A, const int LDA, - const float* B,const int LDB, const float beta,float* C, const int LDC) ; - -int verify(float* v_res, float* v_ref, int len) { - int passed = 1; - int i; - for (i = 0; i < len; ++i) { - if (fabs(v_res[i] - v_ref[i]) > 0.001*v_ref[i]) { - __builtin_abort (); - } - } - return passed; -} - - -int main(int argc, char* argv[]){ - - Matrix A,B,Bt,C,Cref; - int a1,a2,a3,i,j; - struct timespec start_time1, end_time1; - struct timespec start_time2, end_time2; - long int nanosecs,total_ops; - float gflopsTiled,gflopsCPU; - - a1 = 35; - a2 = 28; - a3 = 47; - - A.height = a1; - A.width = a2; - A.stride = (((A.width-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - A.hpad = (((A.height-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - A.elements = (float*)malloc(A.stride * A.hpad* sizeof(float)); - - B.height = a2; - B.width = a3; - B.stride = (((B.width-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - B.hpad = (((B.height-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - B.elements = (float*)malloc(B.stride * B.hpad * sizeof(float)); - - /* Bt is same as B but stored in column-major order */ - Bt.height = B.height; - Bt.width = B.width; - Bt.stride = B.stride; - Bt.hpad = B.hpad; - Bt.elements = (float*)malloc(Bt.stride * Bt.hpad * sizeof(float)); - - C.height = a1; - C.width = a3; - C.stride = (((C.width-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - C.hpad = (((C.height-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - C.elements = (float*)malloc(C.stride * C.hpad * sizeof(float)); - - Cref.height = a1; - Cref.width = a3; - Cref.stride = (((Cref.width-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - Cref.hpad = (((Cref.height-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - Cref.elements = (float*)malloc(Cref.stride * Cref.hpad * sizeof(float)); - - for(i = 0; i < A.hpad ; i++) - for(j = 0; j < A.stride; j++) { - if (( j<A.width ) && (i<A.height)) { - A.elements[i*A.stride + j] = (i % 3); - } else { - A.elements[i*A.stride + j] = 0.0; - } - } - - /* Initialize B and Bt */ - for(i = 0; i < B.hpad ; i++) - for(j = 0; j < B.stride; j++) { - if (( j<B.width ) && (i<B.height)) { - B.elements[i*B.stride+j] = (j % 2); - Bt.elements[j*Bt.stride+i] = B.elements[i*B.stride+j] ; - } else { - B.elements[i*B.stride+j] = 0.0; - Bt.elements[j*Bt.stride+i] = 0.0; - } - } - - /* zero C, and Cref */ - for(i = 0; i < C.hpad; i++) - for(j = 0; j < C.stride; j++) { - C.elements[i*C.stride+j] = 0.0; - Cref.elements[i*Cref.stride+j] = 0.0; - } - - simple_sgemm_tt(A.height,B.width,B.height,1.0,A.elements,A.stride,B.elements,B.stride,1.0,Cref.elements,Cref.stride); - tiled_sgemm_tt(A.height,B.width,B.height,1.0,A.elements,A.stride,B.elements,B.stride,1.0,C.elements,C.stride); - - verify(C.elements, Cref.elements, C.height * C.stride); - return 0; -} - -void simple_sgemm_tt(const int M,const int N,const int K,const float alpha, const float* A,const int LDA, -const float* B,const int LDB, const float beta,float* C, const int LDC) { - /* A,B, and C are in row-major order */ - int c_row,c_col,inner; - float sum; - for (c_col = 0 ; c_col<N; c_col++ ) { - for (c_row = 0 ; c_row<M; c_row++ ) { - sum = 0.0 ; - for (inner = 0 ; inner<K; inner++ ) { - sum += A[c_row*LDA + inner] * B[inner*LDB + c_col] ; - } - C[c_row*LDC + c_col] = alpha*sum + beta*C[ c_row*LDC + c_col] ; - } - } -} - -/*************************** - - tiled_sgemm_tt: Tiled matrix multiplication: - -***************************/ - -void tiled_sgemm_tt(const int M, const int N, const int K, const float alpha, const float*A, const int LDA, - const float*B, const int LDB, const float beta, float*C, const int LDC){ - -#pragma omp target teams map(to:A[M*K],B[K*N]) map(from:C[M*N]) -#pragma omp distribute collapse(2) - for (int C_row_start=0 ; C_row_start < M ; C_row_start+=BLOCK_SIZE) - for (int C_col_start=0 ; C_col_start < N ; C_col_start+=BLOCK_SIZE) - { -// Each team has a local copy of these mini matrices - float As[BLOCK_SIZE][BLOCK_SIZE]; - float Bs[BLOCK_SIZE][BLOCK_SIZE]; -#pragma omp parallel - { - int C_row, C_col; - float Cval = 0.0; - - for (int kblock = 0; kblock < K ; kblock += BLOCK_SIZE ) - { -#pragma omp for collapse(2) - for (int row=0 ; row < BLOCK_SIZE ; row++) - for (int col=0 ; col < BLOCK_SIZE ; col++) - { - C_row = C_row_start + row; - C_col = C_col_start + col; - if ((C_row < M) && (kblock + col < K)) - As[row][col] = A[(C_row*LDA)+ kblock + col]; - else - As[row][col] = 0; - if ((kblock + row < K) && C_col < N) - Bs[row][col] = B[((kblock+row)*LDB)+ C_col]; - else - Bs[row][col] = 0; - } - -#pragma omp for collapse(2) - for (int row=0 ; row < BLOCK_SIZE ; row++) - for (int col=0 ; col < BLOCK_SIZE ; col++) - { - for (int e = 0; e < BLOCK_SIZE; ++e) - Cval += As[row][e] * Bs[e][col]; - } - } /* End for kblock .. */ - - -#pragma omp for collapse(2) - for (int row=0 ; row < BLOCK_SIZE ; row++) - for (int col=0 ; col < BLOCK_SIZE ; col++) - { - C_row = C_row_start + row; - C_col = C_col_start + col; - if ((C_row < M) && (C_col < N)) - C[(C_row*LDC)+C_col] = alpha*Cval + beta*C[(C_row*LDC)+C_col]; - - } - } /* end parallel */ - } /* end target teams distribute */ -} diff --git a/libgomp/testsuite/libgomp.hsa.c/tiling-2.c b/libgomp/testsuite/libgomp.hsa.c/tiling-2.c deleted file mode 100644 index 2756d14..0000000 --- a/libgomp/testsuite/libgomp.hsa.c/tiling-2.c +++ /dev/null @@ -1,258 +0,0 @@ -/* - - matmul.c : Matrix Multiplication with tiling for openmp4 example - -*/ - -#include <stdlib.h> -#include <math.h> - -#define BLOCK_SIZE 16 -/* - #define BLOCK_SIZE 32 -*/ -#define NSECPERSEC 1000000000L - -typedef struct { - int width; - int height; - int stride; - int hpad; - float* elements; -} Matrix; - -/* Correctly extract the number of nanoseconds from the two time structures */ -long int get_nanosecs( struct timespec start_time, struct timespec end_time) { - long int nanosecs; - if ((end_time.tv_nsec-start_time.tv_nsec)<0) nanosecs = - ((((long int) end_time.tv_sec- (long int) start_time.tv_sec )-1)*NSECPERSEC ) + - ( NSECPERSEC + (long int) end_time.tv_nsec - (long int) start_time.tv_nsec) ; - else nanosecs = - (((long int) end_time.tv_sec- (long int) start_time.tv_sec )*NSECPERSEC ) + - ( (long int) end_time.tv_nsec - (long int) start_time.tv_nsec ); - return nanosecs; -} - -void simple_sgemm_tt(const int M,const int N,const int K,const float alpha, const float* A,const int LDA, - const float* B,const int LDB, const float beta,float* C, const int LDC) ; -void simple_sgemm_tn(const int M,const int N,const int K,const float alpha, const float* A,const int LDA, - const float* B,const int LDB, const float beta,float* C, const int LDC) ; -void tiled_sgemm_tt(const int M,const int N,const int K,const float alpha, const float*A, const int LDA, - const float* B,const int LDB, const float beta,float* C, const int LDC) ; - -int verify(float* v_res, float* v_ref, int len) { - int passed = 1; - int i; - for (i = 0; i < len; ++i) { - if (fabs(v_res[i] - v_ref[i]) > 0.001*v_ref[i]) { - __builtin_abort (); - } - } - return passed; -} - - -int main(int argc, char* argv[]){ - - Matrix A,B,Bt,C,Cref; - int a1,a2,a3,i,j; - struct timespec start_time1, end_time1; - struct timespec start_time2, end_time2; - long int nanosecs,total_ops; - float gflopsTiled,gflopsCPU; - - a1 = 35; - a2 = 28; - a3 = 47; - - A.height = a1; - A.width = a2; - A.stride = (((A.width-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - A.hpad = (((A.height-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - A.elements = (float*)malloc(A.stride * A.hpad* sizeof(float)); - - B.height = a2; - B.width = a3; - B.stride = (((B.width-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - B.hpad = (((B.height-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - B.elements = (float*)malloc(B.stride * B.hpad * sizeof(float)); - - /* Bt is same as B but stored in column-major order */ - Bt.height = B.height; - Bt.width = B.width; - Bt.stride = B.stride; - Bt.hpad = B.hpad; - Bt.elements = (float*)malloc(Bt.stride * Bt.hpad * sizeof(float)); - - C.height = a1; - C.width = a3; - C.stride = (((C.width-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - C.hpad = (((C.height-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - C.elements = (float*)malloc(C.stride * C.hpad * sizeof(float)); - - Cref.height = a1; - Cref.width = a3; - Cref.stride = (((Cref.width-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - Cref.hpad = (((Cref.height-1)/BLOCK_SIZE)+1) * BLOCK_SIZE; - Cref.elements = (float*)malloc(Cref.stride * Cref.hpad * sizeof(float)); - - for(i = 0; i < A.hpad ; i++) - for(j = 0; j < A.stride; j++) { - if (( j<A.width ) && (i<A.height)) { - A.elements[i*A.stride + j] = (i % 3); - } else { - A.elements[i*A.stride + j] = 0.0; - } - } - - /* Initialize B and Bt */ - for(i = 0; i < B.hpad ; i++) - for(j = 0; j < B.stride; j++) { - if (( j<B.width ) && (i<B.height)) { - B.elements[i*B.stride+j] = (j % 2); - Bt.elements[j*Bt.stride+i] = B.elements[i*B.stride+j] ; - } else { - B.elements[i*B.stride+j] = 0.0; - Bt.elements[j*Bt.stride+i] = 0.0; - } - } - - /* zero C, and Cref */ - for(i = 0; i < C.hpad; i++) - for(j = 0; j < C.stride; j++) { - C.elements[i*C.stride+j] = 0.0; - Cref.elements[i*Cref.stride+j] = 0.0; - } - - simple_sgemm_tt(A.height,B.width,B.height,1.0,A.elements,A.stride,B.elements,B.stride,1.0,Cref.elements,Cref.stride); - tiled_sgemm_tt(A.height,B.width,B.height,1.0,A.elements,A.stride,B.elements,B.stride,1.0,C.elements,C.stride); - - verify(C.elements, Cref.elements, C.height * C.stride); - return 0; -} - -void simple_sgemm_tt(const int M,const int N,const int K,const float alpha, const float* A,const int LDA, -const float* B,const int LDB, const float beta,float* C, const int LDC) { - /* A,B, and C are in row-major order */ - int c_row,c_col,inner; - float sum; - for (c_col = 0 ; c_col<N; c_col++ ) { - for (c_row = 0 ; c_row<M; c_row++ ) { - sum = 0.0 ; - for (inner = 0 ; inner<K; inner++ ) { - sum += A[c_row*LDA + inner] * B[inner*LDB + c_col] ; - } - C[c_row*LDC + c_col] = alpha*sum + beta*C[ c_row*LDC + c_col] ; - } - } -} - -/*************************** - - tiled_sgemm_tt: Tiled matrix multiplication: - -***************************/ - -void tiled_sgemm_tt(const int M, const int N, const int K, const float alpha, const float*A, const int LDA, - const float*B, const int LDB, const float beta, float*C, const int LDC){ - -#pragma omp target teams map(to:A[M*K],B[K*N]) map(from:C[M*N]) -#pragma omp distribute collapse(2) - for (int C_row_start=0 ; C_row_start < M ; C_row_start+=BLOCK_SIZE) { - for (int C_col_start=0 ; C_col_start < N ; C_col_start+=BLOCK_SIZE) { - -// We now have M/BLOCK_SIZE * N/BLOCK_SIZE teams = (M*N)/(BLOCK_SIZE*BLOCK_SIZE) -// The grid global dimensions are M,N,1 -// The grid local dimensions are BLOCK_SIZE,BLOCK_SIZE,1 - -// ------------------------------------------------------------------- -// The rest of this code forms the HSAIL kernel with the -// pairs of "parallel for collapse(2)" loops replaced with a barrier. -// The kernel initializes these values -// C_row_start = get_group_id(0) * BLOCK_SIZE -// C_col_start = get_group_id(1) * BLOCK_SIZE -// row=get_local_id(0) -// col=get_local_id(1) -// ------------------------------------------------------------------- - -// Each team has a local copy of these mini matrices - float As[BLOCK_SIZE][BLOCK_SIZE]; - float Bs[BLOCK_SIZE][BLOCK_SIZE]; - float Cs[BLOCK_SIZE][BLOCK_SIZE]; - int C_row, C_col; - - /* Zero Cs for this BLOCK */ -// - - - - - - - - - - - - - - - - - - - - -// REPLACE NEXT THREE LINES WITH A BARRIER -#pragma omp parallel for collapse(2) - for (int row=0 ; row < BLOCK_SIZE ; row++) { - for (int col=0 ; col < BLOCK_SIZE ; col++) { -// END BARRIER -// - - - - - - - - - - - - - - - - - - - - - Cs[row][col] = 0.0; - } - } - - // This kblock loop is run on the master thread of each team - for (int kblock = 0; kblock < K ; kblock += BLOCK_SIZE ) { - - // Copy global memory values to local memory -// - - - - - - - - - - - - - - - - - - - - -// REPLACE NEXT THREE LINES WITH A BARRIER -#pragma omp parallel for collapse(2) - for (int row=0 ; row < BLOCK_SIZE ; row++) { - for (int col=0 ; col < BLOCK_SIZE ; col++) { -// END BARRIER -// - - - - - - - - - - - - - - - - - - - - - C_row = C_row_start + row; - C_col = C_col_start + col; - if ((C_row < M) && (kblock + col < K)) - As[row][col] = A[(C_row*LDA)+ kblock + col]; - else - As[row][col] = 0; - if ((kblock + row < K) && C_col < N) - Bs[row][col] = B[((kblock+row)*LDB)+ C_col]; - else - Bs[row][col] = 0; - } - } - - // Calculate Cs <- Sum(As X Bs) across all kblocks -// - - - - - - - - - - - - - - - - - - - - -// REPLACE NEXT THREE LINES WITH A BARRIER -#pragma omp parallel for collapse(2) - for (int row=0 ; row < BLOCK_SIZE ; row++) { - for (int col=0 ; col < BLOCK_SIZE ; col++) { -// END BARRIER -// - - - - - - - - - - - - - - - - - - - - - for (int e = 0; e < BLOCK_SIZE; ++e) - Cs[row][col] += As[row][e] * Bs[e][col]; - } - } - - } /* End for kblock .. */ - - - // Scale Update actual C from Cs -// - - - - - - - - - - - - - - - - - - - - -// REPLACE NEXT THREE LINES WITH A BARRIER -#pragma omp parallel for collapse(2) - for (int row=0 ; row < BLOCK_SIZE ; row++) { - for (int col=0 ; col < BLOCK_SIZE ; col++) { -// END BARRIER -// - - - - - - - - - - - - - - - - - - - - - C_row = C_row_start + row; - C_col = C_col_start + col; - if ((C_row < M) && (C_col < N)) { - C[(C_row*LDC)+C_col] = alpha*Cs[row][col] + beta*C[(C_row*LDC)+C_col]; - } - } - } - -// ------------------------------------------------------------------- -// This is the end of the kernel - - } - } - -} |