This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [ptx] partitioning optimization
- From: Nathan Sidwell <nathan at acm dot org>
- To: Bernd Schmidt <bschmidt at redhat dot com>, GCC Patches <gcc-patches at gcc dot gnu dot org>
- Date: Fri, 13 Nov 2015 15:06:56 -0500
- Subject: Re: [ptx] partitioning optimization
- Authentication-results: sourceware.org; auth=none
- References: <564270D6 dot 6090303 at acm dot org> <56432F50 dot 6000208 at redhat dot com> <564349B5 dot 5050202 at acm dot org> <56434E7D dot 4070803 at redhat dot com>
On 11/11/15 09:19, Bernd Schmidt wrote:
On 11/11/2015 02:59 PM, Nathan Sidwell wrote:
That's not the problem. How to conditionally enable the test is the
difficulty. I suspect porting something concerning accel_compiler from
the libgomp testsuite is needed?
Maybe a check_effective_target_offload_nvptx which tries to see if
-foffload=nvptx gives an error (I would hope it does if it's unsupported).
This patch seems to do the trick. tested on an offload-aware build (passes),
and a regular x86_64-linux build (skipped as unsupported).
ok?
nathan
2015-11-13 Nathan Sidwell <nathan@codesourcery.com>
* lib/target-supports.exp (check_effective_target_offload_nvptx): New.
* gcc.dg/goacc/nvptx-merged-loop.c: New.
Index: testsuite/gcc.dg/goacc/nvptx-merged-loop.c
===================================================================
--- testsuite/gcc.dg/goacc/nvptx-merged-loop.c (revision 0)
+++ testsuite/gcc.dg/goacc/nvptx-merged-loop.c (working copy)
@@ -0,0 +1,30 @@
+/* { dg-do link } */
+/* { dg-require-effective-target offload_nvptx } */
+/* { dg-options "-fopenacc -O2 -foffload=-fdump-rtl-mach\\ -dumpbase\\ nvptx-merged-loop.c\\ -Wa,--no-verify" } */
+
+#define N (32*32*32+17)
+void __attribute__ ((noinline)) Foo (int *ary)
+{
+ int ix;
+
+#pragma acc parallel num_workers(32) vector_length(32) copyout(ary[0:N])
+ {
+ /* Loop partitioning should be merged. */
+#pragma acc loop worker vector
+ for (unsigned ix = 0; ix < N; ix++)
+ {
+ ary[ix] = ix;
+ }
+ }
+}
+
+int main ()
+{
+ int ary[N];
+
+ Foo (ary);
+
+ return 0;
+}
+
+/* { dg-final { scan-rtl-dump "Merging loop .* into " "mach" } } */
Index: testsuite/lib/target-supports.exp
===================================================================
--- testsuite/lib/target-supports.exp (revision 230324)
+++ testsuite/lib/target-supports.exp (working copy)
@@ -6716,3 +6716,11 @@ proc check_effective_target_vect_max_red
}
return 0
}
+
+# Return 1 if there is an nvptx offload compiler.
+
+proc check_effective_target_offload_nvptx { } {
+ return [check_no_compiler_messages offload_nvptx object {
+ int main () {return 0;}
+ } "-foffload=nvptx-none" ]
+}