OffTarget

Observed issues › hecbench/gcc-c-float16-array-in-target-lto-ice

Appears to behave incorrectlyCompiler crash at link time · C · OpenMP · Open

GCC: compiler crash on _Float16 data in a target region

Legality: contested. An independent adjudicating reading finds the result not required: _Float16 is not part of the base languages OpenMP names; section 1.6 says an implementation need not support a construct the base language implementation does not support. Whether the device compiler handles _Float16 conversions is not specified, so the ICE is not a violation of a required result. Both readings are right. How this is checked

GCC 16.2 stops with an internal compiler error in lto1 when a target region reads an array of _Float16 values; the same program with float elements compiles and runs.

Affected variants

CompilerMI300A gfx942
unified memory (XNACK 1)
MI300A gfx942
copy (XNACK 0)
MI250X gfx90a
copy (XNACK 0)
MI300X gfx942
copy (XNACK 0)
MI350X gfx950
copy (XNACK 0)
A100 sm_80
separate memory
A100 sm_80
managed memory
AFAR 24.3.0correctcorrectunknown
not run
correctcorrectn/an/a
CCE 19.0.0correctcorrectunknown
not run
unknown
not run
unknown
not run
n/an/a
CCE 20.0.2correctcorrectunknown
not run
unknown
not run
unknown
not run
n/an/a
CCE 21.0.3correctcorrectunknown
not run
unknown
not run
unknown
not run
n/an/a
CCE 21.0.4correctcorrectunknown
not run
unknown
not run
unknown
not run
n/an/a
NVHPC 25.11n/an/an/an/an/acorrectcorrect
NVHPC 26.9n/an/an/an/an/acorrectcorrect
GCC 16.2.0behaves incorrectlybehaves incorrectlybehaves incorrectlybehaves incorrectlybehaves incorrectlybehaves incorrectlyn/a

"behaves incorrectly" / "not supported": this issue, confirmed in up to four environments (another build directory, another node, a repeat). "fails another way": a different failure, not counted as this issue. "unknown": not run there, or the runs disagree.

Messages

ResultCompiler and GPUMessage
behaves incorrectlyGCC 16.2.0, MI300A unified
GCC 16.2.0, MI300A copy
show
In function ‘main._omp_fn.1’:
lto1: internal compiler error: in make_ssa_name_fn, at tree-ssanames.cc:355
0x1a963dd internal_error(char const*, ...)
	<dir>:787
0x7b22c3 fancy_abort(char const*, int, char const*)
	<dir>:1813
0x79cecc make_ssa_name_fn(function*, tree_node*, gimple*, unsigned int)
	<dir>:355
behaves incorrectlyGCC 16.2.0, MI250X
show
In function ‘main._omp_fn.1’:
lto1: internal compiler error: in make_ssa_name_fn, at tree-ssanames.cc:355
0x1a963dd internal_error(char const*, ...)
	<dir><host>-gcc-src/gcc-16.2.0/gcc/diagnostic-global-context.cc:787
0x7b22c3 fancy_abort(char const*, int, char const*)
	<dir><host>-gcc-src/gcc-16.2.0/gcc/diagnostics/context.cc:1813
0x79cecc make_ssa_name_fn(function*, tree_node*, gimple*, unsigned int)
	<dir><host>-gcc-src/gcc-16.2.0/gcc/tree-ssanames.cc:355
behaves incorrectlyGCC 16.2.0, MI300X
GCC 16.2.0, MI350X
show
In function ‘main._omp_fn.1’:
lto1: internal compiler error: in make_ssa_name_fn, at tree-ssanames.cc:355
0x1a78aed internal_error(char const*, ...)
	<home>/<user>/gcc-work/gcc-16.2.0/gcc/diagnostic-global-context.cc:787
0x7a88fe fancy_abort(char const*, int, char const*)
	<home>/<user>/gcc-work/gcc-16.2.0/gcc/diagnostics/context.cc:1813
0x79332c make_ssa_name_fn(function*, tree_node*, gimple*, unsigned int)
	<home>/<user>/gcc-work/gcc-16.2.0/gcc/tree-ssanames.cc:355
behaves incorrectlyGCC 16.2.0, A100 separate
show
In function ‘main._omp_fn.1’:
lto1: internal compiler error: in make_ssa_name_fn, at tree-ssanames.cc:355
0x17b7f6b internal_error(char const*, ...)
	<home>/gcc-build/gcc-16.2.0/gcc/diagnostic-global-context.cc:787
0x6ba0dd fancy_abort(char const*, int, char const*)
	<home>/gcc-build/gcc-16.2.0/gcc/diagnostics/context.cc:1813
0x6a7358 make_ssa_name_fn(function*, tree_node*, gimple*, unsigned int)
	<home>/gcc-build/gcc-16.2.0/gcc/tree-ssanames.cc:355

Reproduce

GCC 16.2.0, MI300A unified

export LD_LIBRARY_PATH=$GCC_16_2/lib64:/opt/rocm-7.2.1/lib
$GCC_16_2/bin/gcc -fopenmp -foffload=amdgcn-amdhsa -foffload-options=amdgcn-amdhsa=-march=gfx942 -O2 repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro

GCC 16.2.0, A100 separate

export PATH=/usr/local/cuda/bin:$PATH
export LD_LIBRARY_PATH=$GCC_16_2/lib64
$GCC_16_2/bin/gcc -fopenmp -foffload=nvptx-none -foffload-options=nvptx-none=-march=sm_80 -foffload-options=nvptx-none=-latomic -O2 repro.c -o repro
OMP_TARGET_OFFLOAD=mandatory ./repro

AFAR 24.3.0, MI300A unified (runs correctly)

$AFAR/bin/amdclang -fopenmp --offload-arch=gfx942 -O2 repro.c -o repro
HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
Program (59 lines)
// # Finding (hecbench): GCC: a target region that reads an array of _Float16 values crashes the compiler (lto1)
//
// - toolchains and modes that fail: gcc_g90 (GCC 16.2.0), gcc_sm80 (GCC 16.2.0); modes: copy; -O2
// - build: gcc -fopenmp -foffload=amdgcn-amdhsa -foffload-options=amdgcn-amdhsa=-march=gfx90a -O2 (gcc_g90); gcc -fopenmp
//   -foffload=nvptx-none -foffload-options=nvptx-none=-march=sm_80 -foffload-options=nvptx-none=-latomic -O2 (gcc_sm80)
// - failure: ice
// - observed: `Please submit a full bug report, with preprocessed source (by using -freport-bug).`; identical signature on 2 of 2
//   failing cells
// - expected: compiles and prints `total= 2048.000000` with `check=pass`
// - rule: "The functionality added to the target directive is the inclusion of an executable region to be executed on a device"
//   (5.2 §13.8 p.284 l.12-13): the region here is a loop that converts each element of a mapped array of _Float16 to float, and
//   the specification names no base type that a target region may not use; with float elements instead the compile succeeds
// - same-compiler control: controls/f16-ctl-float-array.c (one change from this file; correct on gcc_g90 (GCC 16.2.0), gcc_sm80
//   (GCC 16.2.0); rows results/hecbench/red/tio_g90.jsonl, results/hecbench/red/tuo.jsonl, results/hecbench/red/wt1.jsonl)
// - other families: afar (AFAR 24.3.0), cce19 (CCE 19.0.0), cce20 (CCE 20.0.2), cce21 (CCE 21.0.3), cce21_4 (CCE 21.0.4), nvhpc
//   (NVHPC 25.11), nvhpc269 (NVHPC 26.9) run the file correctly (rows results/hecbench/red/tio_g90.jsonl,
//   results/hecbench/red/tuo.jsonl, results/hecbench/red/wt1.jsonl)
// - gate and confirmer: campaign gate of the programs below: 1 cells, 1 confirmed in the four environments (a-d; (c) on another
//   node or GPU; rows results/hecbench/dev/gate_reruns.part.<host>.2090963.jsonl,
//   results/hecbench/dev/gate_reruns.part.<host>.2185748.jsonl); 2 of 2 cells reproduce in all four environments (a-d; (c) on
//   another node or GPU than (a); rows results/hecbench/red/tio_g90.jsonl, results/hecbench/red/tio_gate_abd.jsonl,
//   results/hecbench/red/tio_gate_c.jsonl, results/hecbench/red/wt1.jsonl, results/hecbench/red/wt_gate_abd.jsonl,
//   results/hecbench/red/wt_gate_c.jsonl)
// - origin: HeCBench `src/f16sp-omp/main.cpp` (https://github.com/ORNL/HeCBench, commit
//   8f44ff3e69fa003f870963045988190a275534cb), whose gcc_g90 and gcc_sm80 cells fail with the same compiler line, `lto1: internal
//   compiler error: in make_ssa_name_fn, at tree-ssanames.cc:355`. This file is a small program written for this finding and
//   reduced with tools/reduce.py; it holds no benchmark source.
// - programs (1): hecbench-f16sp-omp
// - legality: offload.engine.prerun.check (the rules, and AFAR's front end on Tuolumne, 2026-11-06: two witnesses) accepts this file
//   and the control; CCE and NVHPC compile and run it correctly (front ends of two other families). GCC is a compiler under test
//   only and never a legality witness.
#include <stdio.h>
#include <stdlib.h>
#include <omp.h>

typedef _Float16 half_t;

int main(void)
{
  const int size = 1024;
  int hf = 1;
  half_t *ax = malloc(size * sizeof(half_t));
  float *out = malloc(size * sizeof(float));
  float total = 0.f;
  for (int i = 0; i < size; i++) { ax[i] = 2; out[i] = 0; }
  #pragma omp target map(from: hf)
  hf = omp_is_initial_device() ? 1 : 0;

  #pragma omp target map(to: ax[0:size]) map(from: out[0:size])
  for (int i = 0; i < size; i++)
    out[i] = (float)ax[i];

  for (int i = 0; i < size; i++) total += out[i];
  printf("total= %f\n", total);
  printf("RESULT ran_on_host=%d check=%s\n", hf, total == 2048.f ? "pass" : "fail");
  free(ax);
  free(out);
  return 0;
}

Source in the repository

Something to add? Dispute its legality · Report it fixed in a newer release · Report a mistake on this page

Expected and observed

Observed where it behaves incorrectly

In function ‘main._omp_fn.1’:
lto1: internal compiler error: in make_ssa_name_fn, at tree-ssanames.cc:355
0x1a963dd internal_error(char const*, ...)
	<dir>:787
0x7b22c3 fancy_abort(char const*, int, char const*)
	<dir>:1813
0x79cecc make_ssa_name_fn(function*, tree_node*, gimple*, unsigned int)
	<dir>:355

What the specification requires

The specification text this program relies on and the compiler is expected to support.

OpenMP 5.2 §13.8 p.284 l.12-13

The functionality added to the target directive is the inclusion of an executable region to be executed on a device

Evidence

Same compiler, one change

  • one change from the finding; correct on gcc_g90, gcc_sm80: AFAR 24.3.0: correct; CCE 19.0.0: correct; CCE 20.0.2: correct; CCE 21.0.3: correct; CCE 21.0.4: correct; NVHPC 25.11: correct; NVHPC 26.9: correct; GCC 16.2.0: correct
    Control program
    #include <stdio.h>
    #include <stdlib.h>
    #include <omp.h>
    
    typedef float half_t;
    
    int main(void)
    {
      const int size = 1024;
      int hf = 1;
      half_t *ax = malloc(size * sizeof(half_t));
      float *out = malloc(size * sizeof(float));
      float total = 0.f;
      for (int i = 0; i < size; i++) { ax[i] = 2; out[i] = 0; }
      #pragma omp target map(from: hf)
      hf = omp_is_initial_device() ? 1 : 0;
    
      #pragma omp target map(to: ax[0:size]) map(from: out[0:size])
      for (int i = 0; i < size; i++)
        out[i] = (float)ax[i];
    
      for (int i = 0; i < size; i++) total += out[i];
      printf("total= %f\n", total);
      printf("RESULT ran_on_host=%d check=%s\n", hf, total == 2048.f ? "pass" : "fail");
      free(ax);
      free(out);
      return 0;
    }
    

Other compilers run it correctly

AFAR 24.3.0, CCE 19.0.0, CCE 20.0.2, CCE 21.0.3, CCE 21.0.4, NVHPC 25.11, NVHPC 26.9