Observed issues › catalog/cce20-c-defaultmap-private-storage-rejected
craycc rejects defaultmap(private) (OpenMP 6.0)
Legality: independently confirmed. One independent reading was uncertain; 2 adjudicating readings that weighed both find the program legal and the result required: OpenMP 6.0 lists private among the implicit-behavior keywords of defaultmap, with variable-category defaulting to all. The program is therefore conforming and a 6.0 compiler must accept the keyword. a, b and host_flag are explicitly mapped. c and w are implicitly private on the target, so the host w stays 0 and b gets 2a+1, as expected. The refusal is a compiler error. 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 | not supported | not supported | not supported | 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 | fails another way | fails another way | fails another way | fails another way | fails another way | fails another way | 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 CCE 20.0.2, MI300A unified CCE 20.0.2, MI300A copy CCE 20.0.2, MI250X | showcce20-c-defaultmap-private-storage-rejected.c:24:80: error: expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'
24 | #pragma omp target map(to: a) map(from: b) map(tofrom: host_flag) defaultmap(private)
| ^
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"cce20-c-defaultmap-private-storage-rejected.c", line 24: error: invalid text in pragma
#pragma omp target map(to: a) map(from: b) map(tofrom: host_flag) defaultmap(private)
^
1 error detected in the compilation of "cce20-c-defaultmap-private-storage-rejected.c". |
| fails another way | GCC 16.2.0, MI300A unified GCC 16.2.0, MI300A copy GCC 16.2.0, MI250X GCC 16.2.0, MI300X GCC 16.2.0, MI350X GCC 16.2.0, A100 separate | showcce20-c-defaultmap-private-storage-rejected.c: In function ‘main’:
cce20-c-defaultmap-private-storage-rejected.c:24:80: error: expected ‘alloc’, ‘to’, ‘from’, ‘tofrom’, ‘firstprivate’, ‘none’ or ‘default’ before ‘private’
24 | #pragma omp target map(to: a) map(from: b) map(tofrom: host_flag) defaultmap(private)
| ^~~~~~~ |
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=60 -O2 repro.c -o repro HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
CCE 20.0.2, MI300A unified
module load cce/20.0.2 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=60 -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=60 -fopenmp-version=60 -O2 repro.c -o repro HSA_XNACK=1 OMP_TARGET_OFFLOAD=mandatory ./repro
Program (80 lines)
// # Draft finding: omp60-c-clause-defaultmap-private on cce20
//
// - compiler: cce20
// - toolchain: cce20; arch: gfx942; machine: tuolumne; memory mode: unified; opt: -O2
// - program: catalog omp60-c-clause-defaultmap-private (c, version 60, sha256 09a36af3c3918990)
// - expected: None (None; )
// - actual: rejected (diagnostic: <tmpdir>/omp60-c-clause-defaultmap-private.c:34:80: error: expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap')
// - 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 6.0, C11): the C port (offload.engine.f2c) of omp60-clause-defaultmap-private. Do not edit.
// id: omp60-c-clause-defaultmap-private
// primary atom: omp.clause.defaultmap.private
// legal: yes
// observable: yes
// ran_on_host: inside the target region of interest
// evidence: construct
// needs: none
// std: c11
// C: the check also reads w after the region: the defaultmap category keeps the region's writes to the aggregate w private; without the clause w is mapped tofrom and they come back (the Fortran program checks b only)
#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 c;
int w[n];
for (int i = 1; i <= n; i++) {
a[i - 1] = i;
b[i - 1] = 0;
ref[i - 1] = 2 * i + 1;
}
host_flag = 1;
c = 0;
for (int k1_ = 0; k1_ < n; k1_++)
w[k1_] = 0;
#pragma omp target map(to: a) map(from: b) map(tofrom: host_flag) defaultmap(private)
{
host_flag = (omp_is_initial_device() ? 1 : 0);
c = 1;
for (int i = 1; i <= n; i++) {
w[i - 1] = 2 * a[i - 1];
b[i - 1] = w[i - 1] + c;
}
}
bool wz = true;
for (int k_ = 0; k_ < n; k_++)
wz = wz && (w[k_] == 0);
bool t2_ = true;
for (int k3_ = 0; k3_ < n; k3_++)
t2_ = t2_ && (b[k3_] == ref[k3_]);
ok = t2_ && wz;
printf("RESULT ran_on_host=%d check=%s\n", host_flag, ok ? "pass" : "fail");
return 0;
}
// spec (requirement): In OpenMP 6.0 the defaultmap clause sets the implicit data-mapping or data-sharing attribute of
// the variables referenced in a target construct, and its implicit-behavior keywords include private.
// "The defaultmap clause controls the implicitly determined data-mapping attributes or implicitly determined
// data-sharing attributes of certain variables that are referenced in a target construct" (6.0 §7.9.9 p.292 l.2-4)
// spec (requirement): A defaultmap clause with an implicit behavior other than default determines the attribute of
// the variables of its category.
// "If a defaultmap clause is present for the category of the variable and specifies an implicit behavior other than
// default, the data-mapping attribute or data-sharing attribute is determined by that clause." (6.0 §7.9.3 p.276
// l.21-23)
// spec (this program): defaultmap(private) is a 6.0 form: the 6.0 argument table lists private among the
// implicit-behavior keywords (a table row without line numbers, so it is not quoted), but the 6.0 semantics
// paragraph of the clause does not describe it, so no sentence states what private does and the cited text is the
// closest normative rule. The compiler rejects the keyword with a diagnostic that lists alloc, from, to, tofrom,
// firstprivate, none, default and present.
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
cce20-c-defaultmap-private-storage-rejected.c:24:80: error: expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'
24 | #pragma omp target map(to: a) map(from: b) map(tofrom: host_flag) defaultmap(private)
| ^
1 error generated.What the specification requires
The specification text this program relies on and the compiler is expected to support.
OpenMP 6.0 §7.9.9 p.292 l.2-4
The defaultmap clause controls the implicitly determined data-mapping attributes or implicitly determined data-sharing attributes of certain variables that are referenced in a target construct
OpenMP 6.0 §7.9.3 p.276 l.21-23
If a defaultmap clause is present for the category of the variable and specifies an implicit behavior other than default, the data-mapping attribute or data-sharing attribute is determined by that clause.
Evidence
Same compiler, one change
- defaultmap(firstprivate) instead of defaultmap(private): 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 c; int w[n]; for (int i = 1; i <= n; i++) { a[i - 1] = i; b[i - 1] = 0; ref[i - 1] = 2 * i + 1; } host_flag = 1; c = 0; for (int k1_ = 0; k1_ < n; k1_++) w[k1_] = 0; #pragma omp target map(to: a) map(from: b) map(tofrom: host_flag) defaultmap(firstprivate) { host_flag = (omp_is_initial_device() ? 1 : 0); c = 1; for (int i = 1; i <= n; i++) { w[i - 1] = 2 * a[i - 1]; b[i - 1] = w[i - 1] + c; } } bool wz = true; for (int k_ = 0; k_ < n; k_++) wz = wz && (w[k_] == 0); bool t2_ = true; for (int k3_ = 0; k3_ < n; k3_++) t2_ = t2_ && (b[k3_] == ref[k3_]); ok = t2_ && wz; 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