Author: Yaxun (Sam) Liu Date: 2021-02-04T18:04:54-05:00 New Revision: e355110040d17dcbf0fcfcba386c4d96fe6fe7ec
URL: https://github.com/llvm/llvm-project/commit/e355110040d17dcbf0fcfcba386c4d96fe6fe7ec DIFF: https://github.com/llvm/llvm-project/commit/e355110040d17dcbf0fcfcba386c4d96fe6fe7ec.diff LOG: [CUDA][HIP] Fix checking dependent initalizer Defer constant checking of dependent initializer to template instantiation since it cannot be done for dependent values. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D95840 Added: clang/test/SemaCUDA/dependent-device-var.cu Modified: clang/lib/Sema/SemaCUDA.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index ee91eb4c5deb..b907374b466f 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -530,9 +530,12 @@ void Sema::checkAllowedCUDAInitializer(VarDecl *VD) { if (!AllowedInit && (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>())) { auto *Init = VD->getInit(); + // isConstantInitializer cannot be called with dependent value, therefore + // we skip checking dependent value here. This is OK since + // checkAllowedCUDAInitializer is called again when the template is + // instantiated. AllowedInit = - ((VD->getType()->isDependentType() || Init->isValueDependent()) && - VD->isConstexpr()) || + VD->getType()->isDependentType() || Init->isValueDependent() || Init->isConstantInitializer(Context, VD->getType()->isReferenceType()); } diff --git a/clang/test/SemaCUDA/dependent-device-var.cu b/clang/test/SemaCUDA/dependent-device-var.cu new file mode 100644 index 000000000000..41139c8e64c1 --- /dev/null +++ b/clang/test/SemaCUDA/dependent-device-var.cu @@ -0,0 +1,18 @@ +// RUN: %clang_cc1 -fsyntax-only -verify=host,com -x hip %s +// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify=dev,com -x hip %s + +#include "Inputs/cuda.h" + +template<typename T> +__device__ int fun1(T x) { + // Check type-dependent constant is allowed in initializer. + static __device__ int a = sizeof(x); + static __device__ int b = x; + // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + return a + b; +} + +__device__ int fun1_caller() { + return fun1(1); + // com-note@-1 {{in instantiation of function template specialization 'fun1<int>' requested here}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
