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

[Artem Belevich via cfe-commits]

Author: Artem Belevich
Date: 2020-12-14T11:53:18-08:00
New Revision: 0936655bac78f6e9cb84dc3feb30c32012100839

URL: https://github.com/llvm/llvm-project/commit/0936655bac78f6e9cb84dc3feb30c32012100839
DIFF: https://github.com/llvm/llvm-project/commit/0936655bac78f6e9cb84dc3feb30c32012100839.diff

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

`isCUDADeviceBuiltinSurfaceType()`/`isCUDADeviceBuiltinTextureType()` do not
work on dependent types as they rely on specific type attributes.

Differential Revision: https://reviews.llvm.org/D92893

Added: 
    

Modified: 
    clang/include/clang/Basic/Attr.td
    clang/test/SemaCUDA/device-use-host-var.cu

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 51f654fc7613..79902c8f5b89 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1079,6 +1079,7 @@ def CUDADeviceBuiltinSurfaceType : InheritableAttr {
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[CXXRecord]>;
   let Documentation = [CUDADeviceBuiltinSurfaceTypeDocs];
+  let MeaningfulToClassTemplateDefinition = 1;
 }
 
 def CUDADeviceBuiltinTextureType : InheritableAttr {
@@ -1087,6 +1088,7 @@ def CUDADeviceBuiltinTextureType : InheritableAttr {
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[CXXRecord]>;
   let Documentation = [CUDADeviceBuiltinTextureTypeDocs];
+  let MeaningfulToClassTemplateDefinition = 1;
 }
 
 def CUDAGlobal : InheritableAttr {

diff  --git a/clang/test/SemaCUDA/device-use-host-var.cu b/clang/test/SemaCUDA/device-use-host-var.cu
index cf5514610a42..c8ef7dbbb18d 100644
--- a/clang/test/SemaCUDA/device-use-host-var.cu
+++ b/clang/test/SemaCUDA/device-use-host-var.cu
@@ -158,3 +158,23 @@ void dev_lambda_capture_by_copy(int *out) {
   });
 }
 
+// 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 
diff erently 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}}
+  }
+};