[OpenMP] Static class members not implicitly captured from derived classes

[jhuber6]

When a derived class is present in a target region it will try to automatically declare members on the device. This is not correctly handled when done through a derived class that references the parent. See the example on godbolt https://godbolt.org/z/zvG5nr6KW and below.

class C {
public:
  C() { asm volatile(""); }
  virtual ~C() { asm volatile(""); }
  static int x;
};

int C::x = 0;
class D : public C {
public:
  D() { x = 1; }
};

void foo() {
#pragma omp target
  {
    D d;
  }
}

Here, the variable x is not captured on the device and thus shows up as an extern variable. This is then never resolved and fails at link time.

[jhuber6]

@alexey-bataev Any clue where this logic lives?

[alexey-bataev]

The variable C::x is just declared, not defined, see https://godbolt.org/z/Mo3aK578h for example. So, it is correct that it is just marked extern.

[jhuber6]

The variable C::x is just declared, not defined, see https://godbolt.org/z/Mo3aK578h for example. So, it is correct that it is just marked extern.

Whoops, forgot to copy that bit when I reduced this. I've updated the issue. If you explicitly declare it for the device it works, see https://godbolt.org/z/bPMx3h3o8. My guess is that it doesn't automatically add variables outside of the class definition.

[jhuber6]

Turns out you don't even need the inheritance https://godbolt.org/z/3ErrG1YPf.

[alexey-bataev]

You need explicitly mark definition as declare target

[jhuber6]

You need explicitly mark definition as declare target

Okay, so this is a wontfix? I don't know the expected OpenMP standards behavior for this.

[alexey-bataev]

Need to check the standard, previously it was a feature, not a bug :)

[jhuber6]

Need to check the standard, previously it was a feature, not a bug :)

Alright, this is related to a V&V failure I reduced https://github.com/OpenMP-Validation-and-Verification/OpenMP_VV/blob/master/tests/5.2/unified_shared_mem/test_target_VirDestr.cpp.

[llvmbot]

@llvm/issue-subscribers-openmp

Author: Joseph Huber (jhuber6)

When a derived class is present in a target region it will try to automatically declare members on the device. This is not correctly handled when done through a derived class that references the parent. See the example on godbolt https://godbolt.org/z/zvG5nr6KW and below. ```c++ class C { public: C() { asm volatile(""); } virtual ~C() { asm volatile(""); } static int x; };

int C::x = 0;
class D : public C {
public:
D() { x = 1; }
};

void foo() {
#pragma omp target
{
D d;
}
}

Here, the variable `x` is not captured on the device and thus shows up as an `extern` variable. This is then never resolved and fails at link time.
</details>

[alexey-bataev]

This is for unified shared memory, I suppose definition on the host must be visible from the device in this case.

[tschwinge]

Not sure to which extent it's related, but https://gcc.gnu.org/PR118484 "ICE during IPA pass: cp, segfault in determine_versionability ipa-cp.cc:467" also discusses this test case. In particular, with the known issues worked around, we get:

ptxas /tmp/ccFhejUv.o, line 94; error   : Unknown symbol '_ZTV4base'

..., which is "vtable for base".