uses_allocators with allocator traits or a memory space: "sorry, unimplemented"

// # Finding (GCC 16.2): uses_allocators with traits or a memory space is not implemented (C); the reproducer is the catalog program
//   omp52-c-clause-uses_allocators-traits
//
// - toolchains: gcc_g90 (GCC 16.2.0, AMD Instinct MI250X gfx90a, Tioga) and gcc_sm80 (GCC 16.2.0,
//   NVIDIA A100 sm_80, wingtip); separate memory only; -O2; OMP_TARGET_OFFLOAD=mandatory
// - build: gcc -fopenmp -foffload=amdgcn-amdhsa -foffload-options=amdgcn-amdhsa=-march=gfx90a -O2 (gcc_g90); -foffload=nvptx-none
//   -foffload-options=nvptx-none=-march=sm_80 -foffload-options=nvptx-none=-latomic -O2 (gcc_sm80)
// - failure: rejected: sorry, unimplemented: .uses_allocators. clause
// - observed: `sorry, unimplemented: 'uses_allocators' clause` on both GPUs.
// - expected: compiles; `RESULT ran_on_host=0 check=pass`
// - GCC 16.2 documentation (libgomp.texi of the GCC 16.2.0 release tarball, "OpenMP Implementation Status", OpenMP 5.2 table): "New
//   memspace and traits modifiers for uses_allocators": P (partial), "Only predefined allocators". The message is GCC's own "sorry,
//   unimplemented".
// - spec: the uses_allocators clause and its traits and memspace modifiers (5.2 §6.8 p.181).
// - other families: AFAR reports uses_allocators as not yet implemented; CCE rejects the 5.2 traits form
//   (catalog/cce-c-uses-allocators-traits-rejected.c).
// - gate and confirmer: 10 cells (5 gcc_g90, 5 gcc_sm80), each confirmed in the gate's four environments (a-d;
//   results/gcc/gate-tioga/ and gate-wingtip/ verdicts.jsonl); confirmer verdicts (results/gcc/confirm/batch.gcc_*.jsonl):
//   unsupported (R05-unsupported) 10
// - catalog programs (5): omp50-c-clause-uses_allocators-traits, omp52-c-clause-uses_allocators-memspace,
//   omp52-c-clause-uses_allocators-traits, omp60-c-clause-uses_allocators-memspace, omp60-c-clause-uses_allocators-traits
// - legality: the catalog programs are legal by construction (offload.engine.gen, c_gen and acc_gen
//   check every program against the catalog rules and the offload scope as it is built; no
//   compiler). GCC is a compiler under test only: never a legality witness, here or for another
//   family's case (catalog/machines.yaml under_test_only).
// Generated by compliance.c_gen (OpenMP 5.2, C11): the C port (offload.engine.f2c) of omp52-clause-uses_allocators-traits. Do not edit.
// id: omp52-c-clause-uses_allocators-traits
// primary atom: omp.clause.uses_allocators.traits
// legal: yes
// observable: no (unverified: the check covers the device kernel, not the atom)
// ran_on_host: inside the target region of interest
// evidence: unverified
// needs: none
// std: c11
// unverified (oracle_not_written): re-categorised (was implementation_defined): memory from the allocator 'will be byte
// aligned to at least the value specified for the alignment trait' (5.2 §6.2 p.173, 6.0 §8.2 p.306), so a check that t's address is
// a multiple of 64 is a deterministic oracle of the original; it is not written. A mutant without the trait could still be 64-byte
// aligned (any alignment conforms then), AFAR rejects uses_allocators and CCE 19's mutant without the trait is an internal compiler
// error (final-review re-audit)
#include <stdio.h>
#include <stdbool.h>
#include <omp.h>

enum { n = 64 };

int main(void) {
  int a[n];
  int b[n];
  int ref[n];
  int host_flag;
  bool ok;
  int t;
  omp_allocator_handle_t my_alloc;
  const const omp_alloctrait_t my_traits[1] = {{omp_atk_alignment, 64}};
  for (int i = 1; i <= n; i++) {
    a[i - 1] = i;
    b[i - 1] = 0;
    ref[i - 1] = 2 * i + 1;
  }
  host_flag = 1;
  #pragma omp target map(to: a) map(from: b) map(tofrom: host_flag) \
      uses_allocators(traits(my_traits): my_alloc)
  {
    host_flag = (omp_is_initial_device() ? 1 : 0);
    #pragma omp parallel for default(none) shared(a, b, my_alloc) private(t) \
        allocate(my_alloc: t)
    for (int i = 1; i <= n; i++) {
      t = 2 * a[i - 1];
      b[i - 1] = t + 1;
    }
  }
  bool t1_ = true;
  for (int k2_ = 0; k2_ < n; k2_++)
    t1_ = t1_ && (b[k2_] == ref[k2_]);
  ok = t1_;
  printf("RESULT ran_on_host=%d check=%s\n", host_flag, ok ? "pass" : "fail");
  return 0;
}