mirror of
https://github.com/llvm/llvm-project.git
synced 2025-04-16 12:26:44 +00:00
2b76e20ea7
Treat ctor/dtor in device var init as host device function so that they can be used to initialize file-scope device variables to match nvcc behavior. If they are non-trivial they will be diagnosed. We cannot add implicit host device attrs to non-trivial ctor/dtor since determining whether they are non-trivial needs to know whether they have a trivial body and all their member and base classes' ctor/dtor have trivial body, which is affected by where their bodies are defined or instantiated. Fixes: #72261 Fixes: SWDEV-432412
58 lines
1.2 KiB
Plaintext
58 lines
1.2 KiB
Plaintext
// RUN: %clang_cc1 -isystem %S/Inputs -fsyntax-only -verify %s
|
|
// RUN: %clang_cc1 -isystem %S/Inputs -fcuda-is-device -fsyntax-only -verify %s
|
|
|
|
#include <cuda.h>
|
|
|
|
// Check trivial ctor/dtor
|
|
struct A {
|
|
int x;
|
|
A() {}
|
|
~A() {}
|
|
};
|
|
|
|
__device__ A a;
|
|
|
|
// Check trivial ctor/dtor of template class
|
|
template<typename T>
|
|
struct TA {
|
|
T x;
|
|
TA() {}
|
|
~TA() {}
|
|
};
|
|
|
|
__device__ TA<int> ta;
|
|
|
|
// Check non-trivial ctor/dtor in parent template class
|
|
template<typename T>
|
|
struct TB {
|
|
T x;
|
|
TB() { static int nontrivial_ctor = 1; }
|
|
~TB() {}
|
|
};
|
|
|
|
template<typename T>
|
|
struct TC : TB<T> {
|
|
T x;
|
|
TC() {}
|
|
~TC() {}
|
|
};
|
|
|
|
template class TC<int>;
|
|
|
|
__device__ TC<int> tc; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
|
|
|
|
// Check trivial ctor specialization
|
|
template <typename T>
|
|
struct C {
|
|
explicit C() {};
|
|
};
|
|
|
|
template <> C<int>::C() {};
|
|
__device__ C<int> ci_d;
|
|
C<int> ci_h;
|
|
|
|
// Check non-trivial ctor specialization
|
|
template <> C<float>::C() { static int nontrivial_ctor = 1; }
|
|
__device__ C<float> cf_d; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
|
|
C<float> cf_h;
|