This revision was automatically updated to reflect the committed changes. yaxunl marked 2 inline comments as done. Closed by commit rGe355110040d1: [CUDA][HIP] Fix checking dependent initalizer (authored by yaxunl). Herald added a project: clang.
Changed prior to commit: https://reviews.llvm.org/D95840?vs=320642&id=321577#toc Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D95840/new/ https://reviews.llvm.org/D95840 Files: clang/lib/Sema/SemaCUDA.cpp clang/test/SemaCUDA/dependent-device-var.cu Index: clang/test/SemaCUDA/dependent-device-var.cu =================================================================== --- /dev/null +++ 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}} +} Index: clang/lib/Sema/SemaCUDA.cpp =================================================================== --- clang/lib/Sema/SemaCUDA.cpp +++ clang/lib/Sema/SemaCUDA.cpp @@ -530,9 +530,12 @@ 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()); }
Index: clang/test/SemaCUDA/dependent-device-var.cu =================================================================== --- /dev/null +++ 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}} +} Index: clang/lib/Sema/SemaCUDA.cpp =================================================================== --- clang/lib/Sema/SemaCUDA.cpp +++ clang/lib/Sema/SemaCUDA.cpp @@ -530,9 +530,12 @@ 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()); }
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits