Observed issues › catalog/cce-c-allocate-allocator-modifier-rejected
Appears not to supportRefused at compile time · C · OpenMP · Fixed in 20.0.2
craycc rejects the allocator modifier of the allocate clause
Legality: independently confirmed. 2 independent readings find the program legal and the result required. How this is checked
Affected variants
| Compiler | MI300A 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.0 | correct | correct | correct | correct | correct | n/a | n/a |
| CCE 19.0.0 | not supported | not supported | not supported | unknown not run | unknown not run | n/a | n/a |
| CCE 20.0.2 | correct | correct | correct | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.3 | correct | correct | correct | unknown not run | unknown not run | n/a | n/a |
| CCE 21.0.4 | correct | correct | correct | unknown not run | unknown not run | n/a | n/a |
| NVHPC 25.11 | n/a | n/a | n/a | n/a | n/a | fails another way | fails another way |
| NVHPC 26.9 | n/a | n/a | n/a | n/a | n/a | fails another way | fails another way |
| GCC 16.2.0 | correct | correct | correct | correct | correct | correct | n/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
| Result | Compiler and GPU | Message |
|---|---|---|
| not supported | CCE 19.0.0, MI300A unified CCE 19.0.0, MI300A copy CCE 19.0.0, MI250X | showcce-c-allocate-allocator-modifier-rejected.c:25:18: error: call to undeclared function 'allocator'; ISO C99 and later do not support implicit function declarations [-Wimplicit-function-declaration]
25 | allocate(allocator(omp_default_mem_alloc): t)
| ^
cce-c-allocate-allocator-modifier-rejected.c:25:18: note: did you mean 'alloca'?
/usr/include/alloca.h:32:14: note: 'alloca' declared here
32 | extern void *alloca (size_t __size) __THROW;
| ^
1 error generated. |
| fails another way | NVHPC 25.11, A100 separate NVHPC 25.11, A100 managed NVHPC 26.9, A100 separate NVHPC 26.9, A100 managed | show"cce-c-allocate-allocator-modifier-rejected.c", line 21: error: invalid text in pragma
uses_allocators(omp_default_mem_alloc)
^
"cce-c-allocate-allocator-modifier-rejected.c", line 21: error: extra text after expected end of preprocessing directive
uses_allocators(omp_default_mem_alloc)
^
"cce-c-allocate-allocator-modifier-rejected.c", line 25: error: invalid text in pragma
allocate(allocator(omp_default_mem_alloc): t) |
Reproduce
CCE 19.0.0, MI300A unified
module load cce/19.0.0 craype-accel-amd-gfx942 rocm/6.4.3 module unload cray-mpich craycc -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx942 -fopenmp-version=52 -O2 repro.c -o repro HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
AFAR 24.3.0, MI300A unified (runs correctly)
$AFAR/bin/amdclang -fopenmp --offload-arch=gfx942 -fopenmp-version=52 -fopenmp-version=52 -O2 repro.c -o repro HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
Program (72 lines)
// # Draft finding: omp52-c-clause-allocate-allocator on cce19
//
// - compiler: cce19
// - toolchain: cce19; arch: gfx942; machine: tuolumne; memory mode: unified; opt: -O2
// - program: catalog omp52-c-clause-allocate-allocator (c, version 52, sha256 f3fb2637e316cfa9)
// - expected: None (None; )
// - actual: rejected (diagnostic: <tmpdir>/omp52-c-clause-allocate-allocator.c:35:18: error: call to undeclared function 'allocator'; ISO C99 and later do not support implicit function declarations [-Wimplicit-function-dec)
// - verdict: compiler_bug by rule R11-compile-rejection: a legal program the compiler rejects with a diagnostic (or whose link fails on symbols the compiler's own code needs, or whose compile does not finish) while another compiler family compiles it
// - reason:
// - citations: RUL-COMPILE-TIME (rulings.md)
// - legality: legal (legal by construction: offload.engine.gen / c_gen / acc_gen check every program against the catalog rules and the offload scope as it is built (no compiler))
// - witnesses: none
// - withdrawn: none
//
// Generated by compliance.c_gen (OpenMP 5.2, C11): the C port (offload.engine.f2c) of omp52-clause-allocate-allocator. Do not edit.
// id: omp52-c-clause-allocate-allocator
// primary atom: omp.clause.allocate.allocator
// 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 (needs_other_infrastructure): the program cannot observe which allocator provides the private copy of t (no OpenMP routine reports a variable's allocator; omp_default_mem_alloc and the default allocator both return ordinary memory; C mutant survived on AFAR)
#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;
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(omp_default_mem_alloc)
{
host_flag = (omp_is_initial_device() ? 1 : 0);
#pragma omp parallel for default(none) shared(a, b) private(t) \
allocate(allocator(omp_default_mem_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;
}
// spec (requirement): The allocate clause of OpenMP 5.2 accepts the allocator-complex-modifier form allocator(...),
// with the syntax and semantics of the allocator clause.
// "The allocator-complex-modifier and align-modifier have the same syntax and semantics for the allocate clause as
// the allocator and align clauses have for the allocate directive." (5.2 §6.6 p.178 l.16-18)
// spec (requirement): The allocator clause names the memory allocator used for allocations associated with the
// construct.
// "The allocator clause specifies the memory allocator to be used for allocations associated with the construct on
// which the clause appears." (5.2 §6.4 p.175 l.21-23)
// spec (this program): allocate(allocator(omp_default_mem_alloc): t) uses the complex-modifier spelling; craycc reads
// allocator as a call to an undeclared function and rejects the clause, while the 5.0 spelling
// allocate(omp_default_mem_alloc: t) is accepted.
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
cce-c-allocate-allocator-modifier-rejected.c:25:18: error: call to undeclared function 'allocator'; ISO C99 and later do not support implicit function declarations [-Wimplicit-function-declaration]
25 | allocate(allocator(omp_default_mem_alloc): t)
| ^
cce-c-allocate-allocator-modifier-rejected.c:25:18: note: did you mean 'alloca'?
/usr/include/alloca.h:32:14: note: 'alloca' declared here
32 | extern void *alloca (size_t __size) __THROW;
| ^
1 error generated.What the specification requires
The specification text this program relies on and the compiler is expected to support.
OpenMP 5.2 §6.6 p.178 l.16-18
The allocator-complex-modifier and align-modifier have the same syntax and semantics for the allocate clause as the allocator and align clauses have for the allocate directive.
OpenMP 5.2 §6.4 p.175 l.21-23
The allocator clause specifies the memory allocator to be used for allocations associated with the construct on which the clause appears.
Evidence
Same compiler, one change
- allocate(omp_default_mem_alloc: t), the 5.0 spelling: 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
Control program
#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; 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(omp_default_mem_alloc) { host_flag = (omp_is_initial_device() ? 1 : 0); #pragma omp parallel for default(none) shared(a, b) private(t) \ allocate(omp_default_mem_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; }
Other compilers run it correctly
AFAR 24.3.0, GCC 16.2.0