d37a39207b
Currently if CUDA/HIP users use template class with virtual dtor and std::string data member with C++20 and MSVC. When the template class is explicitly instantiated, there is error about host function called by host device function (used to be undefined symbols in linking stage before member destructors were checked by deferred diagnostics). It was caused by clang inferring host/device attributes for default dtors. Since all dtors of member and parent classes have implicit host device attrs, clang infers the virtual dtor have implicit host and device attrs. Since virtual dtor of explicitly instantiated template class must be emitted, this causes constexpr dtor of std::string emitted, which calls a host function which was not emitted on device side. This is a serious issue since it prevents users from using std::string with C++20 on Windows. When inferring host device attr of virtual dtor of explicit template class instantiation, clang should be conservative since it is sure to be emitted. Since an implicit host device function may call a host function, clang cannot assume it is always available on device. This guarantees dtors that may call host functions not to have implicit device attr, therefore will not be emitted on device side. Fixes: https://github.com/llvm/llvm-project/issues/108548 Fixes: SWDEV-517435
105 lines
2.2 KiB
Plaintext
105 lines
2.2 KiB
Plaintext
// RUN: %clang_cc1 %s -std=c++20 -fsyntax-only -verify=host
|
|
// RUN: %clang_cc1 %s -std=c++20 -fcuda-is-device -fsyntax-only -verify=dev
|
|
|
|
// host-no-diagnostics
|
|
|
|
#include "Inputs/cuda.h"
|
|
|
|
// Virtual dtor ~B() of explicit instantiation B<float> must
|
|
// be emitted, which causes host_fun() called.
|
|
namespace ExplicitInstantiationExplicitDevDtor {
|
|
void host_fun() // dev-note {{'host_fun' declared here}}
|
|
{}
|
|
|
|
template <unsigned>
|
|
constexpr void hd_fun() {
|
|
host_fun(); // dev-error {{reference to __host__ function 'host_fun' in __host__ __device__ function}}
|
|
}
|
|
|
|
struct A {
|
|
constexpr ~A() { // dev-note {{called by '~B'}}
|
|
hd_fun<8>(); // dev-note {{called by '~A'}}
|
|
}
|
|
};
|
|
|
|
template <typename T>
|
|
struct B {
|
|
public:
|
|
virtual __device__ ~B() = default;
|
|
A _a;
|
|
};
|
|
|
|
template class B<float>;
|
|
}
|
|
|
|
// The implicit host/device attrs of virtual dtor ~B() should be
|
|
// conservatively inferred, where constexpr member dtor's should
|
|
// not be considered device since they may call host functions.
|
|
// Therefore B<float>::~B() should not have implicit device attr.
|
|
// However C<float>::~C() should have implicit device attr since
|
|
// it is trivial.
|
|
namespace ExplicitInstantiationDtorNoAttr {
|
|
void host_fun()
|
|
{}
|
|
|
|
template <unsigned>
|
|
constexpr void hd_fun() {
|
|
host_fun();
|
|
}
|
|
|
|
struct A {
|
|
constexpr ~A() {
|
|
hd_fun<8>();
|
|
}
|
|
};
|
|
|
|
template <typename T>
|
|
struct B {
|
|
public:
|
|
virtual ~B() = default;
|
|
A _a;
|
|
};
|
|
|
|
template <typename T>
|
|
struct C {
|
|
public:
|
|
virtual ~C() = default;
|
|
};
|
|
|
|
template class B<float>;
|
|
template class C<float>;
|
|
__device__ void foo() {
|
|
C<float> x;
|
|
}
|
|
}
|
|
|
|
// Dtors of implicit template class instantiation are not
|
|
// conservatively inferred because the invalid usage can
|
|
// be diagnosed.
|
|
namespace ImplicitInstantiation {
|
|
void host_fun() // dev-note {{'host_fun' declared here}}
|
|
{}
|
|
|
|
template <unsigned>
|
|
constexpr void hd_fun() {
|
|
host_fun(); // dev-error {{reference to __host__ function 'host_fun' in __host__ __device__ function}}
|
|
}
|
|
|
|
struct A {
|
|
constexpr ~A() { // dev-note {{called by '~B'}}
|
|
hd_fun<8>(); // dev-note {{called by '~A'}}
|
|
}
|
|
};
|
|
|
|
template <typename T>
|
|
struct B {
|
|
public:
|
|
~B() = default; // dev-note {{called by 'foo'}}
|
|
A _a;
|
|
};
|
|
|
|
__device__ void foo() {
|
|
B<float> x;
|
|
}
|
|
}
|