A firstprivate scalar allocated with omp_const_mem_alloc has the wrong value on the device

// # Draft finding: omp52-c-w-allocate-const-firstprivate on afar
//
// - compiler: afar
// - toolchain: afar; arch: gfx942; machine: tuolumne; memory mode: unified; opt: -O2
// - program: catalog omp52-c-w-allocate-const-firstprivate (c, version 52, sha256 2f1a38cbd60a463c)
// - expected: check=pass (trusted; elision + model)
// - actual: wrong_answer check=fail (check_fail: ran_on_host=0 check=fail)
// - oracle audit: run on AFAR, CCE 19, 20, 21 (copy mode) and NVHPC: no toolchain passes the original (NVHPC rejects the clauses), so the mutant cannot be credited (unaudited, not weak); in sequential elision the original prints check=pass and the mutant (x's contribution removed) check=fail, so the check depends on the firstprivate value
// - verdict: compiler_bug by rule R20-runtime-bug: a legal program fails at run time, deterministically, against an expected value from two or more independent sources with an executed witness, and another toolchain, mode, arch or the elision gives the expected value
// - reason: expected trusted check=pass (elision + model): stored main verdict: elision + model agree; witnesses: elision
// - citations: RUL-DERIVED-NEEDS-EXECUTION (rulings.md); RUL-WITNESS-PER-MODE (rulings.md); RUL-SAME-COMPILER-RUNS (rulings.md); RUL-PRESENCE-WITNESS (rulings.md); RUL-METAMORPHIC-OTHER-FAMILY (rulings.md)
// - legality: legal: offload.engine.prerun.check accepts it (rerun: offload.engine.rules with the
//   race rule and the restriction checks, and the AFAR clang front end, not under test here); no
//   compiler under test, gfortran or gcc is a legality witness
// - witnesses: elision
// - withdrawn: none
//
// Generated by compliance.widen (OpenMP 5.2, C11; catalog widening). Do not edit.
// id: omp52-c-w-allocate-const-firstprivate
// primary atom: pair.omp.clause.allocate+omp.clause.firstprivate
// atoms: pair.omp.clause.allocate+omp.clause.firstprivate, omp.clause.allocate, omp.clause.firstprivate
// legal: yes
// observable: yes
// ran_on_host: inside the target region of interest
// evidence: construct
// needs: none
// std: c11
// feature: a firstprivate scalar allocated with omp_const_mem_alloc (uses_allocators) on a target loop
// spec: the parts' own programs
// pair of: omp.clause.allocate + omp.clause.firstprivate (each has a program of its own)
#include <stdio.h>
#include <stdbool.h>
#include <omp.h>

enum { n = 64 };

int main(void) {
  int a[n], b[n], ref[n];
  int host_flag = 1;
  int x = 5;
  for (int i = 0; i < n; i++) {
    a[i] = i + 1;
    b[i] = 0;
    ref[i] = i + 1 + 5;
  }
  #pragma omp target teams distribute parallel for map(to: a) map(from: b) map(tofrom: \
      host_flag) firstprivate(x) allocate(omp_const_mem_alloc: x) \
      uses_allocators(omp_const_mem_alloc)
  for (int i = 0; i < n; i++) {
    b[i] = a[i] + x;
    if (i == 0) host_flag = omp_is_initial_device();
  }
  bool ok = true;
  for (int i = 0; i < n; i++)
    ok = ok && (b[i] == ref[i]);
  printf("RESULT ran_on_host=%d check=%s\n", host_flag, ok ? "pass" : "fail");
  return 0;
}
// spec (requirement): Variables in the constant memory space may be initialized through the firstprivate clause.
// "Variables allocated in the omp_const_mem_space memory space may be initialized through the firstprivate clause"
//   (5.2 §6.1 p.171 l.9-10)
// spec (requirement): A firstprivate list item that also appears in allocate gets its storage from that allocator and
//   is initialized from the original list item.
// "If a list item in the clause also appears in a data-sharing attribute clause on the same directive that privatizes
//   the list item, allocations that arise from that list item in the clause will be provided by the memory
//   allocator." (5.2 §6.6 p.178 l.12-14)
// spec (requirement): The new list item of a firstprivate clause is initialized from the original list item.
// "In addition, the new list item is initialized from the original list item that exists before the construct." (5.2
//   §5.4.4 p.112 l.11-12)
// spec (this program): firstprivate(x) allocate(omp_const_mem_alloc: x) with uses_allocators(omp_const_mem_alloc) is
//   the form 5.2 names for constant memory, so every device copy of x must hold 5; craycc yields the wrong value on
//   the device.