diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h --- a/clang/include/clang/AST/Type.h +++ b/clang/include/clang/AST/Type.h @@ -488,7 +488,12 @@ /// Returns true if the address space in these qualifiers is equal to or /// a superset of the address space in the argument qualifiers. bool isAddressSpaceSupersetOf(Qualifiers other) const { - return isAddressSpaceSupersetOf(getAddressSpace(), other.getAddressSpace()); + return isAddressSpaceSupersetOf(getAddressSpace(), + other.getAddressSpace()) || + (!hasAddressSpace() && + (other.getAddressSpace() == LangAS::opencl_private || + other.getAddressSpace() == LangAS::opencl_local || + other.getAddressSpace() == LangAS::opencl_global)); } /// Determines if these qualifiers compatibly include another set. diff --git a/clang/test/SemaOpenCLCXX/address-space-lambda.cl b/clang/test/SemaOpenCLCXX/address-space-lambda.cl --- a/clang/test/SemaOpenCLCXX/address-space-lambda.cl +++ b/clang/test/SemaOpenCLCXX/address-space-lambda.cl @@ -31,8 +31,6 @@ //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () {{.*}}const __generic' auto priv2 = []() __generic {}; priv2(); - auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}} - priv3(); //expected-error{{no matching function for call to object of type}} __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}} //expected-note{{conversion candidate of type 'void (*)()'}} const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}} diff --git a/clang/test/SemaOpenCLCXX/pr-45472.cpp b/clang/test/SemaOpenCLCXX/pr-45472.cpp new file mode 100644 --- /dev/null +++ b/clang/test/SemaOpenCLCXX/pr-45472.cpp @@ -0,0 +1,9 @@ +// RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s +// FIXME: expected error is not emitted due to wrong address space inference +// see https://bugs.llvm.org/show_bug.cgi?id=45472 for more details. +// XFAIL: * + +__kernel void test_qual() { + auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}} + priv3(); //expected-error{{no matching function for call to object of type}} +} \ No newline at end of file diff --git a/clang/test/SemaSYCL/address-space-parameter-conversions.cpp b/clang/test/SemaSYCL/address-space-parameter-conversions.cpp new file mode 100644 --- /dev/null +++ b/clang/test/SemaSYCL/address-space-parameter-conversions.cpp @@ -0,0 +1,57 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s + +void bar(int &Data) {} +void bar2(int &Data) {} +void bar(__attribute__((opencl_private)) int &Data) {} +void foo(int *Data) {} +void foo2(int *Data) {} +void foo(__attribute__((opencl_private)) int *Data) {} + +template +void tmpl(T *t) {} + +void usages() { + __attribute__((opencl_global)) int *GLOB; + __attribute__((opencl_private)) int *PRIV; + __attribute__((opencl_local)) int *LOC; + int *NoAS; + + bar(*GLOB); + bar2(*GLOB); + + bar(*PRIV); + bar2(*PRIV); + + bar(*NoAS); + bar2(*NoAS); + + bar(*LOC); + bar2(*LOC); + + foo(GLOB); + foo2(GLOB); + foo(PRIV); + foo2(PRIV); + foo(NoAS); + foo2(NoAS); + foo(LOC); + foo2(LOC); + + tmpl(GLOB); + tmpl(PRIV); + tmpl(NoAS); + tmpl(LOC); + + (void)static_cast(GLOB); + (void)static_cast(GLOB); + // FIXME: determine if we can warn on the below conversions. + int *i = GLOB; + void *v = GLOB; + (void)i; + (void)v; + + // expected-error@+1{{address space is negative}} + __attribute__((address_space(-1))) int *TooLow; + // expected-error@+1{{unknown type name '__generic'}} + __generic int *IsGeneric; +}