[SYCL] Re-use OpenCL address space attributes for SYCL

Today we re-use OpenCL parsed attributes, but have separate SYCL address
space semantic attributes as current implementation of OpenCL semantics
breaks valid C++. This patch enables re-use of OpenCL semantic
attributes by allowing conversions between types qualified with OpenCL
address spaces and type w/o address space qualifiers. Clang compiler
(almost) always adds address space qualifiers in OpenCL mode, so it
should not affect OpenCL mode.

NOTE: this change also disables implicit conversion between the
unqualified types and types qualified with
`__attribute__((address_space(N)))`, enabled by one of the previous SYCL
patches.

Signed-off-by: Alexey Bader <alexey.bader@intel.com>

clang/include/clang/AST/Type.h

@@ -486,15 +486,12 @@ class Qualifiers {
   /// 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()) ||
-        (!hasAddressSpace() &&
-         (other.getAddressSpace() == LangAS::sycl_private ||
-          other.getAddressSpace() == LangAS::sycl_local ||
-          other.getAddressSpace() == LangAS::sycl_global ||
-          other.getAddressSpace() == LangAS::sycl_constant ||
-          other.getAddressSpace() == LangAS::sycl_generic));
+    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.

clang/include/clang/Basic/AddressSpaces.h

@@ -42,14 +42,6 @@ enum class LangAS : unsigned {
   cuda_constant,
   cuda_shared,
 
-  sycl_global,
-  sycl_local,
-  sycl_constant,
-  sycl_private,
-  // Likely never used, but useful in the future to reserve the spot in the
-  // enum.
-  sycl_generic,
-
   // Pointer size and extension address spaces.
   ptr32_sptr,
   ptr32_uptr,

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10292,8 +10292,6 @@ def err_builtin_launder_invalid_arg : Error<
   "'__builtin_launder' is not allowed">;
 
 // SYCL-specific diagnostics
-def err_sycl_attribute_address_space_invalid : Error<
-  "address space is outside the valid range of values">;
 def err_sycl_kernel_name_class_not_top_level : Error<
  "kernel name class and its template argument classes' declarations can only "
  "nest in a namespace: %0">;

clang/include/clang/Sema/ParsedAttr.h

@@ -534,24 +534,6 @@ class ParsedAttr final
     }
   }
 
-  /// If this is an OpenCL addr space attribute returns its SYCL representation
-  /// in LangAS, otherwise returns default addr space.
-  LangAS asSYCLLangAS() const {
-    switch (getKind()) {
-    case ParsedAttr::AT_OpenCLConstantAddressSpace:
-      return LangAS::sycl_constant;
-    case ParsedAttr::AT_OpenCLGlobalAddressSpace:
-      return LangAS::sycl_global;
-    case ParsedAttr::AT_OpenCLLocalAddressSpace:
-      return LangAS::sycl_local;
-    case ParsedAttr::AT_OpenCLPrivateAddressSpace:
-      return LangAS::sycl_private;
-    case ParsedAttr::AT_OpenCLGenericAddressSpace:
-    default:
-      return LangAS::Default;
-    }
-  }
-
   AttributeCommonInfo::Kind getKind() const { return getParsedKind(); }
 };
 

clang/lib/AST/ASTContext.cpp

@@ -836,11 +836,6 @@ static const LangASMap *getAddressSpaceMap(const TargetInfo &T,
         5, // cuda_device
         6, // cuda_constant
         7, // cuda_shared
-        1, // sycl_global
-        3, // sycl_local
-        2, // sycl_constant
-        0, // sycl_private
-        4, // sycl_generic
         8, // ptr32_sptr
         9, // ptr32_uptr
         10 // ptr64

clang/lib/AST/TypePrinter.cpp

@@ -1792,16 +1792,12 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
   case LangAS::Default:
     return "";
   case LangAS::opencl_global:
-  case LangAS::sycl_global:
     return "__global";
   case LangAS::opencl_local:
-  case LangAS::sycl_local:
     return "__local";
   case LangAS::opencl_private:
-  case LangAS::sycl_private:
     return "__private";
   case LangAS::opencl_constant:
-  case LangAS::sycl_constant:
     return "__constant";
   case LangAS::opencl_generic:
     return "__generic";

clang/lib/Basic/Targets/AMDGPU.cpp

@@ -48,11 +48,6 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsGenMap = {
     Global,   // cuda_device
     Constant, // cuda_constant
     Local,    // cuda_shared
-    Global,   // sycl_global
-    Local,    // sycl_local
-    Constant, // sycl_constant
-    Private,  // sycl_private
-    Generic,  // sycl_generic
     Generic,  // ptr32_sptr
     Generic,  // ptr32_uptr
     Generic   // ptr64
@@ -68,11 +63,6 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsPrivMap = {
     Global,   // cuda_device
     Constant, // cuda_constant
     Local,    // cuda_shared
-    Global,   // sycl_global
-    Local,    // sycl_local
-    Constant, // sycl_constant
-    Private,  // sycl_private
-    Generic,  // sycl_generic
     Generic,  // ptr32_sptr
     Generic,  // ptr32_uptr
     Generic   // ptr64

clang/lib/Basic/Targets/NVPTX.h

@@ -33,12 +33,6 @@ static const unsigned NVPTXAddrSpaceMap[] = {
     1, // cuda_device
     4, // cuda_constant
     3, // cuda_shared
-    1, // sycl_global
-    3, // sycl_local
-    4, // sycl_constant
-    0, // sycl_private
-    // FIXME: generic has to be added to the target
-    0, // sycl_generic
     0, // ptr32_sptr
     0, // ptr32_uptr
     0  // ptr64

clang/lib/Basic/Targets/SPIR.h

@@ -33,11 +33,6 @@ static const unsigned SPIRAddrSpaceMap[] = {
     0, // cuda_device
     0, // cuda_constant
     0, // cuda_shared
-    1, // sycl_global
-    3, // sycl_local
-    2, // sycl_constant
-    0, // sycl_private
-    4, // sycl_generic
     0, // ptr32_sptr
     0, // ptr32_uptr
     0  // ptr64
@@ -53,11 +48,6 @@ static const unsigned SYCLAddrSpaceMap[] = {
     0, // cuda_device
     0, // cuda_constant
     0, // cuda_shared
-    1, // sycl_global
-    3, // sycl_local
-    2, // sycl_constant
-    0, // sycl_private
-    4, // sycl_generic
     0, // ptr32_sptr
     0, // ptr32_uptr
     0  // ptr64
@@ -70,11 +60,9 @@ class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo {
     TLSSupported = false;
     VLASupported = false;
     LongWidth = LongAlign = 64;
-    if (Triple.getEnvironment() == llvm::Triple::SYCLDevice) {
-      AddrSpaceMap = &SYCLAddrSpaceMap;
-    } else {
-      AddrSpaceMap = &SPIRAddrSpaceMap;
-    }
+    AddrSpaceMap = (Triple.getEnvironment() == llvm::Triple::SYCLDevice)
+                       ? &SYCLAddrSpaceMap
+                       : &SPIRAddrSpaceMap;
     UseAddrSpaceMapMangling = true;
     HasLegalHalfType = true;
     HasFloat16 = true;

clang/lib/Basic/Targets/TCE.h

@@ -40,12 +40,6 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = {
     0, // cuda_device
     0, // cuda_constant
     0, // cuda_shared
-    3, // sycl_global
-    4, // sycl_local
-    5, // sycl_constant
-    0, // sycl_private
-    // FIXME: generic has to be added to the target
-    0, // sycl_generic
     0, // ptr32_sptr
     0, // ptr32_uptr
     0, // ptr64

clang/lib/Basic/Targets/X86.h

@@ -32,11 +32,6 @@ static const unsigned X86AddrSpaceMap[] = {
     0,   // cuda_device
     0,   // cuda_constant
     0,   // cuda_shared
-    0,   // sycl_global
-    0,   // sycl_local
-    0,   // sycl_constant
-    0,   // sycl_private
-    0,   // sycl_generic
     270, // ptr32_sptr
     271, // ptr32_uptr
     272  // ptr64

clang/lib/Sema/SemaType.cpp

@@ -5963,35 +5963,14 @@ static bool BuildAddressSpaceIndex(Sema &S, LangAS &ASIdx,
     llvm::APSInt max(addrSpace.getBitWidth());
     max =
         Qualifiers::MaxAddressSpace - (unsigned)LangAS::FirstTargetAddressSpace;
-
     if (addrSpace > max) {
       S.Diag(AttrLoc, diag::err_attribute_address_space_too_high)
           << (unsigned)max.getZExtValue() << AddrSpace->getSourceRange();
       return false;
     }
 
-    if (S.LangOpts.SYCLIsDevice && (addrSpace >= 4)) {
-      S.Diag(AttrLoc, diag::err_sycl_attribute_address_space_invalid)
-          << AddrSpace->getSourceRange();
-      return false;
-    }
-
-    ASIdx = getLangASFromTargetAS(
-                             static_cast<unsigned>(addrSpace.getZExtValue()));
-
-    if (S.LangOpts.SYCLIsDevice) {
-      ASIdx =
-          [](unsigned AS) {
-            switch (AS) {
-            case 0: return LangAS::sycl_private;
-            case 1: return LangAS::sycl_global;
-            case 2: return LangAS::sycl_constant;
-            case 3: return LangAS::sycl_local;
-            case 4: default: llvm_unreachable("Invalid SYCL AS");
-            }
-          }(static_cast<unsigned>(ASIdx) -
-            static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
-    }
+    ASIdx =
+        getLangASFromTargetAS(static_cast<unsigned>(addrSpace.getZExtValue()));
     return true;
   }
 
@@ -6117,8 +6096,7 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
       Attr.setInvalid();
   } else {
     // The keyword-based type attributes imply which address space to use.
-    ASIdx = S.getLangOpts().SYCLIsDevice ? 
-                Attr.asSYCLLangAS() : Attr.asOpenCLLangAS();
+    ASIdx = Attr.asOpenCLLangAS();
     if (ASIdx == LangAS::Default)
       llvm_unreachable("Invalid address space");
 

clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp

@@ -9,7 +9,7 @@ void foo(int * Data) {}
 // CHECK-DAG: define spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
 void foo2(int * Data) {}
 // CHECK-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
-void foo(__attribute__((address_space(3))) int * Data) {}
+void foo(__attribute__((opencl_local)) int * Data) {}
 // CHECK-DAG: define spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](i32 addrspace(3)* %
 
 template<typename T>
@@ -18,12 +18,11 @@ void tmpl(T t){}
 
 void usages() {
   // CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca i32 addrspace(1)*
-  __attribute__((address_space(1))) int *GLOB;
+  __attribute__((opencl_global)) int *GLOB;
   // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)*
   __attribute__((opencl_local)) int *LOC;
   // CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32 addrspace(4)*
   int *NoAS;
-
   // CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca i32*
   __attribute__((opencl_private)) int *PRIV;
 
@@ -94,57 +93,23 @@ void usages() {
 // CHECK-DAG: define linkonce_odr spir_func void [[GEN_TMPL]](i32 addrspace(4)* %
 
 void usages2() {
-  __attribute__((address_space(0))) int *PRIV_NUM;
-  // CHECK-DAG: [[PRIV_NUM:%[a-zA-Z0-9_]+]] = alloca i32*
-  __attribute__((address_space(0))) int *PRIV_NUM2;
-  // CHECK-DAG: [[PRIV_NUM2:%[a-zA-Z0-9_]+]] = alloca i32*
   __attribute__((opencl_private)) int *PRIV;
   // CHECK-DAG: [[PRIV:%[a-zA-Z0-9_]+]] = alloca i32*
-  __attribute__((address_space(1))) int *GLOB_NUM;
-  // CHECK-DAG: [[GLOB_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)*
   __attribute__((opencl_global)) int *GLOB;
   // CHECK-DAG: [[GLOB:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)*
-  __attribute__((address_space(2))) int *CONST_NUM;
-  // CHECK-DAG: [[CONST_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)*
   __attribute__((opencl_constant)) int *CONST;
   // CHECK-DAG: [[CONST:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)*
-  __attribute__((address_space(3))) int *LOCAL_NUM;
-  // CHECK-DAG: [[LOCAL_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)*
   __attribute__((opencl_local)) int *LOCAL;
   // CHECK-DAG: [[LOCAL:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)*
 
-  bar(*PRIV_NUM);
-  // CHECK-DAG: [[PRIV_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM]]
-  // CHECK-DAG: [[PRIV_NUM_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM_LOAD]] to i32 addrspace(4)*
-  // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM_ASCAST]])
-  bar(*PRIV_NUM2);
-  // CHECK-DAG: [[PRIV_NUM2_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM2]]
-  // CHECK-DAG: [[PRIV_NUM2_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM2_LOAD]] to i32 addrspace(4)*
-  // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM2_ASCAST]])
   bar(*PRIV);
   // CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]]
   // CHECK-DAG: [[PRIV_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_LOAD]] to i32 addrspace(4)*
   // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_ASCAST]])
-  bar(*GLOB_NUM);
-  // CHECK-DAG: [[GLOB_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB_NUM]]
-  // CHECK-DAG: [[GLOB_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_NUM_LOAD]] to i32 addrspace(4)*
-  // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_NUM_CAST]])
   bar(*GLOB);
   // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]]
   // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)*
   // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST]])
-  bar(*CONST_NUM);
-  // CHECK-DAG: [[CONST_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST_NUM]]
-  // CHECK-DAG: [[CONST_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_NUM_LOAD]] to i32 addrspace(4)*
-  // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_NUM_CAST]])
-  bar(*CONST);
-  // CHECK-DAG: [[CONST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST]]
-  // CHECK-DAG: [[CONST_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_LOAD]] to i32 addrspace(4)*
-  // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_CAST]])
-  bar2(*LOCAL_NUM);
-  // CHECK-DAG: [[LOCAL_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL_NUM]]
-  // CHECK-DAG: [[LOCAL_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_NUM_LOAD]] to i32 addrspace(4)*
-  // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOCAL_NUM_CAST]])
   bar2(*LOCAL);
   // CHECK-DAG: [[LOCAL_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL]]
   // CHECK-DAG: [[LOCAL_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_LOAD]] to i32 addrspace(4)*

clang/test/SemaOpenCLCXX/address-space-lambda.cl

@@ -31,8 +31,8 @@ __kernel void test_qual() {
 //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}}
+  auto priv3 = []() __global {}; //ex pected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //ex pected-note{{conversion candidate of type 'void (*)()'}}
+  priv3(); //ex pected-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}}

clang/test/SemaSYCL/address-space-parameter-conversions.cpp

@@ -13,7 +13,7 @@ void tmpl(T *t){}
 void usages() {
   __attribute__((opencl_global)) int *GLOB;
   __attribute__((opencl_private)) int *PRIV;
-  __attribute__((address_space(3))) int *LOC;
+  __attribute__((opencl_local)) int *LOC;
   int *NoAS;
 
   bar(*GLOB);
@@ -53,10 +53,6 @@ void usages() {
 
   // expected-error@+1{{address space is negative}}
   __attribute__((address_space(-1))) int *TooLow;
-  // expected-error@+1{{address space is outside the valid range of values}}
-  __attribute__((address_space(6))) int *TooHigh;
-  // expected-error@+1{{address space is outside the valid range of values}}
-  __attribute__((address_space(4))) int *TriedGeneric;
   // expected-error@+1{{unknown type name '__generic'}}
   __generic int *IsGeneric;
 

clang/test/SemaTemplate/address_space-dependent.cpp

@@ -43,7 +43,7 @@ void neg() {
 
 template <long int I>
 void tooBig() {
-  __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388590)}}
+  __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388595)}}
 }
 
 template <long int I>