From: Tobias Burnus <tobias@codesourcery.com>
To: Tom de Vries <tdevries@suse.de>, gcc-patches <gcc-patches@gcc.gnu.org>
Cc: Jakub Jelinek <jakub@redhat.com>
Subject: Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70
Date: Tue, 22 Feb 2022 17:03:07 +0100 [thread overview]
Message-ID: <221e5215-3b8b-fa9f-4e7b-2faa86b8b991@codesourcery.com> (raw)
In-Reply-To: <6b0f4fe5-06d4-5ede-5bd1-a53ed82f6d36@suse.de>
Hi Tom,
On 22.02.22 15:43, Tom de Vries wrote:
> On 2/17/22 18:24, Tobias Burnus wrote:
>> --- a/gcc/config/nvptx/t-omp-device
>> +++ b/gcc/config/nvptx/t-omp-device
>> @@ -1,4 +1,4 @@
>> echo kind: gpu > $@
>> echo arch: nvptx >> $@
>> - echo isa: sm_30 sm_35 >> $@
>> + echo isa: sm_30 sm_35 sm_53 sm_70 sm_75 sm_80 >> $@
>
> I'm not sure I understand how this is used. Is this user-visible? Is
> there a libgomp test-case where we can observe a difference?
That's used for OpenMP context selectors like; that way, one can generate,
e.g. one code used with nvptx and one with gcn as with:
#pragma omp declare variant (on_nvptx) match(construct={target},device={arch(nvptx)})
#pragma omp declare variant (on_gcn) match(construct={target},device={arch(gcn)})
...
#pragma omp target map(from:v)
v = on ();
which then either calls 'on' or 'on_nvptx' or 'on_gcn'
(from libgomp/testsuite/libgomp.c/target-42.c)
The following testcases use 'arch(nvptx)':
libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
libgomp/testsuite/libgomp.c/target-42.c
libgomp/testsuite/libgomp.c/usleep.h
libgomp/testsuite/libgomp.fortran/declare-variant-1.f90
For ISA, there is only one run-time test:
libgomp/testsuite/libgomp.c/declare-variant-1.c
but only for x86-64: match (device={isa("avx512f")})
The sm_35 also appears, but only in the compile-time tests:
gcc/testsuite/{c-c++-common,gfortran.dg}/gomp/declare-variant-{9,10}.*
Tobias
-----------------
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
next prev parent reply other threads:[~2022-02-22 16:03 UTC|newest]
Thread overview: 14+ messages / expand[flat|nested] mbox.gz Atom feed top
2022-02-17 17:24 Tobias Burnus
2022-02-22 14:26 ` Tom de Vries
2022-02-22 14:39 ` Tom de Vries
2022-02-22 14:43 ` Tom de Vries
2022-02-22 16:03 ` Tobias Burnus [this message]
2022-02-24 8:29 ` Tom de Vries
2022-02-24 10:01 ` [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c Tom de Vries
2022-02-24 10:09 ` Jakub Jelinek
2022-02-24 10:32 ` Tom de Vries
2022-02-24 10:38 ` Jakub Jelinek
2023-11-30 14:48 ` Fix 'libgomp.c/declare-variant-3-*.c' compilation for configurations where GCN offloading is enabled in addition to nvptx (was: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c) Thomas Schwinge
2023-11-30 14:49 ` Thomas Schwinge
2023-11-30 14:57 ` In 'libgomp.c/declare-variant-{3,4}-*.c', restrict 'scan-offload-tree-dump's to 'only_for_offload_target [...]' " Thomas Schwinge
2023-11-30 14:59 ` Spin 'dg-do run' part of 'libgomp.c/declare-variant-3-sm30.c' off into new 'libgomp.c/declare-variant-3.c' " Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=221e5215-3b8b-fa9f-4e7b-2faa86b8b991@codesourcery.com \
--to=tobias@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=tdevries@suse.de \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).