[Bug target/96835] New: Constructor in offload template class

[Bug target/96835] New: Constructor in offload template class

[tobias.weinzierl at durham dot ac.uk]

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

            Bug ID: 96835
           Summary: Constructor in offload template class
           Product: gcc
           Version: 10.1.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: tobias.weinzierl at durham dot ac.uk
  Target Milestone: ---

I have a simple templated vector class with multiple constructors:

#ifdef GPUOffloading
#pragma omp declare target
#endif
template<int Size, typename Scalar>
struct tarch::la::Vector {
  private:
    Scalar _values[Size];

  public:
    Vector();

    Vector (const Scalar& initialValue);

[...]

This code translates fine and I can use the vector class as long as I do not
use one of the constructors (all the other non-overloaded routines work fine).
I either get

lto1: internal compiler error: Segmentation fault

or an undefined symbol error.

[Bug target/96835] Constructor in offload template class

[jakub at gcc dot gnu.org]

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

Jakub Jelinek <jakub at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |jakub at gcc dot gnu.org
   Last reconfirmed|                            |2020-08-28
             Status|UNCONFIRMED                 |WAITING
     Ever confirmed|0                           |1

--- Comment #1 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
Same problems as in PR96833.

[Bug target/96835] Constructor in offload template class

[tobias.weinzierl at durham dot ac.uk]

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

--- Comment #2 from Tobias Weinzierl <tobias.weinzierl at durham dot ac.uk> ---
#include <stdio.h>

#pragma omp declare target
template<int sz>
struct vector {
  int values_[sz];
  vector();
  vector(int const& init_val);
  int dot(vector o) {
    int res = 0;
    for (int i = 0; i < sz; ++ i)
      res += values_[i] * o.values_[i];
    return res;
  }
};

template<int sz>
vector<sz>::vector(int const& init_val) {
  for (int i = 0; i < sz; ++ i) values_[i] = init_val;
}
template<int sz>
vector<sz>::vector() : vector(0) {
}

#pragma omp end declare target

int main() {
  int res;
  #pragma omp target map(from:res)
  {
    vector<4> v1(1);
    vector<4> v2(2);
    res = v1.dot(v2);
  }
  printf("%d\n", res);
  return 0;
}

[Bug target/96835] Constructor in offload template class

[tobias.weinzierl at durham dot ac.uk]

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

--- Comment #3 from Tobias Weinzierl <tobias.weinzierl at durham dot ac.uk> ---
The full compilation error is 

+ g++-10 -fopenmp -foffload=nvptx-none bug.cpp -o bug
ptxas /tmp/cc1XobxJ.o, line 253; error   : Illegal operand type to instruction
'ld'
ptxas /tmp/cc1XobxJ.o, line 266; error   : Label expected for argument 0 of
instruction 'call'
ptxas /tmp/cc1XobxJ.o, line 266; error   : Function '_ZN6vectorILi4EEC1ERKi'
not declared in this scope
ptxas /tmp/cc1XobxJ.o, line 266; fatal   : Call target not recognized
ptxas fatal   : Ptx assembly aborted due to errors
nvptx-as: ptxas returned 255 exit status
mkoffload: fatal error: x86_64-linux-gnu-accel-nvptx-none-gcc-10 returned 1
exit status
compilation terminated.
lto-wrapper: fatal error:
/usr/lib/gcc/x86_64-linux-gnu/10//accel/nvptx-none/mkoffload returned 1 exit
status
compilation terminated.
/usr/bin/ld: error: lto-wrapper failed
collect2: error: ld returned 1 exit status

As long as the constructor is not used within the map region, the code
translates fine.

[Bug target/96835] Constructor in offload template class

[tobias.weinzierl at durham dot ac.uk]

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

--- Comment #4 from Tobias Weinzierl <tobias.weinzierl at durham dot ac.uk> ---
Created attachment 49339
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=49339&action=edit
Reproducer

[Bug target/96835] Constructor in offload template class

[redi at gcc dot gnu.org]

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

Jonathan Wakely <redi at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|WAITING                     |UNCONFIRMED
     Ever confirmed|1                           |0

[Bug target/96835] Constructor in offload template class

[burnus at gcc dot gnu.org]

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

--- Comment #5 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tobias Weinzierl from comment #4)
> Created attachment 49339 [details]
> Reproducer

Compiles here with mainline (11.0.0 20201104) and nvptx offloading (-O0).
I wonder whether that was fixed by:
  OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)
r11-3505-g2a10a2c0689db280ee3a94164504b7196b8370f4

Thus, I think the issue of comment 0 (the internal compiler error) is fixed and
the automatic 'declare target' works now on mainline.


However, the run-time JIT fails here (with and without 'omp declare target) –
as in comment 3:

libgomp: Link error log ptxas application ptx input, line 255; error   : Label
expected for argument 0 of instruction 'call'
ptxas application ptx input, line 255; error   : Function 
'_ZN6vectorILi4EEC1ERKi' not declared in this scope

 ^-- this is "vector<4>::vector(int const&)"
 ^-- and 'C1' means that it is the "complete object constructor"


My bet is that this is the nvptx alias issue PR 97106 (see also PR 97102 and,
regarding a newer PTX ISA PR96005).

Solution: Either resolving the alias in GCC (in mkoffload) or (conditionally)
bumping the minimal CUDA version to 10.0 such that .alias is available (PTX ISA
>= 6.3).

[Bug target/96835] Constructor in offload template class

[tobias.weinzierl at durham dot ac.uk]

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

--- Comment #6 from Tobias Weinzierl <tobias.weinzierl at durham dot ac.uk> ---
We've found some more stuff. This works:

#include <iostream>

#define mydt double

#pragma omp declare target
struct vector {  
  vector(mydt x, mydt y);
  mydt dot(vector o);
  mydt v[2];
}; 

vector::vector(mydt x, mydt y) {
  v[0]=x;
  v[1]=y;
}

mydt vector::dot(vector o) { return v[0] * o.v[0] + v[1]*o.v[1]; }  
#pragma omp end declare target

int main() {
  mydt res;
  vector v1(1,1);
  vector v2(1,3);
  #pragma omp target map(from:res)
  {
    res = v1.dot(v2); 
  }
  std::cout << res << "\n";
}

As soon as you instantiate the vector within the target region, it does not
work anymore:


int main() {
  mydt res;
  #pragma omp target map(from:res)
  {
    vector v1(1,1);
    vector v2(1,3);
    res = v1.dot(v2); 
  }
  std::cout << res << "\n";
}

[Bug target/96835] Constructor in offload template class

[tobias.weinzierl at durham dot ac.uk]

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

--- Comment #7 from Tobias Weinzierl <tobias.weinzierl at durham dot ac.uk> ---
Adding a default constructor to the vector class still does not allow us to
create the object on the target:

#include <iostream>

#define mydt double

#pragma omp declare target
struct vector {  
  vector() {};
  vector(mydt x, mydt y);
  mydt dot(vector o);
  mydt v[2];
}; 

vector::vector(mydt x, mydt y) {
  v[0]=x;
  v[1]=y;
}

mydt vector::dot(vector o) { return v[0] * o.v[0] + v[1]*o.v[1]; }  
#pragma omp end declare target

int main() {
  mydt res;
  vector v1(1,1);
  vector v2(1,3);
  #pragma omp target map(from:res)
  {
    vector v3;
    res = v1.dot(v2); 
  }
  std::cout << res << "\n";
}


However, if we replace the default constructor with a default clause, then we
ca finally create a vector on the GPU, too:

#pragma omp declare target
struct vector {  
  vector() = default;
  [...]