diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -530,9 +530,12 @@ if (!AllowedInit && (VD->hasAttr() || VD->hasAttr())) { 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 --- /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 +__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' requested here}} +}