[libgomp,testsuite,nvptx] Add libgomp.c/declare-variant-3-sm*.c

Message ID c276d17e-184d-b1c3-8e35-85f51cc6e7b4@suse.de
State New
Headers
Series [libgomp,testsuite,nvptx] Add libgomp.c/declare-variant-3-sm*.c |

Commit Message

Tom de Vries Feb. 24, 2022, 10:01 a.m. UTC
  [ was: Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70 ]

On 2/24/22 09:29, Tom de Vries wrote:
> I'll try to submit a patch with one or more test-cases.

Hi,

These test-cases exercise the omp declare variant construct using the 
available nvptx isas.

OK for trunk?

Thanks,
- Tom
  

Comments

Jakub Jelinek Feb. 24, 2022, 10:09 a.m. UTC | #1
On Thu, Feb 24, 2022 at 11:01:22AM +0100, Tom de Vries wrote:
> [ was: Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70 ]
> 
> On 2/24/22 09:29, Tom de Vries wrote:
> > I'll try to submit a patch with one or more test-cases.
> 
> Hi,
> 
> These test-cases exercise the omp declare variant construct using the
> available nvptx isas.
> 
> OK for trunk?
> 
> Thanks,
> - Tom

> [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c
> 
> Add openmp test-cases that test the omp declare variant construct:
> ...
>   #pragma omp declare variant (f30) match (device={isa("sm_30")})
> ...
> using the available nvptx isas.
> 
> On a Pascal board GT 1030 with sm_61, we have these unsupported:
> ...
> UNSUPPORTED: libgomp.c/declare-variant-3-sm70.c
> UNSUPPORTED: libgomp.c/declare-variant-3-sm75.c
> UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
> ...
> and on a Turing board T400 with sm_75, we have this only this one:
> ...
> UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
> ...
> 
> Tested on x86_64 with nvptx accelerator.

I think testing it through dg-do link tests with -fdump-tree-optimized
or so would be better, you wouldn't need access to actual hardware level
and checking in the dump what function is actually called for each case is
easy.

	Jakub
  
Tom de Vries Feb. 24, 2022, 10:32 a.m. UTC | #2
On 2/24/22 11:09, Jakub Jelinek wrote:
> On Thu, Feb 24, 2022 at 11:01:22AM +0100, Tom de Vries wrote:
>> [ was: Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70 ]
>>
>> On 2/24/22 09:29, Tom de Vries wrote:
>>> I'll try to submit a patch with one or more test-cases.
>>
>> Hi,
>>
>> These test-cases exercise the omp declare variant construct using the
>> available nvptx isas.
>>
>> OK for trunk?
>>
>> Thanks,
>> - Tom
> 
>> [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c
>>
>> Add openmp test-cases that test the omp declare variant construct:
>> ...
>>    #pragma omp declare variant (f30) match (device={isa("sm_30")})
>> ...
>> using the available nvptx isas.
>>
>> On a Pascal board GT 1030 with sm_61, we have these unsupported:
>> ...
>> UNSUPPORTED: libgomp.c/declare-variant-3-sm70.c
>> UNSUPPORTED: libgomp.c/declare-variant-3-sm75.c
>> UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
>> ...
>> and on a Turing board T400 with sm_75, we have this only this one:
>> ...
>> UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
>> ...
>>
>> Tested on x86_64 with nvptx accelerator.
> 
> I think testing it through dg-do link tests with -fdump-tree-optimized
> or so would be better, you wouldn't need access to actual hardware level
> and checking in the dump what function is actually called for each case is
> easy.
> 

Done, expect for the sm_30 test which is still dg-do run (although I've 
added the compile time test) which should pass on all boards (since we 
don't support below sm_30).

OK for trunk?

Thanks,
- Tom
  
Jakub Jelinek Feb. 24, 2022, 10:38 a.m. UTC | #3
On Thu, Feb 24, 2022 at 11:32:53AM +0100, Tom de Vries wrote:
> libgomp/ChangeLog:
> 
> 2022-02-24  Tom de Vries  <tdevries@suse.de>
> 
> 	* testsuite/libgomp.c/declare-variant-3-sm30.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm35.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm53.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm70.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm75.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm80.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3.h: New header file.

LGTM, thanks.

	Jakub
  
Thomas Schwinge Nov. 30, 2023, 2:48 p.m. UTC | #4
Hi!

On 2022-02-24T11:32:53+0100, Tom de Vries via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
> [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

\o/ Yay for test cases!

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
> @@ -0,0 +1,7 @@
> +/* { dg-do run { target { offload_target_nvptx } } } */
> +/* { dg-additional-options "-foffload=-misa=sm_30" } */
> +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */

Etc.

Pushed to master branch commit 3f5a3b7539e066b539e81b901687facdea4e1bac
"Fix 'libgomp.c/declare-variant-3-*.c' compilation for configurations where GCN offloading is enabled in addition to nvptx",
see attached.


Grüße
 Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  
Thomas Schwinge Nov. 30, 2023, 2:49 p.m. UTC | #5
Hi!

On 2023-11-30T15:48:25+0100, I wrote:
> On 2022-02-24T11:32:53+0100, Tom de Vries via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
>> [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c
>
> \o/ Yay for test cases!
>
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
>> @@ -0,0 +1,7 @@
>> +/* { dg-do run { target { offload_target_nvptx } } } */
>> +/* { dg-additional-options "-foffload=-misa=sm_30" } */
>> +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
>
> Etc.
>
> Pushed to master branch commit 3f5a3b7539e066b539e81b901687facdea4e1bac
> "Fix 'libgomp.c/declare-variant-3-*.c' compilation for configurations where GCN offloading is enabled in addition to nvptx",
> see attached.

..., now.


Grüße
 Thomas


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  

Patch

[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

Add openmp test-cases that test the omp declare variant construct:
...
  #pragma omp declare variant (f30) match (device={isa("sm_30")})
...
using the available nvptx isas.

On a Pascal board GT 1030 with sm_61, we have these unsupported:
...
UNSUPPORTED: libgomp.c/declare-variant-3-sm70.c
UNSUPPORTED: libgomp.c/declare-variant-3-sm75.c
UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
...
and on a Turing board T400 with sm_75, we have this only this one:
...
UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
...

Tested on x86_64 with nvptx accelerator.

libgomp/ChangeLog:

2022-02-24  Tom de Vries  <tdevries@suse.de>

	* testsuite/lib/libgomp.exp
	(check_effective_target_offload_device_nvptx_sm_xx)
	(check_effective_target_offload_device_nvptx_sm_30)
	(check_effective_target_offload_device_nvptx_sm_35)
	(check_effective_target_offload_device_nvptx_sm_53)
	(check_effective_target_offload_device_nvptx_sm_70)
	(check_effective_target_offload_device_nvptx_sm_75)
	(check_effective_target_offload_device_nvptx_sm_80): New proc.
	* testsuite/libgomp.c/declare-variant-3-sm30.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm35.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm53.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm70.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm75.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm80.c: New test.
	* testsuite/libgomp.c/declare-variant-3.h: New header file.

---
 libgomp/testsuite/lib/libgomp.exp                  | 46 +++++++++++++++
 .../testsuite/libgomp.c/declare-variant-3-sm30.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm35.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm53.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm70.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm75.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm80.c   |  5 ++
 libgomp/testsuite/libgomp.c/declare-variant-3.h    | 66 ++++++++++++++++++++++
 8 files changed, 142 insertions(+)

diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 8c5ecfff0ac..d664863b15c 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -426,6 +426,52 @@  proc check_effective_target_offload_device_nvptx { } {
     } ]
 }
 
+# Return 1 if using nvptx offload device which supports -misa=sm_$SM.
+proc check_effective_target_offload_device_nvptx_sm_xx { sm } {
+    if { ![check_effective_target_offload_device_nvptx] } {
+	return 0
+    }
+    return [check_runtime_nocache offload_device_nvptx_sm_$sm {
+      int main ()
+	{
+	  int x = 1;
+	  #pragma omp target map(tofrom: x)
+	    x--;
+	  return x;
+	}
+    } "-foffload=-misa=sm_$sm" ]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_30 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 30]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_35 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 35]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_53 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 53]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_70 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 70]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_75 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 75]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_80 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 80]
+}
+
 # Return 1 if at least one Nvidia GPU is accessible.
 
 proc check_effective_target_openacc_nvidia_accel_present { } {
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
new file mode 100644
index 00000000000..7c680b07a94
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
@@ -0,0 +1,5 @@ 
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_30 } */
+/* { dg-additional-options "-foffload=-misa=sm_30" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
new file mode 100644
index 00000000000..b8b2a714248
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
@@ -0,0 +1,5 @@ 
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_35 } */
+/* { dg-additional-options "-foffload=-misa=sm_35" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
new file mode 100644
index 00000000000..cfc7ee6f137
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
@@ -0,0 +1,5 @@ 
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_53 } */
+/* { dg-additional-options "-foffload=-misa=sm_53" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
new file mode 100644
index 00000000000..4527cc1376c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
@@ -0,0 +1,5 @@ 
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_70 } */
+/* { dg-additional-options "-foffload=-misa=sm_70" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
new file mode 100644
index 00000000000..8e7da369fc5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
@@ -0,0 +1,5 @@ 
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_75 } */
+/* { dg-additional-options "-foffload=-misa=sm_75" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
new file mode 100644
index 00000000000..5cf5b2867a2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
@@ -0,0 +1,5 @@ 
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_80 } */
+/* { dg-additional-options "-foffload=-misa=sm_80" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3.h b/libgomp/testsuite/libgomp.c/declare-variant-3.h
new file mode 100644
index 00000000000..772fc20a519
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3.h
@@ -0,0 +1,66 @@ 
+#pragma omp declare target
+int
+f30 (void)
+{
+  return 30;
+}
+
+int
+f35 (void)
+{
+  return 35;
+}
+
+int
+f53 (void)
+{
+  return 53;
+}
+
+int
+f70 (void)
+{
+  return 70;
+}
+
+int
+f75 (void)
+{
+  return 75;
+}
+
+int
+f80 (void)
+{
+  return 80;
+}
+
+#pragma omp declare variant (f30) match (device={isa("sm_30")})
+#pragma omp declare variant (f35) match (device={isa("sm_35")})
+#pragma omp declare variant (f53) match (device={isa("sm_53")})
+#pragma omp declare variant (f70) match (device={isa("sm_70")})
+#pragma omp declare variant (f75) match (device={isa("sm_75")})
+#pragma omp declare variant (f80) match (device={isa("sm_80")})
+int
+f (void)
+{
+  return 0;
+}
+
+#pragma omp end declare target
+
+int
+main (void)
+{
+  int v = 0;
+
+  #pragma omp target map(from:v)
+  v = f ();
+
+  if (v == 0)
+    __builtin_abort ();
+
+  __builtin_printf ("Nvptx accelerator: sm_%d\n", v);
+
+  return 0;
+}