Author: Yaxun (Sam) Liu Date: 2023-10-17T10:00:32-04:00 New Revision: fc53b1abf7d5e54012ea77a9bc8f6ccb7b487f13
URL: https://github.com/llvm/llvm-project/commit/fc53b1abf7d5e54012ea77a9bc8f6ccb7b487f13 DIFF: https://github.com/llvm/llvm-project/commit/fc53b1abf7d5e54012ea77a9bc8f6ccb7b487f13.diff LOG: [CUDA][HIP] Fix init var diag in temmplate (#69081) Currently clang diagnoses the following code: (https://godbolt.org/z/s8zK3E5P5) but nvcc does not. ` struct A { constexpr A(){} }; struct B { A a; int b; }; template<typename T> __global__ void kernel( ) { __shared__ B x; } ` Clang generates an implicit trivial ctor for struct B, which should be allowed for initializing a shared variable. However, the body of the ctor is defined only if the template kernel is instantiated. Clang checks the initialization of variable in non-instantiated templates, where it cannot find the body of the ctor, therefore diagnoses it. This patch skips the check for non-instantiated templates. Added: Modified: clang/lib/Sema/SemaCUDA.cpp clang/test/SemaCUDA/Inputs/cuda-initializers.h clang/test/SemaCUDA/device-var-init.cu Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 7c4083e4ec4d4bb..d993499cf4a6e6e 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -632,6 +632,13 @@ bool HasAllowedCUDADeviceStaticInitializer(Sema &S, VarDecl *VD, } // namespace void Sema::checkAllowedCUDAInitializer(VarDecl *VD) { + // Return early if VD is inside a non-instantiated template function since + // the implicit constructor is not defined yet. + if (const FunctionDecl *FD = + dyn_cast_or_null<FunctionDecl>(VD->getDeclContext())) + if (FD->isDependentContext()) + return; + // Do not check dependent variables since the ctor/dtor/initializer are not // determined. Do it after instantiation. if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage() || diff --git a/clang/test/SemaCUDA/Inputs/cuda-initializers.h b/clang/test/SemaCUDA/Inputs/cuda-initializers.h index 837b726a13e0f4b..b1e7a1bd48fb576 100644 --- a/clang/test/SemaCUDA/Inputs/cuda-initializers.h +++ b/clang/test/SemaCUDA/Inputs/cuda-initializers.h @@ -143,3 +143,14 @@ struct T_F_NED { struct T_FA_NED { NED ned[2]; }; + +// contexpr empty ctor -- allowed +struct CEEC { + constexpr CEEC() {} +}; + +// Compiler generated trivial ctor -- allowed +struct CGTC { + CEEC ceec; + int a; +}; diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu index 9d499bddbe1b31a..ee7a9e2276f2df0 100644 --- a/clang/test/SemaCUDA/device-var-init.cu +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -31,6 +31,14 @@ __device__ ECD d_ecd_i{}; __shared__ ECD s_ecd_i{}; __constant__ ECD c_ecd_i{}; +__device__ CEEC d_ceec; +__shared__ CEEC s_ceec; +__constant__ CEEC c_ceec; + +__device__ CGTC d_cgtc; +__shared__ CGTC s_cgtc; +__constant__ CGTC c_cgtc; + __device__ EC d_ec_i(3); // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}} __shared__ EC s_ec_i(3); @@ -213,6 +221,17 @@ __device__ void df_sema() { static const __device__ int cds = 1; static const __constant__ int cdc = 1; + for (int i = 0; i < 10; i++) { + static __device__ CEEC sd_ceec; + static __shared__ CEEC ss_ceec; + static __constant__ CEEC sc_ceec; + __shared__ CEEC s_ceec; + + static __device__ CGTC sd_cgtc; + static __shared__ CGTC ss_cgtc; + static __constant__ CGTC sc_cgtc; + __shared__ CGTC s_cgtc; + } // __shared__ does not need to be explicitly static. __shared__ int lsi; @@ -431,6 +450,35 @@ template <typename T> __global__ void bar() { __shared__ T bad; // expected-error@-1 {{initialization is not supported for __shared__ variables.}} + for (int i = 0; i < 10; i++) { + static __device__ CEEC sd_ceec; + static __shared__ CEEC ss_ceec; + static __constant__ CEEC sc_ceec; + __shared__ CEEC s_ceec; + + static __device__ CGTC sd_cgtc; + static __shared__ CGTC ss_cgtc; + static __constant__ CGTC sc_cgtc; + __shared__ CGTC s_cgtc; + } +} + +// Check specialization of template function. +template <> +__global__ void bar<int>() { + __shared__ NontrivialInitializer bad; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} + for (int i = 0; i < 10; i++) { + static __device__ CEEC sd_ceec; + static __shared__ CEEC ss_ceec; + static __constant__ CEEC sc_ceec; + __shared__ CEEC s_ceec; + + static __device__ CGTC sd_cgtc; + static __shared__ CGTC ss_cgtc; + static __constant__ CGTC sc_cgtc; + __shared__ CGTC s_cgtc; + } } void instantiate() { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
