https://gcc.gnu.org/g:fb0dad07be32276b35bae92a7a3c51fcd3daa972
commit fb0dad07be32276b35bae92a7a3c51fcd3daa972
Author: Paul-Antoine Arras <par...@baylibre.com>
Date: Wed Nov 20 15:28:58 2024 +0100
OpenMP: common C/C++ testcases for dispatch + adjust_args
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/declare-variant-2.c: Adjust dg-error directives.
* c-c++-common/gomp/adjust-args-1.c: New test.
* c-c++-common/gomp/adjust-args-2.c: New test.
* c-c++-common/gomp/declare-variant-dup-match-clause.c: New test.
* c-c++-common/gomp/dispatch-1.c: New test.
* c-c++-common/gomp/dispatch-2.c: New test.
* c-c++-common/gomp/dispatch-3.c: New test.
* c-c++-common/gomp/dispatch-4.c: New test.
* c-c++-common/gomp/dispatch-5.c: New test.
* c-c++-common/gomp/dispatch-6.c: New test.
* c-c++-common/gomp/dispatch-7.c: New test.
* c-c++-common/gomp/dispatch-8.c: New test.
* c-c++-common/gomp/dispatch-9.c: New test.
* c-c++-common/gomp/dispatch-10.c: New test.
libgomp/ChangeLog:
* testsuite/libgomp.c-c++-common/dispatch-1.c: New test.
* testsuite/libgomp.c-c++-common/dispatch-2.c: New test.
(cherry picked from commit 377eff7c38e8df1388b0a1eeb6afdbcaf12c3551)
Diff:
---
gcc/testsuite/ChangeLog.omp | 20 ++++++
gcc/testsuite/c-c++-common/gomp/adjust-args-1.c | 30 ++++++++
gcc/testsuite/c-c++-common/gomp/adjust-args-2.c | 31 ++++++++
.../c-c++-common/gomp/declare-variant-2.c | 4 +-
.../gomp/declare-variant-dup-match-clause.c | 6 ++
gcc/testsuite/c-c++-common/gomp/dispatch-1.c | 71 ++++++++++++++++++
gcc/testsuite/c-c++-common/gomp/dispatch-10.c | 36 ++++++++++
gcc/testsuite/c-c++-common/gomp/dispatch-2.c | 28 ++++++++
gcc/testsuite/c-c++-common/gomp/dispatch-3.c | 12 ++++
gcc/testsuite/c-c++-common/gomp/dispatch-4.c | 18 +++++
gcc/testsuite/c-c++-common/gomp/dispatch-5.c | 34 +++++++++
gcc/testsuite/c-c++-common/gomp/dispatch-6.c | 18 +++++
gcc/testsuite/c-c++-common/gomp/dispatch-7.c | 21 ++++++
gcc/testsuite/c-c++-common/gomp/dispatch-8.c | 63 ++++++++++++++++
gcc/testsuite/c-c++-common/gomp/dispatch-9.c | 17 +++++
libgomp/ChangeLog.omp | 8 +++
.../testsuite/libgomp.c-c++-common/dispatch-1.c | 76 ++++++++++++++++++++
.../testsuite/libgomp.c-c++-common/dispatch-2.c | 84 ++++++++++++++++++++++
18 files changed, 575 insertions(+), 2 deletions(-)
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 2162805e5708..ccec4308bf22 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,23 @@
+2025-01-27 Paul-Antoine Arras <par...@baylibre.com>
+
+ Backported from master:
+ 2024-11-20 Paul-Antoine Arras <par...@baylibre.com>
+
+ * c-c++-common/gomp/declare-variant-2.c: Adjust dg-error directives.
+ * c-c++-common/gomp/adjust-args-1.c: New test.
+ * c-c++-common/gomp/adjust-args-2.c: New test.
+ * c-c++-common/gomp/declare-variant-dup-match-clause.c: New test.
+ * c-c++-common/gomp/dispatch-1.c: New test.
+ * c-c++-common/gomp/dispatch-2.c: New test.
+ * c-c++-common/gomp/dispatch-3.c: New test.
+ * c-c++-common/gomp/dispatch-4.c: New test.
+ * c-c++-common/gomp/dispatch-5.c: New test.
+ * c-c++-common/gomp/dispatch-6.c: New test.
+ * c-c++-common/gomp/dispatch-7.c: New test.
+ * c-c++-common/gomp/dispatch-8.c: New test.
+ * c-c++-common/gomp/dispatch-9.c: New test.
+ * c-c++-common/gomp/dispatch-10.c: New test.
+
2025-01-27 Paul-Antoine Arras <par...@baylibre.com>
Backported from master:
diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c
b/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c
new file mode 100644
index 000000000000..728abe620922
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c
@@ -0,0 +1,30 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+int f (int a, void *b, float c[2]);
+
+#pragma omp declare variant (f) match (construct={dispatch}) adjust_args
(nothing: a) adjust_args (need_device_ptr: b, c)
+int f0 (int a, void *b, float c[2]);
+#pragma omp declare variant (f) match (construct={dispatch}) adjust_args
(nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c)
+int f1 (int a, void *b, float c[2]);
+
+int test () {
+ int a;
+ void *b;
+ float c[2];
+ struct {int a;} s;
+
+ s.a = f0 (a, b, c);
+ #pragma omp dispatch
+ s.a = f0 (a, b, c);
+
+ f1 (a, b, c);
+ #pragma omp dispatch
+ s.a = f1 (a, b, c);
+
+ return s.a;
+}
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device
\\(\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ =
__builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ =
__builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c
b/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c
new file mode 100644
index 000000000000..d2a4a5f4ec45
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c
@@ -0,0 +1,31 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+int f (int a, void *b, float c[2]);
+
+#pragma omp declare variant (f) match (construct={dispatch}) adjust_args
(nothing: a) adjust_args (need_device_ptr: b, c)
+int f0 (int a, void *b, float c[2]);
+#pragma omp declare variant (f) adjust_args (need_device_ptr: b, c) match
(construct={dispatch}) adjust_args (nothing: a)
+int f1 (int a, void *b, float c[2]);
+
+void test () {
+ int a;
+ void *b;
+ float c[2];
+
+ #pragma omp dispatch
+ f0 (a, b, c);
+
+ #pragma omp dispatch device (-4852)
+ f0 (a, b, c);
+
+ #pragma omp dispatch device (a + a)
+ f0 (a, b, c);
+}
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device
\\(\\);" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ =
__builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ =
__builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ =
__builtin_omp_get_mapped_ptr \\(&c, -4852\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ =
__builtin_omp_get_mapped_ptr \\(b, -4852\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-not "#pragma omp dispatch device" "gimple" } }
*/
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c
b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c
index 43657486c996..27c1cdff0443 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c
@@ -8,9 +8,9 @@ void f3 (void);
void f4 (void);
#pragma omp declare variant match(user={condition(0)}) /* { dg-error "expected
'\\(' before 'match'" } */
void f5 (void);
-#pragma omp declare variant (f1) /* { dg-error "expected 'match' before
end of line" } */
+#pragma omp declare variant (f1) /* { dg-error "expected 'match' clause
before end of line" } */
void f6 (void);
-#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' before
'simd'" } */
+#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' clause
before 'simd'" } */
void f7 (void);
#pragma omp declare variant (f1) match /* { dg-error "expected '\\(' before
end of line" } */
void f8 (void);
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-dup-match-clause.c
b/gcc/testsuite/c-c++-common/gomp/declare-variant-dup-match-clause.c
new file mode 100644
index 000000000000..14dc681b59de
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-dup-match-clause.c
@@ -0,0 +1,6 @@
+/* Check that errors are detected if a match clause appears more than once,
+ even if compatible. */
+
+void f();
+#pragma omp declare variant(f) match(construct={target})
match(construct={target}) /* { dg-error "too many 'match' clauses" } */
+void g();
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-1.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-1.c
new file mode 100644
index 000000000000..2a4e939d0002
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-1.c
@@ -0,0 +1,71 @@
+#include <stdint.h>
+
+int f0 (int, long, double);
+void f2 (void);
+int f3 (void);
+void (*f4) (void);
+void f5 (int*, int);
+
+int *pp (void);
+int dd (void);
+
+void f1 (void)
+{
+ int a, c;
+ long b;
+ double x;
+ struct {int a; float b; short c;} s, *sp;
+ int arr[3];
+
+#pragma omp dispatch
+ c = f0 (a, b, x);
+#pragma omp dispatch
+ x = f0 (a * 4, 2 - b, x * x);
+#pragma omp dispatch
+ s.a = f0 (a, sp->c, x);
+#pragma omp dispatch
+ sp->c = f0 (s.a - 2, b / 3, x * 5);
+#pragma omp dispatch
+ arr[0] = f0 (arr[1], !b, arr[2]);
+#pragma omp dispatch
+ (*sp).c = f0 (s.a, b, x);
+#pragma omp dispatch
+ sp->b = f0 (s.a++, b % 4, --x);
+#pragma omp dispatch
+ f0 (f3(), b, s.b);
+#pragma omp dispatch
+ f2 ();
+#pragma omp dispatch
+ f4 ();
+#pragma omp dispatch
+ f5 (pp(), pp()[dd()]);
+
+#pragma omp dispatch nocontext(sp->a * x + arr[2])
+ f2 ();
+#pragma omp dispatch nocontext(arr - (intptr_t)(x / s.b))
+ f2 ();
+#pragma omp dispatch nocontext(x == s.c || b != c)
+ f2 ();
+#pragma omp dispatch novariants(b << sp->c)
+ f2 ();
+#pragma omp dispatch novariants(!arr | s.a)
+ f2 ();
+#pragma omp dispatch novariants(s.c ? f3() : a & c)
+ f2 ();
+#pragma omp dispatch nowait
+ f2 ();
+#pragma omp dispatch device(-25373654)
+ f2 ();
+#pragma omp dispatch device(b * (int)(x - sp->b))
+ f2 ();
+#pragma omp dispatch is_device_ptr(arr)
+ f2 ();
+#pragma omp dispatch is_device_ptr(sp)
+ f2 ();
+#pragma omp dispatch depend(inout: sp)
+ f2 ();
+#pragma omp dispatch depend(inoutset: arr[:2])
+ f2 ();
+#pragma omp dispatch depend(out: arr)
+ f2 ();
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-10.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-10.c
new file mode 100644
index 000000000000..40b3fc89a123
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-10.c
@@ -0,0 +1,36 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+int *f();
+
+struct t {
+ int *a, *b;
+};
+
+void construct(int *x, int *y);
+void noconstruct(int *x, int *y);
+
+#pragma omp declare variant(construct) match(construct={dispatch})
adjust_args(need_device_ptr: x,y)
+#pragma omp declare variant(noconstruct) match(implementation={vendor(gnu)})
+void bar(int *x, int *y);
+
+int nocontext, novariant;
+
+void sub(struct t *s, void *y)
+{
+ bar( f(), s->b);
+ #pragma omp dispatch device(0) is_device_ptr(s)
+ bar( f(), s->b);
+
+ bar( (int *) y, s->b);
+ #pragma omp dispatch device(0) is_device_ptr(y)
+ bar( (int *) y, s->b);
+}
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device" 2
"gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device" 4
"gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device
\\(0\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device
\\(D\\.\[0-9\]+\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_get_mapped_ptr" 3 "gimple"
} } */
+/* { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = s->b;\[\r\n\]*
*D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(D\\.\[0-9\]+, 0\\);" 2 "gimple"
} } */
+/* { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = f \\(\\);\[\r\n\]*
*D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(D\\.\[0-9\]+, 0\\);" 1 "gimple"
} } */
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-2.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-2.c
new file mode 100644
index 000000000000..24ab9545b73e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-2.c
@@ -0,0 +1,28 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+int f0 (void);
+int f1 (void);
+#pragma omp declare variant (f0) match (construct={dispatch})
+#pragma omp declare variant (f1) match (implementation={vendor(gnu)})
+int f2 (void);
+
+int test (void)
+{
+ int a;
+#pragma omp dispatch
+ a = f2 ();
+#pragma omp dispatch novariants(1)
+ a = f2 ();
+#pragma omp dispatch novariants(0)
+ a = f2 ();
+#pragma omp dispatch nocontext(1)
+ a = f2 ();
+#pragma omp dispatch nocontext(0)
+ a = f2 ();
+ return a;
+}
+
+/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-3.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c
new file mode 100644
index 000000000000..dba1d4ffdc8c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c
@@ -0,0 +1,12 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f2 (void);
+
+void test (void)
+{
+#pragma omp dispatch /* { dg-final { scan-tree-dump-not "#pragma omp task"
"gimple" } } */
+ f2 ();
+#pragma omp dispatch nowait /* { dg-final { scan-tree-dump-not "nowait"
"gimple" } } */
+ f2 ();
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-4.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-4.c
new file mode 100644
index 000000000000..ac746f755beb
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-4.c
@@ -0,0 +1,18 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f2 (int a);
+
+void test (void)
+{
+ int a;
+
+#pragma omp dispatch device(-25373654)
+/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device
\\(-25373654\\);" 1 "gimple" } } */
+ f2 (a);
+#pragma omp dispatch device(a + a)
+/* { dg-final { scan-tree-dump-times "(D\.\[0-9]+) = a \\* 2;.*#pragma omp
dispatch.*__builtin_omp_set_default_device \\(\\1\\);.*f2 \\(a\\)" 2 "gimple" }
} */
+ f2 (a);
+}
+
+/* { dg-final { scan-tree-dump-times "(D\.\[0-9]+) =
__builtin_omp_get_default_device \\(\\);.*__builtin_omp_set_default_device
\\(\\1\\);" 4 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-5.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c
new file mode 100644
index 000000000000..608c6d9f154e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c
@@ -0,0 +1,34 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f1 (void* p, int arr[]);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args
(need_device_ptr: p, arr)
+void f2 (void* p, int arr[]);
+
+void test (void)
+{
+ void *p;
+ int arr[2];
+
+/* Note there are multiple matches because every variable capturing matches in
addition,
+ i.e. scan-tree-dump-times = 1 plus number of captures used for backward
references.
+
+ For the first scan-tree-dump, on some targets the
__builtin_omp_get_mapped_ptr get
+ swapped. As the counting does not work well with variable-name capturing,
it only
+ uses 'scan-tree-dump' and not 'scan-tree-dump-times'. */
+
+#pragma omp dispatch
+/* { dg-final { scan-tree-dump "#pragma omp dispatch\[ \t\n\r\{]*int
(D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*void \\*
(D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[
\t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&arr, \\1\\);\[ \t\n\r]*\\3 =
__builtin_omp_get_mapped_ptr \\(p, \\1\\);\[ \t\n\r]*f1 \\(\\3, \\2\\);|#pragma
omp dispatch\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\* (D\.\[0-9]+);\[
\t\n\r]*void \\* (D\.\[0-9]+);\[ \t\n\r]*\\4 = __builtin_omp_get_default_device
\\(\\);\[ \t\n\r]*\\5 = __builtin_omp_get_mapped_ptr \\(p, \\4\\);\[
\t\n\r]*\\6 = __builtin_omp_get_mapped_ptr \\(&arr, \\4\\);\[ \t\n\r]*f1
\\(\\5, \\6\\);" "gimple" } } */
+ f2 (p, arr);
+#pragma omp dispatch is_device_ptr(p)
+/* { dg-final { scan-tree-dump-times "#pragma omp dispatch
is_device_ptr\\(p\\)\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\*
(D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[
\t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(&arr, \\1\\);\[ \t\n\r]*f1 \\(p,
\\2\\);" 3 "gimple" } } */
+ f2 (p, arr);
+#pragma omp dispatch is_device_ptr(arr)
+/* { dg-final { scan-tree-dump-times "#pragma omp dispatch
is_device_ptr\\(arr\\)\[ \t\n\r\{]*int (D\.\[0-9]+);\[ \t\n\r]*void \\*
(D\.\[0-9]+);\[ \t\n\r]*\\1 = __builtin_omp_get_default_device \\(\\);\[
\t\n\r]*\\2 = __builtin_omp_get_mapped_ptr \\(p, \\1\\);\[ \t\n\r]*f1 \\(\\2,
&arr\\);" 3 "gimple" } } */
+ f2 (p, arr);
+#pragma omp dispatch is_device_ptr(p, arr)
+/* { dg-final { scan-tree-dump-times "#pragma omp dispatch
is_device_ptr\\(arr\\) is_device_ptr\\(p\\)\[ \t\n\r\{]*f1 \\(p, &arr\\);" 1
"gimple" } } */
+ f2 (p, arr);
+}
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-6.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c
new file mode 100644
index 000000000000..4e900bb5e8da
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c
@@ -0,0 +1,18 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-ompexp" } */
+
+void f2 (void* p);
+
+void test (void)
+{
+ void *p;
+
+#pragma omp dispatch
+/* { dg-final { scan-tree-dump-not "__builtin_GOMP_task " "ompexp" } } */
+ f2 (p);
+#pragma omp dispatch depend(inout: p)
+/* { dg-final { scan-tree-dump-times "(D\.\[0-9]+)\\\[2] = &p;\[
\n]*__builtin_GOMP_taskwait_depend \\(&\\1\\);" 2 "ompexp" } } */
+ f2 (p);
+}
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-7.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c
new file mode 100644
index 000000000000..44a62608eb0b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+int f0 (void);
+int f1 (void);
+#pragma omp declare variant (f0) match (construct={dispatch})
+#pragma omp declare variant (f1) match (implementation={vendor(gnu)})
+int f2 (void);
+
+int test (void)
+{
+ int a, n;
+#pragma omp dispatch novariants(n < 1024) nocontext(n > 1024)
+ a = f2 ();
+ return a;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma omp dispatch nocontext\\(0\\)
novariants\\(0\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-8.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-8.c
new file mode 100644
index 000000000000..11713554a826
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-8.c
@@ -0,0 +1,63 @@
+int g1() { return 100; }
+int g2() { return 200; }
+
+#pragma omp declare variant (g1) match
(construct={dispatch},user={condition(1)})
+#pragma omp declare variant (g2) match (user={condition(1)})
+int g() { return 300; }
+
+
+int f1(int x) { return 1 + x; }
+int f2(int x) { return 2 + x; }
+
+#pragma omp declare variant (f1) match
(construct={dispatch},user={condition(1)})
+#pragma omp declare variant (f2) match (user={condition(1)})
+int f(int x) { return 3 + x; }
+
+
+void
+test ()
+{
+ int res;
+
+ // Call f2(g2()) due to user condition(1)
+ res = f (g ());
+ __builtin_printf("%d\n", res);
+ if (res != 202)
+ __builtin_abort ();
+
+ // Call 'f1(g1())' due to construct 'dispatch' and user conditional(1)
+ #pragma omp dispatch
+ res = f (g ());
+ __builtin_printf("%d\n", res);
+ if (res != 101)
+ __builtin_abort ();
+
+ // Call 'f2(g2())' due to nocontext (i.e. ignore construct 'dispatch') and
user conditional(1)
+ #pragma omp dispatch nocontext(1)
+ res = f (g ());
+ __builtin_printf("%d\n", res);
+ if (res != 202)
+ __builtin_abort ();
+
+ // Call 'f' due to 'novariants'
+ // Call 'g1' due to construct 'dispatch' and user conditional(1)
+ #pragma omp dispatch novariants(1)
+ res = f (g ());
+ __builtin_printf("%d\n", res); /* ACTUAL RESULT: 303, i.e. 'g' and not
'g1' was called */
+ if (res != 103)
+ __builtin_abort ();
+
+ // Call 'f' due to 'novariants'
+ // Call 'g2' due to nocontext (i.e. ignore construct 'dispatch') and user
conditional(1)
+ #pragma omp dispatch novariants(1) nocontext(1)
+ res = f (g ());
+ __builtin_printf("%d\n", res); /* ACTUAL RESULT: 303, i.e. 'g' and not
'g2' was called */
+ if (res != 203)
+ __builtin_abort ();
+}
+
+int
+main ()
+{
+ test ();
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-9.c
b/gcc/testsuite/c-c++-common/gomp/dispatch-9.c
new file mode 100644
index 000000000000..b4b46b946fdc
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-9.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
+/* Check that the right call to f is wrapped in a GOMP_DISPATCH internal
function
+ before translation and that it is stripped during gimplification. */
+
+int f(int);
+void g(int *x)
+{
+ #pragma omp dispatch
+ x[f(1)] = f(f(2));
+ // ^ only this call to f is a dispatch call
+}
+
+/* { dg-final { scan-tree-dump "\.GOMP_DISPATCH \\(f \\(f \\(2\\)\\)\\)"
"original" } } */
+/* { dg-final { scan-tree-dump-times "\.GOMP_DISPATCH" 1 "original" } } */
+/* { dg-final { scan-tree-dump-not "\.GOMP_DISPATCH" "gimple" } } */
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 04bdec462c29..65232a1b8612 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2025-01-27 Paul-Antoine Arras <par...@baylibre.com>
+
+ Backported from master:
+ 2024-11-20 Paul-Antoine Arras <par...@baylibre.com>
+
+ * testsuite/libgomp.c-c++-common/dispatch-1.c: New test.
+ * testsuite/libgomp.c-c++-common/dispatch-2.c: New test.
+
2025-01-23 Tobias Burnus <tbur...@baylibre.com>
Backported from master:
diff --git a/libgomp/testsuite/libgomp.c-c++-common/dispatch-1.c
b/libgomp/testsuite/libgomp.c-c++-common/dispatch-1.c
new file mode 100644
index 000000000000..0efc075a8596
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/dispatch-1.c
@@ -0,0 +1,76 @@
+// Adapted from OpenMP examples
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int baz (double *d_bv, const double *d_av, int n)
+{
+#pragma omp distribute parallel for
+ for (int i = 0; i < n; i++)
+ d_bv[i] = d_av[i] * i;
+ return -3;
+}
+
+int bar (double *d_bv, const double *d_av, int n)
+{
+#pragma omp target is_device_ptr(d_bv, d_av)
+ for (int i = 0; i < n; i++)
+ d_bv[i] = d_av[i] * i;
+ return -2;
+}
+
+#pragma omp declare variant(bar) match(construct={dispatch})
adjust_args(need_device_ptr: bv, av)
+#pragma omp declare variant(baz) match(implementation={vendor(gnu)})
+int foo (double *bv, const double *av, int n)
+{
+ for (int i = 0; i < n; i++)
+ bv[i] = av[i] * i;
+ return -1;
+}
+
+int test (int n)
+{
+ const double e = 2.71828;
+
+ double *av = (double *) malloc (n * sizeof (*av));
+ double *bv = (double *) malloc (n * sizeof (*bv));
+ double *d_bv = (double *) malloc (n * sizeof (*d_bv));
+
+ for (int i = 0; i < n; i++)
+ {
+ av[i] = e * i;
+ bv[i] = 0.0;
+ d_bv[i] = 0.0;
+ }
+
+ int f, last_dev = omp_get_num_devices () - 1;
+#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev)
if (n == 1024)
+ {
+ #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024)
device(last_dev)
+ f = foo (d_bv, av, n);
+ }
+
+ foo (bv, av, n);
+ for (int i = 0; i < n; i++)
+ {
+ if (d_bv[i] != bv[i])
+ {
+ fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i],
bv[i]);
+ return 1;
+ }
+ }
+ return f;
+}
+
+int
+main (void)
+{
+ int ret = test(1023);
+ if (ret != -1) return 1;
+ ret = test(1024);
+ if (ret != -2) return 1;
+ ret = test(1025);
+ if (ret != -3) return 1;
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c
b/libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c
new file mode 100644
index 000000000000..faa0d8a1d1c1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/dispatch-2.c
@@ -0,0 +1,84 @@
+// Adapted from OpenMP examples
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int baz (double *d_bv, const double *d_av, int n);
+int bar (double *d_bv, const double *d_av, int n);
+
+#pragma omp declare variant(bar) match(construct={dispatch})
adjust_args(need_device_ptr: f_bv, f_av)
+#pragma omp declare variant(baz) match(implementation={vendor(gnu)})
+int foo (double *f_bv, const double *f_av, int n);
+
+int baz (double *bv, const double *av, int n);
+int bar (double *bv, const double *av, int n);
+
+int foo (double *bv, const double *av, int n)
+{
+ for (int i = 0; i < n; i++)
+ bv[i] = av[i] * i;
+ return -1;
+}
+
+int baz (double *d_bv, const double *d_av, int n)
+{
+#pragma omp distribute parallel for
+ for (int i = 0; i < n; i++)
+ d_bv[i] = d_av[i] * i;
+ return -3;
+}
+
+int bar (double *d_bv, const double *d_av, int n)
+{
+#pragma omp target is_device_ptr(d_bv, d_av)
+ for (int i = 0; i < n; i++)
+ d_bv[i] = d_av[i] * i;
+ return -2;
+}
+
+int test (int n)
+{
+ const double e = 2.71828;
+
+ double *av = (double *) malloc (n * sizeof (*av));
+ double *bv = (double *) malloc (n * sizeof (*bv));
+ double *d_bv = (double *) malloc (n * sizeof (*d_bv));
+
+ for (int i = 0; i < n; i++)
+ {
+ av[i] = e * i;
+ bv[i] = 0.0;
+ d_bv[i] = 0.0;
+ }
+
+ int f, last_dev = omp_get_num_devices () - 1;
+#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev)
if (n == 1024)
+ {
+ #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024)
device(last_dev)
+ f = foo (d_bv, av, n);
+ }
+
+ foo (bv, av, n);
+ for (int i = 0; i < n; i++)
+ {
+ if (d_bv[i] != bv[i])
+ {
+ fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i],
bv[i]);
+ return 1;
+ }
+ }
+ return f;
+}
+
+int
+main (void)
+{
+ int ret = test(1023);
+ if (ret != -1) return 1;
+ ret = test(1024);
+ if (ret != -2) return 1;
+ ret = test(1025);
+ if (ret != -3) return 1;
+ return 0;
+}
