[Bug lto/127589] New: [OpenMP] Offload: ICE in device LTO if the host code uses '_Float16', unsupported by gcn and nvptx

[burnus at gcc dot gnu.org via Gcc-bugs]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=127589

            Bug ID: 127589
           Summary: [OpenMP] Offload: ICE in device LTO if the host code
                    uses '_Float16', unsupported by gcn and nvptx
           Product: gcc
           Version: 17.0
            Status: UNCONFIRMED
          Keywords: ice-on-valid-code, openmp
          Severity: normal
          Priority: P3
         Component: lto
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
  Target Milestone: ---
            Target: gcn,nvptx

While reducing the following testcase, a common pattern emerged:

  All of them use _Float16

As one of the reduced test cases shows, even:

using __half = _Float16;
void f16AtomicOnGlobalMem(__half) {
#pragma omp target
  ;
}

is enough, i.e. the device side does not even use _Float16 in the code gen!

* * *

The following examples are reduced from https://github.com/ORNL/HeCBench/

The error message of the first three and the later ones differs with regards to
where the message is printed:

Both start with "lto1: internal compiler error: tree check: expected class
'type', have 'exceptional' (error_mark)" but then continue with

* in lto_fixup_prevailing_type, at lto/lto-common.cc:1006
* in type_with_linkage_p, at ipa-utils.h:219

but otherwise is identical.

And all are to be compiled with 'g++ -fopenmp' – albeit the reduced tests might
also work with 'gcc -fopenmp'.

NOTE: For _Float128, it fails with:
  lto1: fatal error: nvptx-none - 128-bit-precision floating-point numbers
  unsupported (mode 'TF')





Reduced src/addBiasResidualLayerNorm-omp testcase:
--------------------------------------------------
template <int> void layer(int, int, int) {
  _Float16 h_output;
#pragma omp target
  h_output = 0;
}
int main_m, main_repeat;
main() {
  for (int n;;)
    layer<1>(main_m, n, main_repeat);

Reduced src/bgmv-omp testcase:
------------------------------


Reduced src/attentionMergeState-omp testcase:
---------------------------------------------
void merge_attn_states_kernel2(_Float16 *) {
#pragma omp target
  ;
}

Reduced src/f16atomic-omp testcase:
-----------------------------------
using __half = _Float16;
void f16AtomicOnGlobalMem(__half) {
#pragma omp target
  ;
}

Reduced src/gelu-omp testcase:
------------------------------
using __half = _Float16;
int gelu_bias_loop_base_batch_size, gelu_bias_loop_base_height;
void gelu_bias_loop_base(__half) {
#pragma omp target teams distribute
  for (int batch = 0; batch < gelu_bias_loop_base_batch_size; batch++)
    for (int x;;)
      ;
}