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}} + } +}; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
