aboutsummaryrefslogtreecommitdiff
path: root/clang/lib/Sema/SemaCUDA.cpp
diff options
context:
space:
mode:
authorYaxun (Sam) Liu <yaxun.liu@amd.com>2023-10-17 10:00:32 -0400
committerGitHub <noreply@github.com>2023-10-17 10:00:32 -0400
commitfc53b1abf7d5e54012ea77a9bc8f6ccb7b487f13 (patch)
tree4d337f3d8a5d137fc11b9d20c6fc17fa38dfddfd /clang/lib/Sema/SemaCUDA.cpp
parent096eba148df7dcddf9872544fbf510a2c1a9785c (diff)
downloadllvm-fc53b1abf7d5e54012ea77a9bc8f6ccb7b487f13.zip
llvm-fc53b1abf7d5e54012ea77a9bc8f6ccb7b487f13.tar.gz
llvm-fc53b1abf7d5e54012ea77a9bc8f6ccb7b487f13.tar.bz2
[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.
Diffstat (limited to 'clang/lib/Sema/SemaCUDA.cpp')
-rw-r--r--clang/lib/Sema/SemaCUDA.cpp7
1 files changed, 7 insertions, 0 deletions
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 7c4083e..d993499 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() ||