Nvfortran OpenMP offload segfaults when updating derived type with allocatable member

Accelerated Computing HPC Compilers Legacy PGI Compilers
1

Hi,

We’ve been trying to figure out how to minimize data transfers of large derived types in our code. These derived types have a mix of scalar and allocatable array components. A pattern we’re trying to achieve:

  1. allocate member arrays and initialize with data on GPU
  2. update member scalars using user inputs
  3. update the scalars using !$omp target update

However, doing anything with the member arrays on the GPU afterward result in a cuda segfault. We have an MRE here: GitHub - JorgeG94/bug-copy-nvhpc-
which compiles bork and nobork binaries with make with nvfortran in PATH. The problem goes away with -gpu=mem:managed, but we’re using separate so we can manually manage memory more efficiently. We’ve tried compiler versions 24.9, 25.5, 25.7 which all give the below error:

Failing in Thread:1
Accelerator Fatal Error: call to cuStreamSynchronize returned error 700 (CUDA_ERROR_ILLEGAL_ADDRESS): Illegal address during kernel execution
 File: /home/edwardy/test2.f90
 Function: init:12
 Line: 31

I found a previous response on this forum which I think suggests that this problem is arising because of shallow copies of derived type members. So the host address of the allocatable members is being sent to the device, hence the segfault. I guess this would mean we would have to avoid updates of derived types with allocatable members altogether?

Is this understanding correct?

PS I gather from old PGI notes that shallow copies of updates is expected behaviour of OpenACC 2.6. But is that also true of OpenMP updates?

2

We’ve narrowed it down and made a reproducible of the real things in a C++ example. Had to do in C++ because the gnu compiler with GPU offloading that I have at hand is 14.2.0 and it errors out if you put a struct in the map() clause in Fortran. The output of the code below compiling with : nvc++ -mp=gpu bug.cpp is:

[jlv900@gadi-gpu-v100-0106 MOM6-examples]$ nvc++ -mp=gpu ed-2.cpp
[jlv900@gadi-gpu-v100-0106 MOM6-examples]$ ./a.out
Allocating c.a on host...
  Address of c.a on host: 0x20dd0d0
Mapping c and c.a to device...
  Address of c.a on device: 0x149c83afa200
Checking device address again, to confirm it doesn't change between target regions
  Address of c.a on device: 0x149c83afa200
!$omp target update to(c)...
  Address of c.a on device (after update): 0x20dd0d0

with g++ -fopenmp bug.cpp:

[jlv900@gadi-gpu-v100-0106 MOM6-examples]$ g++ -fopenmp ed-2.cpp
[jlv900@gadi-gpu-v100-0106 MOM6-examples]$ ./a.out
Allocating c.a on host...
  Address of c.a on host: 0xa6f7e0
Mapping c and c.a to device...
  Address of c.a on device: 0x14fa17400200
Checking device address again, to confirm it doesn't change between target regions
  Address of c.a on device: 0x14fa17400200
!$omp target update to(c)...
  Address of c.a on device (after update): 0x14fa17400200

So the device address with nvc++ is being updated to the host? while in G++ it is not.

The code is:

#include <iostream>
#include <cstdint>
#include <cstdlib>
#include <cstring>
#include <omp.h>

struct container {
    double* a;
};

int main() {
    container c;
    uintptr_t addr_a_host = 0;
    uintptr_t addr_a_dev = 0;

    std::cout << "Allocating c.a on host..." << std::endl;

    // Allocate 10x20 array
    c.a = static_cast<double*>(std::calloc(10 * 20, sizeof(double)));

    addr_a_host = reinterpret_cast<uintptr_t>(c.a);
    std::cout << "  Address of c.a on host: 0x" << std::hex << addr_a_host << std::dec << std::endl;

    std::cout << "Mapping c and c.a to device..." << std::endl;

    #pragma omp target enter data map(to: c, c.a[0:10*20])

    // Query address of c.a on device
    #pragma omp target map(from: addr_a_dev)
    {
        addr_a_dev = reinterpret_cast<uintptr_t>(c.a);
    }

    std::cout << "  Address of c.a on device: 0x" << std::hex << addr_a_dev << std::dec << std::endl;

    // Confirm address doesn't change across target regions
    #pragma omp target map(from: addr_a_dev)
    {
        addr_a_dev = reinterpret_cast<uintptr_t>(c.a);
    }

    std::cout << "Checking device address again, to confirm it doesn't change between target regions" << std::endl;
    std::cout << "  Address of c.a on device: 0x" << std::hex << addr_a_dev << std::dec << std::endl;

    std::cout << "!$omp target update to(c)..." << std::endl;

    #pragma omp target update to(c)

    #pragma omp target map(from: addr_a_dev)
    {
        addr_a_dev = reinterpret_cast<uintptr_t>(c.a);
    }

    std::cout << "  Address of c.a on device (after update): 0x" << std::hex << addr_a_dev << std::dec << std::endl;

    std::free(c.a);
    return 0;
}

which is a carbon copy of our Fortran example:

program test

    use iso_fortran_env, only: int64

    implicit none
    type container
        double precision, allocatable :: a(:, :)
    end type container
    type(container) :: c
    integer(int64) :: addr_a, addr_a_d

    print '(A)', "Allocating c%a on host..."

    allocate( c%a(10, 20), source=0.d0)

    addr_a = loc(c%a)

    print '(A, Z)', "  Address of c%a on host: ", addr_a
    print '(A)', "Mapping c, c%a to device..."
    !$omp target enter data map(to: c, c%a)

    !$omp target map(from: addr_a_d)
    addr_a_d = loc(c%a)
    !$omp end target

    print '(A, Z)', "  Address of c%a on device: ", addr_a_d

    !$omp target map(from: addr_a_d)
    addr_a_d = loc(c%a)
    !$omp end target

    print '(A)', "Checking device address again, to confirm it doesn't change between target regions"
    print '(A, Z)', "  Address of c%a on device: ", addr_a_d

    print '(A)', "!$omp target update to(c)..."
    !$omp target update to(c)

    !$omp target map(from: addr_a_d)
    addr_a_d = loc(c%a)
    !$omp end target

    print '(A, Z)', "  Address of c%a on device (after update):", addr_a_d

end program test
3

Yes, you identified the issue in that updating an aggregate type with dynamic data members will perform a shallow copy of the type. In this case, the host addresses of a dynamic member would get copied over if the aggregate type’s object gets updated to the device.

The OpenMP standard defines “target update” behavior as:

The target_update directive makes the corresponding list items in the device data environment consistent with their original list items, according to the specified data-motion clauses.

Our implementation conforms to this given a pointer member is part of the struct.

I wont go as far to say that the GNU implementation is non-conforming, there could be other parts of the standard that allow it or give them the freedom to implement how they choose, but I do see it as problematic. Besides the issue of not allowing strucs in map clause, another example would be pointer swapping as the following code illustrates.

#include <iostream>
#include <cstdio>
#include <cstdint>
#include <cstdlib>
#include <cstring>
#include <omp.h>

struct container {
    double* a;
    double * b;
};

int main() {
    container c;
    double * tmp;
    uintptr_t addr_a_dev = 0;
    uintptr_t addr_b_dev = 0;

    // Allocate 10x20 array
    c.a = static_cast<double*>(std::calloc(10 * 20, sizeof(double)));
    c.b = static_cast<double*>(std::calloc(10 * 20, sizeof(double)));

    std::cout << "Mapping c and c.a to device..." << std::endl;
    #pragma omp target enter data map(to: c, c.a[0:10*20],c.b[0:10*20])
    #pragma omp target map(from: addr_a_dev,addr_b_dev)
    {
        printf("Device pointers are %p %p\n",c.a,c.b);
        addr_a_dev = reinterpret_cast<uintptr_t>(c.a);
        addr_b_dev = reinterpret_cast<uintptr_t>(c.b);
    }
    std::cout << "  Address of c.a on device: " << std::hex << addr_a_dev << std::dec << std::endl;
    std::cout << "  Address of c.b on device: " << std::hex << addr_b_dev << std::dec << std::endl;
    std::cout << "Swap pointers" << std::endl;
    tmp=c.a;
    c.a=c.b;
    c.b=tmp;

    //#pragma omp target update to(c)

    #pragma omp target map(from: addr_a_dev,addr_b_dev) map(to:c.a,c.b)
    {
        printf("Pointers should be swapped %p %p\n",c.a,c.b);
        addr_a_dev = reinterpret_cast<uintptr_t>(c.a);
        addr_b_dev = reinterpret_cast<uintptr_t>(c.b);
    }
    std::cout << "  Address of c.a on device: " << std::hex << addr_a_dev << std::dec << std::endl;
    std::cout << "  Address of c.b on device: " << std::hex << addr_b_dev << std::dec << std::endl;


    std::free(c.a);
    return 0;
}

nvc++ can correctly map the swapped pointers to the device side addresses, which GNU does not.

% nvc++ -mp=gpu ptrswp.cpp -w
% a.out
Mapping c and c.a to device...
  Address of c.a on device: 14e06a6fa200
  Address of c.b on device: 14e06a6faa00
Swap pointers
  Address of c.a on device: 14e06a6faa00
  Address of c.b on device: 14e06a6fa200

% g++ -fopenmp ptrswp.cpp
% a.out
Mapping c and c.a to device...
  Address of c.a on device: 146f2be00a00
  Address of c.b on device: 146f2be00200
Swap pointers
  Address of c.a on device: 146f2be00a00
  Address of c.b on device: 146f2be00200
4

Thanks for confirming Mat!
Do you have a suggestion for an alternative to !$omp target update to(c) for nvfortran? In C/C++, I guess we could write a custom update function that manages the pointers, but how could this be done in Fortran? Would it be better to map(delete: c, c%...) then map(to: c, c%...) (c being the aggregate type with dynamic members)?

5

You’d just update “c%a”. Though to be more correct, you’d use triplet notation, “c%a(1:10,1:20)”. It’s Fortran so the shape and size of the array is known, so it’s not necessary,

In your original question you were talking about the scalar members, in which case, you’d do the same, “c%scalar”. You just don’t want to update the parent object itself since it will overwrite the dynamic data members.

I’m not sure what to do about gfortran since they don’t support using dynamic members.

% gfortran -w test.F90 -fopenmp
test.F90:20:36:

   20 |     !$omp target enter data map(to: c, c%a)
      |                                    1
Error: List item ‘c’ with allocatable components is not permitted in map clause at (1)

In C/C++, I guess we could write a custom update function that manages the pointers, but how could this be done in Fortran? Would it be better to map(delete: c, c%...) then map(to: c, c%...) (c being the aggregate type with dynamic members)?

Is “a” changing during the run? i.e. allocating/deallocate/allocating?

In that case, you want “map(delete: c%a)” before the deallocate, then “map(to: c%a(1:size))” after the next allocation.

Or is “a” a pointer that dynamically changes what device memory it points at during the run?

In OpenACC, you’d do an “attach” in this case, but I’m not sure of the equivalent in OpenMP. Though I can do some research if this is what you’re doing.

Another possibility is to mix in CUDA Fortran device array. Of course it will make your code non-portable to other compilers, but is could work.

6

In our case, c%a doesn’t change during the run, but the scalar members of c might. Individually updating scalar members of c is an option, but there are many of them (20 or so), so is a bit tedious. We were hopingupdate to(c) could be a concise way of updating the c member scalars on the GPU without affecting c%a. I guess it would be hard to make it neat if we also had to “reattach” member allocatables after every update…

Anyway, I guess it’s enough to know what update is doing in NVHPC so we can work around it.