[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

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/include/clang/Basic/Attr.td

 }
 
 def CUDADeviceBuiltinSurfaceType : InheritableAttr {
   let Spellings = [GNU<"device_builtin_surface_type">,
                    Declspec<"__device_builtin_surface_type__">];
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[CXXRecord]>;
   let Documentation = [CUDADeviceBuiltinSurfaceTypeDocs];
+  let MeaningfulToClassTemplateDefinition = 1;
 }
 
 def CUDADeviceBuiltinTextureType : InheritableAttr {
   let Spellings = [GNU<"device_builtin_texture_type">,
                    Declspec<"__device_builtin_texture_type__">];
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[CXXRecord]>;
   let Documentation = [CUDADeviceBuiltinTextureTypeDocs];
+  let MeaningfulToClassTemplateDefinition = 1;
 }
 
 def CUDAGlobal : InheritableAttr {
   let Spellings = [GNU<"global">, Declspec<"__global__">];
   let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }

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

void dev_lambda_capture_by_copy(int *out) {
   *out = ref_dev_var;
   *out = ref_constant_var;
   *out = ref_shared_var;
   *out = ref_constexpr_var;
   *out = ref_const_var;
   });
 }
 
+// Texture references are special. As far as C++ is concerned they are host
+// variables that are referenced from device code. However, they are handled
+// very differently by the compiler under the hood and such references are
+// allowed. Compiler should produce no warning here, but it should diagnose the
+// same case without the device_builtin_texture_type attribute.
+template <class, int = 1, int = 1>
+struct __attribute__((device_builtin_texture_type)) texture {
+  static texture<int> ref;
+  __device__ int c() {
+    auto &x = ref;
+  }
+};
+
+template <class, int = 1, int = 1>
+struct  not_a_texture {
+  static not_a_texture<int> ref;
+  __device__ int c() {
+    auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}}
+  }
+};