[Bug libgomp/121573] New: #pragma acc declare device_resident does not work in c/c++ with gcc 15.2

[schulz.benjamin at googlemail dot com via Gcc-bugs]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=121573

            Bug ID: 121573
           Summary: #pragma acc declare device_resident does not work in
                    c/c++ with gcc 15.2
           Product: gcc
           Version: 15.2.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: libgomp
          Assignee: unassigned at gcc dot gnu.org
          Reporter: schulz.benjamin at googlemail dot com
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

the following snipped works with nvc++ with options 

  -acc -std=c++20  -Msafeptr=all  -Minfo=all

and produces correct code.

/opt/nvidia/hpc_sdk/Linux_x86_64/25.5/compilers/bin/nvc++ 
-I/home/benni/projects/openmptestnew/openmpoffloatest/gpu_compiler_test
-I/home/benni/projects/openmptestnew/openmpoffloatest -acc -std=c++20 
-Msafeptr=all  -Minfo=all -std=c++23 -MD -MT
CMakeFiles/gpu_compiler_test.dir/main.cpp.o -MF
CMakeFiles/gpu_compiler_test.dir/main.cpp.o.d -o
CMakeFiles/gpu_compiler_test.dir/main.cpp.o -c
/home/benni/projects/openmptestnew/openmpoffloatest/main.cpp
in:
     22, Accelerator serial kernel generated
         CUDA shared memory used for _T4_1185
         Generating NVIDIA GPU code
     31, Generating copyout(out) [if not already present]
         Generating NVIDIA GPU code



It fails to compile with gcc 15.2, with options 
-fopenacc -foffload=nvptx-none  -fcf-protection=none -fno-stack-protector  
-no-pie  -U_FORTIFY_SOURCE


#include <openacc.h>

#include <stdio.h>


typedef struct
{
    int a;
    float b;
} mystruct;

#pragma acc declare device_resident (dev_var)
mystruct dev_var;


int main(void)
{

//serial kernel to initialize the device only variable

#pragma acc serial
   {
        dev_var = {42,3.14f};   // assign inside device code
    }

    // Try to read back into a host variable 

mystruct out;

#pragma acc parallel copyout(out)
    {
        if (acc_on_device(acc_device_nvidia)!=0)
            printf("on nvidia\n");
        else
            printf("not on nvidia\n");

      printf("on parallel region, should be 42; dev_var.a=%d \n", dev_var.a);
      printf("on parallel region, should be 3.14: dev_var.b=%f \n", dev_var.b);

      out = dev_var;
    }


    printf("Device-resident struct copied back, should be 42 and 3.14: a=%d
b=%f\n", out.a, out.b);

    return 0;
}






here is the error from gcc:

/usr/bin/gmake  -f CMakeFiles/Makefile2 CMakeFiles/gpu_compiler_test.dir/all
cc1plus: Warnung:
/home/benni/projects/openmptestnew/openmpoffloatest/gpu_compiler_test: not a
directory
/home/benni/projects/openmptestnew/openmpoffloatest/main.cpp:12:38: error:
»dev_var« was not declared
   12 | #pragma acc declare device_resident (dev_var)
      |                                      ^~~~~~~
/home/benni/projects/openmptestnew/openmpoffloatest/main.cpp:12:9: error: no
valid clause in #pragma acc declare« 
   12 | #pragma acc declare device_resident (dev_var)
      |         ^~~
gmake[3]: *** [CMakeFiles/gpu_compiler_test.dir/build.make:79:
CMakeFiles/gpu_compiler_test.dir/main.cpp.o] error 1
gmake[2]: *** [CMakeFiles/Makefile2:87: CMakeFiles/gpu_compiler_test.dir/all]
Fehler 2
gmake[1]: *** [CMakeFiles/Makefile2:94: CMakeFiles/gpu_compiler_test.dir/rule]
Fehler 2
gmake: *** [/home/benni/projects/openmptestnew/openmpoffloatest/Makefile:124:
gpu_compiler_test] error 2