yaxunl created this revision. yaxunl added a reviewer: tra. https://reviews.llvm.org/D68753
Files: test/SemaCUDA/constexpr-ctor.cu Index: test/SemaCUDA/constexpr-ctor.cu =================================================================== --- /dev/null +++ test/SemaCUDA/constexpr-ctor.cu @@ -0,0 +1,43 @@ +// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only \ +// RUN: -fcuda-is-device -verify=dev -verify-ignore-unexpected=note %s +// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only \ +// RUN: -verify=host -verify-ignore-unexpected=note %s + +// host-no-diagnostics + +#include "Inputs/cuda.h" + +struct A { + A(); +}; + +template <class T> struct B { + T a; + constexpr B() = default; +}; + +template <class T> struct C { + T a; + constexpr C() = default; // dev-error{{reference to __host__ function 'A' in __host__ __device__ function}} +}; + +template <class T> struct D { + T a; + constexpr D() = default; // dev-error{{reference to __host__ function 'A' in __host__ __device__ function}} +}; + +__host__ void f() { B<A> x; } +__device__ void f() { C<A> x; } // causes C<A>::C emitted on device + +struct foo { + __host__ foo() { B<A> x; } + __device__ foo() { D<A> x; } +}; + +__host__ void g() { foo x; } +__device__ void g() { foo x; } // causes D<A>::D emitted on device + +struct bar { + __host__ bar() { B<A> x; } + __device__ bar() { B<A> x; } // no error since no instantiation of bar +};
Index: test/SemaCUDA/constexpr-ctor.cu =================================================================== --- /dev/null +++ test/SemaCUDA/constexpr-ctor.cu @@ -0,0 +1,43 @@ +// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only \ +// RUN: -fcuda-is-device -verify=dev -verify-ignore-unexpected=note %s +// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only \ +// RUN: -verify=host -verify-ignore-unexpected=note %s + +// host-no-diagnostics + +#include "Inputs/cuda.h" + +struct A { + A(); +}; + +template <class T> struct B { + T a; + constexpr B() = default; +}; + +template <class T> struct C { + T a; + constexpr C() = default; // dev-error{{reference to __host__ function 'A' in __host__ __device__ function}} +}; + +template <class T> struct D { + T a; + constexpr D() = default; // dev-error{{reference to __host__ function 'A' in __host__ __device__ function}} +}; + +__host__ void f() { B<A> x; } +__device__ void f() { C<A> x; } // causes C<A>::C emitted on device + +struct foo { + __host__ foo() { B<A> x; } + __device__ foo() { D<A> x; } +}; + +__host__ void g() { foo x; } +__device__ void g() { foo x; } // causes D<A>::D emitted on device + +struct bar { + __host__ bar() { B<A> x; } + __device__ bar() { B<A> x; } // no error since no instantiation of bar +};
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits