[CUDA] Do not diagnose host/device variable access in dependent types.

[tra]

Summary

isCUDADeviceBuiltinSurfaceType()/isCUDADeviceBuiltinTextureType() do not
work on dependent types as they rely on specific type attributes and that
results in mis-diagnosing access when the check happens on class templates used as texture references.
https://godbolt.org/z/B7rtxR

[yaxunl]

LGTM. Can we have a test?

[hliao]

LGTM if there's a regression test available.

[tra]

Found another corner case (reference within a template with the surface/texture
attibute.) and figured out a better fix.

Added a test case.

[tra]

My first variant of the patch only helped with some cases when the surface/texture attribute type was used.
Trying to reduce real-world failure resulted in an example that I've added as the test case which was still failing with this patch applied.

Thanks to @rsmith, who suggested applying MeaningfulToClassTemplateDefinition, the attributes should work for templates now, too.

[yaxunl]

clang/test/SemaCUDA/device-use-host-var.cu

169 ↗

(On Diff #311065)

use __device__ will make it more readable. same as below.

[hliao]

LGTM if you revise the test based on Sam's suggestion on the test case.

[tra]

Use device in the test case.

clang/lib/AST/Type.cpp

if (!Pointer)
     return false;
 
   QualType Pointee = Pointer->getPointeeType();
   return Pointee->isVoidType() || Pointee->isRecordType();
 }
 
 /// Check if the specified type is the CUDA device builtin surface type.
 bool Type::isCUDADeviceBuiltinSurfaceType() const {
+  assert(!isDependentType());
   if (const auto *RT = getAs<RecordType>())
     return RT->getDecl()->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>();
   return false;
 }
 
 /// Check if the specified type is the CUDA device builtin texture type.
 bool Type::isCUDADeviceBuiltinTextureType() const {
+  assert(!isDependentType());
   if (const auto *RT = getAs<RecordType>())
     return RT->getDecl()->hasAttr<CUDADeviceBuiltinTextureTypeAttr>();
   return false;
 }
 
 bool Type::hasSizedVLAType() const {
   if (!isVariablyModifiedType()) return false;
 

clang/lib/Sema/SemaExpr.cpp

bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
   // CUDA/HIP: Diagnose invalid references of host global variables in device
   // functions. Reference of device global variables in host functions is
   // allowed through shadow variables therefore it is not diagnosed.
   if (LangOpts.CUDAIsDevice) {
     auto *FD = dyn_cast_or_null<FunctionDecl>(CurContext);
     auto Target = IdentifyCUDATarget(FD);
     if (FD && Target != CFT_Host) {
       const auto *VD = dyn_cast<VarDecl>(D);
-      if (VD && VD->hasGlobalStorage() && !VD->hasAttr<CUDADeviceAttr>() &&
-          !VD->hasAttr<CUDAConstantAttr>() && !VD->hasAttr<CUDASharedAttr>() &&
+      if (VD && VD->hasGlobalStorage() && !VD->getType()->isDependentType() &&
+          !VD->hasAttr<CUDADeviceAttr>() && !VD->hasAttr<CUDAConstantAttr>() &&
+          !VD->hasAttr<CUDASharedAttr>() &&
           !VD->getType()->isCUDADeviceBuiltinSurfaceType() &&
           !VD->getType()->isCUDADeviceBuiltinTextureType() &&
           !VD->isConstexpr() && !VD->getType().isConstQualified())
         targetDiag(*Locs.begin(), diag::err_ref_bad_target)
             << /*host*/ 2 << /*variable*/ 1 << VD << Target;
     }
   }