[nvptx] Pass -m32/-m64 to host_compiler if it has multilib support

Hi Richard,
After differing NUM_POLY_INT_COEFFS fix for AArch64/nvptx offloading, the following minimal test:

int main()
{
  int x;
  #pragma omp target map(x)
    x = 5;
  return x;
}

compiled with -fopenmp -foffload=nvptx-none now fails with:
gcc: error: unrecognized command-line option '-m64'
nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit status compilation terminated.

As mentioned in RFC email, this happens because nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host compiler depending on whether
offload_abi is OFFLOAD_ABI_LP64 or OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these options.

Based on your suggestion in: https://gcc.gnu.org/pipermail/gcc/2024-July/244470.html,
The attached patch generates new macro HOST_MULTILIB derived from $enable_as_accelerator_for, and in mkoffload.cc it gates passing -m32/-m64
to host_compiler on HOST_MULTILIB. I verified that the macro is set to 0 for aarch64 host (and thus avoids above unrecognized command line option error),
and is set to 1 for x86_64 host.

Does the patch look OK ?

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

Thanks,
Prathamesh
[nvptx] Pass -m32/-m64 to host_compiler if it has multilib support.

gcc/ChangeLog:
	* configure.ac: Generate new macro HOST_MULTILIB.
	* config.in: Regenerate.
	* configure: Likewise.
	* config/nvptx/mkoffload.cc (compile_native): Gate appending
	"-m32"/"-m64" to argv_obstack on HOST_MULTILIB.
	(main): Likewise.

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

On Thu, Aug 8, 2024 at 6:11 AM Prathamesh Kulkarni
<prathameshk@nvidia.com> wrote:
>
> Hi Richard,
> After differing NUM_POLY_INT_COEFFS fix for AArch64/nvptx offloading, the following minimal test:
>
> int main()
> {
>   int x;
>   #pragma omp target map(x)
>     x = 5;
>   return x;
> }
>
> compiled with -fopenmp -foffload=nvptx-none now fails with:
> gcc: error: unrecognized command-line option '-m64'
> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit status compilation terminated.
>
> As mentioned in RFC email, this happens because nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host compiler depending on whether
> offload_abi is OFFLOAD_ABI_LP64 or OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these options.
>
> Based on your suggestion in: https://gcc.gnu.org/pipermail/gcc/2024-July/244470.html,
> The attached patch generates new macro HOST_MULTILIB derived from $enable_as_accelerator_for, and in mkoffload.cc it gates passing -m32/-m64
> to host_compiler on HOST_MULTILIB. I verified that the macro is set to 0 for aarch64 host (and thus avoids above unrecognized command line option error),
> and is set to 1 for x86_64 host.
>
> Does the patch look OK ?

Note I think the usage of the name MULTILIB here is wrong because
aarch64 (and riscv) could have MUTLILIB support just the options are
different. For aarch64, it would be -mabi=ilp32/-mabi=lp64 (riscv it
is more complex).

This most likely should be something more complex due to the above.
Maybe call it HOST_64_32 but even that seems wrong due to Aarch64
having ILP32 support and such.
What about HOST_64ABI_OPTS="-mabi=lp64"/HOST_32ABI_OPTS="-mabi=ilp32"
but  I am not sure if that would be enough to support RISCV which
requires two options.

Thanks,
Andrew Pinski

>
> Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>
>
> Thanks,
> Prathamesh

Hi Prathamesh!

On 2024-08-08T06:46:25-0700, Andrew Pinski <pinskia@gmail.com> wrote:
> On Thu, Aug 8, 2024 at 6:11 AM Prathamesh Kulkarni
> <prathameshk@nvidia.com> wrote:
>> After differing NUM_POLY_INT_COEFFS fix for AArch64/nvptx offloading, the following minimal test:

First, thanks for your work on enabling this!  I will say that I had the
plan to re-engage with Nvidia to hire us (as initial implementors of
GCC/nvptx offloading) to make AArch64/nvptx offloading work, but now that
Nvidia has its own GCC team, that's great that you're able to work on
this yourself!  :-)

Please CC me for GCC/nvptx issues for (at least potentially...) faster
response times.

>> compiled with -fopenmp -foffload=nvptx-none now fails with:
>> gcc: error: unrecognized command-line option '-m64'
>> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit status compilation terminated.

Heh.  Yeah...

>> As mentioned in RFC email, this happens because nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host compiler depending on whether
>> offload_abi is OFFLOAD_ABI_LP64 or OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these options.
>>
>> Based on your suggestion in: https://gcc.gnu.org/pipermail/gcc/2024-July/244470.html,
>> The attached patch generates new macro HOST_MULTILIB derived from $enable_as_accelerator_for, and in mkoffload.cc it gates passing -m32/-m64
>> to host_compiler on HOST_MULTILIB. I verified that the macro is set to 0 for aarch64 host (and thus avoids above unrecognized command line option error),
>> and is set to 1 for x86_64 host.
>>
>> Does the patch look OK ?
>
> Note I think the usage of the name MULTILIB here is wrong because
> aarch64 (and riscv) could have MUTLILIB support just the options are
> different.

I also think the proposed patch is not quite the right hammer for the
issue at hand.

> For aarch64, it would be -mabi=ilp32/-mabi=lp64 (riscv it
> is more complex).
>
> This most likely should be something more complex due to the above.

Right.

> Maybe call it HOST_64_32 but even that seems wrong due to Aarch64
> having ILP32 support and such.

Right.

> What about HOST_64ABI_OPTS="-mabi=lp64"/HOST_32ABI_OPTS="-mabi=ilp32"
> but  I am not sure if that would be enough to support RISCV which
> requires two options.

So, my idea is: instead of the current strategy that the host
'TARGET_OFFLOAD_OPTIONS' synthesizes '-foffload-abi=lp64' etc., which the
'mkoffload's then interpret and re-synthesize '-m64' etc. -- how about we
instead directly tell the 'mkoffload's the relevant ABI options?  That
is, 'TARGET_OFFLOAD_OPTIONS' instead synthesizes '-foffload-abi=-m64'
etc., which the 'mkoffload's can then readily use.  Could you please give
that a try, and/or does anyone see any issues with that approach?

And use something like '-foffload-abi=disable' to replace the current:

    /* PR libgomp/65099: Currently, we only support offloading in 64-bit
       configurations.  */
    if (offload_abi == OFFLOAD_ABI_LP64)
      {

(As discussed before, this should be done differently altogether, but
that's for another day.)


Grüße
 Thomas

> -----Original Message-----
> From: Thomas Schwinge <tschwinge@baylibre.com>
> Sent: Friday, August 9, 2024 12:55 AM
> To: Prathamesh Kulkarni <prathameshk@nvidia.com>
> Cc: Andrew Pinski <pinskia@gmail.com>; Richard Biener
> <rguenther@suse.de>; gcc-patches@gcc.gnu.org; Jakub Jelinek
> <jakub@redhat.com>
> Subject: Re: [nvptx] Pass -m32/-m64 to host_compiler if it has
> multilib support
> 
> External email: Use caution opening links or attachments
> 
> 
> Hi Prathamesh!
Hi Thomas,
> 
> On 2024-08-08T06:46:25-0700, Andrew Pinski <pinskia@gmail.com> wrote:
> > On Thu, Aug 8, 2024 at 6:11 AM Prathamesh Kulkarni
> > <prathameshk@nvidia.com> wrote:
> >> After differing NUM_POLY_INT_COEFFS fix for AArch64/nvptx
> offloading, the following minimal test:
> 
> First, thanks for your work on enabling this!  I will say that I had
> the plan to re-engage with Nvidia to hire us (as initial implementors
> of GCC/nvptx offloading) to make AArch64/nvptx offloading work, but
> now that Nvidia has its own GCC team, that's great that you're able to
> work on this yourself!  :-)
> 
> Please CC me for GCC/nvptx issues for (at least potentially...) faster
> response times.
Thanks, will do 😊
> 
> >> compiled with -fopenmp -foffload=nvptx-none now fails with:
> >> gcc: error: unrecognized command-line option '-m64'
> >> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit
> status compilation terminated.
> 
> Heh.  Yeah...
> 
> >> As mentioned in RFC email, this happens because
> >> nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host compiler
> depending on whether offload_abi is OFFLOAD_ABI_LP64 or
> OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these
> options.
> >>
> >> Based on your suggestion in:
> >> https://gcc.gnu.org/pipermail/gcc/2024-July/244470.html,
> >> The attached patch generates new macro HOST_MULTILIB derived from
> >> $enable_as_accelerator_for, and in mkoffload.cc it gates passing
> >> -m32/-m64 to host_compiler on HOST_MULTILIB. I verified that the
> macro is set to 0 for aarch64 host (and thus avoids above unrecognized
> command line option error), and is set to 1 for x86_64 host.
> >>
> >> Does the patch look OK ?
> >
> > Note I think the usage of the name MULTILIB here is wrong because
> > aarch64 (and riscv) could have MUTLILIB support just the options are
> > different.
> 
> I also think the proposed patch is not quite the right hammer for the
> issue at hand.
> 
> > For aarch64, it would be -mabi=ilp32/-mabi=lp64 (riscv it is more
> > complex).
> >
> > This most likely should be something more complex due to the above.
> 
> Right.
> 
> > Maybe call it HOST_64_32 but even that seems wrong due to Aarch64
> > having ILP32 support and such.
> 
> Right.
> 
> > What about HOST_64ABI_OPTS="-mabi=lp64"/HOST_32ABI_OPTS="-
> mabi=ilp32"
> > but  I am not sure if that would be enough to support RISCV which
> > requires two options.
> 
> So, my idea is: instead of the current strategy that the host
> 'TARGET_OFFLOAD_OPTIONS' synthesizes '-foffload-abi=lp64' etc., which
> the 'mkoffload's then interpret and re-synthesize '-m64' etc. -- how
> about we instead directly tell the 'mkoffload's the relevant ABI
> options?  That is, 'TARGET_OFFLOAD_OPTIONS' instead synthesizes '-
> foffload-abi=-m64'
> etc., which the 'mkoffload's can then readily use.  Could you please
> give that a try, and/or does anyone see any issues with that approach?
> 
> And use something like '-foffload-abi=disable' to replace the current:
> 
>     /* PR libgomp/65099: Currently, we only support offloading in 64-
> bit
>        configurations.  */
>     if (offload_abi == OFFLOAD_ABI_LP64)
>       {
> 
> (As discussed before, this should be done differently altogether, but
> that's for another day.)
Sorry, I don't quite follow. Currently we enable offloading if offload_abi == OFFLOAD_ABI_LP64,
which is synthesized from -foffload-abi=lp64. If we change -foffload-abi to instead specify
host-specific ABI opts, I guess mkoffload will still need to somehow figure out which ABI is used,
so it can disable offloading for 32-bit ? I suppose we could adjust TARGET_OFFLOAD_OPTIONS for each
host to pass -foffload-abi=disable if TARGET_ILP32 is set and offload target is nvptx, but not sure
if that'd be correct ?

In the attached patch, I added another option -foffload-abi-host-opts to specify host abi
opts, and leave -foffload-abi to specify if ABI is 32/64 bit which mkoffload can use to
enable/disable offloading (as before).
Does that look OK ?

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

Thanks,
Prathamesh
> 
> 
> Grüße
>  Thomas
[nvptx] Pass host specific ABI opts from mkoffload.

The patch adds an option -foffload-abi-host-opts, which
is set by host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's value
to host_compiler.

gcc/ChangeLog:
	* common.opt (foffload-abi-host-opts): New option.
	* config/aarch64/aarch64.cc (aarch64_offload_options): Set
	-foffload-abi-host-opts.
	* config/i386/i386-opts.cc (ix86_offload_options): Likewise.
	* config/rs6000/rs6000.cc (rs6000_offload_options): Likewise.
	* config/nvptx/mkoffload.cc (host_abi_opts): Define.
	(compile_native): Append host_abi_opts to argv_obstack.
	(main): Handle option -foffload-abi-host-opts.
	* lto-wrapper.cc (append_compiler_options): Handle
	-foffload-abi-host-opts.
	* opts.cc (common_handle_option): Likewise.

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

diff --git a/gcc/common.opt b/gcc/common.opt
index ea39f87ae71..d1a9efb9513 100644
--- a/gcc/common.opt
+++ b/gcc/common.opt
@@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
 EnumValue
 Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
 
+foffload-abi-host-opts=
+Common Driver Joined MissingArgError(option or option=abi missing after %qs)
+-foffload-abi-host-opts=<options>=<abi> Specify host abi options.
+
 fomit-frame-pointer
 Common Var(flag_omit_frame_pointer) Optimization
 When possible do not generate stack frames.
diff --git a/gcc/config/aarch64/aarch64.cc b/gcc/config/aarch64/aarch64.cc
index 2ac5a22c848..7418cb1fb69 100644
--- a/gcc/config/aarch64/aarch64.cc
+++ b/gcc/config/aarch64/aarch64.cc
@@ -18999,9 +18999,9 @@ static char *
 aarch64_offload_options (void)
 {
   if (TARGET_ILP32)
-    return xstrdup ("-foffload-abi=ilp32");
+    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-mabi=ilp32");
   else
-    return xstrdup ("-foffload-abi=lp64");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-mabi=lp64");
 }
 
 static struct machine_function *
diff --git a/gcc/config/i386/i386-options.cc b/gcc/config/i386/i386-options.cc
index 1c8f7835af2..bd960674e5d 100644
--- a/gcc/config/i386/i386-options.cc
+++ b/gcc/config/i386/i386-options.cc
@@ -3669,8 +3669,8 @@ char *
 ix86_offload_options (void)
 {
   if (TARGET_LP64)
-    return xstrdup ("-foffload-abi=lp64");
-  return xstrdup ("-foffload-abi=ilp32");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
+  return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
 }
 
 /* Handle "cdecl", "stdcall", "fastcall", "regparm", "thiscall",
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index 503b1abcefd..d5ca2386641 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -61,6 +61,7 @@ static const char *omp_requires_file;
 static const char *ptx_dumpbase;
 
 enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
+const char *host_abi_opts = NULL;
 
 /* Delete tempfiles.  */
 
@@ -607,17 +608,9 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
   obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
   obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
   obstack_ptr_grow (&argv_obstack, ".c");
-  switch (offload_abi)
-    {
-    case OFFLOAD_ABI_LP64:
-      obstack_ptr_grow (&argv_obstack, "-m64");
-      break;
-    case OFFLOAD_ABI_ILP32:
-      obstack_ptr_grow (&argv_obstack, "-m32");
-      break;
-    default:
-      gcc_unreachable ();
-    }
+  if (!host_abi_opts)
+    fatal_error (input_location, "-foffload-abi-host-opts not specified.");
+  obstack_ptr_grow (&argv_obstack, host_abi_opts);
   obstack_ptr_grow (&argv_obstack, infile);
   obstack_ptr_grow (&argv_obstack, "-c");
   obstack_ptr_grow (&argv_obstack, "-o");
@@ -721,6 +714,8 @@ main (int argc, char **argv)
 			 "unrecognizable argument of option " STR);
 	}
 #undef STR
+      else if (startswith (argv[i], "-foffload-abi-host-opts="))
+	host_abi_opts = argv[i] + strlen ("-foffload-abi-host-opts=");
       else if (strcmp (argv[i], "-fopenmp") == 0)
 	fopenmp = true;
       else if (strcmp (argv[i], "-fopenacc") == 0)
diff --git a/gcc/config/rs6000/rs6000.cc b/gcc/config/rs6000/rs6000.cc
index 0bcc6a2d0ab..decdf49a1f5 100644
--- a/gcc/config/rs6000/rs6000.cc
+++ b/gcc/config/rs6000/rs6000.cc
@@ -17333,9 +17333,9 @@ static char *
 rs6000_offload_options (void)
 {
   if (TARGET_64BIT)
-    return xstrdup ("-foffload-abi=lp64");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
   else
-    return xstrdup ("-foffload-abi=ilp32");
+    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
 }
 
 
diff --git a/gcc/lto-wrapper.cc b/gcc/lto-wrapper.cc
index 6bfc96590a5..1ecc4997e5a 100644
--- a/gcc/lto-wrapper.cc
+++ b/gcc/lto-wrapper.cc
@@ -745,6 +745,7 @@ append_compiler_options (obstack *argv_obstack, vec<cl_decoded_option> opts)
 	case OPT_fopenacc:
 	case OPT_fopenacc_dim_:
 	case OPT_foffload_abi_:
+	case OPT_foffload_abi_host_opts_:
 	case OPT_fcf_protection_:
 	case OPT_fasynchronous_unwind_tables:
 	case OPT_funwind_tables:
diff --git a/gcc/opts.cc b/gcc/opts.cc
index 0b7b137c376..79118237ce4 100644
--- a/gcc/opts.cc
+++ b/gcc/opts.cc
@@ -3069,6 +3069,7 @@ common_handle_option (struct gcc_options *opts,
       break;
 
     case OPT_foffload_abi_:
+    case OPT_foffload_abi_host_opts_:
 #ifdef ACCEL_COMPILER
       /* Handled in the 'mkoffload's.  */
 #else

Hi Prathamesh!

On 2024-08-12T07:50:07+0000, Prathamesh Kulkarni <prathameshk@nvidia.com> wrote:
>> From: Thomas Schwinge <tschwinge@baylibre.com>
>> Sent: Friday, August 9, 2024 12:55 AM

>> On 2024-08-08T06:46:25-0700, Andrew Pinski <pinskia@gmail.com> wrote:
>> > On Thu, Aug 8, 2024 at 6:11 AM Prathamesh Kulkarni
>> > <prathameshk@nvidia.com> wrote:
>> >> After differing NUM_POLY_INT_COEFFS fix for AArch64/nvptx
>> offloading, the following minimal test:
>> 
>> First, thanks for your work on enabling this!  I will say that I had
>> the plan to re-engage with Nvidia to hire us (as initial implementors
>> of GCC/nvptx offloading) to make AArch64/nvptx offloading work, but
>> now that Nvidia has its own GCC team, that's great that you're able to
>> work on this yourself!  :-)
>> 
>> Please CC me for GCC/nvptx issues for (at least potentially...) faster
>> response times.
> Thanks, will do 😊

Heh, so much for "potentially": I'm not able to spend a lot of time on
this right now, as I shall soon be out of office.  Quickly:

>> >> compiled with -fopenmp -foffload=nvptx-none now fails with:
>> >> gcc: error: unrecognized command-line option '-m64'
>> >> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit
>> status compilation terminated.
>> 
>> Heh.  Yeah...
>> 
>> >> As mentioned in RFC email, this happens because
>> >> nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host compiler
>> depending on whether offload_abi is OFFLOAD_ABI_LP64 or
>> OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these
>> options.

>> So, my idea is: instead of the current strategy that the host
>> 'TARGET_OFFLOAD_OPTIONS' synthesizes '-foffload-abi=lp64' etc., which
>> the 'mkoffload's then interpret and re-synthesize '-m64' etc. -- how
>> about we instead directly tell the 'mkoffload's the relevant ABI
>> options?  That is, 'TARGET_OFFLOAD_OPTIONS' instead synthesizes '-
>> foffload-abi=-m64'
>> etc., which the 'mkoffload's can then readily use.  Could you please
>> give that a try, and/or does anyone see any issues with that approach?
>> 
>> And use something like '-foffload-abi=disable' to replace the current:
>> 
>>     /* PR libgomp/65099: Currently, we only support offloading in 64-
>> bit
>>        configurations.  */
>>     if (offload_abi == OFFLOAD_ABI_LP64)
>>       {
>> 
>> (As discussed before, this should be done differently altogether, but
>> that's for another day.)
> Sorry, I don't quite follow. Currently we enable offloading if offload_abi == OFFLOAD_ABI_LP64,
> which is synthesized from -foffload-abi=lp64. If we change -foffload-abi to instead specify
> host-specific ABI opts, I guess mkoffload will still need to somehow figure out which ABI is used,
> so it can disable offloading for 32-bit ? I suppose we could adjust TARGET_OFFLOAD_OPTIONS for each
> host to pass -foffload-abi=disable if TARGET_ILP32 is set and offload target is nvptx, but not sure
> if that'd be correct ?

Basically, yes.  My idea was that all 'TARGET_OFFLOAD_OPTIONS'
implementations return either the correct host flags to be used by the
'mkoffload's (the case that offloading is supported for the current host
flags/ABI configuration), or otherwise return '-foffload-abi=disable'.
For example (untested):

>  char *
>  ix86_offload_options (void)
>  {
>    if (TARGET_LP64)
> -    return xstrdup ("-foffload-abi=lp64");
> +    return xstrdup ("-foffload-abi=-m64");
> -  return xstrdup ("-foffload-abi=ilp32");
> +  return xstrdup ("-foffload-abi=disable");
>  }

That is, only for 'TARGET_LP64' offloading is supported, and via
'-foffload-abi=-m64' the 'mkoffload's know that they need to specify
'-m64'.  For other host flags/ABI configuration, the 'mkoffload's see
'-foffload-abi=disable' and thus disable offload code generation
(replacing the current 'if (offload_abi == OFFLOAD_ABI_LP64)' in
'mkoffload').

> In the attached patch

Yes, that's going in the right direction, thanks!

> I added another option -foffload-abi-host-opts to specify host abi
> opts, and leave -foffload-abi to specify if ABI is 32/64 bit which mkoffload can use to
> enable/disable offloading (as before).

I'm not sure however, if this additional option is really necessary?

In case we're not happy to re-purpose the flag name
'-foffload-abi=[...]', we could also rename that one to
'-foffload-abi-host-opts=[...]'; the former is not user-exposed, so we
may change it as necessary.  (Or, in other words, go with your proposed
'-foffload-abi-host-opts=[...]', but also remove '-foffload-abi=[...]' at
the same time.)


I'll be able to spend more time on this in two weeks.


Grüße
 Thomas


> [nvptx] Pass host specific ABI opts from mkoffload.
>
> The patch adds an option -foffload-abi-host-opts, which
> is set by host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's value
> to host_compiler.
>
> gcc/ChangeLog:
> 	* common.opt (foffload-abi-host-opts): New option.
> 	* config/aarch64/aarch64.cc (aarch64_offload_options): Set
> 	-foffload-abi-host-opts.
> 	* config/i386/i386-opts.cc (ix86_offload_options): Likewise.
> 	* config/rs6000/rs6000.cc (rs6000_offload_options): Likewise.
> 	* config/nvptx/mkoffload.cc (host_abi_opts): Define.
> 	(compile_native): Append host_abi_opts to argv_obstack.
> 	(main): Handle option -foffload-abi-host-opts.
> 	* lto-wrapper.cc (append_compiler_options): Handle
> 	-foffload-abi-host-opts.
> 	* opts.cc (common_handle_option): Likewise.
>
> Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>
>
> diff --git a/gcc/common.opt b/gcc/common.opt
> index ea39f87ae71..d1a9efb9513 100644
> --- a/gcc/common.opt
> +++ b/gcc/common.opt
> @@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
>  EnumValue
>  Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
>  
> +foffload-abi-host-opts=
> +Common Driver Joined MissingArgError(option or option=abi missing after %qs)
> +-foffload-abi-host-opts=<options>=<abi> Specify host abi options.
> +
>  fomit-frame-pointer
>  Common Var(flag_omit_frame_pointer) Optimization
>  When possible do not generate stack frames.
> diff --git a/gcc/config/aarch64/aarch64.cc b/gcc/config/aarch64/aarch64.cc
> index 2ac5a22c848..7418cb1fb69 100644
> --- a/gcc/config/aarch64/aarch64.cc
> +++ b/gcc/config/aarch64/aarch64.cc
> @@ -18999,9 +18999,9 @@ static char *
>  aarch64_offload_options (void)
>  {
>    if (TARGET_ILP32)
> -    return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-mabi=ilp32");
>    else
> -    return xstrdup ("-foffload-abi=lp64");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-mabi=lp64");
>  }
>  
>  static struct machine_function *
> diff --git a/gcc/config/i386/i386-options.cc b/gcc/config/i386/i386-options.cc
> index 1c8f7835af2..bd960674e5d 100644
> --- a/gcc/config/i386/i386-options.cc
> +++ b/gcc/config/i386/i386-options.cc
> @@ -3669,8 +3669,8 @@ char *
>  ix86_offload_options (void)
>  {
>    if (TARGET_LP64)
> -    return xstrdup ("-foffload-abi=lp64");
> -  return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
> +  return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
>  }
>  
>  /* Handle "cdecl", "stdcall", "fastcall", "regparm", "thiscall",
> diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
> index 503b1abcefd..d5ca2386641 100644
> --- a/gcc/config/nvptx/mkoffload.cc
> +++ b/gcc/config/nvptx/mkoffload.cc
> @@ -61,6 +61,7 @@ static const char *omp_requires_file;
>  static const char *ptx_dumpbase;
>  
>  enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
> +const char *host_abi_opts = NULL;
>  
>  /* Delete tempfiles.  */
>  
> @@ -607,17 +608,9 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
>    obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
>    obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
>    obstack_ptr_grow (&argv_obstack, ".c");
> -  switch (offload_abi)
> -    {
> -    case OFFLOAD_ABI_LP64:
> -      obstack_ptr_grow (&argv_obstack, "-m64");
> -      break;
> -    case OFFLOAD_ABI_ILP32:
> -      obstack_ptr_grow (&argv_obstack, "-m32");
> -      break;
> -    default:
> -      gcc_unreachable ();
> -    }
> +  if (!host_abi_opts)
> +    fatal_error (input_location, "-foffload-abi-host-opts not specified.");
> +  obstack_ptr_grow (&argv_obstack, host_abi_opts);
>    obstack_ptr_grow (&argv_obstack, infile);
>    obstack_ptr_grow (&argv_obstack, "-c");
>    obstack_ptr_grow (&argv_obstack, "-o");
> @@ -721,6 +714,8 @@ main (int argc, char **argv)
>  			 "unrecognizable argument of option " STR);
>  	}
>  #undef STR
> +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> +	host_abi_opts = argv[i] + strlen ("-foffload-abi-host-opts=");
>        else if (strcmp (argv[i], "-fopenmp") == 0)
>  	fopenmp = true;
>        else if (strcmp (argv[i], "-fopenacc") == 0)
> diff --git a/gcc/config/rs6000/rs6000.cc b/gcc/config/rs6000/rs6000.cc
> index 0bcc6a2d0ab..decdf49a1f5 100644
> --- a/gcc/config/rs6000/rs6000.cc
> +++ b/gcc/config/rs6000/rs6000.cc
> @@ -17333,9 +17333,9 @@ static char *
>  rs6000_offload_options (void)
>  {
>    if (TARGET_64BIT)
> -    return xstrdup ("-foffload-abi=lp64");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
>    else
> -    return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
>  }
>  
>  
> diff --git a/gcc/lto-wrapper.cc b/gcc/lto-wrapper.cc
> index 6bfc96590a5..1ecc4997e5a 100644
> --- a/gcc/lto-wrapper.cc
> +++ b/gcc/lto-wrapper.cc
> @@ -745,6 +745,7 @@ append_compiler_options (obstack *argv_obstack, vec<cl_decoded_option> opts)
>  	case OPT_fopenacc:
>  	case OPT_fopenacc_dim_:
>  	case OPT_foffload_abi_:
> +	case OPT_foffload_abi_host_opts_:
>  	case OPT_fcf_protection_:
>  	case OPT_fasynchronous_unwind_tables:
>  	case OPT_funwind_tables:
> diff --git a/gcc/opts.cc b/gcc/opts.cc
> index 0b7b137c376..79118237ce4 100644
> --- a/gcc/opts.cc
> +++ b/gcc/opts.cc
> @@ -3069,6 +3069,7 @@ common_handle_option (struct gcc_options *opts,
>        break;
>  
>      case OPT_foffload_abi_:
> +    case OPT_foffload_abi_host_opts_:
>  #ifdef ACCEL_COMPILER
>        /* Handled in the 'mkoffload's.  */
>  #else

> Am 13.08.2024 um 17:48 schrieb Thomas Schwinge <tschwinge@baylibre.com>:
> 
> Hi Prathamesh!
> 
> On 2024-08-12T07:50:07+0000, Prathamesh Kulkarni <prathameshk@nvidia.com> wrote:
>>> From: Thomas Schwinge <tschwinge@baylibre.com>
>>> Sent: Friday, August 9, 2024 12:55 AM
> 
>>> On 2024-08-08T06:46:25-0700, Andrew Pinski <pinskia@gmail.com> wrote:
>>>> On Thu, Aug 8, 2024 at 6:11 AM Prathamesh Kulkarni
>>>> <prathameshk@nvidia.com> wrote:
>>>>> After differing NUM_POLY_INT_COEFFS fix for AArch64/nvptx
>>> offloading, the following minimal test:
>>> 
>>> First, thanks for your work on enabling this!  I will say that I had
>>> the plan to re-engage with Nvidia to hire us (as initial implementors
>>> of GCC/nvptx offloading) to make AArch64/nvptx offloading work, but
>>> now that Nvidia has its own GCC team, that's great that you're able to
>>> work on this yourself!  :-)
>>> 
>>> Please CC me for GCC/nvptx issues for (at least potentially...) faster
>>> response times.
>> Thanks, will do 😊
> 
> Heh, so much for "potentially": I'm not able to spend a lot of time on
> this right now, as I shall soon be out of office.  Quickly:
> 
>>>>> compiled with -fopenmp -foffload=nvptx-none now fails with:
>>>>> gcc: error: unrecognized command-line option '-m64'
>>>>> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit
>>> status compilation terminated.
>>> 
>>> Heh.  Yeah...
>>> 
>>>>> As mentioned in RFC email, this happens because
>>>>> nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host compiler
>>> depending on whether offload_abi is OFFLOAD_ABI_LP64 or
>>> OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these
>>> options.
> 
>>> So, my idea is: instead of the current strategy that the host
>>> 'TARGET_OFFLOAD_OPTIONS' synthesizes '-foffload-abi=lp64' etc., which
>>> the 'mkoffload's then interpret and re-synthesize '-m64' etc. -- how
>>> about we instead directly tell the 'mkoffload's the relevant ABI
>>> options?  That is, 'TARGET_OFFLOAD_OPTIONS' instead synthesizes '-
>>> foffload-abi=-m64'
>>> etc., which the 'mkoffload's can then readily use.  Could you please
>>> give that a try, and/or does anyone see any issues with that approach?
>>> 
>>> And use something like '-foffload-abi=disable' to replace the current:
>>> 
>>>    /* PR libgomp/65099: Currently, we only support offloading in 64-
>>> bit
>>>       configurations.  */
>>>    if (offload_abi == OFFLOAD_ABI_LP64)
>>>      {
>>> 
>>> (As discussed before, this should be done differently altogether, but
>>> that's for another day.)
>> Sorry, I don't quite follow. Currently we enable offloading if offload_abi == OFFLOAD_ABI_LP64,
>> which is synthesized from -foffload-abi=lp64. If we change -foffload-abi to instead specify
>> host-specific ABI opts, I guess mkoffload will still need to somehow figure out which ABI is used,
>> so it can disable offloading for 32-bit ? I suppose we could adjust TARGET_OFFLOAD_OPTIONS for each
>> host to pass -foffload-abi=disable if TARGET_ILP32 is set and offload target is nvptx, but not sure
>> if that'd be correct ?
> 
> Basically, yes.  My idea was that all 'TARGET_OFFLOAD_OPTIONS'
> implementations return either the correct host flags to be used by the
> 'mkoffload's (the case that offloading is supported for the current host
> flags/ABI configuration), or otherwise return '-foffload-abi=disable'.
> For example (untested):
> 
>> char *
>> ix86_offload_options (void)
>> {
>>   if (TARGET_LP64)
>> -    return xstrdup ("-foffload-abi=lp64");
>> +    return xstrdup ("-foffload-abi=-m64");
>> -  return xstrdup ("-foffload-abi=ilp32");
>> +  return xstrdup ("-foffload-abi=disable");
>> }
> 
> That is, only for 'TARGET_LP64' offloading is supported, and via
> '-foffload-abi=-m64' the 'mkoffload's know that they need to specify
> '-m64'.  For other host flags/ABI configuration, the 'mkoffload's see
> '-foffload-abi=disable' and thus disable offload code generation
> (replacing the current 'if (offload_abi == OFFLOAD_ABI_LP64)' in
> 'mkoffload').
> 
>> In the attached patch
> 
> Yes, that's going in the right direction, thanks!
> 
>> I added another option -foffload-abi-host-opts to specify host abi
>> opts, and leave -foffload-abi to specify if ABI is 32/64 bit which mkoffload can use to
>> enable/disable offloading (as before).
> 
> I'm not sure however, if this additional option is really necessary?
> 
> In case we're not happy to re-purpose the flag name
> '-foffload-abi=[...]', we could also rename that one to
> '-foffload-abi-host-opts=[...]'; the former is not user-exposed, so we
> may change it as necessary.  (Or, in other words, go with your proposed
> '-foffload-abi-host-opts=[...]', but also remove '-foffload-abi=[...]' at
> the same time.)
> 
> 
> I'll be able to spend more time on this in two weeks.

Since we do not support 32 -> 64 bit offload wouldn’t the most pragmatic fix be to recognize -m64 in the nvptx backend (and ignore it)?

Richard 


> 
> Grüße
> Thomas
> 
> 
>> [nvptx] Pass host specific ABI opts from mkoffload.
>> 
>> The patch adds an option -foffload-abi-host-opts, which
>> is set by host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's value
>> to host_compiler.
>> 
>> gcc/ChangeLog:
>>    * common.opt (foffload-abi-host-opts): New option.
>>    * config/aarch64/aarch64.cc (aarch64_offload_options): Set
>>    -foffload-abi-host-opts.
>>    * config/i386/i386-opts.cc (ix86_offload_options): Likewise.
>>    * config/rs6000/rs6000.cc (rs6000_offload_options): Likewise.
>>    * config/nvptx/mkoffload.cc (host_abi_opts): Define.
>>    (compile_native): Append host_abi_opts to argv_obstack.
>>    (main): Handle option -foffload-abi-host-opts.
>>    * lto-wrapper.cc (append_compiler_options): Handle
>>    -foffload-abi-host-opts.
>>    * opts.cc (common_handle_option): Likewise.
>> 
>> Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>
>> 
>> diff --git a/gcc/common.opt b/gcc/common.opt
>> index ea39f87ae71..d1a9efb9513 100644
>> --- a/gcc/common.opt
>> +++ b/gcc/common.opt
>> @@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
>> EnumValue
>> Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
>> 
>> +foffload-abi-host-opts=
>> +Common Driver Joined MissingArgError(option or option=abi missing after %qs)
>> +-foffload-abi-host-opts=<options>=<abi> Specify host abi options.
>> +
>> fomit-frame-pointer
>> Common Var(flag_omit_frame_pointer) Optimization
>> When possible do not generate stack frames.
>> diff --git a/gcc/config/aarch64/aarch64.cc b/gcc/config/aarch64/aarch64.cc
>> index 2ac5a22c848..7418cb1fb69 100644
>> --- a/gcc/config/aarch64/aarch64.cc
>> +++ b/gcc/config/aarch64/aarch64.cc
>> @@ -18999,9 +18999,9 @@ static char *
>> aarch64_offload_options (void)
>> {
>>   if (TARGET_ILP32)
>> -    return xstrdup ("-foffload-abi=ilp32");
>> +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-mabi=ilp32");
>>   else
>> -    return xstrdup ("-foffload-abi=lp64");
>> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-mabi=lp64");
>> }
>> 
>> static struct machine_function *
>> diff --git a/gcc/config/i386/i386-options.cc b/gcc/config/i386/i386-options.cc
>> index 1c8f7835af2..bd960674e5d 100644
>> --- a/gcc/config/i386/i386-options.cc
>> +++ b/gcc/config/i386/i386-options.cc
>> @@ -3669,8 +3669,8 @@ char *
>> ix86_offload_options (void)
>> {
>>   if (TARGET_LP64)
>> -    return xstrdup ("-foffload-abi=lp64");
>> -  return xstrdup ("-foffload-abi=ilp32");
>> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
>> +  return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
>> }
>> 
>> /* Handle "cdecl", "stdcall", "fastcall", "regparm", "thiscall",
>> diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
>> index 503b1abcefd..d5ca2386641 100644
>> --- a/gcc/config/nvptx/mkoffload.cc
>> +++ b/gcc/config/nvptx/mkoffload.cc
>> @@ -61,6 +61,7 @@ static const char *omp_requires_file;
>> static const char *ptx_dumpbase;
>> 
>> enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
>> +const char *host_abi_opts = NULL;
>> 
>> /* Delete tempfiles.  */
>> 
>> @@ -607,17 +608,9 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
>>   obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
>>   obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
>>   obstack_ptr_grow (&argv_obstack, ".c");
>> -  switch (offload_abi)
>> -    {
>> -    case OFFLOAD_ABI_LP64:
>> -      obstack_ptr_grow (&argv_obstack, "-m64");
>> -      break;
>> -    case OFFLOAD_ABI_ILP32:
>> -      obstack_ptr_grow (&argv_obstack, "-m32");
>> -      break;
>> -    default:
>> -      gcc_unreachable ();
>> -    }
>> +  if (!host_abi_opts)
>> +    fatal_error (input_location, "-foffload-abi-host-opts not specified.");
>> +  obstack_ptr_grow (&argv_obstack, host_abi_opts);
>>   obstack_ptr_grow (&argv_obstack, infile);
>>   obstack_ptr_grow (&argv_obstack, "-c");
>>   obstack_ptr_grow (&argv_obstack, "-o");
>> @@ -721,6 +714,8 @@ main (int argc, char **argv)
>>             "unrecognizable argument of option " STR);
>>    }
>> #undef STR
>> +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
>> +    host_abi_opts = argv[i] + strlen ("-foffload-abi-host-opts=");
>>       else if (strcmp (argv[i], "-fopenmp") == 0)
>>    fopenmp = true;
>>       else if (strcmp (argv[i], "-fopenacc") == 0)
>> diff --git a/gcc/config/rs6000/rs6000.cc b/gcc/config/rs6000/rs6000.cc
>> index 0bcc6a2d0ab..decdf49a1f5 100644
>> --- a/gcc/config/rs6000/rs6000.cc
>> +++ b/gcc/config/rs6000/rs6000.cc
>> @@ -17333,9 +17333,9 @@ static char *
>> rs6000_offload_options (void)
>> {
>>   if (TARGET_64BIT)
>> -    return xstrdup ("-foffload-abi=lp64");
>> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
>>   else
>> -    return xstrdup ("-foffload-abi=ilp32");
>> +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
>> }
>> 
>> 
>> diff --git a/gcc/lto-wrapper.cc b/gcc/lto-wrapper.cc
>> index 6bfc96590a5..1ecc4997e5a 100644
>> --- a/gcc/lto-wrapper.cc
>> +++ b/gcc/lto-wrapper.cc
>> @@ -745,6 +745,7 @@ append_compiler_options (obstack *argv_obstack, vec<cl_decoded_option> opts)
>>    case OPT_fopenacc:
>>    case OPT_fopenacc_dim_:
>>    case OPT_foffload_abi_:
>> +    case OPT_foffload_abi_host_opts_:
>>    case OPT_fcf_protection_:
>>    case OPT_fasynchronous_unwind_tables:
>>    case OPT_funwind_tables:
>> diff --git a/gcc/opts.cc b/gcc/opts.cc
>> index 0b7b137c376..79118237ce4 100644
>> --- a/gcc/opts.cc
>> +++ b/gcc/opts.cc
>> @@ -3069,6 +3069,7 @@ common_handle_option (struct gcc_options *opts,
>>       break;
>> 
>>     case OPT_foffload_abi_:
>> +    case OPT_foffload_abi_host_opts_:
>> #ifdef ACCEL_COMPILER
>>       /* Handled in the 'mkoffload's.  */
>> #else

> -----Original Message-----
> From: Richard Biener <rguenther@suse.de>
> Sent: Tuesday, August 13, 2024 10:06 PM
> To: Thomas Schwinge <tschwinge@baylibre.com>
> Cc: Prathamesh Kulkarni <prathameshk@nvidia.com>; Andrew Pinski
> <pinskia@gmail.com>; gcc-patches@gcc.gnu.org; Jakub Jelinek
> <jakub@redhat.com>
> Subject: Re: [nvptx] Pass -m32/-m64 to host_compiler if it has
> multilib support
> 
> External email: Use caution opening links or attachments
> 
> 
> > Am 13.08.2024 um 17:48 schrieb Thomas Schwinge
> <tschwinge@baylibre.com>:
> >
> > Hi Prathamesh!
> >
> > On 2024-08-12T07:50:07+0000, Prathamesh Kulkarni
> <prathameshk@nvidia.com> wrote:
> >>> From: Thomas Schwinge <tschwinge@baylibre.com>
> >>> Sent: Friday, August 9, 2024 12:55 AM
> >
> >>> On 2024-08-08T06:46:25-0700, Andrew Pinski <pinskia@gmail.com>
> wrote:
> >>>> On Thu, Aug 8, 2024 at 6:11 AM Prathamesh Kulkarni
> >>>> <prathameshk@nvidia.com> wrote:
> >>>>> After differing NUM_POLY_INT_COEFFS fix for AArch64/nvptx
> >>> offloading, the following minimal test:
> >>>
> >>> First, thanks for your work on enabling this!  I will say that I
> had
> >>> the plan to re-engage with Nvidia to hire us (as initial
> >>> implementors of GCC/nvptx offloading) to make AArch64/nvptx
> >>> offloading work, but now that Nvidia has its own GCC team, that's
> >>> great that you're able to work on this yourself!  :-)
> >>>
> >>> Please CC me for GCC/nvptx issues for (at least potentially...)
> >>> faster response times.
> >> Thanks, will do 😊
> >
> > Heh, so much for "potentially": I'm not able to spend a lot of time
> on
> > this right now, as I shall soon be out of office.  Quickly:
> >
> >>>>> compiled with -fopenmp -foffload=nvptx-none now fails with:
> >>>>> gcc: error: unrecognized command-line option '-m64'
> >>>>> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit
> >>> status compilation terminated.
> >>>
> >>> Heh.  Yeah...
> >>>
> >>>>> As mentioned in RFC email, this happens because
> >>>>> nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host
> >>>>> compiler
> >>> depending on whether offload_abi is OFFLOAD_ABI_LP64 or
> >>> OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these
> >>> options.
> >
> >>> So, my idea is: instead of the current strategy that the host
> >>> 'TARGET_OFFLOAD_OPTIONS' synthesizes '-foffload-abi=lp64' etc.,
> >>> which the 'mkoffload's then interpret and re-synthesize '-m64'
> etc.
> >>> -- how about we instead directly tell the 'mkoffload's the
> relevant
> >>> ABI options?  That is, 'TARGET_OFFLOAD_OPTIONS' instead
> synthesizes
> >>> '- foffload-abi=-m64'
> >>> etc., which the 'mkoffload's can then readily use.  Could you
> please
> >>> give that a try, and/or does anyone see any issues with that
> approach?
> >>>
> >>> And use something like '-foffload-abi=disable' to replace the
> current:
> >>>
> >>>    /* PR libgomp/65099: Currently, we only support offloading in
> 64-
> >>> bit
> >>>       configurations.  */
> >>>    if (offload_abi == OFFLOAD_ABI_LP64)
> >>>      {
> >>>
> >>> (As discussed before, this should be done differently altogether,
> >>> but that's for another day.)
> >> Sorry, I don't quite follow. Currently we enable offloading if
> >> offload_abi == OFFLOAD_ABI_LP64, which is synthesized from
> >> -foffload-abi=lp64. If we change -foffload-abi to instead specify
> >> host-specific ABI opts, I guess mkoffload will still need to
> somehow
> >> figure out which ABI is used, so it can disable offloading for 32-
> bit
> >> ? I suppose we could adjust TARGET_OFFLOAD_OPTIONS for each host to
> pass -foffload-abi=disable if TARGET_ILP32 is set and offload target
> is nvptx, but not sure if that'd be correct ?
> >
> > Basically, yes.  My idea was that all 'TARGET_OFFLOAD_OPTIONS'
> > implementations return either the correct host flags to be used by
> the
> > 'mkoffload's (the case that offloading is supported for the current
> > host flags/ABI configuration), or otherwise return '-foffload-
> abi=disable'.
> > For example (untested):
> >
> >> char *
> >> ix86_offload_options (void)
> >> {
> >>   if (TARGET_LP64)
> >> -    return xstrdup ("-foffload-abi=lp64");
> >> +    return xstrdup ("-foffload-abi=-m64");
> >> -  return xstrdup ("-foffload-abi=ilp32");
> >> +  return xstrdup ("-foffload-abi=disable");
> >> }
> >
> > That is, only for 'TARGET_LP64' offloading is supported, and via
> > '-foffload-abi=-m64' the 'mkoffload's know that they need to specify
> > '-m64'.  For other host flags/ABI configuration, the 'mkoffload's
> see
> > '-foffload-abi=disable' and thus disable offload code generation
> > (replacing the current 'if (offload_abi == OFFLOAD_ABI_LP64)' in
> > 'mkoffload').
> >
> >> In the attached patch
> >
> > Yes, that's going in the right direction, thanks!
> >
> >> I added another option -foffload-abi-host-opts to specify host abi
> >> opts, and leave -foffload-abi to specify if ABI is 32/64 bit which
> >> mkoffload can use to enable/disable offloading (as before).
> >
> > I'm not sure however, if this additional option is really necessary?
Well, my concern was if that'd change the behavior for TARGET_ILP32 ?
IIUC, currently for -foffload-abi=ilp32, mkoffload will create empty C file
for ptx_cfile_name (instead of munged ptx assembly since offloading will be disabled),
and pass that to host compiler with -m32 option (in compile_native).

If we change -foffload-abi to specify ABI host opts, and pass -foffload-abi=disable 
for TARGET_ILP32 in TARGET_OFFLOAD_OPTIONS, mkoffload will no longer be able to
pass 32-bit ABI opts to host compiler, which may result in linker error (arch mismatch?)
if the host object files are 32-bit ABI and xnvptx-none.o is 64-bit (assuming the host
compiler is configured to generate 64-bit code-gen by default) ?

So, I thought to add another option -foffload-abi-host-opts to pass host-specific ABI opts,
and keep -foffload-abi as-is to infer ABI type for enabling/disabling offloading.
> >
> > In case we're not happy to re-purpose the flag name
> > '-foffload-abi=[...]', we could also rename that one to
> > '-foffload-abi-host-opts=[...]'; the former is not user-exposed, so
> we
> > may change it as necessary.  (Or, in other words, go with your
> > proposed '-foffload-abi-host-opts=[...]', but also remove
> > '-foffload-abi=[...]' at the same time.)
> >
> >
> > I'll be able to spend more time on this in two weeks.
> 
> Since we do not support 32 -> 64 bit offload wouldn’t the most
> pragmatic fix be to recognize -m64 in the nvptx backend (and ignore
> it)?
I think nvptx already supports m64 and ignores it.
From nvptx.opt:

m64
Target RejectNegative Mask(ABI64)
Ignored, but preserved for backward compatibility.  Only 64-bit ABI is
supported.

Thanks,
Prathamesh
> 
> Richard
> 
> 
> >
> > Grüße
> > Thomas
> >
> >
> >> [nvptx] Pass host specific ABI opts from mkoffload.
> >>
> >> The patch adds an option -foffload-abi-host-opts, which is set by
> >> host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's
> value
> >> to host_compiler.
> >>
> >> gcc/ChangeLog:
> >>    * common.opt (foffload-abi-host-opts): New option.
> >>    * config/aarch64/aarch64.cc (aarch64_offload_options): Set
> >>    -foffload-abi-host-opts.
> >>    * config/i386/i386-opts.cc (ix86_offload_options): Likewise.
> >>    * config/rs6000/rs6000.cc (rs6000_offload_options): Likewise.
> >>    * config/nvptx/mkoffload.cc (host_abi_opts): Define.
> >>    (compile_native): Append host_abi_opts to argv_obstack.
> >>    (main): Handle option -foffload-abi-host-opts.
> >>    * lto-wrapper.cc (append_compiler_options): Handle
> >>    -foffload-abi-host-opts.
> >>    * opts.cc (common_handle_option): Likewise.
> >>
> >> Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>
> >>
> >> diff --git a/gcc/common.opt b/gcc/common.opt index
> >> ea39f87ae71..d1a9efb9513 100644
> >> --- a/gcc/common.opt
> >> +++ b/gcc/common.opt
> >> @@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32)
> >> Value(OFFLOAD_ABI_ILP32) EnumValue
> >> Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
> >>
> >> +foffload-abi-host-opts=
> >> +Common Driver Joined MissingArgError(option or option=abi missing
> >> +after %qs) -foffload-abi-host-opts=<options>=<abi> Specify host
> abi options.
> >> +
> >> fomit-frame-pointer
> >> Common Var(flag_omit_frame_pointer) Optimization When possible do
> not
> >> generate stack frames.
> >> diff --git a/gcc/config/aarch64/aarch64.cc
> >> b/gcc/config/aarch64/aarch64.cc index 2ac5a22c848..7418cb1fb69
> 100644
> >> --- a/gcc/config/aarch64/aarch64.cc
> >> +++ b/gcc/config/aarch64/aarch64.cc
> >> @@ -18999,9 +18999,9 @@ static char * aarch64_offload_options
> (void)
> >> {
> >>   if (TARGET_ILP32)
> >> -    return xstrdup ("-foffload-abi=ilp32");
> >> +    return xstrdup ("-foffload-abi=ilp32
> >> + -foffload-abi-host-opts=-mabi=ilp32");
> >>   else
> >> -    return xstrdup ("-foffload-abi=lp64");
> >> +    return xstrdup ("-foffload-abi=lp64
> >> + -foffload-abi-host-opts=-mabi=lp64");
> >> }
> >>
> >> static struct machine_function *
> >> diff --git a/gcc/config/i386/i386-options.cc
> >> b/gcc/config/i386/i386-options.cc index 1c8f7835af2..bd960674e5d
> >> 100644
> >> --- a/gcc/config/i386/i386-options.cc
> >> +++ b/gcc/config/i386/i386-options.cc
> >> @@ -3669,8 +3669,8 @@ char *
> >> ix86_offload_options (void)
> >> {
> >>   if (TARGET_LP64)
> >> -    return xstrdup ("-foffload-abi=lp64");
> >> -  return xstrdup ("-foffload-abi=ilp32");
> >> +    return xstrdup ("-foffload-abi=lp64
> >> + -foffload-abi-host-opts=-m64");  return xstrdup
> >> + ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
> >> }
> >>
> >> /* Handle "cdecl", "stdcall", "fastcall", "regparm", "thiscall",
> diff
> >> --git a/gcc/config/nvptx/mkoffload.cc
> b/gcc/config/nvptx/mkoffload.cc
> >> index 503b1abcefd..d5ca2386641 100644
> >> --- a/gcc/config/nvptx/mkoffload.cc
> >> +++ b/gcc/config/nvptx/mkoffload.cc
> >> @@ -61,6 +61,7 @@ static const char *omp_requires_file; static
> const
> >> char *ptx_dumpbase;
> >>
> >> enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
> >> +const char *host_abi_opts = NULL;
> >>
> >> /* Delete tempfiles.  */
> >>
> >> @@ -607,17 +608,9 @@ compile_native (const char *infile, const char
> *outfile, const char *compiler,
> >>   obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
> >>   obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
> >>   obstack_ptr_grow (&argv_obstack, ".c");
> >> -  switch (offload_abi)
> >> -    {
> >> -    case OFFLOAD_ABI_LP64:
> >> -      obstack_ptr_grow (&argv_obstack, "-m64");
> >> -      break;
> >> -    case OFFLOAD_ABI_ILP32:
> >> -      obstack_ptr_grow (&argv_obstack, "-m32");
> >> -      break;
> >> -    default:
> >> -      gcc_unreachable ();
> >> -    }
> >> +  if (!host_abi_opts)
> >> +    fatal_error (input_location, "-foffload-abi-host-opts not
> >> + specified.");  obstack_ptr_grow (&argv_obstack, host_abi_opts);
> >>   obstack_ptr_grow (&argv_obstack, infile);
> >>   obstack_ptr_grow (&argv_obstack, "-c");
> >>   obstack_ptr_grow (&argv_obstack, "-o"); @@ -721,6 +714,8 @@ main
> >> (int argc, char **argv)
> >>             "unrecognizable argument of option " STR);
> >>    }
> >> #undef STR
> >> +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> >> +    host_abi_opts = argv[i] + strlen ("-foffload-abi-host-opts=");
> >>       else if (strcmp (argv[i], "-fopenmp") == 0)
> >>    fopenmp = true;
> >>       else if (strcmp (argv[i], "-fopenacc") == 0) diff --git
> >> a/gcc/config/rs6000/rs6000.cc b/gcc/config/rs6000/rs6000.cc index
> >> 0bcc6a2d0ab..decdf49a1f5 100644
> >> --- a/gcc/config/rs6000/rs6000.cc
> >> +++ b/gcc/config/rs6000/rs6000.cc
> >> @@ -17333,9 +17333,9 @@ static char * rs6000_offload_options (void)
> {
> >>   if (TARGET_64BIT)
> >> -    return xstrdup ("-foffload-abi=lp64");
> >> +    return xstrdup ("-foffload-abi=lp64
> >> + -foffload-abi-host-opts=-m64");
> >>   else
> >> -    return xstrdup ("-foffload-abi=ilp32");
> >> +    return xstrdup ("-foffload-abi=ilp32
> >> + -foffload-abi-host-opts=-m32");
> >> }
> >>
> >>
> >> diff --git a/gcc/lto-wrapper.cc b/gcc/lto-wrapper.cc index
> >> 6bfc96590a5..1ecc4997e5a 100644
> >> --- a/gcc/lto-wrapper.cc
> >> +++ b/gcc/lto-wrapper.cc
> >> @@ -745,6 +745,7 @@ append_compiler_options (obstack *argv_obstack,
> vec<cl_decoded_option> opts)
> >>    case OPT_fopenacc:
> >>    case OPT_fopenacc_dim_:
> >>    case OPT_foffload_abi_:
> >> +    case OPT_foffload_abi_host_opts_:
> >>    case OPT_fcf_protection_:
> >>    case OPT_fasynchronous_unwind_tables:
> >>    case OPT_funwind_tables:
> >> diff --git a/gcc/opts.cc b/gcc/opts.cc index
> 0b7b137c376..79118237ce4
> >> 100644
> >> --- a/gcc/opts.cc
> >> +++ b/gcc/opts.cc
> >> @@ -3069,6 +3069,7 @@ common_handle_option (struct gcc_options
> *opts,
> >>       break;
> >>
> >>     case OPT_foffload_abi_:
> >> +    case OPT_foffload_abi_host_opts_:
> >> #ifdef ACCEL_COMPILER
> >>       /* Handled in the 'mkoffload's.  */ #else

On Thu, 8 Aug 2024, Prathamesh Kulkarni wrote:

> Hi Richard,
> After differing NUM_POLY_INT_COEFFS fix for AArch64/nvptx offloading, the following minimal test:
> 
> int main()
> {
>   int x;
>   #pragma omp target map(x)
>     x = 5;
>   return x;
> }
> 
> compiled with -fopenmp -foffload=nvptx-none now fails with:
> gcc: error: unrecognized command-line option '-m64'
> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit status compilation terminated.
> 
> As mentioned in RFC email, this happens because nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host compiler depending on whether
> offload_abi is OFFLOAD_ABI_LP64 or OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these options.
> 
> Based on your suggestion in: https://gcc.gnu.org/pipermail/gcc/2024-July/244470.html,
> The attached patch generates new macro HOST_MULTILIB derived from $enable_as_accelerator_for, and in mkoffload.cc it gates passing -m32/-m64
> to host_compiler on HOST_MULTILIB. I verified that the macro is set to 0 for aarch64 host (and thus avoids above unrecognized command line option error),
> and is set to 1 for x86_64 host.
> 
> Does the patch look OK ?

The patch looks reasonable to me.

Thanks,
Richard.

> Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>
> 
> Thanks,
> Prathamesh
>

Hi!

On 2024-08-16T15:36:29+0000, Prathamesh Kulkarni <prathameshk@nvidia.com> wrote:
>> > Am 13.08.2024 um 17:48 schrieb Thomas Schwinge
>> <tschwinge@baylibre.com>:
>> > On 2024-08-12T07:50:07+0000, Prathamesh Kulkarni
>> <prathameshk@nvidia.com> wrote:
>> >>> From: Thomas Schwinge <tschwinge@baylibre.com>
>> >>> Sent: Friday, August 9, 2024 12:55 AM
>> >
>> >>> On 2024-08-08T06:46:25-0700, Andrew Pinski <pinskia@gmail.com>
>> wrote:
>> >>>> On Thu, Aug 8, 2024 at 6:11 AM Prathamesh Kulkarni
>> >>>> <prathameshk@nvidia.com> wrote:
>> >>>>> compiled with -fopenmp -foffload=nvptx-none now fails with:
>> >>>>> gcc: error: unrecognized command-line option '-m64'
>> >>>>> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit
>> >>> status compilation terminated.
>> >>>
>> >>> Heh.  Yeah...
>> >>>
>> >>>>> As mentioned in RFC email, this happens because
>> >>>>> nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host
>> >>>>> compiler
>> >>> depending on whether offload_abi is OFFLOAD_ABI_LP64 or
>> >>> OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these
>> >>> options.
>> >
>> >>> So, my idea is: instead of the current strategy that the host
>> >>> 'TARGET_OFFLOAD_OPTIONS' synthesizes '-foffload-abi=lp64' etc.,
>> >>> which the 'mkoffload's then interpret and re-synthesize '-m64' etc.
>> >>> -- how about we instead directly tell the 'mkoffload's the relevant
>> >>> ABI options?  That is, 'TARGET_OFFLOAD_OPTIONS' instead synthesizes
>> >>> '-foffload-abi=-m64'
>> >>> etc., which the 'mkoffload's can then readily use.  Could you please
>> >>> give that a try, and/or does anyone see any issues with that approach?
>> >>>
>> >>> And use something like '-foffload-abi=disable' to replace the current:
>> >>>
>> >>>    /* PR libgomp/65099: Currently, we only support offloading in 64-bit
>> >>>       configurations.  */
>> >>>    if (offload_abi == OFFLOAD_ABI_LP64)
>> >>>      {
>> >>>
>> >>> (As discussed before, this should be done differently altogether,
>> >>> but that's for another day.)
>> >> Sorry, I don't quite follow. Currently we enable offloading if
>> >> offload_abi == OFFLOAD_ABI_LP64, which is synthesized from
>> >> -foffload-abi=lp64. If we change -foffload-abi to instead specify
>> >> host-specific ABI opts, I guess mkoffload will still need to somehow
>> >> figure out which ABI is used, so it can disable offloading for 32-bit
>> >> ? I suppose we could adjust TARGET_OFFLOAD_OPTIONS for each host to
>> pass -foffload-abi=disable if TARGET_ILP32 is set and offload target
>> is nvptx, but not sure if that'd be correct ?
>> >
>> > Basically, yes.  My idea was that all 'TARGET_OFFLOAD_OPTIONS'
>> > implementations return either the correct host flags to be used by the
>> > 'mkoffload's (the case that offloading is supported for the current
>> > host flags/ABI configuration), or otherwise return '-foffload-abi=disable'.

Oh..., you're right of course: we do need to continue to tell the
'mkoffload's which kind of offload code to generate!  My bad...

>> >> I added another option -foffload-abi-host-opts to specify host abi
>> >> opts, and leave -foffload-abi to specify if ABI is 32/64 bit which
>> >> mkoffload can use to enable/disable offloading (as before).
>> >
>> > I'm not sure however, if this additional option is really necessary?
> Well, my concern was if that'd change the behavior for TARGET_ILP32 ?
> IIUC, currently for -foffload-abi=ilp32, mkoffload will create empty C file
> for ptx_cfile_name (instead of munged ptx assembly since offloading will be disabled),
> and pass that to host compiler with -m32 option (in compile_native).
>
> If we change -foffload-abi to specify ABI host opts, and pass -foffload-abi=disable 
> for TARGET_ILP32 in TARGET_OFFLOAD_OPTIONS, mkoffload will no longer be able to
> pass 32-bit ABI opts to host compiler, which may result in linker error (arch mismatch?)
> if the host object files are 32-bit ABI and xnvptx-none.o is 64-bit (assuming the host
> compiler is configured to generate 64-bit code-gen by default) ?
>
> So, I thought to add another option -foffload-abi-host-opts to pass host-specific ABI opts,
> and keep -foffload-abi as-is to infer ABI type for enabling/disabling offloading.

Quite right, yes.

>> -----Original Message-----
>> From: Richard Biener <rguenther@suse.de>
>> Sent: Tuesday, August 13, 2024 10:06 PM

>> Since we do not support 32 -> 64 bit offload

We don't -- but it's generally possible.  As Tobias recently educated
me, the OpenMP specification explicitly does *not* require matching
host 'sizeof (void *)' and device 'sizeof (void *)'.

At the LLVM workshop at ISC High Performance 2024 there was a (short)
presentation of someone who did LLVM offloading from host to a different
architecture, and from there again to a yet different architecture.  Heh!

Anyway:

>> wouldn’t the most
>> pragmatic fix be to recognize -m64 in the nvptx backend (and ignore
>> it)?

> I think nvptx already supports m64 and ignores it.
> From nvptx.opt:
>
> m64
> Target RejectNegative Mask(ABI64)
> Ignored, but preserved for backward compatibility.  Only 64-bit ABI is
> supported.

Right, but that's also not the problem here: the problem is that
'mkoffload' puts '-m64' onto the *host* compiler command line (for
embedding the offload image), which in case of aarch64 isn't the right
thing to do; just happened to do the right thing for x86_64 and
powerpc64le.


Prathamesh's proposed patch:

> [nvptx] Pass host specific ABI opts from mkoffload.
>
> The patch adds an option -foffload-abi-host-opts, which
> is set by host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's value
> to host_compiler.

ACK, conceptually.

> --- a/gcc/common.opt
> +++ b/gcc/common.opt
> @@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
>  EnumValue
>  Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
>  
> +foffload-abi-host-opts=
> +Common Driver Joined MissingArgError(option or option=abi missing after %qs)
> +-foffload-abi-host-opts=<options>=<abi> Specify host abi options.
> +

Here, 'option or option=abi' and '<options>=<abi>' should be just
'options' and '<options>', right?  And, TAB between
'-foffload-abi-host-opts=<options>' and its help text.  And upper-case
ABI.

> --- a/gcc/config/aarch64/aarch64.cc
> +++ b/gcc/config/aarch64/aarch64.cc
> @@ -18999,9 +18999,9 @@ static char *
>  aarch64_offload_options (void)
>  {
>    if (TARGET_ILP32)
> -    return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-mabi=ilp32");
>    else
> -    return xstrdup ("-foffload-abi=lp64");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-mabi=lp64");
>  }

As none of the current offload compilers is set up of ILP32, I suggest we
continue to pass '-foffload-abi=ilp32' without
'-foffload-abi-host-opts=[...]' -- the 'mkoffload's in that case should
get to the point where the latter is used.

> --- a/gcc/config/i386/i386-options.cc
> +++ b/gcc/config/i386/i386-options.cc
> @@ -3669,8 +3669,8 @@ char *
>  ix86_offload_options (void)
>  {
>    if (TARGET_LP64)
> -    return xstrdup ("-foffload-abi=lp64");
> -  return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
> +  return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
>  }

Likewise.

> --- a/gcc/config/rs6000/rs6000.cc
> +++ b/gcc/config/rs6000/rs6000.cc
> @@ -17333,9 +17333,9 @@ static char *
>  rs6000_offload_options (void)
>  {
>    if (TARGET_64BIT)
> -    return xstrdup ("-foffload-abi=lp64");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
>    else
> -    return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
>  }

Likewise.

> --- a/gcc/config/nvptx/mkoffload.cc
> +++ b/gcc/config/nvptx/mkoffload.cc
> @@ -61,6 +61,7 @@ static const char *omp_requires_file;
>  static const char *ptx_dumpbase;
>  
>  enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
> +const char *host_abi_opts = NULL;

Should this be 'offload_abi_host_opts' for similarity with the
command-line option?

> @@ -607,17 +608,9 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
>    obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
>    obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
>    obstack_ptr_grow (&argv_obstack, ".c");
> -  switch (offload_abi)
> -    {
> -    case OFFLOAD_ABI_LP64:
> -      obstack_ptr_grow (&argv_obstack, "-m64");
> -      break;
> -    case OFFLOAD_ABI_ILP32:
> -      obstack_ptr_grow (&argv_obstack, "-m32");
> -      break;
> -    default:
> -      gcc_unreachable ();
> -    }
> +  if (!host_abi_opts)
> +    fatal_error (input_location, "-foffload-abi-host-opts not specified.");

I know we're not doing that consistently, but please use
'%<-foffload-abi-host-opts%>'.

> +  obstack_ptr_grow (&argv_obstack, host_abi_opts);
>    obstack_ptr_grow (&argv_obstack, infile);
>    obstack_ptr_grow (&argv_obstack, "-c");
>    obstack_ptr_grow (&argv_obstack, "-o");
> @@ -721,6 +714,8 @@ main (int argc, char **argv)
>  			 "unrecognizable argument of option " STR);
>  	}
>  #undef STR
> +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> +	host_abi_opts = argv[i] + strlen ("-foffload-abi-host-opts=");

The option parsing in the 'mkoffload's is ad-hoc (not using the proepr
GCC infrastructure; which I'd like to change at some point in time...),
but let's please catch the case that '-foffload-abi-host-opts=[...]'
appears more than once (which could be necessary in certain
configurations, to produce ABI-compatible host code?).  Not necessary to
implement that right now: for now, it'll be fine to 'fatal_error' if
running into a second '-foffload-abi-host-opts=[...]'.

Generally, likewise need to adjust 'gcc/config/gcn/mkoffload.cc'.  I can
test this, or co-author, if you'd like.

> --- a/gcc/lto-wrapper.cc
> +++ b/gcc/lto-wrapper.cc

Don't we also need to adjust 'merge_and_complain':

    case OPT_foffload_abi_:
      if (existing_opt == -1)
        decoded_options.safe_push (*foption);
      else if (foption->value != decoded_options[existing_opt].value)
        fatal_error (input_location,
                     "option %s not used consistently in all LTO input"
                     " files", foption->orig_option_with_args_text);
      break;

> @@ -745,6 +745,7 @@ append_compiler_options (obstack *argv_obstack, vec<cl_decoded_option> opts)
>  	case OPT_fopenacc:
>  	case OPT_fopenacc_dim_:
>  	case OPT_foffload_abi_:
> +	case OPT_foffload_abi_host_opts_:
>  	case OPT_fcf_protection_:
>  	case OPT_fasynchronous_unwind_tables:
>  	case OPT_funwind_tables:

Per my quick reading of the code, that should be correct.

> --- a/gcc/opts.cc
> +++ b/gcc/opts.cc
> @@ -3069,6 +3069,7 @@ common_handle_option (struct gcc_options *opts,
>        break;
>  
>      case OPT_foffload_abi_:
> +    case OPT_foffload_abi_host_opts_:
>  #ifdef ACCEL_COMPILER
>        /* Handled in the 'mkoffload's.  */
>  #else
|        error_at (loc, "%<-foffload-abi%> option can be specified only for "
|              "offload compiler");
|  #endif

Please adjust the diagnostic.  Surely the original option string will be
available for use with '%qs'.


Grüße
 Thomas

> -----Original Message-----
> From: Thomas Schwinge <tschwinge@baylibre.com>
> Sent: Friday, September 6, 2024 2:31 PM
> To: Prathamesh Kulkarni <prathameshk@nvidia.com>; Richard Biener
> <rguenther@suse.de>
> Cc: Andrew Pinski <pinskia@gmail.com>; gcc-patches@gcc.gnu.org; Jakub
> Jelinek <jakub@redhat.com>
> Subject: RE: [nvptx] Pass -m32/-m64 to host_compiler if it has
> multilib support
> 
> External email: Use caution opening links or attachments
> 
> 
> Hi!
Hi Thomas,
Thanks for the review!
> 
> On 2024-08-16T15:36:29+0000, Prathamesh Kulkarni
> <prathameshk@nvidia.com> wrote:
> >> > Am 13.08.2024 um 17:48 schrieb Thomas Schwinge
> >> <tschwinge@baylibre.com>:
> >> > On 2024-08-12T07:50:07+0000, Prathamesh Kulkarni
> >> <prathameshk@nvidia.com> wrote:
> >> >>> From: Thomas Schwinge <tschwinge@baylibre.com>
> >> >>> Sent: Friday, August 9, 2024 12:55 AM
> >> >
> >> >>> On 2024-08-08T06:46:25-0700, Andrew Pinski <pinskia@gmail.com>
> >> wrote:
> >> >>>> On Thu, Aug 8, 2024 at 6:11 AM Prathamesh Kulkarni
> >> >>>> <prathameshk@nvidia.com> wrote:
> >> >>>>> compiled with -fopenmp -foffload=nvptx-none now fails with:
> >> >>>>> gcc: error: unrecognized command-line option '-m64'
> >> >>>>> nvptx mkoffload: fatal error: ../install/bin/gcc returned 1
> >> >>>>> exit
> >> >>> status compilation terminated.
> >> >>>
> >> >>> Heh.  Yeah...
> >> >>>
> >> >>>>> As mentioned in RFC email, this happens because
> >> >>>>> nvptx/mkoffload.cc:compile_native passes -m64/-m32 to host
> >> >>>>> compiler
> >> >>> depending on whether offload_abi is OFFLOAD_ABI_LP64 or
> >> >>> OFFLOAD_ABI_ILP32, and aarch64 backend doesn't recognize these
> >> >>> options.
> >> >
> >> >>> So, my idea is: instead of the current strategy that the host
> >> >>> 'TARGET_OFFLOAD_OPTIONS' synthesizes '-foffload-abi=lp64' etc.,
> >> >>> which the 'mkoffload's then interpret and re-synthesize '-m64'
> etc.
> >> >>> -- how about we instead directly tell the 'mkoffload's the
> >> >>> relevant ABI options?  That is, 'TARGET_OFFLOAD_OPTIONS'
> instead
> >> >>> synthesizes '-foffload-abi=-m64'
> >> >>> etc., which the 'mkoffload's can then readily use.  Could you
> >> >>> please give that a try, and/or does anyone see any issues with
> that approach?
> >> >>>
> >> >>> And use something like '-foffload-abi=disable' to replace the
> current:
> >> >>>
> >> >>>    /* PR libgomp/65099: Currently, we only support offloading
> in 64-bit
> >> >>>       configurations.  */
> >> >>>    if (offload_abi == OFFLOAD_ABI_LP64)
> >> >>>      {
> >> >>>
> >> >>> (As discussed before, this should be done differently
> altogether,
> >> >>> but that's for another day.)
> >> >> Sorry, I don't quite follow. Currently we enable offloading if
> >> >> offload_abi == OFFLOAD_ABI_LP64, which is synthesized from
> >> >> -foffload-abi=lp64. If we change -foffload-abi to instead
> specify
> >> >> host-specific ABI opts, I guess mkoffload will still need to
> >> >> somehow figure out which ABI is used, so it can disable
> offloading
> >> >> for 32-bit ? I suppose we could adjust TARGET_OFFLOAD_OPTIONS
> for
> >> >> each host to
> >> pass -foffload-abi=disable if TARGET_ILP32 is set and offload
> target
> >> is nvptx, but not sure if that'd be correct ?
> >> >
> >> > Basically, yes.  My idea was that all 'TARGET_OFFLOAD_OPTIONS'
> >> > implementations return either the correct host flags to be used
> by
> >> > the 'mkoffload's (the case that offloading is supported for the
> >> > current host flags/ABI configuration), or otherwise return '-
> foffload-abi=disable'.
> 
> Oh..., you're right of course: we do need to continue to tell the
> 'mkoffload's which kind of offload code to generate!  My bad...
> 
> >> >> I added another option -foffload-abi-host-opts to specify host
> abi
> >> >> opts, and leave -foffload-abi to specify if ABI is 32/64 bit
> which
> >> >> mkoffload can use to enable/disable offloading (as before).
> >> >
> >> > I'm not sure however, if this additional option is really
> necessary?
> > Well, my concern was if that'd change the behavior for TARGET_ILP32
> ?
> > IIUC, currently for -foffload-abi=ilp32, mkoffload will create empty
> C
> > file for ptx_cfile_name (instead of munged ptx assembly since
> > offloading will be disabled), and pass that to host compiler with -
> m32 option (in compile_native).
> >
> > If we change -foffload-abi to specify ABI host opts, and pass
> > -foffload-abi=disable for TARGET_ILP32 in TARGET_OFFLOAD_OPTIONS,
> > mkoffload will no longer be able to pass 32-bit ABI opts to host
> > compiler, which may result in linker error (arch mismatch?) if the
> > host object files are 32-bit ABI and xnvptx-none.o is 64-bit
> (assuming the host compiler is configured to generate 64-bit code-gen
> by default) ?
> >
> > So, I thought to add another option -foffload-abi-host-opts to pass
> > host-specific ABI opts, and keep -foffload-abi as-is to infer ABI
> type for enabling/disabling offloading.
> 
> Quite right, yes.
> 
> >> -----Original Message-----
> >> From: Richard Biener <rguenther@suse.de>
> >> Sent: Tuesday, August 13, 2024 10:06 PM
> 
> >> Since we do not support 32 -> 64 bit offload
> 
> We don't -- but it's generally possible.  As Tobias recently educated
> me, the OpenMP specification explicitly does *not* require matching
> host 'sizeof (void *)' and device 'sizeof (void *)'.
> 
> At the LLVM workshop at ISC High Performance 2024 there was a (short)
> presentation of someone who did LLVM offloading from host to a
> different architecture, and from there again to a yet different
> architecture.  Heh!
> 
> Anyway:
> 
> >> wouldn’t the most
> >> pragmatic fix be to recognize -m64 in the nvptx backend (and ignore
> >> it)?
> 
> > I think nvptx already supports m64 and ignores it.
> > From nvptx.opt:
> >
> > m64
> > Target RejectNegative Mask(ABI64)
> > Ignored, but preserved for backward compatibility.  Only 64-bit ABI
> is
> > supported.
> 
> Right, but that's also not the problem here: the problem is that
> 'mkoffload' puts '-m64' onto the *host* compiler command line (for
> embedding the offload image), which in case of aarch64 isn't the right
> thing to do; just happened to do the right thing for x86_64 and
> powerpc64le.
> 
> 
> Prathamesh's proposed patch:
> 
> > [nvptx] Pass host specific ABI opts from mkoffload.
> >
> > The patch adds an option -foffload-abi-host-opts, which is set by
> host
> > in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's value to
> > host_compiler.
> 
> ACK, conceptually.
> 
> > --- a/gcc/common.opt
> > +++ b/gcc/common.opt
> > @@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32)
> > Value(OFFLOAD_ABI_ILP32)  EnumValue
> >  Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
> >
> > +foffload-abi-host-opts=
> > +Common Driver Joined MissingArgError(option or option=abi missing
> > +after %qs) -foffload-abi-host-opts=<options>=<abi> Specify host abi
> options.
> > +
> 
> Here, 'option or option=abi' and '<options>=<abi>' should be just
> 'options' and '<options>', right?  And, TAB between '-foffload-abi-
> host-opts=<options>' and its help text.  And upper-case ABI.
Yes right, sorry. Fixed in the attached patch.
> 
> > --- a/gcc/config/aarch64/aarch64.cc
> > +++ b/gcc/config/aarch64/aarch64.cc
> > @@ -18999,9 +18999,9 @@ static char *
> >  aarch64_offload_options (void)
> >  {
> >    if (TARGET_ILP32)
> > -    return xstrdup ("-foffload-abi=ilp32");
> > +    return xstrdup ("-foffload-abi=ilp32
> > + -foffload-abi-host-opts=-mabi=ilp32");
> >    else
> > -    return xstrdup ("-foffload-abi=lp64");
> > +    return xstrdup ("-foffload-abi=lp64
> > + -foffload-abi-host-opts=-mabi=lp64");
> >  }
> 
> As none of the current offload compilers is set up of ILP32, I suggest
> we continue to pass '-foffload-abi=ilp32' without '-foffload-abi-host-
> opts=[...]' -- the 'mkoffload's in that case should get to the point
> where the latter is used.
Um, would that still possibly result in arch mismatch for host objects and xnvptx-none.o if we don't pass host ABI opts for ILP32 ?
For eg, if the host compiler defaults to 64-bit code-gen (and user requests for 32-bit code gen on host), and we avoid passing host ABI opts for -foffload-abi=ilp32,
it will generate 64-bit xnvptx-none.o (corresponding to empty ptx_cfile_name), while rest of the host objects will be 32-bit, or am I misunderstanding ?

The attached patch avoids passing -foffload-abi-host-opts if -foffload-abi=ilp32.
> 
> > --- a/gcc/config/i386/i386-options.cc
> > +++ b/gcc/config/i386/i386-options.cc
> > @@ -3669,8 +3669,8 @@ char *
> >  ix86_offload_options (void)
> >  {
> >    if (TARGET_LP64)
> > -    return xstrdup ("-foffload-abi=lp64");
> > -  return xstrdup ("-foffload-abi=ilp32");
> > +    return xstrdup ("-foffload-abi=lp64
> > + -foffload-abi-host-opts=-m64");  return xstrdup
> > + ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
> >  }
> 
> Likewise.
> 
> > --- a/gcc/config/rs6000/rs6000.cc
> > +++ b/gcc/config/rs6000/rs6000.cc
> > @@ -17333,9 +17333,9 @@ static char *
> >  rs6000_offload_options (void)
> >  {
> >    if (TARGET_64BIT)
> > -    return xstrdup ("-foffload-abi=lp64");
> > +    return xstrdup ("-foffload-abi=lp64
> > + -foffload-abi-host-opts=-m64");
> >    else
> > -    return xstrdup ("-foffload-abi=ilp32");
> > +    return xstrdup ("-foffload-abi=ilp32
> > + -foffload-abi-host-opts=-m32");
> >  }
> 
> Likewise.
> 
> > --- a/gcc/config/nvptx/mkoffload.cc
> > +++ b/gcc/config/nvptx/mkoffload.cc
> > @@ -61,6 +61,7 @@ static const char *omp_requires_file;  static
> const
> > char *ptx_dumpbase;
> >
> >  enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
> > +const char *host_abi_opts = NULL;
> 
> Should this be 'offload_abi_host_opts' for similarity with the
> command-line option?
Fixed, thanks.
> 
> > @@ -607,17 +608,9 @@ compile_native (const char *infile, const char
> *outfile, const char *compiler,
> >    obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
> >    obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
> >    obstack_ptr_grow (&argv_obstack, ".c");
> > -  switch (offload_abi)
> > -    {
> > -    case OFFLOAD_ABI_LP64:
> > -      obstack_ptr_grow (&argv_obstack, "-m64");
> > -      break;
> > -    case OFFLOAD_ABI_ILP32:
> > -      obstack_ptr_grow (&argv_obstack, "-m32");
> > -      break;
> > -    default:
> > -      gcc_unreachable ();
> > -    }
> > +  if (!host_abi_opts)
> > +    fatal_error (input_location, "-foffload-abi-host-opts not
> > + specified.");
> 
> I know we're not doing that consistently, but please use '%<-foffload-
> abi-host-opts%>'.
Done.
> 
> > +  obstack_ptr_grow (&argv_obstack, host_abi_opts);
> >    obstack_ptr_grow (&argv_obstack, infile);
> >    obstack_ptr_grow (&argv_obstack, "-c");
> >    obstack_ptr_grow (&argv_obstack, "-o"); @@ -721,6 +714,8 @@ main
> > (int argc, char **argv)
> >                        "unrecognizable argument of option " STR);
> >       }
> >  #undef STR
> > +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> > +     host_abi_opts = argv[i] + strlen ("-foffload-abi-host-opts=");
> 
> The option parsing in the 'mkoffload's is ad-hoc (not using the proepr
> GCC infrastructure; which I'd like to change at some point in
> time...), but let's please catch the case that '-foffload-abi-host-
> opts=[...]'
> appears more than once (which could be necessary in certain
> configurations, to produce ABI-compatible host code?).  Not necessary
> to implement that right now: for now, it'll be fine to 'fatal_error'
> if running into a second '-foffload-abi-host-opts=[...]'.
Done.
> 
> Generally, likewise need to adjust 'gcc/config/gcn/mkoffload.cc'.  I
> can test this, or co-author, if you'd like.
Done.
> 
> > --- a/gcc/lto-wrapper.cc
> > +++ b/gcc/lto-wrapper.cc
> 
> Don't we also need to adjust 'merge_and_complain':
Done.
> 
>     case OPT_foffload_abi_:
>       if (existing_opt == -1)
>         decoded_options.safe_push (*foption);
>       else if (foption->value != decoded_options[existing_opt].value)
>         fatal_error (input_location,
>                      "option %s not used consistently in all LTO
> input"
>                      " files", foption->orig_option_with_args_text);
>       break;
> 
> > @@ -745,6 +745,7 @@ append_compiler_options (obstack *argv_obstack,
> vec<cl_decoded_option> opts)
> >       case OPT_fopenacc:
> >       case OPT_fopenacc_dim_:
> >       case OPT_foffload_abi_:
> > +     case OPT_foffload_abi_host_opts_:
> >       case OPT_fcf_protection_:
> >       case OPT_fasynchronous_unwind_tables:
> >       case OPT_funwind_tables:
> 
> Per my quick reading of the code, that should be correct.
> 
> > --- a/gcc/opts.cc
> > +++ b/gcc/opts.cc
> > @@ -3069,6 +3069,7 @@ common_handle_option (struct gcc_options
> *opts,
> >        break;
> >
> >      case OPT_foffload_abi_:
> > +    case OPT_foffload_abi_host_opts_:
> >  #ifdef ACCEL_COMPILER
> >        /* Handled in the 'mkoffload's.  */  #else
> |        error_at (loc, "%<-foffload-abi%> option can be specified
> only for "
> |              "offload compiler");
> |  #endif
> 
> Please adjust the diagnostic.  Surely the original option string will
> be available for use with '%qs'.
Done, thanks.

I verified the patch survives libgomp testing for Aarch64/nvptx offloading.
Could you please test the patch for gcn backend ?
Thanks!

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

Thanks,
Prathamesh
> 
> 
> Grüße
>  Thomas
[nvptx] Pass host specific ABI opts from mkoffload.

The patch adds an option -foffload-abi-host-opts, which
is set by host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's value
to host_compiler.

gcc/ChangeLog:
	* common.opt (foffload-abi-host-opts): New option.
	* config/aarch64/aarch64.cc (aarch64_offload_options): Set
	-foffload-abi-host-opts if -foffload-abi=lp64.
	* config/i386/i386-opts.cc (ix86_offload_options): Likewise.
	* config/rs6000/rs6000.cc (rs6000_offload_options): Likewise.
	* config/nvptx/mkoffload.cc (offload_abi_host_opts): Define.
	(compile_native): Append host_abi_opts to argv_obstack if
	offload_abi == OFFLOAD_ABI_LP64.
	* config/gcn/mkoffload.cc (offload_abi_host_opts): Define.
	(compile_native): Append host_abi_opts to argv_obstack if
	offload_abi == OFFLOAD_ABI_LP64.
	(main): Handle option -foffload-abi-host-opts.
	* lto-wrapper.cc (merge_and_complain): Handle
	-foffload-abi-host-opts.
	(append_compiler_options): Likewise.
	* opts.cc (common_handle_option): Likewise.

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

diff --git a/gcc/common.opt b/gcc/common.opt
index ea39f87ae71..8e14787d2e8 100644
--- a/gcc/common.opt
+++ b/gcc/common.opt
@@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
 EnumValue
 Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
 
+foffload-abi-host-opts=
+Common Driver Joined MissingArgError(option missing after %qs)
+-foffload-abi-host-opts=<options> Specify host ABI options.
+
 fomit-frame-pointer
 Common Var(flag_omit_frame_pointer) Optimization
 When possible do not generate stack frames.
diff --git a/gcc/config/aarch64/aarch64.cc b/gcc/config/aarch64/aarch64.cc
index 6a3f1a23a9f..57198df8044 100644
--- a/gcc/config/aarch64/aarch64.cc
+++ b/gcc/config/aarch64/aarch64.cc
@@ -19002,7 +19002,7 @@ aarch64_offload_options (void)
   if (TARGET_ILP32)
     return xstrdup ("-foffload-abi=ilp32");
   else
-    return xstrdup ("-foffload-abi=lp64");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-mabi=lp64");
 }
 
 static struct machine_function *
diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index b8d981878ed..d608580a1e1 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -133,6 +133,8 @@ static const char *gcn_dumpbase;
 static struct obstack files_to_cleanup;
 
 enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
+const char *offload_abi_host_opts = NULL;
+
 uint32_t elf_arch = EF_AMDGPU_MACH_AMDGCN_GFX900;  // Default GPU architecture.
 uint32_t elf_flags = EF_AMDGPU_FEATURE_SRAMECC_UNSUPPORTED_V4;
 
@@ -819,16 +821,12 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
   obstack_ptr_grow (&argv_obstack, gcn_dumpbase);
   obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
   obstack_ptr_grow (&argv_obstack, ".c");
-  switch (offload_abi)
+  if (offload_abi == OFFLOAD_ABI_LP64)
     {
-    case OFFLOAD_ABI_LP64:
-      obstack_ptr_grow (&argv_obstack, "-m64");
-      break;
-    case OFFLOAD_ABI_ILP32:
-      obstack_ptr_grow (&argv_obstack, "-m32");
-      break;
-    default:
-      gcc_unreachable ();
+      if (!offload_abi_host_opts)
+	fatal_error (input_location,
+		     "%<-foffload-abi-host-opts%> not specified.");
+      obstack_ptr_grow (&argv_obstack, offload_abi_host_opts);
     }
   obstack_ptr_grow (&argv_obstack, infile);
   obstack_ptr_grow (&argv_obstack, "-c");
@@ -998,6 +996,14 @@ main (int argc, char **argv)
 			 "unrecognizable argument of option %<" STR "%>");
 	}
 #undef STR
+      else if (startswith (argv[i], "-foffload-abi-host-opts="))
+	{
+	  if (offload_abi_host_opts)
+	    fatal_error (input_location,
+			 "-foffload-abi-host-opts specified multiple times");
+	  offload_abi_host_opts
+	    = argv[i] + strlen ("-foffload-abi-host-opts=");
+	}
       else if (strcmp (argv[i], "-fopenmp") == 0)
 	fopenmp = true;
       else if (strcmp (argv[i], "-fopenacc") == 0)
diff --git a/gcc/config/i386/i386-options.cc b/gcc/config/i386/i386-options.cc
index f79257cc764..3c5443ab67f 100644
--- a/gcc/config/i386/i386-options.cc
+++ b/gcc/config/i386/i386-options.cc
@@ -3680,7 +3680,7 @@ char *
 ix86_offload_options (void)
 {
   if (TARGET_LP64)
-    return xstrdup ("-foffload-abi=lp64");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
   return xstrdup ("-foffload-abi=ilp32");
 }
 
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index 503b1abcefd..95ef54b85e9 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -61,6 +61,7 @@ static const char *omp_requires_file;
 static const char *ptx_dumpbase;
 
 enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
+const char *offload_abi_host_opts = NULL;
 
 /* Delete tempfiles.  */
 
@@ -607,16 +608,12 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
   obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
   obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
   obstack_ptr_grow (&argv_obstack, ".c");
-  switch (offload_abi)
+  if (offload_abi == OFFLOAD_ABI_LP64)
     {
-    case OFFLOAD_ABI_LP64:
-      obstack_ptr_grow (&argv_obstack, "-m64");
-      break;
-    case OFFLOAD_ABI_ILP32:
-      obstack_ptr_grow (&argv_obstack, "-m32");
-      break;
-    default:
-      gcc_unreachable ();
+      if (!offload_abi_host_opts)
+	fatal_error (input_location,
+		     "%<-foffload-abi-host-opts%> not specified.");
+      obstack_ptr_grow (&argv_obstack, offload_abi_host_opts);
     }
   obstack_ptr_grow (&argv_obstack, infile);
   obstack_ptr_grow (&argv_obstack, "-c");
@@ -721,6 +718,14 @@ main (int argc, char **argv)
 			 "unrecognizable argument of option " STR);
 	}
 #undef STR
+      else if (startswith (argv[i], "-foffload-abi-host-opts="))
+	{
+	  if (offload_abi_host_opts)
+	    fatal_error (input_location,
+			 "-foffload-abi-host-opts specified multiple times");
+	  offload_abi_host_opts
+	    = argv[i] + strlen ("-foffload-abi-host-opts=");
+	}
       else if (strcmp (argv[i], "-fopenmp") == 0)
 	fopenmp = true;
       else if (strcmp (argv[i], "-fopenacc") == 0)
diff --git a/gcc/config/rs6000/rs6000.cc b/gcc/config/rs6000/rs6000.cc
index 08579bc83e6..fe03fd30fd7 100644
--- a/gcc/config/rs6000/rs6000.cc
+++ b/gcc/config/rs6000/rs6000.cc
@@ -17330,7 +17330,7 @@ static char *
 rs6000_offload_options (void)
 {
   if (TARGET_64BIT)
-    return xstrdup ("-foffload-abi=lp64");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
   else
     return xstrdup ("-foffload-abi=ilp32");
 }
diff --git a/gcc/lto-wrapper.cc b/gcc/lto-wrapper.cc
index c07765b37a2..7de045da9b9 100644
--- a/gcc/lto-wrapper.cc
+++ b/gcc/lto-wrapper.cc
@@ -484,6 +484,7 @@ merge_and_complain (vec<cl_decoded_option> &decoded_options,
  
 
 	case OPT_foffload_abi_:
+	case OPT_foffload_abi_host_opts_:
 	  if (existing_opt == -1)
 	    decoded_options.safe_push (*foption);
 	  else if (foption->value != decoded_options[existing_opt].value)
@@ -745,6 +746,7 @@ append_compiler_options (obstack *argv_obstack, vec<cl_decoded_option> opts)
 	case OPT_fopenacc:
 	case OPT_fopenacc_dim_:
 	case OPT_foffload_abi_:
+	case OPT_foffload_abi_host_opts_:
 	case OPT_fcf_protection_:
 	case OPT_fasynchronous_unwind_tables:
 	case OPT_funwind_tables:
diff --git a/gcc/opts.cc b/gcc/opts.cc
index fc6abf6f582..c554e2fdd2c 100644
--- a/gcc/opts.cc
+++ b/gcc/opts.cc
@@ -3070,11 +3070,12 @@ common_handle_option (struct gcc_options *opts,
       break;
 
     case OPT_foffload_abi_:
+    case OPT_foffload_abi_host_opts_:
 #ifdef ACCEL_COMPILER
       /* Handled in the 'mkoffload's.  */
 #else
-      error_at (loc, "%<-foffload-abi%> option can be specified only for "
-		"offload compiler");
+      error_at (loc, "%qs option can be specified only for "
+		"offload compiler", arg);
 #endif
       break;

Hi Prathamesh!

On 2024-09-09T06:31:18+0000, Prathamesh Kulkarni <prathameshk@nvidia.com> wrote:
>> -----Original Message-----
>> From: Thomas Schwinge <tschwinge@baylibre.com>
>> Sent: Friday, September 6, 2024 2:31 PM
>> On 2024-08-16T15:36:29+0000, Prathamesh Kulkarni
>> <prathameshk@nvidia.com> wrote:
>> >> > Am 13.08.2024 um 17:48 schrieb Thomas Schwinge
>> >> <tschwinge@baylibre.com>:
>> >> > On 2024-08-12T07:50:07+0000, Prathamesh Kulkarni
>> >> <prathameshk@nvidia.com> wrote:
>> >> >> I added another option -foffload-abi-host-opts to specify host
>> abi
>> >> >> opts, and leave -foffload-abi to specify if ABI is 32/64 bit
>> which
>> >> >> mkoffload can use to enable/disable offloading (as before).

>> > --- a/gcc/config/aarch64/aarch64.cc
>> > +++ b/gcc/config/aarch64/aarch64.cc
>> > @@ -18999,9 +18999,9 @@ static char *
>> >  aarch64_offload_options (void)
>> >  {
>> >    if (TARGET_ILP32)
>> > -    return xstrdup ("-foffload-abi=ilp32");
>> > +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-mabi=ilp32");
>> >    else
>> > -    return xstrdup ("-foffload-abi=lp64");
>> > +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-mabi=lp64");
>> >  }
>> 
>> As none of the current offload compilers is set up of ILP32, I suggest
>> we continue to pass '-foffload-abi=ilp32' without '-foffload-abi-host-
>> opts=[...]' -- the 'mkoffload's in that case should get to the point
>> where the latter is used.

Oh...  I was wrong with the latter item: I failed to see that the
'mkoffload's still do 'compile_native' even if they don't create an
actual offload image, sorry!

> Um, would that still possibly result in arch mismatch for host objects and xnvptx-none.o if we don't pass host ABI opts for ILP32 ?
> For eg, if the host compiler defaults to 64-bit code-gen (and user requests for 32-bit code gen on host), and we avoid passing host ABI opts for -foffload-abi=ilp32,
> it will generate 64-bit xnvptx-none.o (corresponding to empty ptx_cfile_name), while rest of the host objects will be 32-bit, or am I misunderstanding ?

You're quite right -- my fault.

> The attached patch avoids passing -foffload-abi-host-opts if -foffload-abi=ilp32.

So, sorry for the back and forth.  I think we now agree that we do need
'-foffload-abi-host-opts=[...]' specified in call cases (as you
originally had), and then again unconditionally use
'offload_abi_host_opts' in the 'mkoffload's' 'compile_native' functions.

> Could you please test the patch for gcn backend ?

I'll do that.

> [nvptx] Pass host specific ABI opts from mkoffload.
>
> The patch adds an option -foffload-abi-host-opts, which
> is set by host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's value

"its", by the way.  ;-)

> to host_compiler.

> --- a/gcc/common.opt
> +++ b/gcc/common.opt

> +foffload-abi-host-opts=
> +Common Driver Joined MissingArgError(option missing after %qs)
> +-foffload-abi-host-opts=<options> Specify host ABI options.
> +

Still need TAB between '-foffload-abi-host-opts=<options>' and its help
text.

> --- a/gcc/config/gcn/mkoffload.cc
> +++ b/gcc/config/gcn/mkoffload.cc

> @@ -998,6 +996,14 @@ main (int argc, char **argv)
>  			 "unrecognizable argument of option %<" STR "%>");
>  	}
>  #undef STR
> +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> +	{
> +	  if (offload_abi_host_opts)
> +	    fatal_error (input_location,
> +			 "-foffload-abi-host-opts specified multiple times");

ACK, but again '%<-foffload-abi-host-opts%>', please.  (May also use
another '#define STR "[...]"' for the duplicated string, but I don't
care.)

> --- a/gcc/config/nvptx/mkoffload.cc
> +++ b/gcc/config/nvptx/mkoffload.cc

> @@ -721,6 +718,14 @@ main (int argc, char **argv)
>  			 "unrecognizable argument of option " STR);
>  	}
>  #undef STR
> +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> +	{
> +	  if (offload_abi_host_opts)
> +	    fatal_error (input_location,
> +			 "-foffload-abi-host-opts specified multiple times");

Likewise.

> --- a/gcc/lto-wrapper.cc
> +++ b/gcc/lto-wrapper.cc
> @@ -484,6 +484,7 @@ merge_and_complain (vec<cl_decoded_option> &decoded_options,
>   
>  
>  	case OPT_foffload_abi_:
> +	case OPT_foffload_abi_host_opts_:
>  	  if (existing_opt == -1)
>  	    decoded_options.safe_push (*foption);
>  	  else if (foption->value != decoded_options[existing_opt].value)
> @@ -745,6 +746,7 @@ append_compiler_options (obstack *argv_obstack, vec<cl_decoded_option> opts)
>  	case OPT_fopenacc:
>  	case OPT_fopenacc_dim_:
>  	case OPT_foffload_abi_:
> +	case OPT_foffload_abi_host_opts_:
>  	case OPT_fcf_protection_:
>  	case OPT_fasynchronous_unwind_tables:
>  	case OPT_funwind_tables:

I'm not too familiar with this code, but that now looks right to me.

> --- a/gcc/opts.cc
> +++ b/gcc/opts.cc
> @@ -3070,11 +3070,12 @@ common_handle_option (struct gcc_options *opts,
>        break;
>  
>      case OPT_foffload_abi_:
> +    case OPT_foffload_abi_host_opts_:
>  #ifdef ACCEL_COMPILER
>        /* Handled in the 'mkoffload's.  */
>  #else
> -      error_at (loc, "%<-foffload-abi%> option can be specified only for "
> -		"offload compiler");
> +      error_at (loc, "%qs option can be specified only for "
> +		"offload compiler", arg);
>  #endif
>        break;

With this, using '-foffload-abi=ilp32' with the host compiler results in:

    cc1: error: ‘ilp32’ option can be specified only for offload compiler

..., or for '-foffload-abi-host-opts=-m64' in:

    xgcc: error: ‘-m64’ option can be specified only for offload compiler

..., so 'arg' is only the option argument, not the whole string.

And, incidentally, 'cc1' vs. 'xgcc' means without vs. with 'Driver'
option property (re your 'gcc/common.opt' change).  Which should it be?
'-foffload-abi=[...]' currently doesn't have 'Driver', so probably
'-foffload-abi-host-opts=[...]' also shouldn't?

With those small items addressed, the patch looks good to me, thanks!
(..., and I'll still test GCN offloading.)


Grüße
 Thomas

> -----Original Message-----
> From: Thomas Schwinge <tschwinge@baylibre.com>
> Sent: Monday, September 9, 2024 8:50 PM
> To: Prathamesh Kulkarni <prathameshk@nvidia.com>; Richard Biener
> <rguenther@suse.de>
> Cc: Andrew Pinski <pinskia@gmail.com>; gcc-patches@gcc.gnu.org; Jakub
> Jelinek <jakub@redhat.com>
> Subject: RE: [nvptx] Pass -m32/-m64 to host_compiler if it has
> multilib support
> 
> External email: Use caution opening links or attachments
> 
> 
> Hi Prathamesh!
Hi Thomas,
> 
> On 2024-09-09T06:31:18+0000, Prathamesh Kulkarni
> <prathameshk@nvidia.com> wrote:
> >> -----Original Message-----
> >> From: Thomas Schwinge <tschwinge@baylibre.com>
> >> Sent: Friday, September 6, 2024 2:31 PM On 2024-08-
> 16T15:36:29+0000,
> >> Prathamesh Kulkarni <prathameshk@nvidia.com> wrote:
> >> >> > Am 13.08.2024 um 17:48 schrieb Thomas Schwinge
> >> >> <tschwinge@baylibre.com>:
> >> >> > On 2024-08-12T07:50:07+0000, Prathamesh Kulkarni
> >> >> <prathameshk@nvidia.com> wrote:
> >> >> >> I added another option -foffload-abi-host-opts to specify
> host
> >> abi
> >> >> >> opts, and leave -foffload-abi to specify if ABI is 32/64 bit
> >> which
> >> >> >> mkoffload can use to enable/disable offloading (as before).
> 
> >> > --- a/gcc/config/aarch64/aarch64.cc
> >> > +++ b/gcc/config/aarch64/aarch64.cc
> >> > @@ -18999,9 +18999,9 @@ static char *  aarch64_offload_options
> >> > (void)  {
> >> >    if (TARGET_ILP32)
> >> > -    return xstrdup ("-foffload-abi=ilp32");
> >> > +    return xstrdup ("-foffload-abi=ilp32
> >> > + -foffload-abi-host-opts=-mabi=ilp32");
> >> >    else
> >> > -    return xstrdup ("-foffload-abi=lp64");
> >> > +    return xstrdup ("-foffload-abi=lp64
> >> > + -foffload-abi-host-opts=-mabi=lp64");
> >> >  }
> >>
> >> As none of the current offload compilers is set up of ILP32, I
> >> suggest we continue to pass '-foffload-abi=ilp32' without
> >> '-foffload-abi-host- opts=[...]' -- the 'mkoffload's in that case
> >> should get to the point where the latter is used.
> 
> Oh...  I was wrong with the latter item: I failed to see that the
> 'mkoffload's still do 'compile_native' even if they don't create an
> actual offload image, sorry!
> 
> > Um, would that still possibly result in arch mismatch for host
> objects and xnvptx-none.o if we don't pass host ABI opts for ILP32 ?
> > For eg, if the host compiler defaults to 64-bit code-gen (and user
> > requests for 32-bit code gen on host), and we avoid passing host ABI
> opts for -foffload-abi=ilp32, it will generate 64-bit xnvptx-none.o
> (corresponding to empty ptx_cfile_name), while rest of the host
> objects will be 32-bit, or am I misunderstanding ?
> 
> You're quite right -- my fault.
> 
> > The attached patch avoids passing -foffload-abi-host-opts if -
> foffload-abi=ilp32.
> 
> So, sorry for the back and forth.  I think we now agree that we do
> need '-foffload-abi-host-opts=[...]' specified in call cases (as you
> originally had), and then again unconditionally use
> 'offload_abi_host_opts' in the 'mkoffload's' 'compile_native'
> functions.
Done in the attached patch, thanks.
> 
> > Could you please test the patch for gcn backend ?
> 
> I'll do that.
> 
> > [nvptx] Pass host specific ABI opts from mkoffload.
> >
> > The patch adds an option -foffload-abi-host-opts, which is set by
> host
> > in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes it's value
> 
> "its", by the way.  ;-)
Fixed 😊
> 
> > to host_compiler.
> 
> > --- a/gcc/common.opt
> > +++ b/gcc/common.opt
> 
> > +foffload-abi-host-opts=
> > +Common Driver Joined MissingArgError(option missing after %qs)
> > +-foffload-abi-host-opts=<options> Specify host ABI options.
> > +
> 
> Still need TAB between '-foffload-abi-host-opts=<options>' and its
> help text.
Done.
> 
> > --- a/gcc/config/gcn/mkoffload.cc
> > +++ b/gcc/config/gcn/mkoffload.cc
> 
> > @@ -998,6 +996,14 @@ main (int argc, char **argv)
> >                        "unrecognizable argument of option %<" STR
> "%>");
> >       }
> >  #undef STR
> > +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> > +     {
> > +       if (offload_abi_host_opts)
> > +         fatal_error (input_location,
> > +                      "-foffload-abi-host-opts specified multiple
> > + times");
> 
> ACK, but again '%<-foffload-abi-host-opts%>', please.  (May also use
> another '#define STR "[...]"' for the duplicated string, but I don't
> care.)
Sorry, missed this earlier, fixed.
> 
> > --- a/gcc/config/nvptx/mkoffload.cc
> > +++ b/gcc/config/nvptx/mkoffload.cc
> 
> > @@ -721,6 +718,14 @@ main (int argc, char **argv)
> >                        "unrecognizable argument of option " STR);
> >       }
> >  #undef STR
> > +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> > +     {
> > +       if (offload_abi_host_opts)
> > +         fatal_error (input_location,
> > +                      "-foffload-abi-host-opts specified multiple
> > + times");
> 
> Likewise.
> 
> > --- a/gcc/lto-wrapper.cc
> > +++ b/gcc/lto-wrapper.cc
> > @@ -484,6 +484,7 @@ merge_and_complain (vec<cl_decoded_option>
> > &decoded_options,
> >
> >
> >       case OPT_foffload_abi_:
> > +     case OPT_foffload_abi_host_opts_:
> >         if (existing_opt == -1)
> >           decoded_options.safe_push (*foption);
> >         else if (foption->value !=
> > decoded_options[existing_opt].value)
> > @@ -745,6 +746,7 @@ append_compiler_options (obstack *argv_obstack,
> vec<cl_decoded_option> opts)
> >       case OPT_fopenacc:
> >       case OPT_fopenacc_dim_:
> >       case OPT_foffload_abi_:
> > +     case OPT_foffload_abi_host_opts_:
> >       case OPT_fcf_protection_:
> >       case OPT_fasynchronous_unwind_tables:
> >       case OPT_funwind_tables:
> 
> I'm not too familiar with this code, but that now looks right to me.
> 
> > --- a/gcc/opts.cc
> > +++ b/gcc/opts.cc
> > @@ -3070,11 +3070,12 @@ common_handle_option (struct gcc_options
> *opts,
> >        break;
> >
> >      case OPT_foffload_abi_:
> > +    case OPT_foffload_abi_host_opts_:
> >  #ifdef ACCEL_COMPILER
> >        /* Handled in the 'mkoffload's.  */  #else
> > -      error_at (loc, "%<-foffload-abi%> option can be specified
> only for "
> > -             "offload compiler");
> > +      error_at (loc, "%qs option can be specified only for "
> > +             "offload compiler", arg);
> >  #endif
> >        break;
> 
> With this, using '-foffload-abi=ilp32' with the host compiler results
> in:
> 
>     cc1: error: ‘ilp32’ option can be specified only for offload
> compiler
> 
> ..., or for '-foffload-abi-host-opts=-m64' in:
> 
>     xgcc: error: ‘-m64’ option can be specified only for offload
> compiler
> 
> ..., so 'arg' is only the option argument, not the whole string.
Ah, didn't realize that, sorry. Fixed.
> 
> And, incidentally, 'cc1' vs. 'xgcc' means without vs. with 'Driver'
> option property (re your 'gcc/common.opt' change).  Which should it
> be?
> '-foffload-abi=[...]' currently doesn't have 'Driver', so probably '-
> foffload-abi-host-opts=[...]' also shouldn't?
Indeed, removed Driver from -foffload-abi-host-opts, thanks.
> 
> With those small items addressed, the patch looks good to me, thanks!
> (..., and I'll still test GCN offloading.)
Thanks, I have tested libgomp for aarch64/nvptx offloading.
Is it OK to commit (if testing at your end also passes on gcn) ?

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

Thanks,
Prathamesh
> 
> 
> Grüße
>  Thomas
[nvptx] Pass host specific ABI opts from mkoffload.

The patch adds an option -foffload-abi-host-opts, which
is set by host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes its value
to host_compiler.

gcc/ChangeLog:
	* common.opt (foffload-abi-host-opts): New option.
	* config/aarch64/aarch64.cc (aarch64_offload_options): Pass
	-foffload-abi-host-opts.
	* config/i386/i386-opts.cc (ix86_offload_options): Likewise.
	* config/rs6000/rs6000.cc (rs6000_offload_options): Likewise.
	* config/nvptx/mkoffload.cc (offload_abi_host_opts): Define.
	(compile_native): Append offload_abi_host_opts to argv_obstack.
	(main): Handle option -foffload-abi-host-opts.
	* config/gcn/mkoffload.cc (offload_abi_host_opts): Define.
	(compile_native): Append offload_abi_host_opts to argv_obstack.
	(main): Handle option -foffload-abi-host-opts.
	* lto-wrapper.cc (merge_and_complain): Handle
	-foffload-abi-host-opts.
	(append_compiler_options): Likewise.
	* opts.cc (common_handle_option): Likewise.

Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

diff --git a/gcc/common.opt b/gcc/common.opt
index ea39f87ae71..d270e524ff4 100644
--- a/gcc/common.opt
+++ b/gcc/common.opt
@@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
 EnumValue
 Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
 
+foffload-abi-host-opts=
+Common Joined MissingArgError(option missing after %qs)
+-foffload-abi-host-opts=<options>	Specify host ABI options.
+
 fomit-frame-pointer
 Common Var(flag_omit_frame_pointer) Optimization
 When possible do not generate stack frames.
diff --git a/gcc/config/aarch64/aarch64.cc b/gcc/config/aarch64/aarch64.cc
index 6a3f1a23a9f..6ccf08d1cc0 100644
--- a/gcc/config/aarch64/aarch64.cc
+++ b/gcc/config/aarch64/aarch64.cc
@@ -19000,9 +19000,9 @@ static char *
 aarch64_offload_options (void)
 {
   if (TARGET_ILP32)
-    return xstrdup ("-foffload-abi=ilp32");
+    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-mabi=ilp32");
   else
-    return xstrdup ("-foffload-abi=lp64");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-mabi=lp64");
 }
 
 static struct machine_function *
diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index b8d981878ed..345bbf7709c 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -133,6 +133,8 @@ static const char *gcn_dumpbase;
 static struct obstack files_to_cleanup;
 
 enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
+const char *offload_abi_host_opts = NULL;
+
 uint32_t elf_arch = EF_AMDGPU_MACH_AMDGCN_GFX900;  // Default GPU architecture.
 uint32_t elf_flags = EF_AMDGPU_FEATURE_SRAMECC_UNSUPPORTED_V4;
 
@@ -819,17 +821,10 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
   obstack_ptr_grow (&argv_obstack, gcn_dumpbase);
   obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
   obstack_ptr_grow (&argv_obstack, ".c");
-  switch (offload_abi)
-    {
-    case OFFLOAD_ABI_LP64:
-      obstack_ptr_grow (&argv_obstack, "-m64");
-      break;
-    case OFFLOAD_ABI_ILP32:
-      obstack_ptr_grow (&argv_obstack, "-m32");
-      break;
-    default:
-      gcc_unreachable ();
-    }
+  if (!offload_abi_host_opts)
+    fatal_error (input_location,
+		 "%<-foffload-abi-host-opts%> not specified.");
+  obstack_ptr_grow (&argv_obstack, offload_abi_host_opts);
   obstack_ptr_grow (&argv_obstack, infile);
   obstack_ptr_grow (&argv_obstack, "-c");
   obstack_ptr_grow (&argv_obstack, "-o");
@@ -998,6 +993,15 @@ main (int argc, char **argv)
 			 "unrecognizable argument of option %<" STR "%>");
 	}
 #undef STR
+      else if (startswith (argv[i], "-foffload-abi-host-opts="))
+	{
+	  if (offload_abi_host_opts)
+	    fatal_error (input_location,
+			 "%<-foffload-abi-host-opts%> specified "
+			 "multiple times");
+	  offload_abi_host_opts
+	    = argv[i] + strlen ("-foffload-abi-host-opts=");
+	}
       else if (strcmp (argv[i], "-fopenmp") == 0)
 	fopenmp = true;
       else if (strcmp (argv[i], "-fopenacc") == 0)
diff --git a/gcc/config/i386/i386-options.cc b/gcc/config/i386/i386-options.cc
index f79257cc764..55e0210260f 100644
--- a/gcc/config/i386/i386-options.cc
+++ b/gcc/config/i386/i386-options.cc
@@ -3680,8 +3680,8 @@ char *
 ix86_offload_options (void)
 {
   if (TARGET_LP64)
-    return xstrdup ("-foffload-abi=lp64");
-  return xstrdup ("-foffload-abi=ilp32");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
+  return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
 }
 
 /* Handle "cdecl", "stdcall", "fastcall", "regparm", "thiscall",
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index 503b1abcefd..df16ee64736 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -61,6 +61,7 @@ static const char *omp_requires_file;
 static const char *ptx_dumpbase;
 
 enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
+const char *offload_abi_host_opts = NULL;
 
 /* Delete tempfiles.  */
 
@@ -607,17 +608,10 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
   obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
   obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
   obstack_ptr_grow (&argv_obstack, ".c");
-  switch (offload_abi)
-    {
-    case OFFLOAD_ABI_LP64:
-      obstack_ptr_grow (&argv_obstack, "-m64");
-      break;
-    case OFFLOAD_ABI_ILP32:
-      obstack_ptr_grow (&argv_obstack, "-m32");
-      break;
-    default:
-      gcc_unreachable ();
-    }
+  if (!offload_abi_host_opts)
+    fatal_error (input_location,
+		 "%<-foffload-abi-host-opts%> not specified.");
+  obstack_ptr_grow (&argv_obstack, offload_abi_host_opts);
   obstack_ptr_grow (&argv_obstack, infile);
   obstack_ptr_grow (&argv_obstack, "-c");
   obstack_ptr_grow (&argv_obstack, "-o");
@@ -721,6 +715,15 @@ main (int argc, char **argv)
 			 "unrecognizable argument of option " STR);
 	}
 #undef STR
+      else if (startswith (argv[i], "-foffload-abi-host-opts="))
+	{
+	  if (offload_abi_host_opts)
+	    fatal_error (input_location,
+			 "%<-foffload-abi-host-opts%> specified "
+			 "multiple times");
+	  offload_abi_host_opts
+	    = argv[i] + strlen ("-foffload-abi-host-opts=");
+	}
       else if (strcmp (argv[i], "-fopenmp") == 0)
 	fopenmp = true;
       else if (strcmp (argv[i], "-fopenacc") == 0)
diff --git a/gcc/config/rs6000/rs6000.cc b/gcc/config/rs6000/rs6000.cc
index 08579bc83e6..0bf8bae27f5 100644
--- a/gcc/config/rs6000/rs6000.cc
+++ b/gcc/config/rs6000/rs6000.cc
@@ -17330,9 +17330,9 @@ static char *
 rs6000_offload_options (void)
 {
   if (TARGET_64BIT)
-    return xstrdup ("-foffload-abi=lp64");
+    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
   else
-    return xstrdup ("-foffload-abi=ilp32");
+    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
 }
 
 
diff --git a/gcc/lto-wrapper.cc b/gcc/lto-wrapper.cc
index c07765b37a2..7de045da9b9 100644
--- a/gcc/lto-wrapper.cc
+++ b/gcc/lto-wrapper.cc
@@ -484,6 +484,7 @@ merge_and_complain (vec<cl_decoded_option> &decoded_options,
  
 
 	case OPT_foffload_abi_:
+	case OPT_foffload_abi_host_opts_:
 	  if (existing_opt == -1)
 	    decoded_options.safe_push (*foption);
 	  else if (foption->value != decoded_options[existing_opt].value)
@@ -745,6 +746,7 @@ append_compiler_options (obstack *argv_obstack, vec<cl_decoded_option> opts)
 	case OPT_fopenacc:
 	case OPT_fopenacc_dim_:
 	case OPT_foffload_abi_:
+	case OPT_foffload_abi_host_opts_:
 	case OPT_fcf_protection_:
 	case OPT_fasynchronous_unwind_tables:
 	case OPT_funwind_tables:
diff --git a/gcc/opts.cc b/gcc/opts.cc
index fc6abf6f582..a78f73e57e3 100644
--- a/gcc/opts.cc
+++ b/gcc/opts.cc
@@ -3070,11 +3070,14 @@ common_handle_option (struct gcc_options *opts,
       break;
 
     case OPT_foffload_abi_:
+    case OPT_foffload_abi_host_opts_:
 #ifdef ACCEL_COMPILER
       /* Handled in the 'mkoffload's.  */
 #else
-      error_at (loc, "%<-foffload-abi%> option can be specified only for "
-		"offload compiler");
+      error_at (loc,
+		"%qs option can be specified only for offload compiler",
+		(code == OPT_foffload_abi_) ? "-foffload-abi"
+					    : "-foffload-abi-host-opts");
 #endif
       break;

Hi Prathamesh!

On 2024-09-10T13:22:10+0000, Prathamesh Kulkarni <prathameshk@nvidia.com> wrote:
>> -----Original Message-----
>> From: Thomas Schwinge <tschwinge@baylibre.com>
>> Sent: Monday, September 9, 2024 8:50 PM

>> > Could you please test the patch for gcn backend ?

I've successfully tested x86_64 host with GCN as well as nvptx
offloading, and also ppc64le host with nvptx offloading.

I just realized two more minor things:

> [nvptx] Pass host specific ABI opts from mkoffload.
>
> The patch adds an option -foffload-abi-host-opts, which
> is set by host in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes its value
> to host_compiler.
>

Please add here "	PR target/96265".

> gcc/ChangeLog:
> 	* common.opt (foffload-abi-host-opts): New option.
> 	* config/aarch64/aarch64.cc (aarch64_offload_options): Pass
> 	-foffload-abi-host-opts.
> 	* config/i386/i386-opts.cc (ix86_offload_options): Likewise.
> 	* config/rs6000/rs6000.cc (rs6000_offload_options): Likewise.
> 	* config/nvptx/mkoffload.cc (offload_abi_host_opts): Define.
> 	(compile_native): Append offload_abi_host_opts to argv_obstack.
> 	(main): Handle option -foffload-abi-host-opts.
> 	* config/gcn/mkoffload.cc (offload_abi_host_opts): Define.
> 	(compile_native): Append offload_abi_host_opts to argv_obstack.
> 	(main): Handle option -foffload-abi-host-opts.
> 	* lto-wrapper.cc (merge_and_complain): Handle
> 	-foffload-abi-host-opts.
> 	(append_compiler_options): Likewise.
> 	* opts.cc (common_handle_option): Likewise.
>
> Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>

Given that we're adding a new option to 'gcc/common.opt', do we need to
update (regenerate?) 'gcc/common.opt.urls'?  (I've not yet had the need
myself, and therefore not yet looked up how to do that.)  Or maybe not,
given that '-foffload-abi-host-opts=[...]' isn't documented?

Otherwise looks good to me; OK to push (with these minor items addressed,
as necessary), thanks!


Grüße
 Thomas


> diff --git a/gcc/common.opt b/gcc/common.opt
> index ea39f87ae71..d270e524ff4 100644
> --- a/gcc/common.opt
> +++ b/gcc/common.opt
> @@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
>  EnumValue
>  Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
>  
> +foffload-abi-host-opts=
> +Common Joined MissingArgError(option missing after %qs)
> +-foffload-abi-host-opts=<options>	Specify host ABI options.
> +
>  fomit-frame-pointer
>  Common Var(flag_omit_frame_pointer) Optimization
>  When possible do not generate stack frames.
> diff --git a/gcc/config/aarch64/aarch64.cc b/gcc/config/aarch64/aarch64.cc
> index 6a3f1a23a9f..6ccf08d1cc0 100644
> --- a/gcc/config/aarch64/aarch64.cc
> +++ b/gcc/config/aarch64/aarch64.cc
> @@ -19000,9 +19000,9 @@ static char *
>  aarch64_offload_options (void)
>  {
>    if (TARGET_ILP32)
> -    return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-mabi=ilp32");
>    else
> -    return xstrdup ("-foffload-abi=lp64");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-mabi=lp64");
>  }
>  
>  static struct machine_function *
> diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
> index b8d981878ed..345bbf7709c 100644
> --- a/gcc/config/gcn/mkoffload.cc
> +++ b/gcc/config/gcn/mkoffload.cc
> @@ -133,6 +133,8 @@ static const char *gcn_dumpbase;
>  static struct obstack files_to_cleanup;
>  
>  enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
> +const char *offload_abi_host_opts = NULL;
> +
>  uint32_t elf_arch = EF_AMDGPU_MACH_AMDGCN_GFX900;  // Default GPU architecture.
>  uint32_t elf_flags = EF_AMDGPU_FEATURE_SRAMECC_UNSUPPORTED_V4;
>  
> @@ -819,17 +821,10 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
>    obstack_ptr_grow (&argv_obstack, gcn_dumpbase);
>    obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
>    obstack_ptr_grow (&argv_obstack, ".c");
> -  switch (offload_abi)
> -    {
> -    case OFFLOAD_ABI_LP64:
> -      obstack_ptr_grow (&argv_obstack, "-m64");
> -      break;
> -    case OFFLOAD_ABI_ILP32:
> -      obstack_ptr_grow (&argv_obstack, "-m32");
> -      break;
> -    default:
> -      gcc_unreachable ();
> -    }
> +  if (!offload_abi_host_opts)
> +    fatal_error (input_location,
> +		 "%<-foffload-abi-host-opts%> not specified.");
> +  obstack_ptr_grow (&argv_obstack, offload_abi_host_opts);
>    obstack_ptr_grow (&argv_obstack, infile);
>    obstack_ptr_grow (&argv_obstack, "-c");
>    obstack_ptr_grow (&argv_obstack, "-o");
> @@ -998,6 +993,15 @@ main (int argc, char **argv)
>  			 "unrecognizable argument of option %<" STR "%>");
>  	}
>  #undef STR
> +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> +	{
> +	  if (offload_abi_host_opts)
> +	    fatal_error (input_location,
> +			 "%<-foffload-abi-host-opts%> specified "
> +			 "multiple times");
> +	  offload_abi_host_opts
> +	    = argv[i] + strlen ("-foffload-abi-host-opts=");
> +	}
>        else if (strcmp (argv[i], "-fopenmp") == 0)
>  	fopenmp = true;
>        else if (strcmp (argv[i], "-fopenacc") == 0)
> diff --git a/gcc/config/i386/i386-options.cc b/gcc/config/i386/i386-options.cc
> index f79257cc764..55e0210260f 100644
> --- a/gcc/config/i386/i386-options.cc
> +++ b/gcc/config/i386/i386-options.cc
> @@ -3680,8 +3680,8 @@ char *
>  ix86_offload_options (void)
>  {
>    if (TARGET_LP64)
> -    return xstrdup ("-foffload-abi=lp64");
> -  return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
> +  return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
>  }
>  
>  /* Handle "cdecl", "stdcall", "fastcall", "regparm", "thiscall",
> diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
> index 503b1abcefd..df16ee64736 100644
> --- a/gcc/config/nvptx/mkoffload.cc
> +++ b/gcc/config/nvptx/mkoffload.cc
> @@ -61,6 +61,7 @@ static const char *omp_requires_file;
>  static const char *ptx_dumpbase;
>  
>  enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
> +const char *offload_abi_host_opts = NULL;
>  
>  /* Delete tempfiles.  */
>  
> @@ -607,17 +608,10 @@ compile_native (const char *infile, const char *outfile, const char *compiler,
>    obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
>    obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
>    obstack_ptr_grow (&argv_obstack, ".c");
> -  switch (offload_abi)
> -    {
> -    case OFFLOAD_ABI_LP64:
> -      obstack_ptr_grow (&argv_obstack, "-m64");
> -      break;
> -    case OFFLOAD_ABI_ILP32:
> -      obstack_ptr_grow (&argv_obstack, "-m32");
> -      break;
> -    default:
> -      gcc_unreachable ();
> -    }
> +  if (!offload_abi_host_opts)
> +    fatal_error (input_location,
> +		 "%<-foffload-abi-host-opts%> not specified.");
> +  obstack_ptr_grow (&argv_obstack, offload_abi_host_opts);
>    obstack_ptr_grow (&argv_obstack, infile);
>    obstack_ptr_grow (&argv_obstack, "-c");
>    obstack_ptr_grow (&argv_obstack, "-o");
> @@ -721,6 +715,15 @@ main (int argc, char **argv)
>  			 "unrecognizable argument of option " STR);
>  	}
>  #undef STR
> +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> +	{
> +	  if (offload_abi_host_opts)
> +	    fatal_error (input_location,
> +			 "%<-foffload-abi-host-opts%> specified "
> +			 "multiple times");
> +	  offload_abi_host_opts
> +	    = argv[i] + strlen ("-foffload-abi-host-opts=");
> +	}
>        else if (strcmp (argv[i], "-fopenmp") == 0)
>  	fopenmp = true;
>        else if (strcmp (argv[i], "-fopenacc") == 0)
> diff --git a/gcc/config/rs6000/rs6000.cc b/gcc/config/rs6000/rs6000.cc
> index 08579bc83e6..0bf8bae27f5 100644
> --- a/gcc/config/rs6000/rs6000.cc
> +++ b/gcc/config/rs6000/rs6000.cc
> @@ -17330,9 +17330,9 @@ static char *
>  rs6000_offload_options (void)
>  {
>    if (TARGET_64BIT)
> -    return xstrdup ("-foffload-abi=lp64");
> +    return xstrdup ("-foffload-abi=lp64 -foffload-abi-host-opts=-m64");
>    else
> -    return xstrdup ("-foffload-abi=ilp32");
> +    return xstrdup ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
>  }
>  
>  
> diff --git a/gcc/lto-wrapper.cc b/gcc/lto-wrapper.cc
> index c07765b37a2..7de045da9b9 100644
> --- a/gcc/lto-wrapper.cc
> +++ b/gcc/lto-wrapper.cc
> @@ -484,6 +484,7 @@ merge_and_complain (vec<cl_decoded_option> &decoded_options,
>   
>  
>  	case OPT_foffload_abi_:
> +	case OPT_foffload_abi_host_opts_:
>  	  if (existing_opt == -1)
>  	    decoded_options.safe_push (*foption);
>  	  else if (foption->value != decoded_options[existing_opt].value)
> @@ -745,6 +746,7 @@ append_compiler_options (obstack *argv_obstack, vec<cl_decoded_option> opts)
>  	case OPT_fopenacc:
>  	case OPT_fopenacc_dim_:
>  	case OPT_foffload_abi_:
> +	case OPT_foffload_abi_host_opts_:
>  	case OPT_fcf_protection_:
>  	case OPT_fasynchronous_unwind_tables:
>  	case OPT_funwind_tables:
> diff --git a/gcc/opts.cc b/gcc/opts.cc
> index fc6abf6f582..a78f73e57e3 100644
> --- a/gcc/opts.cc
> +++ b/gcc/opts.cc
> @@ -3070,11 +3070,14 @@ common_handle_option (struct gcc_options *opts,
>        break;
>  
>      case OPT_foffload_abi_:
> +    case OPT_foffload_abi_host_opts_:
>  #ifdef ACCEL_COMPILER
>        /* Handled in the 'mkoffload's.  */
>  #else
> -      error_at (loc, "%<-foffload-abi%> option can be specified only for "
> -		"offload compiler");
> +      error_at (loc,
> +		"%qs option can be specified only for offload compiler",
> +		(code == OPT_foffload_abi_) ? "-foffload-abi"
> +					    : "-foffload-abi-host-opts");
>  #endif
>        break;
>

> -----Original Message-----
> From: Thomas Schwinge <tschwinge@baylibre.com>
> Sent: Tuesday, September 10, 2024 8:19 PM
> To: Prathamesh Kulkarni <prathameshk@nvidia.com>; Richard Biener
> <rguenther@suse.de>
> Cc: Andrew Pinski <pinskia@gmail.com>; gcc-patches@gcc.gnu.org; Jakub
> Jelinek <jakub@redhat.com>
> Subject: RE: [nvptx] Pass -m32/-m64 to host_compiler if it has
> multilib support
> 
> External email: Use caution opening links or attachments
> 
> 
> Hi Prathamesh!
> 
> On 2024-09-10T13:22:10+0000, Prathamesh Kulkarni
> <prathameshk@nvidia.com> wrote:
> >> -----Original Message-----
> >> From: Thomas Schwinge <tschwinge@baylibre.com>
> >> Sent: Monday, September 9, 2024 8:50 PM
> 
> >> > Could you please test the patch for gcn backend ?
> 
> I've successfully tested x86_64 host with GCN as well as nvptx
> offloading, and also ppc64le host with nvptx offloading.
Thanks for the thorough testing!
> 
> I just realized two more minor things:
> 
> > [nvptx] Pass host specific ABI opts from mkoffload.
> >
> > The patch adds an option -foffload-abi-host-opts, which is set by
> host
> > in TARGET_OFFLOAD_OPTIONS, and mkoffload then passes its value to
> > host_compiler.
> >
> 
> Please add here "       PR target/96265".
> 
> > gcc/ChangeLog:
> >       * common.opt (foffload-abi-host-opts): New option.
> >       * config/aarch64/aarch64.cc (aarch64_offload_options): Pass
> >       -foffload-abi-host-opts.
> >       * config/i386/i386-opts.cc (ix86_offload_options): Likewise.
> >       * config/rs6000/rs6000.cc (rs6000_offload_options): Likewise.
> >       * config/nvptx/mkoffload.cc (offload_abi_host_opts): Define.
> >       (compile_native): Append offload_abi_host_opts to
> argv_obstack.
> >       (main): Handle option -foffload-abi-host-opts.
> >       * config/gcn/mkoffload.cc (offload_abi_host_opts): Define.
> >       (compile_native): Append offload_abi_host_opts to
> argv_obstack.
> >       (main): Handle option -foffload-abi-host-opts.
> >       * lto-wrapper.cc (merge_and_complain): Handle
> >       -foffload-abi-host-opts.
> >       (append_compiler_options): Likewise.
> >       * opts.cc (common_handle_option): Likewise.
> >
> > Signed-off-by: Prathamesh Kulkarni <prathameshk@nvidia.com>
> 
> Given that we're adding a new option to 'gcc/common.opt', do we need
> to update (regenerate?) 'gcc/common.opt.urls'?  (I've not yet had the
> need myself, and therefore not yet looked up how to do that.)  Or
> maybe not, given that '-foffload-abi-host-opts=[...]' isn't
> documented?
I checked common.opt.urls doesn't seem to have entry for -foffload-abi,
so I guess it's probably not necessary for -foffload-abi-host-opts either ?
Or should we do it for both the options ?
> 
> Otherwise looks good to me; OK to push (with these minor items
> addressed, as necessary), thanks!
Thanks, I have committed the patch to trunk in:
https://gcc.gnu.org/git/?p=gcc.git;a=commit;h=e783a4a683762487cb003ae48235f3d44875de1b
Will post a follow up patch to regenerate common.opt.urls for -foffload-abi and -foffload-abi-host-opts if required.

Thanks,
Prathamesh
> 
> 
> Grüße
>  Thomas
> 
> 
> > diff --git a/gcc/common.opt b/gcc/common.opt index
> > ea39f87ae71..d270e524ff4 100644
> > --- a/gcc/common.opt
> > +++ b/gcc/common.opt
> > @@ -2361,6 +2361,10 @@ Enum(offload_abi) String(ilp32)
> > Value(OFFLOAD_ABI_ILP32)  EnumValue
> >  Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
> >
> > +foffload-abi-host-opts=
> > +Common Joined MissingArgError(option missing after %qs)
> > +-foffload-abi-host-opts=<options>    Specify host ABI options.
> > +
> >  fomit-frame-pointer
> >  Common Var(flag_omit_frame_pointer) Optimization  When possible do
> > not generate stack frames.
> > diff --git a/gcc/config/aarch64/aarch64.cc
> > b/gcc/config/aarch64/aarch64.cc index 6a3f1a23a9f..6ccf08d1cc0
> 100644
> > --- a/gcc/config/aarch64/aarch64.cc
> > +++ b/gcc/config/aarch64/aarch64.cc
> > @@ -19000,9 +19000,9 @@ static char *
> >  aarch64_offload_options (void)
> >  {
> >    if (TARGET_ILP32)
> > -    return xstrdup ("-foffload-abi=ilp32");
> > +    return xstrdup ("-foffload-abi=ilp32
> > + -foffload-abi-host-opts=-mabi=ilp32");
> >    else
> > -    return xstrdup ("-foffload-abi=lp64");
> > +    return xstrdup ("-foffload-abi=lp64
> > + -foffload-abi-host-opts=-mabi=lp64");
> >  }
> >
> >  static struct machine_function *
> > diff --git a/gcc/config/gcn/mkoffload.cc
> b/gcc/config/gcn/mkoffload.cc
> > index b8d981878ed..345bbf7709c 100644
> > --- a/gcc/config/gcn/mkoffload.cc
> > +++ b/gcc/config/gcn/mkoffload.cc
> > @@ -133,6 +133,8 @@ static const char *gcn_dumpbase;  static struct
> > obstack files_to_cleanup;
> >
> >  enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
> > +const char *offload_abi_host_opts = NULL;
> > +
> >  uint32_t elf_arch = EF_AMDGPU_MACH_AMDGCN_GFX900;  // Default GPU
> architecture.
> >  uint32_t elf_flags = EF_AMDGPU_FEATURE_SRAMECC_UNSUPPORTED_V4;
> >
> > @@ -819,17 +821,10 @@ compile_native (const char *infile, const char
> *outfile, const char *compiler,
> >    obstack_ptr_grow (&argv_obstack, gcn_dumpbase);
> >    obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
> >    obstack_ptr_grow (&argv_obstack, ".c");
> > -  switch (offload_abi)
> > -    {
> > -    case OFFLOAD_ABI_LP64:
> > -      obstack_ptr_grow (&argv_obstack, "-m64");
> > -      break;
> > -    case OFFLOAD_ABI_ILP32:
> > -      obstack_ptr_grow (&argv_obstack, "-m32");
> > -      break;
> > -    default:
> > -      gcc_unreachable ();
> > -    }
> > +  if (!offload_abi_host_opts)
> > +    fatal_error (input_location,
> > +              "%<-foffload-abi-host-opts%> not specified.");
> > + obstack_ptr_grow (&argv_obstack, offload_abi_host_opts);
> >    obstack_ptr_grow (&argv_obstack, infile);
> >    obstack_ptr_grow (&argv_obstack, "-c");
> >    obstack_ptr_grow (&argv_obstack, "-o"); @@ -998,6 +993,15 @@ main
> > (int argc, char **argv)
> >                        "unrecognizable argument of option %<" STR
> "%>");
> >       }
> >  #undef STR
> > +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> > +     {
> > +       if (offload_abi_host_opts)
> > +         fatal_error (input_location,
> > +                      "%<-foffload-abi-host-opts%> specified "
> > +                      "multiple times");
> > +       offload_abi_host_opts
> > +         = argv[i] + strlen ("-foffload-abi-host-opts=");
> > +     }
> >        else if (strcmp (argv[i], "-fopenmp") == 0)
> >       fopenmp = true;
> >        else if (strcmp (argv[i], "-fopenacc") == 0) diff --git
> > a/gcc/config/i386/i386-options.cc b/gcc/config/i386/i386-options.cc
> > index f79257cc764..55e0210260f 100644
> > --- a/gcc/config/i386/i386-options.cc
> > +++ b/gcc/config/i386/i386-options.cc
> > @@ -3680,8 +3680,8 @@ char *
> >  ix86_offload_options (void)
> >  {
> >    if (TARGET_LP64)
> > -    return xstrdup ("-foffload-abi=lp64");
> > -  return xstrdup ("-foffload-abi=ilp32");
> > +    return xstrdup ("-foffload-abi=lp64
> > + -foffload-abi-host-opts=-m64");  return xstrdup
> > + ("-foffload-abi=ilp32 -foffload-abi-host-opts=-m32");
> >  }
> >
> >  /* Handle "cdecl", "stdcall", "fastcall", "regparm", "thiscall",
> diff
> > --git a/gcc/config/nvptx/mkoffload.cc
> b/gcc/config/nvptx/mkoffload.cc
> > index 503b1abcefd..df16ee64736 100644
> > --- a/gcc/config/nvptx/mkoffload.cc
> > +++ b/gcc/config/nvptx/mkoffload.cc
> > @@ -61,6 +61,7 @@ static const char *omp_requires_file;  static
> const
> > char *ptx_dumpbase;
> >
> >  enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
> > +const char *offload_abi_host_opts = NULL;
> >
> >  /* Delete tempfiles.  */
> >
> > @@ -607,17 +608,10 @@ compile_native (const char *infile, const char
> *outfile, const char *compiler,
> >    obstack_ptr_grow (&argv_obstack, ptx_dumpbase);
> >    obstack_ptr_grow (&argv_obstack, "-dumpbase-ext");
> >    obstack_ptr_grow (&argv_obstack, ".c");
> > -  switch (offload_abi)
> > -    {
> > -    case OFFLOAD_ABI_LP64:
> > -      obstack_ptr_grow (&argv_obstack, "-m64");
> > -      break;
> > -    case OFFLOAD_ABI_ILP32:
> > -      obstack_ptr_grow (&argv_obstack, "-m32");
> > -      break;
> > -    default:
> > -      gcc_unreachable ();
> > -    }
> > +  if (!offload_abi_host_opts)
> > +    fatal_error (input_location,
> > +              "%<-foffload-abi-host-opts%> not specified.");
> > + obstack_ptr_grow (&argv_obstack, offload_abi_host_opts);
> >    obstack_ptr_grow (&argv_obstack, infile);
> >    obstack_ptr_grow (&argv_obstack, "-c");
> >    obstack_ptr_grow (&argv_obstack, "-o"); @@ -721,6 +715,15 @@ main
> > (int argc, char **argv)
> >                        "unrecognizable argument of option " STR);
> >       }
> >  #undef STR
> > +      else if (startswith (argv[i], "-foffload-abi-host-opts="))
> > +     {
> > +       if (offload_abi_host_opts)
> > +         fatal_error (input_location,
> > +                      "%<-foffload-abi-host-opts%> specified "
> > +                      "multiple times");
> > +       offload_abi_host_opts
> > +         = argv[i] + strlen ("-foffload-abi-host-opts=");
> > +     }
> >        else if (strcmp (argv[i], "-fopenmp") == 0)
> >       fopenmp = true;
> >        else if (strcmp (argv[i], "-fopenacc") == 0) diff --git
> > a/gcc/config/rs6000/rs6000.cc b/gcc/config/rs6000/rs6000.cc index
> > 08579bc83e6..0bf8bae27f5 100644
> > --- a/gcc/config/rs6000/rs6000.cc
> > +++ b/gcc/config/rs6000/rs6000.cc
> > @@ -17330,9 +17330,9 @@ static char *
> >  rs6000_offload_options (void)
> >  {
> >    if (TARGET_64BIT)
> > -    return xstrdup ("-foffload-abi=lp64");
> > +    return xstrdup ("-foffload-abi=lp64
> > + -foffload-abi-host-opts=-m64");
> >    else
> > -    return xstrdup ("-foffload-abi=ilp32");
> > +    return xstrdup ("-foffload-abi=ilp32
> > + -foffload-abi-host-opts=-m32");
> >  }
> >
> >
> > diff --git a/gcc/lto-wrapper.cc b/gcc/lto-wrapper.cc index
> > c07765b37a2..7de045da9b9 100644
> > --- a/gcc/lto-wrapper.cc
> > +++ b/gcc/lto-wrapper.cc
> > @@ -484,6 +484,7 @@ merge_and_complain (vec<cl_decoded_option>
> > &decoded_options,
> >
> >
> >       case OPT_foffload_abi_:
> > +     case OPT_foffload_abi_host_opts_:
> >         if (existing_opt == -1)
> >           decoded_options.safe_push (*foption);
> >         else if (foption->value !=
> > decoded_options[existing_opt].value)
> > @@ -745,6 +746,7 @@ append_compiler_options (obstack *argv_obstack,
> vec<cl_decoded_option> opts)
> >       case OPT_fopenacc:
> >       case OPT_fopenacc_dim_:
> >       case OPT_foffload_abi_:
> > +     case OPT_foffload_abi_host_opts_:
> >       case OPT_fcf_protection_:
> >       case OPT_fasynchronous_unwind_tables:
> >       case OPT_funwind_tables:
> > diff --git a/gcc/opts.cc b/gcc/opts.cc index
> fc6abf6f582..a78f73e57e3
> > 100644
> > --- a/gcc/opts.cc
> > +++ b/gcc/opts.cc
> > @@ -3070,11 +3070,14 @@ common_handle_option (struct gcc_options
> *opts,
> >        break;
> >
> >      case OPT_foffload_abi_:
> > +    case OPT_foffload_abi_host_opts_:
> >  #ifdef ACCEL_COMPILER
> >        /* Handled in the 'mkoffload's.  */  #else
> > -      error_at (loc, "%<-foffload-abi%> option can be specified
> only for "
> > -             "offload compiler");
> > +      error_at (loc,
> > +             "%qs option can be specified only for offload
> compiler",
> > +             (code == OPT_foffload_abi_) ? "-foffload-abi"
> > +                                         :
> > + "-foffload-abi-host-opts");
> >  #endif
> >        break;
> >