tra updated this revision to Diff 311657.
tra added a comment.
Use __device__ in the 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;
+ __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}}
+ }
+};
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 {
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;
+ __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}}
+ }
+};
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 {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
