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

[Artem Belevich via Phabricator via cfe-commits]

tra updated this revision to Diff 311065.
tra edited the summary of this revision.
tra added a comment.
Herald added a reviewer: aaron.ballman.

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

Added a test case.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D92893/new/

https://reviews.llvm.org/D92893

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


Index: clang/test/SemaCUDA/device-use-host-var.cu
===================================================================
--- clang/test/SemaCUDA/device-use-host-var.cu
+++ clang/test/SemaCUDA/device-use-host-var.cu
@@ -158,3 +158,23 @@
   });
 }
 
+// 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;
+  __attribute__((device)) int c() {
+    auto &x = ref;
+  }
+};
+
+template <class, int = 1, int = 1>
+struct  not_a_texture {
+  static not_a_texture<int> ref;
+  __attribute__((device)) int c() {
+    auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}}
+  }
+};
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ clang/include/clang/Basic/Attr.td
@@ -1079,6 +1079,7 @@
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[CXXRecord]>;
   let Documentation = [CUDADeviceBuiltinSurfaceTypeDocs];
+  let MeaningfulToClassTemplateDefinition = 1;
 }
 
 def CUDADeviceBuiltinTextureType : InheritableAttr {
@@ -1087,6 +1088,7 @@
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[CXXRecord]>;
   let Documentation = [CUDADeviceBuiltinTextureTypeDocs];
+  let MeaningfulToClassTemplateDefinition = 1;
 }
 
 def CUDAGlobal : InheritableAttr {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D92893.311065.patch
Type: text/x-patch
Size: 1738 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20201211/50b0c6b9/attachment-0001.bin>