[clang] [clang][SYCL] Add SYCL spelling for AS attributes (PR #200849)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 10 11:49:34 PDT 2026
https://github.com/elizabethandrews updated https://github.com/llvm/llvm-project/pull/200849
>From 8df10139dac2986fe9df25745fe90e99f1663b33 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews at intel.com>
Date: Fri, 8 May 2026 04:10:12 -0700
Subject: [PATCH 1/7] [clang][SYCL] Add SYCL spelling for AS attributes
This change renames OpenCL address space attributes in Attr.td
to generic offload variants to support OpenCL and SYCL spellings.
The following spellings were added - [[clang::sycl_private]],
[[clang::sycl_global]], [[clang::sycl_local]], [[clang::sycl_constant]],
and [[clang::sycl_generic]].
LangAS::sycl_generic was added to support SYCL's generic address space.
---
clang/include/clang/Basic/AddressSpaces.h | 5 ++
clang/include/clang/Basic/Attr.td | 32 ++++----
clang/include/clang/Basic/AttrDocs.td | 77 +++++++++++++------
.../clang/Basic/DiagnosticSemaKinds.td | 3 +
clang/include/clang/Sema/ParsedAttr.h | 25 +++---
clang/lib/AST/ItaniumMangle.cpp | 5 +-
clang/lib/AST/Type.cpp | 7 +-
clang/lib/AST/TypePrinter.cpp | 23 ++++--
clang/lib/Basic/TargetInfo.cpp | 1 +
clang/lib/Basic/Targets/AArch64.h | 1 +
clang/lib/Basic/Targets/AMDGPU.cpp | 2 +
clang/lib/Basic/Targets/DirectX.h | 1 +
clang/lib/Basic/Targets/NVPTX.h | 1 +
clang/lib/Basic/Targets/SPIR.h | 2 +
clang/lib/Basic/Targets/SystemZ.h | 1 +
clang/lib/Basic/Targets/TCE.h | 1 +
clang/lib/Basic/Targets/WebAssembly.h | 1 +
clang/lib/Basic/Targets/X86.h | 1 +
clang/lib/Headers/__clang_spirv_builtins.h | 10 ++-
clang/lib/Sema/ParsedAttr.cpp | 10 +--
clang/lib/Sema/SemaType.cpp | 19 +++--
.../CodeGenSYCL/address-space-conversions.cpp | 22 +++---
.../CodeGenSYCL/address-space-mangling.cpp | 12 +--
.../SemaSYCL/address-space-conversions.cpp | 39 +++++-----
.../address-space-opencl-sycl-compat.cpp | 34 ++++++++
.../SemaTemplate/address_space-dependent.cpp | 4 +-
26 files changed, 223 insertions(+), 116 deletions(-)
create mode 100644 clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h
index a941805423bca..58b04e50c2e3d 100644
--- a/clang/include/clang/Basic/AddressSpaces.h
+++ b/clang/include/clang/Basic/AddressSpaces.h
@@ -36,6 +36,8 @@ enum class LangAS : unsigned {
opencl_constant,
opencl_private,
opencl_generic,
+ // TODO: Remove opencl_global_device and opencl_global_host after corresponding
+ // attributes are deprecated for the required time.
opencl_global_device,
opencl_global_host,
@@ -46,10 +48,13 @@ enum class LangAS : unsigned {
// SYCL specific address spaces.
sycl_global,
+ // TODO: Remove sycl_global_device and sycl_global_host after corresponding attributes
+ // are deprecated for the required time.
sycl_global_device,
sycl_global_host,
sycl_local,
sycl_private,
+ sycl_generic,
// Pointer size and extension address spaces.
ptr32_sptr,
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 328e70b3ed900..aa0fb8d913322 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1791,44 +1791,46 @@ def OpenCLAccess : Attr {
let Documentation = [OpenCLAccessDocs];
}
-def OpenCLPrivateAddressSpace : TypeAttr {
+def OffloadPrivateAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__private">, CustomKeyword<"private">,
- Clang<"opencl_private">];
- let Documentation = [OpenCLAddressSpacePrivateDocs];
+ Clang<"opencl_private">, Clang<"sycl_private">];
+ let Documentation = [OffloadAddressSpacePrivateDocs];
}
-def OpenCLGlobalAddressSpace : TypeAttr {
+def OffloadGlobalAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__global">, CustomKeyword<"global">,
- Clang<"opencl_global">];
- let Documentation = [OpenCLAddressSpaceGlobalDocs];
+ Clang<"opencl_global">, Clang<"sycl_global">];
+ let Documentation = [OffloadAddressSpaceGlobalDocs];
}
+// TODO: Remove OpenCLGlobalDeviceAddressSpace after deprecation.
def OpenCLGlobalDeviceAddressSpace : TypeAttr {
let Spellings = [Clang<"opencl_global_device">];
let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
}
+// TODO: Remove OpenCLGlobalHostAddressSpace after deprecation.
def OpenCLGlobalHostAddressSpace : TypeAttr {
let Spellings = [Clang<"opencl_global_host">];
let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
}
-def OpenCLLocalAddressSpace : TypeAttr {
+def OffloadLocalAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__local">, CustomKeyword<"local">,
- Clang<"opencl_local">];
- let Documentation = [OpenCLAddressSpaceLocalDocs];
+ Clang<"opencl_local">, Clang<"sycl_local">];
+ let Documentation = [OffloadAddressSpaceLocalDocs];
}
-def OpenCLConstantAddressSpace : TypeAttr {
+def OffloadConstantAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__constant">, CustomKeyword<"constant">,
- Clang<"opencl_constant">];
- let Documentation = [OpenCLAddressSpaceConstantDocs];
+ Clang<"opencl_constant">, Clang<"sycl_constant">];
+ let Documentation = [OffloadAddressSpaceConstantDocs];
}
-def OpenCLGenericAddressSpace : TypeAttr {
+def OffloadGenericAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__generic">, CustomKeyword<"generic">,
- Clang<"opencl_generic">];
- let Documentation = [OpenCLAddressSpaceGenericDocs];
+ Clang<"opencl_generic">, Clang<"sycl_generic">];
+ let Documentation = [OffloadAddressSpaceGenericDocs];
}
def OpenCLNoSVM : Attr {
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 502af9b562ef0..74a6d7edaa4a2 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -5003,13 +5003,20 @@ More details can be found in the OpenCL C language Spec v2.0, Section 6.6.
}];
}
-def DocOpenCLAddressSpaces : DocumentationCategory<"OpenCL Address Spaces"> {
+def DocOffloadAddressSpaces : DocumentationCategory<"OpenCL and SYCL Address Spaces"> {
let Content = [{
The address space qualifier may be used to specify the region of memory that is
-used to allocate the object. OpenCL supports the following address spaces:
-__generic(generic), __global(global), __local(local), __private(private),
+used to allocate the object.
+
+OpenCL supports the following address spaces:
+
+__generic(generic), __global(global), __local(local), __private(private) and
__constant(constant).
+More details can be found in the OpenCL C language Spec v2.0, Section 6.5.
+
+Example:
+
.. code-block:: c
__constant int c = ...;
@@ -5021,14 +5028,32 @@ __constant(constant).
return l;
}
-More details can be found in the OpenCL C language Spec v2.0, Section 6.5.
+The memory model for SYCL devices is derived from the OpenCL memory model. Accordingly
+SYCL defines five address spaces: global, local, private, generic and constant. The
+following attributes correspond to these address spaces:
+
+[[clang::sycl_global]], [[clang::sycl_local]], [[clang::sycl_private]],
+[[clang::sycl_generic]] and [[clang::sycl_constant]] (deprecated)
+
+These attributes are intended for use in the implementation of SYCL run-time
+libraries. A direct declaration of pointers with address spaces is discouraged. Users
+should use the sycl::multi_ptr class to handle address space boundaries and
+interoperability.
+
+More details can be found in SYCL 2020 Specification, Section 3.8.2
+"SYCL device memory model" and Section 4.7.7, "Address space classes"
}];
+
}
-def OpenCLAddressSpaceGenericDocs : Documentation {
- let Category = DocOpenCLAddressSpaces;
- let Heading = "__generic, generic, [[clang::opencl_generic]]";
+def OffloadAddressSpaceGenericDocs : Documentation {
+ let Category = DocOffloadAddressSpaces;
+ let Heading = "__generic, generic, [[clang::opencl_generic]], [[clang::sycl_generic]]";
let Content = [{
+The generic address space is a virtual address space which overlaps the global, local
+and private address spaces.
+
+OpenCL:
The generic address space attribute is only available with OpenCL v2.0 and later.
It can be used with pointer types. Variables in global and local scope and
function parameters in non-kernel functions can have the generic address space
@@ -5038,33 +5063,35 @@ spaces.
}];
}
-def OpenCLAddressSpaceConstantDocs : Documentation {
- let Category = DocOpenCLAddressSpaces;
- let Heading = "__constant, constant, [[clang::opencl_constant]]";
+def OffloadAddressSpaceConstantDocs : Documentation {
+ let Category = DocOffloadAddressSpaces;
+ let Heading = "__constant, constant, [[clang::opencl_constant]], [[clang::sycl_constant]]";
let Content = [{
The constant address space attribute signals that an object is located in
a constant (non-modifiable) memory region. It is available to all work items.
Any type can be annotated with the constant address space attribute. Objects
with the constant address space qualifier can be declared in any scope and must
-have an initializer.
+have an initializer. The constant address space is deprecated in SYCL 2020
+specification.
}];
}
-def OpenCLAddressSpaceGlobalDocs : Documentation {
- let Category = DocOpenCLAddressSpaces;
- let Heading = "__global, global, [[clang::opencl_global]]";
+def OffloadAddressSpaceGlobalDocs : Documentation {
+ let Category = DocOffloadAddressSpaces;
+ let Heading = "__global, global, [[clang::opencl_global]], [[clang::sycl_global]]";
let Content = [{
The global address space attribute specifies that an object is allocated in
global memory, which is accessible by all work items. The content stored in this
-memory area persists between kernel executions. Pointer types to the global
-address space are allowed as function parameters or local variables. Starting
-with OpenCL v2.0, the global address space can be used with global (program
-scope) variables and static local variable as well.
+memory area persists between kernel executions.
+
+In OpenCL, pointer types to the global address space are allowed as function parameters or
+local variables. Starting with OpenCL v2.0, the global address space can be used with global
+(program scope) variables and static local variable as well.
}];
}
def OpenCLAddressSpaceGlobalExtDocs : Documentation {
- let Category = DocOpenCLAddressSpaces;
+ let Category = DocOffloadAddressSpaces;
let Heading = "[[clang::opencl_global_device]], [[clang::opencl_global_host]]";
let Content = [{
The ``global_device`` and ``global_host`` address space attributes specify that
@@ -5088,9 +5115,9 @@ As ``global_device`` and ``global_host`` are a subset of
}];
}
-def OpenCLAddressSpaceLocalDocs : Documentation {
- let Category = DocOpenCLAddressSpaces;
- let Heading = "__local, local, [[clang::opencl_local]]";
+def OffloadAddressSpaceLocalDocs : Documentation {
+ let Category = DocOffloadAddressSpaces;
+ let Heading = "__local, local, [[clang::opencl_local]], [[clang::sycl_local]]";
let Content = [{
The local address space specifies that an object is allocated in the local (work
group) memory area, which is accessible to all work items in the same work
@@ -5101,9 +5128,9 @@ space are allowed. Local address space variables cannot have an initializer.
}];
}
-def OpenCLAddressSpacePrivateDocs : Documentation {
- let Category = DocOpenCLAddressSpaces;
- let Heading = "__private, private, [[clang::opencl_private]]";
+def OffloadAddressSpacePrivateDocs : Documentation {
+ let Category = DocOffloadAddressSpaces;
+ let Heading = "__private, private, [[clang::opencl_private]], [[clang::sycl_private]]";
let Content = [{
The private address space specifies that an object is allocated in the private
(work item) memory. Other work items cannot access the same memory area and its
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 4cd4efc55c416..45ab89aa9ebc4 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -3597,6 +3597,9 @@ def err_attribute_address_multiple_qualifiers : Error<
def warn_attribute_address_multiple_identical_qualifiers : Warning<
"multiple identical address spaces specified for type">,
InGroup<DuplicateDeclSpecifier>;
+def warn_deprecated_sycl_constant : Warning<
+ "'sycl_constant' address space attribute is deprecated">,
+ InGroup<DeprecatedAttributes>;
def err_attribute_not_clinkage : Error<
"function type with %0 attribute must have C linkage">;
def err_function_decl_cmse_ns_call : Error<
diff --git a/clang/include/clang/Sema/ParsedAttr.h b/clang/include/clang/Sema/ParsedAttr.h
index 5387f9fad6cd2..9251c8aafdc71 100644
--- a/clang/include/clang/Sema/ParsedAttr.h
+++ b/clang/include/clang/Sema/ParsedAttr.h
@@ -553,44 +553,45 @@ class ParsedAttr final
/// a Spelling enumeration, the value UINT_MAX is returned.
unsigned getSemanticSpelling() const;
- /// If this is an OpenCL address space attribute, returns its representation
- /// in LangAS, otherwise returns default address space.
+ /// If this is a named address space attribute for OpenCL compilation, returns its
+ /// representation in LangAS, otherwise returns default address space.
LangAS asOpenCLLangAS() const {
switch (getParsedKind()) {
- case ParsedAttr::AT_OpenCLConstantAddressSpace:
+ case ParsedAttr::AT_OffloadConstantAddressSpace:
return LangAS::opencl_constant;
- case ParsedAttr::AT_OpenCLGlobalAddressSpace:
+ case ParsedAttr::AT_OffloadGlobalAddressSpace:
return LangAS::opencl_global;
case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
return LangAS::opencl_global_device;
case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
return LangAS::opencl_global_host;
- case ParsedAttr::AT_OpenCLLocalAddressSpace:
+ case ParsedAttr::AT_OffloadLocalAddressSpace:
return LangAS::opencl_local;
- case ParsedAttr::AT_OpenCLPrivateAddressSpace:
+ case ParsedAttr::AT_OffloadPrivateAddressSpace:
return LangAS::opencl_private;
- case ParsedAttr::AT_OpenCLGenericAddressSpace:
+ case ParsedAttr::AT_OffloadGenericAddressSpace:
return LangAS::opencl_generic;
default:
return LangAS::Default;
}
}
- /// If this is an OpenCL address space attribute, returns its SYCL
+ /// If this is a named address space attribute for SYCL compilation, returns its
/// representation in LangAS, otherwise returns default address space.
LangAS asSYCLLangAS() const {
switch (getKind()) {
- case ParsedAttr::AT_OpenCLGlobalAddressSpace:
+ case ParsedAttr::AT_OffloadGlobalAddressSpace:
return LangAS::sycl_global;
case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
return LangAS::sycl_global_device;
case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
return LangAS::sycl_global_host;
- case ParsedAttr::AT_OpenCLLocalAddressSpace:
+ case ParsedAttr::AT_OffloadLocalAddressSpace:
return LangAS::sycl_local;
- case ParsedAttr::AT_OpenCLPrivateAddressSpace:
+ case ParsedAttr::AT_OffloadPrivateAddressSpace:
return LangAS::sycl_private;
- case ParsedAttr::AT_OpenCLGenericAddressSpace:
+ case ParsedAttr::AT_OffloadGenericAddressSpace:
+ return LangAS::sycl_generic;
default:
return LangAS::Default;
}
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index f58faa03bfa8c..6be573966781c 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -2822,7 +2822,7 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals, const DependentAddressSp
ASString = "CLgeneric";
break;
// <SYCL-addrspace> ::= "SY" [ "global" | "local" | "private" |
- // "device" | "host" ]
+ // "generic" | "device" | "host" ]
case LangAS::sycl_global:
ASString = "SYglobal";
break;
@@ -2838,6 +2838,9 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals, const DependentAddressSp
case LangAS::sycl_private:
ASString = "SYprivate";
break;
+ case LangAS::sycl_generic:
+ ASString = "SYgeneric";
+ break;
// <CUDA-addrspace> ::= "CU" [ "device" | "constant" | "shared" ]
case LangAS::cuda_device:
ASString = "CUdevice";
diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp
index 78983fd38410d..f368b75600d58 100644
--- a/clang/lib/AST/Type.cpp
+++ b/clang/lib/AST/Type.cpp
@@ -84,11 +84,14 @@ bool Qualifiers::isTargetAddressSpaceSupersetOf(LangAS A, LangAS B,
// Consider pointer size address spaces to be equivalent to default.
((isPtrSizeAddressSpace(A) || A == LangAS::Default) &&
(isPtrSizeAddressSpace(B) || B == LangAS::Default)) ||
- // Default is a superset of SYCL address spaces.
- (A == LangAS::Default &&
+ // Default and sycl_generic are supersets of SYCL address spaces.
+ ((A == LangAS::Default || A == LangAS::sycl_generic) &&
(B == LangAS::sycl_private || B == LangAS::sycl_local ||
B == LangAS::sycl_global || B == LangAS::sycl_global_device ||
B == LangAS::sycl_global_host)) ||
+ // Consider sycl_generic address space to be equivalent to default.
+ (A == LangAS::Default && B == LangAS::sycl_generic) ||
+ (B == LangAS::Default && A == LangAS::sycl_generic) ||
// In HIP device compilation, any cuda address space is allowed
// to implicitly cast into the default address space.
(A == LangAS::Default &&
diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp
index 80f5b90ba35c4..ed5d95c360303 100644
--- a/clang/lib/AST/TypePrinter.cpp
+++ b/clang/lib/AST/TypePrinter.cpp
@@ -1992,13 +1992,13 @@ void TypePrinter::printAttributedAfter(const AttributedType *T,
case attr::HLSLResourceDimension:
llvm_unreachable("HLSL resource type attributes handled separately");
- case attr::OpenCLPrivateAddressSpace:
- case attr::OpenCLGlobalAddressSpace:
+ case attr::OffloadPrivateAddressSpace:
+ case attr::OffloadGlobalAddressSpace:
case attr::OpenCLGlobalDeviceAddressSpace:
case attr::OpenCLGlobalHostAddressSpace:
- case attr::OpenCLLocalAddressSpace:
- case attr::OpenCLConstantAddressSpace:
- case attr::OpenCLGenericAddressSpace:
+ case attr::OffloadLocalAddressSpace:
+ case attr::OffloadConstantAddressSpace:
+ case attr::OffloadGenericAddressSpace:
case attr::HLSLGroupSharedAddressSpace:
// FIXME: Update printAttributedBefore to print these once we generate
// AttributedType nodes for them.
@@ -2667,24 +2667,31 @@ 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:
return "__constant";
case LangAS::opencl_generic:
return "__generic";
+ // TODO: Remove *_global_device and *_global_host after corresponding
+ // attributes are deprecated for the required time.
case LangAS::opencl_global_device:
case LangAS::sycl_global_device:
return "__global_device";
case LangAS::opencl_global_host:
case LangAS::sycl_global_host:
return "__global_host";
+ case LangAS::sycl_global:
+ return "sycl_global";
+ case LangAS::sycl_local:
+ return "sycl_local";
+ case LangAS::sycl_private:
+ return "sycl_private";
+ case LangAS::sycl_generic:
+ return "sycl_generic";
case LangAS::cuda_device:
return "__device__";
case LangAS::cuda_constant:
diff --git a/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp
index e6ae89e0948c5..1d91bbaee21ef 100644
--- a/clang/lib/Basic/TargetInfo.cpp
+++ b/clang/lib/Basic/TargetInfo.cpp
@@ -44,6 +44,7 @@ static const LangASMap FakeAddrSpaceMap = {
6, // sycl_global_host
3, // sycl_local
0, // sycl_private
+ 4, // sycl_generic
10, // ptr32_sptr
11, // ptr32_uptr
12, // ptr64
diff --git a/clang/lib/Basic/Targets/AArch64.h b/clang/lib/Basic/Targets/AArch64.h
index 0a29bad81939b..90d8401149c37 100644
--- a/clang/lib/Basic/Targets/AArch64.h
+++ b/clang/lib/Basic/Targets/AArch64.h
@@ -41,6 +41,7 @@ static const unsigned ARM64AddrSpaceMap[] = {
0, // sycl_global_host
0, // sycl_local
0, // sycl_private
+ 0, // sycl_generic
static_cast<unsigned>(AArch64AddrSpace::ptr32_sptr),
static_cast<unsigned>(AArch64AddrSpace::ptr32_uptr),
static_cast<unsigned>(AArch64AddrSpace::ptr64),
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index 48ba5d94df581..5d102e015790f 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -43,6 +43,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsGenMap = {
llvm::AMDGPUAS::GLOBAL_ADDRESS, // sycl_global_host
llvm::AMDGPUAS::LOCAL_ADDRESS, // sycl_local
llvm::AMDGPUAS::PRIVATE_ADDRESS, // sycl_private
+ llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_generic
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_sptr
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_uptr
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr64
@@ -75,6 +76,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsPrivMap = {
llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_global_host
llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_local
llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_private
+ llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_generic
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_sptr
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_uptr
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr64
diff --git a/clang/lib/Basic/Targets/DirectX.h b/clang/lib/Basic/Targets/DirectX.h
index 8b21b86bac264..6eb770f4a960e 100644
--- a/clang/lib/Basic/Targets/DirectX.h
+++ b/clang/lib/Basic/Targets/DirectX.h
@@ -38,6 +38,7 @@ static const unsigned DirectXAddrSpaceMap[] = {
0, // sycl_global_host
0, // sycl_local
0, // sycl_private
+ 0, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h
index 69ee20f38343b..62bf6a514d444 100644
--- a/clang/lib/Basic/Targets/NVPTX.h
+++ b/clang/lib/Basic/Targets/NVPTX.h
@@ -42,6 +42,7 @@ static const unsigned NVPTXAddrSpaceMap[] = {
1, // sycl_global_host
3, // sycl_local
0, // sycl_private
+ 0, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 37fe4a970dfef..88a9e5841b5a3 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -44,6 +44,7 @@ static const unsigned SPIRDefIsPrivMap[] = {
0, // sycl_global_host
0, // sycl_local
0, // sycl_private
+ 0, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
@@ -82,6 +83,7 @@ static const unsigned SPIRDefIsGenMap[] = {
6, // sycl_global_host
3, // sycl_local
0, // sycl_private
+ 4, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/SystemZ.h b/clang/lib/Basic/Targets/SystemZ.h
index 00f7d7a055b24..bc597fe30165b 100644
--- a/clang/lib/Basic/Targets/SystemZ.h
+++ b/clang/lib/Basic/Targets/SystemZ.h
@@ -38,6 +38,7 @@ static const unsigned ZOSAddressMap[] = {
0, // sycl_global_host
0, // sycl_local
0, // sycl_private
+ 0, // sycl_generic
0, // ptr32_sptr
1, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/TCE.h b/clang/lib/Basic/Targets/TCE.h
index 2b22f4c4ec724..d7086b23232c4 100644
--- a/clang/lib/Basic/Targets/TCE.h
+++ b/clang/lib/Basic/Targets/TCE.h
@@ -47,6 +47,7 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = {
0, // sycl_global_host
0, // sycl_local
0, // sycl_private
+ 0, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/WebAssembly.h b/clang/lib/Basic/Targets/WebAssembly.h
index 808342485cad0..b0ee4505625c2 100644
--- a/clang/lib/Basic/Targets/WebAssembly.h
+++ b/clang/lib/Basic/Targets/WebAssembly.h
@@ -38,6 +38,7 @@ static const unsigned WebAssemblyAddrSpaceMap[] = {
0, // sycl_global_host
0, // sycl_local
0, // sycl_private
+ 0, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/X86.h b/clang/lib/Basic/Targets/X86.h
index c7afcc7c86053..31fb984caac51 100644
--- a/clang/lib/Basic/Targets/X86.h
+++ b/clang/lib/Basic/Targets/X86.h
@@ -42,6 +42,7 @@ static const unsigned X86AddrSpaceMap[] = {
0, // sycl_global_host
0, // sycl_local
0, // sycl_private
+ 0, // sycl_generic
270, // ptr32_sptr
271, // ptr32_uptr
272, // ptr64
diff --git a/clang/lib/Headers/__clang_spirv_builtins.h b/clang/lib/Headers/__clang_spirv_builtins.h
index 9c7215f506508..9b3ac8c134ab1 100644
--- a/clang/lib/Headers/__clang_spirv_builtins.h
+++ b/clang/lib/Headers/__clang_spirv_builtins.h
@@ -26,13 +26,17 @@
#define __SPIRV_convergent __attribute__((convergent))
#define __SPIRV_inline __attribute__((always_inline))
+#ifdef __SYCL_DEVICE_ONLY__
+#define __global __attribute__((sycl_global))
+#define __local __attribute__((sycl_local))
+#define __private __attribute__((sycl_private))
+#define __constant __attribute__((sycl_constant))
+#define __generic __attribute__((sycl_generic))
+#else
#define __global __attribute__((opencl_global))
#define __local __attribute__((opencl_local))
#define __private __attribute__((opencl_private))
#define __constant __attribute__((opencl_constant))
-#ifdef __SYCL_DEVICE_ONLY__
-#define __generic
-#else
#define __generic __attribute__((opencl_generic))
#endif
diff --git a/clang/lib/Sema/ParsedAttr.cpp b/clang/lib/Sema/ParsedAttr.cpp
index 2b5ad33ad7b29..49dec6188d877 100644
--- a/clang/lib/Sema/ParsedAttr.cpp
+++ b/clang/lib/Sema/ParsedAttr.cpp
@@ -228,13 +228,13 @@ bool ParsedAttr::slidesFromDeclToDeclSpecLegacyBehavior() const {
// possible, we would like this list to go away entirely.
switch (getParsedKind()) {
case AT_AddressSpace:
- case AT_OpenCLPrivateAddressSpace:
- case AT_OpenCLGlobalAddressSpace:
+ case AT_OffloadPrivateAddressSpace:
+ case AT_OffloadGlobalAddressSpace:
case AT_OpenCLGlobalDeviceAddressSpace:
case AT_OpenCLGlobalHostAddressSpace:
- case AT_OpenCLLocalAddressSpace:
- case AT_OpenCLConstantAddressSpace:
- case AT_OpenCLGenericAddressSpace:
+ case AT_OffloadLocalAddressSpace:
+ case AT_OffloadConstantAddressSpace:
+ case AT_OffloadGenericAddressSpace:
case AT_NeonPolyVectorType:
case AT_NeonVectorType:
case AT_ArmMveStrictPolymorphism:
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 846474fe94adf..1d13e632c51d6 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -6684,8 +6684,13 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
if (S.getLangOpts().HLSL)
ASIdx = Attr.asHLSLLangAS();
- if (ASIdx == LangAS::Default)
- llvm_unreachable("Invalid address space");
+ if (ASIdx == LangAS::Default) {
+ if (S.getLangOpts().SYCLIsDevice &&
+ Attr.getKind() == ParsedAttr::AT_OffloadConstantAddressSpace)
+ S.Diag(Attr.getLoc(), diag::warn_deprecated_sycl_constant);
+ else
+ llvm_unreachable("Invalid address space");
+ }
if (DiagnoseMultipleAddrSpaceAttributes(S, Type.getAddressSpace(), ASIdx,
Attr.getLoc())) {
@@ -9092,13 +9097,13 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type,
// it it breaks large amounts of Linux software.
attr.setUsedAsTypeAttr();
break;
- case ParsedAttr::AT_OpenCLPrivateAddressSpace:
- case ParsedAttr::AT_OpenCLGlobalAddressSpace:
+ case ParsedAttr::AT_OffloadPrivateAddressSpace:
+ case ParsedAttr::AT_OffloadGlobalAddressSpace:
case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
- case ParsedAttr::AT_OpenCLLocalAddressSpace:
- case ParsedAttr::AT_OpenCLConstantAddressSpace:
- case ParsedAttr::AT_OpenCLGenericAddressSpace:
+ case ParsedAttr::AT_OffloadLocalAddressSpace:
+ case ParsedAttr::AT_OffloadConstantAddressSpace:
+ case ParsedAttr::AT_OffloadGenericAddressSpace:
case ParsedAttr::AT_AddressSpace:
HandleAddressSpaceTypeAttribute(type, attr, state);
attr.setUsedAsTypeAttr();
diff --git a/clang/test/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CodeGenSYCL/address-space-conversions.cpp
index 506a24fb4a3ba..3eecdded06364 100644
--- a/clang/test/CodeGenSYCL/address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/address-space-conversions.cpp
@@ -3,13 +3,13 @@ void bar(int &Data) {}
// CHECK-DAG: define{{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) %
void bar2(int &Data) {}
// CHECK-DAG: define{{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) %
-void bar(__attribute__((opencl_local)) int &Data) {}
+void bar(__attribute__((sycl_local)) int &Data) {}
// CHECK-DAG: define{{.*}} void [[LOC_REF:@[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) %
void foo(int *Data) {}
// CHECK-DAG: define{{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef %
void foo2(int *Data) {}
// CHECK-DAG: define{{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef %
-void foo(__attribute__((opencl_local)) int *Data) {}
+void foo(__attribute__((sycl_local)) int *Data) {}
// CHECK-DAG: define{{.*}} void [[LOC_PTR:@[a-zA-Z0-9_]+]](ptr addrspace(3) noundef %
template <typename T>
@@ -19,11 +19,11 @@ void tmpl(T t) {}
[[clang::sycl_external]] void usages() {
int *NoAS;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr addrspace(4)
- __attribute__((opencl_global)) int *GLOB;
+ __attribute__((sycl_global)) int *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1)
- __attribute__((opencl_local)) int *LOC;
+ __attribute__((sycl_local)) int *LOC;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3)
- __attribute__((opencl_private)) int *PRIV;
+ __attribute__((sycl_private)) int *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr
__attribute__((opencl_global_device)) int *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5)
@@ -57,24 +57,24 @@ void tmpl(T t) {}
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[NoAS_CAST]], ptr addrspace(4) [[GLOB]].ascast
- GLOB = (__attribute__((opencl_global)) int *)NoAS;
+ GLOB = (__attribute__((sycl_global)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr addrspace(3)
// CHECK-DAG: store ptr addrspace(3) [[NoAS_CAST]], ptr addrspace(4) [[LOC]].ascast
- LOC = (__attribute__((opencl_local)) int *)NoAS;
+ LOC = (__attribute__((sycl_local)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr
// CHECK-DAG: store ptr [[NoAS_CAST]], ptr addrspace(4) [[PRIV]].ascast
- PRIV = (__attribute__((opencl_private)) int *)NoAS;
- // From opencl_global_[host/device] address spaces to opencl_global
+ PRIV = (__attribute__((sycl_private)) int *)NoAS;
+ // From opencl_global_[host/device] address spaces to sycl_global
// CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr addrspace(4) [[GLOB_DEVICE]].ascast
// CHECK-DAG: [[GLOBDEVICE_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[GLOBDEVICE_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_CAST]], ptr addrspace(4) [[GLOB]].ascast
- GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE;
+ GLOB = (__attribute__((sycl_global)) int *)GLOBDEVICE;
// CHECK-DAG: [[GLOBHOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(6), ptr addrspace(4) [[GLOB_HOST]].ascast
// CHECK-DAG: [[GLOBHOST_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(6) [[GLOBHOST_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[GLOBHOST_CAST]], ptr addrspace(4) [[GLOB]].ascast
- GLOB = (__attribute__((opencl_global)) int *)GLOBHOST;
+ GLOB = (__attribute__((sycl_global)) int *)GLOBHOST;
bar(*GLOB);
// CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast
diff --git a/clang/test/CodeGenSYCL/address-space-mangling.cpp b/clang/test/CodeGenSYCL/address-space-mangling.cpp
index ecc2d4b43a159..3e44be24ffc3e 100644
--- a/clang/test/CodeGenSYCL/address-space-mangling.cpp
+++ b/clang/test/CodeGenSYCL/address-space-mangling.cpp
@@ -3,9 +3,9 @@
// REQUIRES: x86-registered-target
-void foo(__attribute__((opencl_global)) int *);
-void foo(__attribute__((opencl_local)) int *);
-void foo(__attribute__((opencl_private)) int *);
+void foo(__attribute__((sycl_global)) int *);
+void foo(__attribute__((sycl_local)) int *);
+void foo(__attribute__((sycl_private)) int *);
void foo(int *);
// SPIR: declare spir_func void @_Z3fooPU3AS1i(ptr addrspace(1) noundef) #1
@@ -19,9 +19,9 @@ void foo(int *);
// X86: declare void @_Z3fooPi(ptr noundef) #1
[[clang::sycl_external]] void test() {
- __attribute__((opencl_global)) int *glob;
- __attribute__((opencl_local)) int *loc;
- __attribute__((opencl_private)) int *priv;
+ __attribute__((sycl_global)) int *glob;
+ __attribute__((sycl_local)) int *loc;
+ __attribute__((sycl_private)) int *priv;
int *def;
foo(glob);
foo(loc);
diff --git a/clang/test/SemaSYCL/address-space-conversions.cpp b/clang/test/SemaSYCL/address-space-conversions.cpp
index d8758248499de..9d209dbe5f8d7 100644
--- a/clang/test/SemaSYCL/address-space-conversions.cpp
+++ b/clang/test/SemaSYCL/address-space-conversions.cpp
@@ -2,28 +2,29 @@
void bar(int &Data) {}
void bar2(int &Data) {}
-void bar(__attribute__((opencl_private)) int &Data) {}
+void bar(__attribute__((sycl_private)) int &Data) {}
void foo(int *Data) {}
void foo2(int *Data) {}
-void foo(__attribute__((opencl_private)) int *Data) {}
-void baz(__attribute__((opencl_private)) int *Data) {} // expected-note {{candidate function not viable: cannot pass pointer to generic address space as a pointer to address space '__private' in 1st argument}}
+void foo(__attribute__((sycl_private)) int *Data) {}
+void baz(__attribute__((sycl_private)) int *Data) {} // expected-note {{candidate function not viable: cannot pass pointer to generic address space as a pointer to address space 'sycl_private' in 1st argument}}
template <typename T>
void tmpl(T *t) {}
void usages() {
- __attribute__((opencl_global)) int *GLOB;
- __attribute__((opencl_private)) int *PRIV;
- __attribute__((opencl_local)) int *LOC;
+ __attribute__((sycl_global)) int *GLOB;
+ __attribute__((sycl_private)) int *PRIV;
+ __attribute__((sycl_local)) int *LOC;
+ __attribute__((sycl_constant)) int *ptr1; // expected-warning {{'sycl_constant' address space attribute is deprecated}}
int *NoAS;
- GLOB = PRIV; // expected-error {{assigning '__private int *' to '__global int *' changes address space of pointer}}
- GLOB = LOC; // expected-error {{assigning '__local int *' to '__global int *' changes address space of pointer}}
- PRIV = static_cast<__attribute__((opencl_private)) int *>(GLOB); // expected-error {{static_cast from '__global int *' to '__private int *' is not allowed}}
- PRIV = static_cast<__attribute__((opencl_private)) int *>(LOC); // expected-error {{static_cast from '__local int *' to '__private int *' is not allowed}}
- NoAS = GLOB + PRIV; // expected-error {{invalid operands to binary expression ('__global int *' and '__private int *')}}
- NoAS = GLOB + LOC; // expected-error {{invalid operands to binary expression ('__global int *' and '__local int *')}}
- NoAS += GLOB; // expected-error {{invalid operands to binary expression ('int *' and '__global int *')}}
+ GLOB = PRIV; // expected-error {{assigning 'sycl_private int *' to 'sycl_global int *' changes address space of pointer}}
+ GLOB = LOC; // expected-error {{assigning 'sycl_local int *' to 'sycl_global int *' changes address space of pointer}}
+ PRIV = static_cast<__attribute__((sycl_private)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to 'sycl_private int *' is not allowed}}
+ PRIV = static_cast<__attribute__((sycl_private)) int *>(LOC); // expected-error {{static_cast from 'sycl_local int *' to 'sycl_private int *' is not allowed}}
+ NoAS = GLOB + PRIV; // expected-error {{invalid operands to binary expression ('sycl_global int *' and 'sycl_private int *')}}
+ NoAS = GLOB + LOC; // expected-error {{invalid operands to binary expression ('sycl_global int *' and 'sycl_local int *')}}
+ NoAS += GLOB; // expected-error {{invalid operands to binary expression ('int *' and 'sycl_global int *')}}
bar(*GLOB);
bar2(*GLOB);
@@ -53,10 +54,10 @@ void usages() {
// Implicit casts to named address space are disallowed
baz(NoAS); // expected-error {{no matching function for call to 'baz'}}
- __attribute__((opencl_local)) int *l = NoAS; // expected-error {{cannot initialize a variable of type '__local int *' with an lvalue of type 'int *'}}
+ __attribute__((sycl_local)) int *l = NoAS; // expected-error {{cannot initialize a variable of type 'sycl_local int *' with an lvalue of type 'int *'}}
// Explicit casts between disjoint address spaces are disallowed
- GLOB = (__attribute__((opencl_global)) int *)PRIV; // expected-error {{C-style cast from '__private int *' to '__global int *' converts between mismatching address spaces}}
+ GLOB = (__attribute__((sycl_global)) int *)PRIV; // expected-error {{C-style cast from 'sycl_private int *' to 'sycl_global int *' converts between mismatching address spaces}}
(void)static_cast<int *>(GLOB);
(void)static_cast<void *>(GLOB);
@@ -69,12 +70,12 @@ void usages() {
bar(*GLOB_HOST);
bar2(*GLOB_HOST);
GLOB = GLOB_HOST;
- GLOB_HOST = GLOB; // expected-error {{assigning '__global int *' to '__global_host int *' changes address space of pointer}}
- GLOB_HOST = static_cast<__attribute__((opencl_global_host)) int *>(GLOB); // expected-error {{static_cast from '__global int *' to '__global_host int *' is not allowed}}
+ GLOB_HOST = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_host int *' changes address space of pointer}}
+ GLOB_HOST = static_cast<__attribute__((opencl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_host int *' is not allowed}}
__attribute__((opencl_global_device)) int *GLOB_DEVICE;
bar(*GLOB_DEVICE);
bar2(*GLOB_DEVICE);
GLOB = GLOB_DEVICE;
- GLOB_DEVICE = GLOB; // expected-error {{assigning '__global int *' to '__global_device int *' changes address space of pointer}}
- GLOB_DEVICE = static_cast<__attribute__((opencl_global_device)) int *>(GLOB); // expected-error {{static_cast from '__global int *' to '__global_device int *' is not allowed}}
+ GLOB_DEVICE = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_device int *' changes address space of pointer}}
+ GLOB_DEVICE = static_cast<__attribute__((opencl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_device int *' is not allowed}}
}
diff --git a/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp b/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
new file mode 100644
index 0000000000000..89c4fa4873086
--- /dev/null
+++ b/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
@@ -0,0 +1,34 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
+
+// Test openCL and SYCL spelling conversions for address space
+// attributes.
+
+void test_incompatible() {
+ __attribute__((opencl_global)) int *opencl_global;
+ int [[clang::sycl_local]] *sycl_local;
+ int [[clang::sycl_private]] *sycl_private;
+
+ // Address space attributes are resolved using mode of compilation and not the spelling itself. This results in the SYCL spelling
+ // being used in both instances of each diagnostic despite openCL spelling being used.
+ opencl_global = sycl_local; // expected-error {{assigning 'sycl_local int *' to 'sycl_global int *' changes address space of pointer}}
+ opencl_global = sycl_private; // expected-error {{assigning 'sycl_private int *' to 'sycl_global int *' changes address space of pointer}}
+ sycl_local = opencl_global; // expected-error {{assigning 'sycl_global int *' to 'sycl_local int *' changes address space of pointer}}
+}
+
+void test_to_generic_mixed() {
+ __attribute__((opencl_generic)) int *opencl_gen;
+ int [[clang::sycl_generic]] *sycl_gen;
+
+ __attribute__((opencl_global)) int *opencl_global;
+ int [[clang::sycl_local]] *sycl_local;
+ int [[clang::sycl_private]] *sycl_private;
+
+ opencl_gen = sycl_local;
+ opencl_gen = sycl_private;
+ sycl_gen = opencl_global;
+
+}
+
+void overload_test(__attribute__((opencl_global)) int *p) { (void)p; } // expected-note {{previous definition is here}}
+void overload_test(__attribute__((sycl_global)) int *p) { (void)p; } // expected-error {{redefinition of 'overload_test'}}
+
diff --git a/clang/test/SemaTemplate/address_space-dependent.cpp b/clang/test/SemaTemplate/address_space-dependent.cpp
index 3fdccb2c71a76..d6f25923b69b5 100644
--- a/clang/test/SemaTemplate/address_space-dependent.cpp
+++ b/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 (8388580)}}
+ __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388579)}}
}
template <long int I>
@@ -101,7 +101,7 @@ int main() {
car<1, 2, 3>(); // expected-note {{in instantiation of function template specialization 'car<1, 2, 3>' requested here}}
HasASTemplateFields<1> HASTF;
neg<-1>(); // expected-note {{in instantiation of function template specialization 'neg<-1>' requested here}}
- correct<0x7FFFE4>();
+ correct<0x7FFFE3>();
tooBig<8388650>(); // expected-note {{in instantiation of function template specialization 'tooBig<8388650L>' requested here}}
__attribute__((address_space(1))) char *x;
>From bfe1c766e0e5a35afb834f3b8f2b787ffc261cb5 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews at intel.com>
Date: Fri, 29 May 2026 12:58:46 -0700
Subject: [PATCH 2/7] Choose address space mapping based on spelling
LangAS is selected based on spelling of attribute as opposed
to mode of compilation. It is unclear to me whether this is the
right move, and I am hoping for review comments to further discuss
this.
This changes behavior for all current sycl code using openCL spelling
for attributes and does not allow conversions between openCL and
sycl spellings even though they technically refer to the same
address space.
---
clang/include/clang/Basic/Attr.td | 74 ++++++++++++++++++-
clang/include/clang/Sema/ParsedAttr.h | 56 +-------------
clang/lib/Parse/ParseDecl.cpp | 2 +-
clang/lib/Sema/ParsedAttr.cpp | 24 ++++++
clang/lib/Sema/SemaType.cpp | 17 ++---
.../Builtins/generic_cast_to_ptr_explicit.c | 12 ++-
.../CodeGenSYCL/address-space-conversions.cpp | 6 +-
.../amd-address-space-conversions.cpp | 24 +++---
.../cuda-address-space-conversions.cpp | 24 +++---
.../SemaSYCL/address-space-conversions.cpp | 8 +-
.../address-space-opencl-sycl-compat.cpp | 13 ++--
11 files changed, 153 insertions(+), 107 deletions(-)
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index aa0fb8d913322..86ba95eef4d7c 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1795,42 +1795,112 @@ def OffloadPrivateAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__private">, CustomKeyword<"private">,
Clang<"opencl_private">, Clang<"sycl_private">];
let Documentation = [OffloadAddressSpacePrivateDocs];
+ let AdditionalMembers = [{
+ static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
+ return A.getAttributeSpellingListIndex() == GNU_sycl_private ||
+ A.getAttributeSpellingListIndex() == CXX11_clang_sycl_private ||
+ A.getAttributeSpellingListIndex() == C23_clang_sycl_private;
+ }
+ static inline LangAS getLangAS(const AttributeCommonInfo& A) {
+ return isSYCLSpelling(A) ? LangAS::sycl_private : LangAS::opencl_private;
+ }
+ }];
}
def OffloadGlobalAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__global">, CustomKeyword<"global">,
Clang<"opencl_global">, Clang<"sycl_global">];
let Documentation = [OffloadAddressSpaceGlobalDocs];
+ let AdditionalMembers = [{
+ static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
+ return A.getAttributeSpellingListIndex() == GNU_sycl_global ||
+ A.getAttributeSpellingListIndex() == CXX11_clang_sycl_global ||
+ A.getAttributeSpellingListIndex() == C23_clang_sycl_global;
+ }
+ static inline LangAS getLangAS(const AttributeCommonInfo& A) {
+ return isSYCLSpelling(A) ? LangAS::sycl_global : LangAS::opencl_global;
+ }
+ }];
}
// TODO: Remove OpenCLGlobalDeviceAddressSpace after deprecation.
def OpenCLGlobalDeviceAddressSpace : TypeAttr {
- let Spellings = [Clang<"opencl_global_device">];
+ let Spellings = [Clang<"opencl_global_device">, Clang<"sycl_global_device">];
let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
+ let AdditionalMembers = [{
+ static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
+ return A.getAttributeSpellingListIndex() == GNU_sycl_global_device ||
+ A.getAttributeSpellingListIndex() == CXX11_clang_sycl_global_device ||
+ A.getAttributeSpellingListIndex() == C23_clang_sycl_global_device;
+ }
+ static inline LangAS getLangAS(const AttributeCommonInfo& A) {
+ return isSYCLSpelling(A) ? LangAS::sycl_global_device : LangAS::opencl_global_device;
+ }
+ }];
}
// TODO: Remove OpenCLGlobalHostAddressSpace after deprecation.
def OpenCLGlobalHostAddressSpace : TypeAttr {
- let Spellings = [Clang<"opencl_global_host">];
+ let Spellings = [Clang<"opencl_global_host">, Clang<"sycl_global_host">];
let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
+ let AdditionalMembers = [{
+ static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
+ return A.getAttributeSpellingListIndex() == GNU_sycl_global_host ||
+ A.getAttributeSpellingListIndex() == CXX11_clang_sycl_global_host ||
+ A.getAttributeSpellingListIndex() == C23_clang_sycl_global_host;
+ }
+ static inline LangAS getLangAS(const AttributeCommonInfo& A) {
+ return isSYCLSpelling(A) ? LangAS::sycl_global_host : LangAS::opencl_global_host;
+ }
+ }];
}
def OffloadLocalAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__local">, CustomKeyword<"local">,
Clang<"opencl_local">, Clang<"sycl_local">];
let Documentation = [OffloadAddressSpaceLocalDocs];
+ let AdditionalMembers = [{
+ static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
+ return A.getAttributeSpellingListIndex() == GNU_sycl_local ||
+ A.getAttributeSpellingListIndex() == CXX11_clang_sycl_local ||
+ A.getAttributeSpellingListIndex() == C23_clang_sycl_local;
+ }
+ static inline LangAS getLangAS(const AttributeCommonInfo& A) {
+ return isSYCLSpelling(A) ? LangAS::sycl_local : LangAS::opencl_local;
+ }
+ }];
}
def OffloadConstantAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__constant">, CustomKeyword<"constant">,
Clang<"opencl_constant">, Clang<"sycl_constant">];
let Documentation = [OffloadAddressSpaceConstantDocs];
+ let AdditionalMembers = [{
+ static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
+ return A.getAttributeSpellingListIndex() == GNU_sycl_constant ||
+ A.getAttributeSpellingListIndex() == CXX11_clang_sycl_constant ||
+ A.getAttributeSpellingListIndex() == C23_clang_sycl_constant;
+ }
+ static inline LangAS getLangAS(const AttributeCommonInfo& A) {
+ return isSYCLSpelling(A) ? LangAS::Default : LangAS::opencl_constant;
+ }
+ }];
}
def OffloadGenericAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__generic">, CustomKeyword<"generic">,
Clang<"opencl_generic">, Clang<"sycl_generic">];
let Documentation = [OffloadAddressSpaceGenericDocs];
+ let AdditionalMembers = [{
+ static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
+ return A.getAttributeSpellingListIndex() == GNU_sycl_generic ||
+ A.getAttributeSpellingListIndex() == CXX11_clang_sycl_generic ||
+ A.getAttributeSpellingListIndex() == C23_clang_sycl_generic;
+ }
+ static inline LangAS getLangAS(const AttributeCommonInfo& A) {
+ return isSYCLSpelling(A) ? LangAS::sycl_generic : LangAS::opencl_generic;
+ }
+ }];
}
def OpenCLNoSVM : Attr {
diff --git a/clang/include/clang/Sema/ParsedAttr.h b/clang/include/clang/Sema/ParsedAttr.h
index 9251c8aafdc71..ddb26e89bdaa6 100644
--- a/clang/include/clang/Sema/ParsedAttr.h
+++ b/clang/include/clang/Sema/ParsedAttr.h
@@ -553,60 +553,8 @@ class ParsedAttr final
/// a Spelling enumeration, the value UINT_MAX is returned.
unsigned getSemanticSpelling() const;
- /// If this is a named address space attribute for OpenCL compilation, returns its
- /// representation in LangAS, otherwise returns default address space.
- LangAS asOpenCLLangAS() const {
- switch (getParsedKind()) {
- case ParsedAttr::AT_OffloadConstantAddressSpace:
- return LangAS::opencl_constant;
- case ParsedAttr::AT_OffloadGlobalAddressSpace:
- return LangAS::opencl_global;
- case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
- return LangAS::opencl_global_device;
- case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
- return LangAS::opencl_global_host;
- case ParsedAttr::AT_OffloadLocalAddressSpace:
- return LangAS::opencl_local;
- case ParsedAttr::AT_OffloadPrivateAddressSpace:
- return LangAS::opencl_private;
- case ParsedAttr::AT_OffloadGenericAddressSpace:
- return LangAS::opencl_generic;
- default:
- return LangAS::Default;
- }
- }
-
- /// If this is a named address space attribute for SYCL compilation, returns its
- /// representation in LangAS, otherwise returns default address space.
- LangAS asSYCLLangAS() const {
- switch (getKind()) {
- case ParsedAttr::AT_OffloadGlobalAddressSpace:
- return LangAS::sycl_global;
- case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
- return LangAS::sycl_global_device;
- case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
- return LangAS::sycl_global_host;
- case ParsedAttr::AT_OffloadLocalAddressSpace:
- return LangAS::sycl_local;
- case ParsedAttr::AT_OffloadPrivateAddressSpace:
- return LangAS::sycl_private;
- case ParsedAttr::AT_OffloadGenericAddressSpace:
- return LangAS::sycl_generic;
- default:
- return LangAS::Default;
- }
- }
-
- /// If this is an HLSL address space attribute, returns its representation
- /// in LangAS, otherwise returns default address space.
- LangAS asHLSLLangAS() const {
- switch (getParsedKind()) {
- case ParsedAttr::AT_HLSLGroupSharedAddressSpace:
- return LangAS::hlsl_groupshared;
- default:
- return LangAS::Default;
- }
- }
+ /// Returns the appropriate LangAS for this address space attribute.
+ LangAS asLangAS() const;
AttributeCommonInfo::Kind getKind() const {
return AttributeCommonInfo::Kind(Info.AttrKind);
diff --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp
index e2ac86bc5e064..1a4b92ca99b99 100644
--- a/clang/lib/Parse/ParseDecl.cpp
+++ b/clang/lib/Parse/ParseDecl.cpp
@@ -7227,7 +7227,7 @@ void Parser::InitCXXThisScopeForDeclaratorIfRelevant(
// prototype for the method.
if (getLangOpts().OpenCLCPlusPlus) {
for (ParsedAttr &attr : DS.getAttributes()) {
- LangAS ASIdx = attr.asOpenCLLangAS();
+ LangAS ASIdx = attr.asLangAS();
if (ASIdx != LangAS::Default) {
Q.addAddressSpace(ASIdx);
break;
diff --git a/clang/lib/Sema/ParsedAttr.cpp b/clang/lib/Sema/ParsedAttr.cpp
index 49dec6188d877..08b9aea3bdb15 100644
--- a/clang/lib/Sema/ParsedAttr.cpp
+++ b/clang/lib/Sema/ParsedAttr.cpp
@@ -12,6 +12,7 @@
#include "clang/Sema/ParsedAttr.h"
#include "clang/AST/ASTContext.h"
+#include "clang/AST/Attr.h"
#include "clang/Basic/AttrSubjectMatchRules.h"
#include "clang/Basic/IdentifierTable.h"
#include "clang/Basic/TargetInfo.h"
@@ -312,3 +313,26 @@ void clang::takeAndConcatenateAttrs(ParsedAttributes &First,
if (Second.Range.getEnd().isValid())
First.Range.setEnd(Second.Range.getEnd());
}
+
+LangAS ParsedAttr::asLangAS() const {
+ switch (getParsedKind()) {
+ case ParsedAttr::AT_OffloadGlobalAddressSpace:
+ return OffloadGlobalAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
+ return OpenCLGlobalDeviceAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
+ return OpenCLGlobalHostAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OffloadLocalAddressSpace:
+ return OffloadLocalAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OffloadPrivateAddressSpace:
+ return OffloadPrivateAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OffloadConstantAddressSpace:
+ return OffloadConstantAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OffloadGenericAddressSpace:
+ return OffloadGenericAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_HLSLGroupSharedAddressSpace:
+ return LangAS::hlsl_groupshared;
+ default:
+ return LangAS::Default;
+ }
+}
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 1d13e632c51d6..326910b564df1 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -5407,7 +5407,7 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
// them later while creating QualType.
if (FTI.MethodQualifiers)
for (ParsedAttr &attr : FTI.MethodQualifiers->getAttributes()) {
- LangAS ASIdxNew = attr.asOpenCLLangAS();
+ LangAS ASIdxNew = attr.asLangAS();
if (DiagnoseMultipleAddrSpaceAttributes(S, ASIdx, ASIdxNew,
attr.getLoc()))
D.setInvalidType(true);
@@ -6678,15 +6678,12 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
else
Attr.setInvalid();
} else {
- // The keyword-based type attributes imply which address space to use.
- ASIdx = S.getLangOpts().SYCLIsDevice ? Attr.asSYCLLangAS()
- : Attr.asOpenCLLangAS();
- if (S.getLangOpts().HLSL)
- ASIdx = Attr.asHLSLLangAS();
-
- if (ASIdx == LangAS::Default) {
- if (S.getLangOpts().SYCLIsDevice &&
- Attr.getKind() == ParsedAttr::AT_OffloadConstantAddressSpace)
+ // Type attributes imply which address space to use.
+ ASIdx = Attr.asLangAS();
+
+ if (ASIdx == LangAS::Default &&
+ Attr.getKind() == ParsedAttr::AT_OffloadConstantAddressSpace) {
+ if (OffloadConstantAddressSpaceAttr::isSYCLSpelling(Attr))
S.Diag(Attr.getLoc(), diag::warn_deprecated_sycl_constant);
else
llvm_unreachable("Invalid address space");
diff --git a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
index 30f4ecb589a5c..b896d76897fbf 100644
--- a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
+++ b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
@@ -4,8 +4,14 @@
#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_EXTERNAL [[clang::sycl_external]]
+#define __global __attribute__((sycl_global))
+#define __local __attribute__((sycl_local))
+#define __private __attribute__((sycl_private))
#else
#define SYCL_EXTERNAL
+#define __global __attribute__((opencl_global))
+#define __local __attribute__((opencl_local))
+#define __private __attribute__((opencl_private))
#endif
// CHECK: spir_func noundef ptr @{{.*}}test_cast_to_private{{.*}}(ptr addrspace(4) noundef readnone [[P:%.*]]
@@ -13,7 +19,7 @@
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr @llvm.spv.generic.cast.to.ptr.explicit.p0(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr [[SPV_CAST]]
//
-SYCL_EXTERNAL __attribute__((opencl_private)) int* test_cast_to_private(int* p) {
+SYCL_EXTERNAL __private int* test_cast_to_private(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 7);
}
@@ -22,7 +28,7 @@ SYCL_EXTERNAL __attribute__((opencl_private)) int* test_cast_to_private(int* p)
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(1) @llvm.spv.generic.cast.to.ptr.explicit.p1(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr addrspace(1) [[SPV_CAST]]
//
-SYCL_EXTERNAL __attribute__((opencl_global)) int* test_cast_to_global(int* p) {
+SYCL_EXTERNAL __global int* test_cast_to_global(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 5);
}
@@ -31,6 +37,6 @@ SYCL_EXTERNAL __attribute__((opencl_global)) int* test_cast_to_global(int* p) {
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(3) @llvm.spv.generic.cast.to.ptr.explicit.p3(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr addrspace(3) [[SPV_CAST]]
//
-SYCL_EXTERNAL __attribute__((opencl_local)) int* test_cast_to_local(int* p) {
+SYCL_EXTERNAL __local int* test_cast_to_local(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 4);
}
diff --git a/clang/test/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CodeGenSYCL/address-space-conversions.cpp
index 3eecdded06364..f331b8367b614 100644
--- a/clang/test/CodeGenSYCL/address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/address-space-conversions.cpp
@@ -25,9 +25,9 @@ void tmpl(T t) {}
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3)
__attribute__((sycl_private)) int *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr
- __attribute__((opencl_global_device)) int *GLOBDEVICE;
+ __attribute__((sycl_global_device)) int *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5)
- __attribute__((opencl_global_host)) int *GLOBHOST;
+ __attribute__((sycl_global_host)) int *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(6)
// CHECK-DAG: [[NoAS]].ascast = addrspacecast ptr [[NoAS]] to ptr addrspace(4)
@@ -66,7 +66,7 @@ void tmpl(T t) {}
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr
// CHECK-DAG: store ptr [[NoAS_CAST]], ptr addrspace(4) [[PRIV]].ascast
PRIV = (__attribute__((sycl_private)) int *)NoAS;
- // From opencl_global_[host/device] address spaces to sycl_global
+ // From sycl_global_[host/device] address spaces to sycl_global
// CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr addrspace(4) [[GLOB_DEVICE]].ascast
// CHECK-DAG: [[GLOBDEVICE_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[GLOBDEVICE_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_CAST]], ptr addrspace(4) [[GLOB]].ascast
diff --git a/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp b/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
index 17a98195318ad..a46b9660b2ef9 100644
--- a/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
@@ -3,13 +3,13 @@ void bar(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
void bar2(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
-void bar(__attribute__((opencl_local)) int &Data) {}
+void bar(__attribute__((sycl_local)) int &Data) {}
// CHECK-DAG: define {{.*}} void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) %
void foo(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef %
void foo2(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef %
-void foo(__attribute__((opencl_local)) int *Data) {}
+void foo(__attribute__((sycl_local)) int *Data) {}
// CHECK-DAG: define {{.*}} void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef %
template <typename T>
@@ -19,15 +19,15 @@ void tmpl(T t) {}
[[clang::sycl_external]] void usages() {
int *NoAS;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8, addrspace(5)
- __attribute__((opencl_global)) int *GLOB;
+ __attribute__((sycl_global)) int *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
- __attribute__((opencl_local)) int *LOC;
+ __attribute__((sycl_local)) int *LOC;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 4, addrspace(5)
- __attribute__((opencl_private)) int *PRIV;
+ __attribute__((sycl_private)) int *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5), align 4, addrspace(5)
- __attribute__((opencl_global_device)) int *GLOBDEVICE;
+ __attribute__((sycl_global_device)) int *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
- __attribute__((opencl_global_host)) int *GLOBHOST;
+ __attribute__((sycl_global_host)) int *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
LOC = nullptr;
// CHECK-DAG: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]].ascast, align 4
@@ -45,22 +45,22 @@ void tmpl(T t) {}
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[NoAS_LOAD]] to ptr
// CHECK-DAG: store ptr %5, ptr [[NoAS]].ascast, align 8
- GLOB = (__attribute__((opencl_global)) int *)NoAS;
+ GLOB = (__attribute__((sycl_global)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr %6 to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) %7, ptr [[GLOB]].ascast, align 8
- LOC = (__attribute__((opencl_local)) int *)NoAS;
+ LOC = (__attribute__((sycl_local)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3)
// CHECK-DAG: store ptr addrspace(3) %9, ptr [[LOC]].ascast, align 4
- PRIV = (__attribute__((opencl_private)) int *)NoAS;
+ PRIV = (__attribute__((sycl_private)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(5)
// CHECK-DAG: store ptr addrspace(5) [[NoAS_CAST]], ptr [[PRIV]].ascast, align 4
- GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE;
+ GLOB = (__attribute__((sycl_global)) int *)GLOBDEVICE;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]DEVICE.ascast, align 8
// CHECK-DAG: store ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8
- GLOB = (__attribute__((opencl_global)) int *)GLOBHOST;
+ GLOB = (__attribute__((sycl_global)) int *)GLOBHOST;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]HOST.ascast, align 8
// CHECK-DAG: tore ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8
bar(*GLOB);
diff --git a/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp b/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
index ffb601e62c118..3427450547fce 100644
--- a/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
@@ -3,13 +3,13 @@ void bar(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
void bar2(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
-void bar(__attribute__((opencl_local)) int &Data) {}
+void bar(__attribute__((sycl_local)) int &Data) {}
// CHECK-DAG: define {{.*}} void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) %
void foo(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef %
void foo2(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef %
-void foo(__attribute__((opencl_local)) int *Data) {}
+void foo(__attribute__((sycl_local)) int *Data) {}
// CHECK-DAG: define {{.*}} void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef %
template <typename T>
@@ -19,15 +19,15 @@ void tmpl(T t);
[[clang::sycl_external]] void usages() {
int *NoAS;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8
- __attribute__((opencl_global)) int *GLOB;
+ __attribute__((sycl_global)) int *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
- __attribute__((opencl_local)) int *LOC;
+ __attribute__((sycl_local)) int *LOC;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 8
- __attribute__((opencl_private)) int *PRIV;
+ __attribute__((sycl_private)) int *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr, align 8
- __attribute__((opencl_global_device)) int *GLOBDEVICE;
+ __attribute__((sycl_global_device)) int *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
- __attribute__((opencl_global_host)) int *GLOBHOST;
+ __attribute__((sycl_global_host)) int *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
LOC = nullptr;
// CHECK-DAG: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]], align 8
@@ -44,21 +44,21 @@ void tmpl(T t);
NoAS = (int *)PRIV;
// CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[PRIV]], align 8
// CHECK-DAG: store ptr [[LOC_LOAD]], ptr [[NoAS]], align 8
- GLOB = (__attribute__((opencl_global)) int *)NoAS;
+ GLOB = (__attribute__((sycl_global)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[NoAS_CAST]], ptr [[GLOB]], align 8
- LOC = (__attribute__((opencl_local)) int *)NoAS;
+ LOC = (__attribute__((sycl_local)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3)
// CHECK-DAG: store ptr addrspace(3) [[NoAS_CAST]], ptr [[LOC]], align 8
- PRIV = (__attribute__((opencl_private)) int *)NoAS;
+ PRIV = (__attribute__((sycl_private)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: store ptr [[NoAS_LOAD]], ptr [[PRIV]], align 8
- GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE;
+ GLOB = (__attribute__((sycl_global)) int *)GLOBDEVICE;
// CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_DEVICE]], align 8
// CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_LOAD]], ptr %GLOB, align 8
- GLOB = (__attribute__((opencl_global)) int *)GLOBHOST;
+ GLOB = (__attribute__((sycl_global)) int *)GLOBHOST;
// CHECK-DAG: [[GLOB_HOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_HOST]], align 8
// CHECK-DAG: store ptr addrspace(1) [[GLOB_HOST_LOAD]], ptr [[GLOB]], align 8
bar(*GLOB);
diff --git a/clang/test/SemaSYCL/address-space-conversions.cpp b/clang/test/SemaSYCL/address-space-conversions.cpp
index 9d209dbe5f8d7..0112ccae4c775 100644
--- a/clang/test/SemaSYCL/address-space-conversions.cpp
+++ b/clang/test/SemaSYCL/address-space-conversions.cpp
@@ -66,16 +66,16 @@ void usages() {
(void)i;
(void)v;
- __attribute__((opencl_global_host)) int *GLOB_HOST;
+ __attribute__((sycl_global_host)) int *GLOB_HOST;
bar(*GLOB_HOST);
bar2(*GLOB_HOST);
GLOB = GLOB_HOST;
GLOB_HOST = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_host int *' changes address space of pointer}}
- GLOB_HOST = static_cast<__attribute__((opencl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_host int *' is not allowed}}
- __attribute__((opencl_global_device)) int *GLOB_DEVICE;
+ GLOB_HOST = static_cast<__attribute__((sycl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_host int *' is not allowed}}
+ __attribute__((sycl_global_device)) int *GLOB_DEVICE;
bar(*GLOB_DEVICE);
bar2(*GLOB_DEVICE);
GLOB = GLOB_DEVICE;
GLOB_DEVICE = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_device int *' changes address space of pointer}}
- GLOB_DEVICE = static_cast<__attribute__((opencl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_device int *' is not allowed}}
+ GLOB_DEVICE = static_cast<__attribute__((sycl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_device int *' is not allowed}}
}
diff --git a/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp b/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
index 89c4fa4873086..7b4575e60c8ad 100644
--- a/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
+++ b/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
@@ -10,9 +10,9 @@ void test_incompatible() {
// Address space attributes are resolved using mode of compilation and not the spelling itself. This results in the SYCL spelling
// being used in both instances of each diagnostic despite openCL spelling being used.
- opencl_global = sycl_local; // expected-error {{assigning 'sycl_local int *' to 'sycl_global int *' changes address space of pointer}}
- opencl_global = sycl_private; // expected-error {{assigning 'sycl_private int *' to 'sycl_global int *' changes address space of pointer}}
- sycl_local = opencl_global; // expected-error {{assigning 'sycl_global int *' to 'sycl_local int *' changes address space of pointer}}
+ opencl_global = sycl_local; // expected-error {{assigning 'sycl_local int *' to '__global int *' changes address space of pointer}}
+ opencl_global = sycl_private; // expected-error {{assigning 'sycl_private int *' to '__global int *' changes address space of pointer}}
+ sycl_local = opencl_global; // expected-error {{assigning '__global int *' to 'sycl_local int *' changes address space of pointer}}
}
void test_to_generic_mixed() {
@@ -23,12 +23,13 @@ void test_to_generic_mixed() {
int [[clang::sycl_local]] *sycl_local;
int [[clang::sycl_private]] *sycl_private;
+ //FIXME: Why don't these throw an error?
opencl_gen = sycl_local;
opencl_gen = sycl_private;
- sycl_gen = opencl_global;
+ sycl_gen = opencl_global; // expected-error {{assigning '__global int *' to 'sycl_generic int *' changes address space of pointer}}
}
-void overload_test(__attribute__((opencl_global)) int *p) { (void)p; } // expected-note {{previous definition is here}}
-void overload_test(__attribute__((sycl_global)) int *p) { (void)p; } // expected-error {{redefinition of 'overload_test'}}
+void overload_test(__attribute__((opencl_global)) int *p) { (void)p; }
+void overload_test(__attribute__((sycl_global)) int *p) { (void)p; }
>From 071ba770ef59dd5c5c8338d2dce360a5cf925feb Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews at intel.com>
Date: Mon, 1 Jun 2026 08:34:55 -0700
Subject: [PATCH 3/7] Update global_host and global_device attributes as well
for uniformity.
---
clang/include/clang/Basic/Attr.td | 12 ++++++------
clang/include/clang/Basic/AttrDocs.td | 4 ++--
clang/lib/AST/TypePrinter.cpp | 12 ++++++------
clang/lib/Sema/ParsedAttr.cpp | 12 ++++++------
clang/lib/Sema/SemaType.cpp | 4 ++--
clang/test/SemaSYCL/address-space-conversions.cpp | 8 ++++----
6 files changed, 26 insertions(+), 26 deletions(-)
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 86ba95eef4d7c..00fb4dc8e24ee 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1823,10 +1823,10 @@ def OffloadGlobalAddressSpace : TypeAttr {
}];
}
-// TODO: Remove OpenCLGlobalDeviceAddressSpace after deprecation.
-def OpenCLGlobalDeviceAddressSpace : TypeAttr {
+// TODO: Remove OffloadGlobalDeviceAddressSpace after deprecation.
+def OffloadGlobalDeviceAddressSpace : TypeAttr {
let Spellings = [Clang<"opencl_global_device">, Clang<"sycl_global_device">];
- let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
+ let Documentation = [OffloadAddressSpaceGlobalExtDocs];
let AdditionalMembers = [{
static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
return A.getAttributeSpellingListIndex() == GNU_sycl_global_device ||
@@ -1839,10 +1839,10 @@ def OpenCLGlobalDeviceAddressSpace : TypeAttr {
}];
}
-// TODO: Remove OpenCLGlobalHostAddressSpace after deprecation.
-def OpenCLGlobalHostAddressSpace : TypeAttr {
+// TODO: Remove OffloadGlobalHostAddressSpace after deprecation.
+def OffloadGlobalHostAddressSpace : TypeAttr {
let Spellings = [Clang<"opencl_global_host">, Clang<"sycl_global_host">];
- let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
+ let Documentation = [OffloadAddressSpaceGlobalExtDocs];
let AdditionalMembers = [{
static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
return A.getAttributeSpellingListIndex() == GNU_sycl_global_host ||
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 74a6d7edaa4a2..d8e009121df44 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -5090,9 +5090,9 @@ local variables. Starting with OpenCL v2.0, the global address space can be used
}];
}
-def OpenCLAddressSpaceGlobalExtDocs : Documentation {
+def OffloadAddressSpaceGlobalExtDocs : Documentation {
let Category = DocOffloadAddressSpaces;
- let Heading = "[[clang::opencl_global_device]], [[clang::opencl_global_host]]";
+ let Heading = "[[clang::opencl_global_device]], [[clang::opencl_global_host]], [[clang::sycl_global_device]], [[clang::sycl_global_host]]";
let Content = [{
The ``global_device`` and ``global_host`` address space attributes specify that
an object is allocated in global memory on the device/host. It helps to
diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp
index ed5d95c360303..ace0963158b84 100644
--- a/clang/lib/AST/TypePrinter.cpp
+++ b/clang/lib/AST/TypePrinter.cpp
@@ -1994,8 +1994,8 @@ void TypePrinter::printAttributedAfter(const AttributedType *T,
case attr::OffloadPrivateAddressSpace:
case attr::OffloadGlobalAddressSpace:
- case attr::OpenCLGlobalDeviceAddressSpace:
- case attr::OpenCLGlobalHostAddressSpace:
+ case attr::OffloadGlobalDeviceAddressSpace:
+ case attr::OffloadGlobalHostAddressSpace:
case attr::OffloadLocalAddressSpace:
case attr::OffloadConstantAddressSpace:
case attr::OffloadGenericAddressSpace:
@@ -2676,13 +2676,9 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
return "__constant";
case LangAS::opencl_generic:
return "__generic";
- // TODO: Remove *_global_device and *_global_host after corresponding
- // attributes are deprecated for the required time.
case LangAS::opencl_global_device:
- case LangAS::sycl_global_device:
return "__global_device";
case LangAS::opencl_global_host:
- case LangAS::sycl_global_host:
return "__global_host";
case LangAS::sycl_global:
return "sycl_global";
@@ -2692,6 +2688,10 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
return "sycl_private";
case LangAS::sycl_generic:
return "sycl_generic";
+ case LangAS::sycl_global_device:
+ return "sycl_global_device";
+ case LangAS::sycl_global_host:
+ return "sycl_global_host";
case LangAS::cuda_device:
return "__device__";
case LangAS::cuda_constant:
diff --git a/clang/lib/Sema/ParsedAttr.cpp b/clang/lib/Sema/ParsedAttr.cpp
index 08b9aea3bdb15..be60e19819431 100644
--- a/clang/lib/Sema/ParsedAttr.cpp
+++ b/clang/lib/Sema/ParsedAttr.cpp
@@ -231,8 +231,8 @@ bool ParsedAttr::slidesFromDeclToDeclSpecLegacyBehavior() const {
case AT_AddressSpace:
case AT_OffloadPrivateAddressSpace:
case AT_OffloadGlobalAddressSpace:
- case AT_OpenCLGlobalDeviceAddressSpace:
- case AT_OpenCLGlobalHostAddressSpace:
+ case AT_OffloadGlobalDeviceAddressSpace:
+ case AT_OffloadGlobalHostAddressSpace:
case AT_OffloadLocalAddressSpace:
case AT_OffloadConstantAddressSpace:
case AT_OffloadGenericAddressSpace:
@@ -318,10 +318,10 @@ LangAS ParsedAttr::asLangAS() const {
switch (getParsedKind()) {
case ParsedAttr::AT_OffloadGlobalAddressSpace:
return OffloadGlobalAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
- return OpenCLGlobalDeviceAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
- return OpenCLGlobalHostAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OffloadGlobalDeviceAddressSpace:
+ return OffloadGlobalDeviceAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OffloadGlobalHostAddressSpace:
+ return OffloadGlobalHostAddressSpaceAttr::getLangAS(*this);
case ParsedAttr::AT_OffloadLocalAddressSpace:
return OffloadLocalAddressSpaceAttr::getLangAS(*this);
case ParsedAttr::AT_OffloadPrivateAddressSpace:
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 326910b564df1..1fdf5c1a183c0 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -9096,8 +9096,8 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type,
break;
case ParsedAttr::AT_OffloadPrivateAddressSpace:
case ParsedAttr::AT_OffloadGlobalAddressSpace:
- case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
- case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
+ case ParsedAttr::AT_OffloadGlobalDeviceAddressSpace:
+ case ParsedAttr::AT_OffloadGlobalHostAddressSpace:
case ParsedAttr::AT_OffloadLocalAddressSpace:
case ParsedAttr::AT_OffloadConstantAddressSpace:
case ParsedAttr::AT_OffloadGenericAddressSpace:
diff --git a/clang/test/SemaSYCL/address-space-conversions.cpp b/clang/test/SemaSYCL/address-space-conversions.cpp
index 0112ccae4c775..0b0ec9fe2f09b 100644
--- a/clang/test/SemaSYCL/address-space-conversions.cpp
+++ b/clang/test/SemaSYCL/address-space-conversions.cpp
@@ -70,12 +70,12 @@ void usages() {
bar(*GLOB_HOST);
bar2(*GLOB_HOST);
GLOB = GLOB_HOST;
- GLOB_HOST = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_host int *' changes address space of pointer}}
- GLOB_HOST = static_cast<__attribute__((sycl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_host int *' is not allowed}}
+ GLOB_HOST = GLOB; // expected-error {{assigning 'sycl_global int *' to 'sycl_global_host int *' changes address space of pointer}}
+ GLOB_HOST = static_cast<__attribute__((sycl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to 'sycl_global_host int *' is not allowed}}
__attribute__((sycl_global_device)) int *GLOB_DEVICE;
bar(*GLOB_DEVICE);
bar2(*GLOB_DEVICE);
GLOB = GLOB_DEVICE;
- GLOB_DEVICE = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_device int *' changes address space of pointer}}
- GLOB_DEVICE = static_cast<__attribute__((sycl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_device int *' is not allowed}}
+ GLOB_DEVICE = GLOB; // expected-error {{assigning 'sycl_global int *' to 'sycl_global_device int *' changes address space of pointer}}
+ GLOB_DEVICE = static_cast<__attribute__((sycl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to 'sycl_global_device int *' is not allowed}}
}
>From dab659e3fd360cb97ff5d422e2e80627ecb9c505 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews at intel.com>
Date: Mon, 8 Jun 2026 07:31:58 -0700
Subject: [PATCH 4/7] Revert "Update global_host and global_device attributes
as well for uniformity."
This reverts commit 071ba770ef59dd5c5c8338d2dce360a5cf925feb.
---
clang/include/clang/Basic/Attr.td | 12 ++++++------
clang/include/clang/Basic/AttrDocs.td | 4 ++--
clang/lib/AST/TypePrinter.cpp | 12 ++++++------
clang/lib/Sema/ParsedAttr.cpp | 12 ++++++------
clang/lib/Sema/SemaType.cpp | 4 ++--
clang/test/SemaSYCL/address-space-conversions.cpp | 8 ++++----
6 files changed, 26 insertions(+), 26 deletions(-)
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 00fb4dc8e24ee..86ba95eef4d7c 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1823,10 +1823,10 @@ def OffloadGlobalAddressSpace : TypeAttr {
}];
}
-// TODO: Remove OffloadGlobalDeviceAddressSpace after deprecation.
-def OffloadGlobalDeviceAddressSpace : TypeAttr {
+// TODO: Remove OpenCLGlobalDeviceAddressSpace after deprecation.
+def OpenCLGlobalDeviceAddressSpace : TypeAttr {
let Spellings = [Clang<"opencl_global_device">, Clang<"sycl_global_device">];
- let Documentation = [OffloadAddressSpaceGlobalExtDocs];
+ let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
let AdditionalMembers = [{
static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
return A.getAttributeSpellingListIndex() == GNU_sycl_global_device ||
@@ -1839,10 +1839,10 @@ def OffloadGlobalDeviceAddressSpace : TypeAttr {
}];
}
-// TODO: Remove OffloadGlobalHostAddressSpace after deprecation.
-def OffloadGlobalHostAddressSpace : TypeAttr {
+// TODO: Remove OpenCLGlobalHostAddressSpace after deprecation.
+def OpenCLGlobalHostAddressSpace : TypeAttr {
let Spellings = [Clang<"opencl_global_host">, Clang<"sycl_global_host">];
- let Documentation = [OffloadAddressSpaceGlobalExtDocs];
+ let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
let AdditionalMembers = [{
static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
return A.getAttributeSpellingListIndex() == GNU_sycl_global_host ||
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index d8e009121df44..74a6d7edaa4a2 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -5090,9 +5090,9 @@ local variables. Starting with OpenCL v2.0, the global address space can be used
}];
}
-def OffloadAddressSpaceGlobalExtDocs : Documentation {
+def OpenCLAddressSpaceGlobalExtDocs : Documentation {
let Category = DocOffloadAddressSpaces;
- let Heading = "[[clang::opencl_global_device]], [[clang::opencl_global_host]], [[clang::sycl_global_device]], [[clang::sycl_global_host]]";
+ let Heading = "[[clang::opencl_global_device]], [[clang::opencl_global_host]]";
let Content = [{
The ``global_device`` and ``global_host`` address space attributes specify that
an object is allocated in global memory on the device/host. It helps to
diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp
index ace0963158b84..ed5d95c360303 100644
--- a/clang/lib/AST/TypePrinter.cpp
+++ b/clang/lib/AST/TypePrinter.cpp
@@ -1994,8 +1994,8 @@ void TypePrinter::printAttributedAfter(const AttributedType *T,
case attr::OffloadPrivateAddressSpace:
case attr::OffloadGlobalAddressSpace:
- case attr::OffloadGlobalDeviceAddressSpace:
- case attr::OffloadGlobalHostAddressSpace:
+ case attr::OpenCLGlobalDeviceAddressSpace:
+ case attr::OpenCLGlobalHostAddressSpace:
case attr::OffloadLocalAddressSpace:
case attr::OffloadConstantAddressSpace:
case attr::OffloadGenericAddressSpace:
@@ -2676,9 +2676,13 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
return "__constant";
case LangAS::opencl_generic:
return "__generic";
+ // TODO: Remove *_global_device and *_global_host after corresponding
+ // attributes are deprecated for the required time.
case LangAS::opencl_global_device:
+ case LangAS::sycl_global_device:
return "__global_device";
case LangAS::opencl_global_host:
+ case LangAS::sycl_global_host:
return "__global_host";
case LangAS::sycl_global:
return "sycl_global";
@@ -2688,10 +2692,6 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
return "sycl_private";
case LangAS::sycl_generic:
return "sycl_generic";
- case LangAS::sycl_global_device:
- return "sycl_global_device";
- case LangAS::sycl_global_host:
- return "sycl_global_host";
case LangAS::cuda_device:
return "__device__";
case LangAS::cuda_constant:
diff --git a/clang/lib/Sema/ParsedAttr.cpp b/clang/lib/Sema/ParsedAttr.cpp
index be60e19819431..08b9aea3bdb15 100644
--- a/clang/lib/Sema/ParsedAttr.cpp
+++ b/clang/lib/Sema/ParsedAttr.cpp
@@ -231,8 +231,8 @@ bool ParsedAttr::slidesFromDeclToDeclSpecLegacyBehavior() const {
case AT_AddressSpace:
case AT_OffloadPrivateAddressSpace:
case AT_OffloadGlobalAddressSpace:
- case AT_OffloadGlobalDeviceAddressSpace:
- case AT_OffloadGlobalHostAddressSpace:
+ case AT_OpenCLGlobalDeviceAddressSpace:
+ case AT_OpenCLGlobalHostAddressSpace:
case AT_OffloadLocalAddressSpace:
case AT_OffloadConstantAddressSpace:
case AT_OffloadGenericAddressSpace:
@@ -318,10 +318,10 @@ LangAS ParsedAttr::asLangAS() const {
switch (getParsedKind()) {
case ParsedAttr::AT_OffloadGlobalAddressSpace:
return OffloadGlobalAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OffloadGlobalDeviceAddressSpace:
- return OffloadGlobalDeviceAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OffloadGlobalHostAddressSpace:
- return OffloadGlobalHostAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
+ return OpenCLGlobalDeviceAddressSpaceAttr::getLangAS(*this);
+ case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
+ return OpenCLGlobalHostAddressSpaceAttr::getLangAS(*this);
case ParsedAttr::AT_OffloadLocalAddressSpace:
return OffloadLocalAddressSpaceAttr::getLangAS(*this);
case ParsedAttr::AT_OffloadPrivateAddressSpace:
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 1fdf5c1a183c0..326910b564df1 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -9096,8 +9096,8 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type,
break;
case ParsedAttr::AT_OffloadPrivateAddressSpace:
case ParsedAttr::AT_OffloadGlobalAddressSpace:
- case ParsedAttr::AT_OffloadGlobalDeviceAddressSpace:
- case ParsedAttr::AT_OffloadGlobalHostAddressSpace:
+ case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
+ case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
case ParsedAttr::AT_OffloadLocalAddressSpace:
case ParsedAttr::AT_OffloadConstantAddressSpace:
case ParsedAttr::AT_OffloadGenericAddressSpace:
diff --git a/clang/test/SemaSYCL/address-space-conversions.cpp b/clang/test/SemaSYCL/address-space-conversions.cpp
index 0b0ec9fe2f09b..0112ccae4c775 100644
--- a/clang/test/SemaSYCL/address-space-conversions.cpp
+++ b/clang/test/SemaSYCL/address-space-conversions.cpp
@@ -70,12 +70,12 @@ void usages() {
bar(*GLOB_HOST);
bar2(*GLOB_HOST);
GLOB = GLOB_HOST;
- GLOB_HOST = GLOB; // expected-error {{assigning 'sycl_global int *' to 'sycl_global_host int *' changes address space of pointer}}
- GLOB_HOST = static_cast<__attribute__((sycl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to 'sycl_global_host int *' is not allowed}}
+ GLOB_HOST = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_host int *' changes address space of pointer}}
+ GLOB_HOST = static_cast<__attribute__((sycl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_host int *' is not allowed}}
__attribute__((sycl_global_device)) int *GLOB_DEVICE;
bar(*GLOB_DEVICE);
bar2(*GLOB_DEVICE);
GLOB = GLOB_DEVICE;
- GLOB_DEVICE = GLOB; // expected-error {{assigning 'sycl_global int *' to 'sycl_global_device int *' changes address space of pointer}}
- GLOB_DEVICE = static_cast<__attribute__((sycl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to 'sycl_global_device int *' is not allowed}}
+ GLOB_DEVICE = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_device int *' changes address space of pointer}}
+ GLOB_DEVICE = static_cast<__attribute__((sycl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_device int *' is not allowed}}
}
>From 3461ac6844a5e2e1f4e5c30910a32d691ebecff7 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews at intel.com>
Date: Mon, 8 Jun 2026 07:32:10 -0700
Subject: [PATCH 5/7] Revert "Choose address space mapping based on spelling"
This reverts commit bfe1c766e0e5a35afb834f3b8f2b787ffc261cb5.
---
clang/include/clang/Basic/Attr.td | 74 +------------------
clang/include/clang/Sema/ParsedAttr.h | 56 +++++++++++++-
clang/lib/Parse/ParseDecl.cpp | 2 +-
clang/lib/Sema/ParsedAttr.cpp | 24 ------
clang/lib/Sema/SemaType.cpp | 17 +++--
.../Builtins/generic_cast_to_ptr_explicit.c | 12 +--
.../CodeGenSYCL/address-space-conversions.cpp | 6 +-
.../amd-address-space-conversions.cpp | 24 +++---
.../cuda-address-space-conversions.cpp | 24 +++---
.../SemaSYCL/address-space-conversions.cpp | 8 +-
.../address-space-opencl-sycl-compat.cpp | 13 ++--
11 files changed, 107 insertions(+), 153 deletions(-)
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 86ba95eef4d7c..aa0fb8d913322 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1795,112 +1795,42 @@ def OffloadPrivateAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__private">, CustomKeyword<"private">,
Clang<"opencl_private">, Clang<"sycl_private">];
let Documentation = [OffloadAddressSpacePrivateDocs];
- let AdditionalMembers = [{
- static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
- return A.getAttributeSpellingListIndex() == GNU_sycl_private ||
- A.getAttributeSpellingListIndex() == CXX11_clang_sycl_private ||
- A.getAttributeSpellingListIndex() == C23_clang_sycl_private;
- }
- static inline LangAS getLangAS(const AttributeCommonInfo& A) {
- return isSYCLSpelling(A) ? LangAS::sycl_private : LangAS::opencl_private;
- }
- }];
}
def OffloadGlobalAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__global">, CustomKeyword<"global">,
Clang<"opencl_global">, Clang<"sycl_global">];
let Documentation = [OffloadAddressSpaceGlobalDocs];
- let AdditionalMembers = [{
- static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
- return A.getAttributeSpellingListIndex() == GNU_sycl_global ||
- A.getAttributeSpellingListIndex() == CXX11_clang_sycl_global ||
- A.getAttributeSpellingListIndex() == C23_clang_sycl_global;
- }
- static inline LangAS getLangAS(const AttributeCommonInfo& A) {
- return isSYCLSpelling(A) ? LangAS::sycl_global : LangAS::opencl_global;
- }
- }];
}
// TODO: Remove OpenCLGlobalDeviceAddressSpace after deprecation.
def OpenCLGlobalDeviceAddressSpace : TypeAttr {
- let Spellings = [Clang<"opencl_global_device">, Clang<"sycl_global_device">];
+ let Spellings = [Clang<"opencl_global_device">];
let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
- let AdditionalMembers = [{
- static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
- return A.getAttributeSpellingListIndex() == GNU_sycl_global_device ||
- A.getAttributeSpellingListIndex() == CXX11_clang_sycl_global_device ||
- A.getAttributeSpellingListIndex() == C23_clang_sycl_global_device;
- }
- static inline LangAS getLangAS(const AttributeCommonInfo& A) {
- return isSYCLSpelling(A) ? LangAS::sycl_global_device : LangAS::opencl_global_device;
- }
- }];
}
// TODO: Remove OpenCLGlobalHostAddressSpace after deprecation.
def OpenCLGlobalHostAddressSpace : TypeAttr {
- let Spellings = [Clang<"opencl_global_host">, Clang<"sycl_global_host">];
+ let Spellings = [Clang<"opencl_global_host">];
let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
- let AdditionalMembers = [{
- static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
- return A.getAttributeSpellingListIndex() == GNU_sycl_global_host ||
- A.getAttributeSpellingListIndex() == CXX11_clang_sycl_global_host ||
- A.getAttributeSpellingListIndex() == C23_clang_sycl_global_host;
- }
- static inline LangAS getLangAS(const AttributeCommonInfo& A) {
- return isSYCLSpelling(A) ? LangAS::sycl_global_host : LangAS::opencl_global_host;
- }
- }];
}
def OffloadLocalAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__local">, CustomKeyword<"local">,
Clang<"opencl_local">, Clang<"sycl_local">];
let Documentation = [OffloadAddressSpaceLocalDocs];
- let AdditionalMembers = [{
- static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
- return A.getAttributeSpellingListIndex() == GNU_sycl_local ||
- A.getAttributeSpellingListIndex() == CXX11_clang_sycl_local ||
- A.getAttributeSpellingListIndex() == C23_clang_sycl_local;
- }
- static inline LangAS getLangAS(const AttributeCommonInfo& A) {
- return isSYCLSpelling(A) ? LangAS::sycl_local : LangAS::opencl_local;
- }
- }];
}
def OffloadConstantAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__constant">, CustomKeyword<"constant">,
Clang<"opencl_constant">, Clang<"sycl_constant">];
let Documentation = [OffloadAddressSpaceConstantDocs];
- let AdditionalMembers = [{
- static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
- return A.getAttributeSpellingListIndex() == GNU_sycl_constant ||
- A.getAttributeSpellingListIndex() == CXX11_clang_sycl_constant ||
- A.getAttributeSpellingListIndex() == C23_clang_sycl_constant;
- }
- static inline LangAS getLangAS(const AttributeCommonInfo& A) {
- return isSYCLSpelling(A) ? LangAS::Default : LangAS::opencl_constant;
- }
- }];
}
def OffloadGenericAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__generic">, CustomKeyword<"generic">,
Clang<"opencl_generic">, Clang<"sycl_generic">];
let Documentation = [OffloadAddressSpaceGenericDocs];
- let AdditionalMembers = [{
- static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
- return A.getAttributeSpellingListIndex() == GNU_sycl_generic ||
- A.getAttributeSpellingListIndex() == CXX11_clang_sycl_generic ||
- A.getAttributeSpellingListIndex() == C23_clang_sycl_generic;
- }
- static inline LangAS getLangAS(const AttributeCommonInfo& A) {
- return isSYCLSpelling(A) ? LangAS::sycl_generic : LangAS::opencl_generic;
- }
- }];
}
def OpenCLNoSVM : Attr {
diff --git a/clang/include/clang/Sema/ParsedAttr.h b/clang/include/clang/Sema/ParsedAttr.h
index ddb26e89bdaa6..9251c8aafdc71 100644
--- a/clang/include/clang/Sema/ParsedAttr.h
+++ b/clang/include/clang/Sema/ParsedAttr.h
@@ -553,8 +553,60 @@ class ParsedAttr final
/// a Spelling enumeration, the value UINT_MAX is returned.
unsigned getSemanticSpelling() const;
- /// Returns the appropriate LangAS for this address space attribute.
- LangAS asLangAS() const;
+ /// If this is a named address space attribute for OpenCL compilation, returns its
+ /// representation in LangAS, otherwise returns default address space.
+ LangAS asOpenCLLangAS() const {
+ switch (getParsedKind()) {
+ case ParsedAttr::AT_OffloadConstantAddressSpace:
+ return LangAS::opencl_constant;
+ case ParsedAttr::AT_OffloadGlobalAddressSpace:
+ return LangAS::opencl_global;
+ case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
+ return LangAS::opencl_global_device;
+ case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
+ return LangAS::opencl_global_host;
+ case ParsedAttr::AT_OffloadLocalAddressSpace:
+ return LangAS::opencl_local;
+ case ParsedAttr::AT_OffloadPrivateAddressSpace:
+ return LangAS::opencl_private;
+ case ParsedAttr::AT_OffloadGenericAddressSpace:
+ return LangAS::opencl_generic;
+ default:
+ return LangAS::Default;
+ }
+ }
+
+ /// If this is a named address space attribute for SYCL compilation, returns its
+ /// representation in LangAS, otherwise returns default address space.
+ LangAS asSYCLLangAS() const {
+ switch (getKind()) {
+ case ParsedAttr::AT_OffloadGlobalAddressSpace:
+ return LangAS::sycl_global;
+ case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
+ return LangAS::sycl_global_device;
+ case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
+ return LangAS::sycl_global_host;
+ case ParsedAttr::AT_OffloadLocalAddressSpace:
+ return LangAS::sycl_local;
+ case ParsedAttr::AT_OffloadPrivateAddressSpace:
+ return LangAS::sycl_private;
+ case ParsedAttr::AT_OffloadGenericAddressSpace:
+ return LangAS::sycl_generic;
+ default:
+ return LangAS::Default;
+ }
+ }
+
+ /// If this is an HLSL address space attribute, returns its representation
+ /// in LangAS, otherwise returns default address space.
+ LangAS asHLSLLangAS() const {
+ switch (getParsedKind()) {
+ case ParsedAttr::AT_HLSLGroupSharedAddressSpace:
+ return LangAS::hlsl_groupshared;
+ default:
+ return LangAS::Default;
+ }
+ }
AttributeCommonInfo::Kind getKind() const {
return AttributeCommonInfo::Kind(Info.AttrKind);
diff --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp
index 1a4b92ca99b99..e2ac86bc5e064 100644
--- a/clang/lib/Parse/ParseDecl.cpp
+++ b/clang/lib/Parse/ParseDecl.cpp
@@ -7227,7 +7227,7 @@ void Parser::InitCXXThisScopeForDeclaratorIfRelevant(
// prototype for the method.
if (getLangOpts().OpenCLCPlusPlus) {
for (ParsedAttr &attr : DS.getAttributes()) {
- LangAS ASIdx = attr.asLangAS();
+ LangAS ASIdx = attr.asOpenCLLangAS();
if (ASIdx != LangAS::Default) {
Q.addAddressSpace(ASIdx);
break;
diff --git a/clang/lib/Sema/ParsedAttr.cpp b/clang/lib/Sema/ParsedAttr.cpp
index 08b9aea3bdb15..49dec6188d877 100644
--- a/clang/lib/Sema/ParsedAttr.cpp
+++ b/clang/lib/Sema/ParsedAttr.cpp
@@ -12,7 +12,6 @@
#include "clang/Sema/ParsedAttr.h"
#include "clang/AST/ASTContext.h"
-#include "clang/AST/Attr.h"
#include "clang/Basic/AttrSubjectMatchRules.h"
#include "clang/Basic/IdentifierTable.h"
#include "clang/Basic/TargetInfo.h"
@@ -313,26 +312,3 @@ void clang::takeAndConcatenateAttrs(ParsedAttributes &First,
if (Second.Range.getEnd().isValid())
First.Range.setEnd(Second.Range.getEnd());
}
-
-LangAS ParsedAttr::asLangAS() const {
- switch (getParsedKind()) {
- case ParsedAttr::AT_OffloadGlobalAddressSpace:
- return OffloadGlobalAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
- return OpenCLGlobalDeviceAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
- return OpenCLGlobalHostAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OffloadLocalAddressSpace:
- return OffloadLocalAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OffloadPrivateAddressSpace:
- return OffloadPrivateAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OffloadConstantAddressSpace:
- return OffloadConstantAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_OffloadGenericAddressSpace:
- return OffloadGenericAddressSpaceAttr::getLangAS(*this);
- case ParsedAttr::AT_HLSLGroupSharedAddressSpace:
- return LangAS::hlsl_groupshared;
- default:
- return LangAS::Default;
- }
-}
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 326910b564df1..1d13e632c51d6 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -5407,7 +5407,7 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
// them later while creating QualType.
if (FTI.MethodQualifiers)
for (ParsedAttr &attr : FTI.MethodQualifiers->getAttributes()) {
- LangAS ASIdxNew = attr.asLangAS();
+ LangAS ASIdxNew = attr.asOpenCLLangAS();
if (DiagnoseMultipleAddrSpaceAttributes(S, ASIdx, ASIdxNew,
attr.getLoc()))
D.setInvalidType(true);
@@ -6678,12 +6678,15 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
else
Attr.setInvalid();
} else {
- // Type attributes imply which address space to use.
- ASIdx = Attr.asLangAS();
-
- if (ASIdx == LangAS::Default &&
- Attr.getKind() == ParsedAttr::AT_OffloadConstantAddressSpace) {
- if (OffloadConstantAddressSpaceAttr::isSYCLSpelling(Attr))
+ // The keyword-based type attributes imply which address space to use.
+ ASIdx = S.getLangOpts().SYCLIsDevice ? Attr.asSYCLLangAS()
+ : Attr.asOpenCLLangAS();
+ if (S.getLangOpts().HLSL)
+ ASIdx = Attr.asHLSLLangAS();
+
+ if (ASIdx == LangAS::Default) {
+ if (S.getLangOpts().SYCLIsDevice &&
+ Attr.getKind() == ParsedAttr::AT_OffloadConstantAddressSpace)
S.Diag(Attr.getLoc(), diag::warn_deprecated_sycl_constant);
else
llvm_unreachable("Invalid address space");
diff --git a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
index b896d76897fbf..30f4ecb589a5c 100644
--- a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
+++ b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
@@ -4,14 +4,8 @@
#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_EXTERNAL [[clang::sycl_external]]
-#define __global __attribute__((sycl_global))
-#define __local __attribute__((sycl_local))
-#define __private __attribute__((sycl_private))
#else
#define SYCL_EXTERNAL
-#define __global __attribute__((opencl_global))
-#define __local __attribute__((opencl_local))
-#define __private __attribute__((opencl_private))
#endif
// CHECK: spir_func noundef ptr @{{.*}}test_cast_to_private{{.*}}(ptr addrspace(4) noundef readnone [[P:%.*]]
@@ -19,7 +13,7 @@
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr @llvm.spv.generic.cast.to.ptr.explicit.p0(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr [[SPV_CAST]]
//
-SYCL_EXTERNAL __private int* test_cast_to_private(int* p) {
+SYCL_EXTERNAL __attribute__((opencl_private)) int* test_cast_to_private(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 7);
}
@@ -28,7 +22,7 @@ SYCL_EXTERNAL __private int* test_cast_to_private(int* p) {
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(1) @llvm.spv.generic.cast.to.ptr.explicit.p1(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr addrspace(1) [[SPV_CAST]]
//
-SYCL_EXTERNAL __global int* test_cast_to_global(int* p) {
+SYCL_EXTERNAL __attribute__((opencl_global)) int* test_cast_to_global(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 5);
}
@@ -37,6 +31,6 @@ SYCL_EXTERNAL __global int* test_cast_to_global(int* p) {
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(3) @llvm.spv.generic.cast.to.ptr.explicit.p3(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr addrspace(3) [[SPV_CAST]]
//
-SYCL_EXTERNAL __local int* test_cast_to_local(int* p) {
+SYCL_EXTERNAL __attribute__((opencl_local)) int* test_cast_to_local(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 4);
}
diff --git a/clang/test/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CodeGenSYCL/address-space-conversions.cpp
index f331b8367b614..3eecdded06364 100644
--- a/clang/test/CodeGenSYCL/address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/address-space-conversions.cpp
@@ -25,9 +25,9 @@ void tmpl(T t) {}
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3)
__attribute__((sycl_private)) int *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr
- __attribute__((sycl_global_device)) int *GLOBDEVICE;
+ __attribute__((opencl_global_device)) int *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5)
- __attribute__((sycl_global_host)) int *GLOBHOST;
+ __attribute__((opencl_global_host)) int *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(6)
// CHECK-DAG: [[NoAS]].ascast = addrspacecast ptr [[NoAS]] to ptr addrspace(4)
@@ -66,7 +66,7 @@ void tmpl(T t) {}
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr
// CHECK-DAG: store ptr [[NoAS_CAST]], ptr addrspace(4) [[PRIV]].ascast
PRIV = (__attribute__((sycl_private)) int *)NoAS;
- // From sycl_global_[host/device] address spaces to sycl_global
+ // From opencl_global_[host/device] address spaces to sycl_global
// CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr addrspace(4) [[GLOB_DEVICE]].ascast
// CHECK-DAG: [[GLOBDEVICE_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[GLOBDEVICE_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_CAST]], ptr addrspace(4) [[GLOB]].ascast
diff --git a/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp b/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
index a46b9660b2ef9..17a98195318ad 100644
--- a/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
@@ -3,13 +3,13 @@ void bar(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
void bar2(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
-void bar(__attribute__((sycl_local)) int &Data) {}
+void bar(__attribute__((opencl_local)) int &Data) {}
// CHECK-DAG: define {{.*}} void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) %
void foo(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef %
void foo2(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef %
-void foo(__attribute__((sycl_local)) int *Data) {}
+void foo(__attribute__((opencl_local)) int *Data) {}
// CHECK-DAG: define {{.*}} void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef %
template <typename T>
@@ -19,15 +19,15 @@ void tmpl(T t) {}
[[clang::sycl_external]] void usages() {
int *NoAS;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8, addrspace(5)
- __attribute__((sycl_global)) int *GLOB;
+ __attribute__((opencl_global)) int *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
- __attribute__((sycl_local)) int *LOC;
+ __attribute__((opencl_local)) int *LOC;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 4, addrspace(5)
- __attribute__((sycl_private)) int *PRIV;
+ __attribute__((opencl_private)) int *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5), align 4, addrspace(5)
- __attribute__((sycl_global_device)) int *GLOBDEVICE;
+ __attribute__((opencl_global_device)) int *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
- __attribute__((sycl_global_host)) int *GLOBHOST;
+ __attribute__((opencl_global_host)) int *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
LOC = nullptr;
// CHECK-DAG: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]].ascast, align 4
@@ -45,22 +45,22 @@ void tmpl(T t) {}
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[NoAS_LOAD]] to ptr
// CHECK-DAG: store ptr %5, ptr [[NoAS]].ascast, align 8
- GLOB = (__attribute__((sycl_global)) int *)NoAS;
+ GLOB = (__attribute__((opencl_global)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr %6 to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) %7, ptr [[GLOB]].ascast, align 8
- LOC = (__attribute__((sycl_local)) int *)NoAS;
+ LOC = (__attribute__((opencl_local)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3)
// CHECK-DAG: store ptr addrspace(3) %9, ptr [[LOC]].ascast, align 4
- PRIV = (__attribute__((sycl_private)) int *)NoAS;
+ PRIV = (__attribute__((opencl_private)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(5)
// CHECK-DAG: store ptr addrspace(5) [[NoAS_CAST]], ptr [[PRIV]].ascast, align 4
- GLOB = (__attribute__((sycl_global)) int *)GLOBDEVICE;
+ GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]DEVICE.ascast, align 8
// CHECK-DAG: store ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8
- GLOB = (__attribute__((sycl_global)) int *)GLOBHOST;
+ GLOB = (__attribute__((opencl_global)) int *)GLOBHOST;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]HOST.ascast, align 8
// CHECK-DAG: tore ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8
bar(*GLOB);
diff --git a/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp b/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
index 3427450547fce..ffb601e62c118 100644
--- a/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
@@ -3,13 +3,13 @@ void bar(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
void bar2(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
-void bar(__attribute__((sycl_local)) int &Data) {}
+void bar(__attribute__((opencl_local)) int &Data) {}
// CHECK-DAG: define {{.*}} void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) %
void foo(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef %
void foo2(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef %
-void foo(__attribute__((sycl_local)) int *Data) {}
+void foo(__attribute__((opencl_local)) int *Data) {}
// CHECK-DAG: define {{.*}} void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef %
template <typename T>
@@ -19,15 +19,15 @@ void tmpl(T t);
[[clang::sycl_external]] void usages() {
int *NoAS;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8
- __attribute__((sycl_global)) int *GLOB;
+ __attribute__((opencl_global)) int *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
- __attribute__((sycl_local)) int *LOC;
+ __attribute__((opencl_local)) int *LOC;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 8
- __attribute__((sycl_private)) int *PRIV;
+ __attribute__((opencl_private)) int *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr, align 8
- __attribute__((sycl_global_device)) int *GLOBDEVICE;
+ __attribute__((opencl_global_device)) int *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
- __attribute__((sycl_global_host)) int *GLOBHOST;
+ __attribute__((opencl_global_host)) int *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
LOC = nullptr;
// CHECK-DAG: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]], align 8
@@ -44,21 +44,21 @@ void tmpl(T t);
NoAS = (int *)PRIV;
// CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[PRIV]], align 8
// CHECK-DAG: store ptr [[LOC_LOAD]], ptr [[NoAS]], align 8
- GLOB = (__attribute__((sycl_global)) int *)NoAS;
+ GLOB = (__attribute__((opencl_global)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[NoAS_CAST]], ptr [[GLOB]], align 8
- LOC = (__attribute__((sycl_local)) int *)NoAS;
+ LOC = (__attribute__((opencl_local)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3)
// CHECK-DAG: store ptr addrspace(3) [[NoAS_CAST]], ptr [[LOC]], align 8
- PRIV = (__attribute__((sycl_private)) int *)NoAS;
+ PRIV = (__attribute__((opencl_private)) int *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: store ptr [[NoAS_LOAD]], ptr [[PRIV]], align 8
- GLOB = (__attribute__((sycl_global)) int *)GLOBDEVICE;
+ GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE;
// CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_DEVICE]], align 8
// CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_LOAD]], ptr %GLOB, align 8
- GLOB = (__attribute__((sycl_global)) int *)GLOBHOST;
+ GLOB = (__attribute__((opencl_global)) int *)GLOBHOST;
// CHECK-DAG: [[GLOB_HOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_HOST]], align 8
// CHECK-DAG: store ptr addrspace(1) [[GLOB_HOST_LOAD]], ptr [[GLOB]], align 8
bar(*GLOB);
diff --git a/clang/test/SemaSYCL/address-space-conversions.cpp b/clang/test/SemaSYCL/address-space-conversions.cpp
index 0112ccae4c775..9d209dbe5f8d7 100644
--- a/clang/test/SemaSYCL/address-space-conversions.cpp
+++ b/clang/test/SemaSYCL/address-space-conversions.cpp
@@ -66,16 +66,16 @@ void usages() {
(void)i;
(void)v;
- __attribute__((sycl_global_host)) int *GLOB_HOST;
+ __attribute__((opencl_global_host)) int *GLOB_HOST;
bar(*GLOB_HOST);
bar2(*GLOB_HOST);
GLOB = GLOB_HOST;
GLOB_HOST = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_host int *' changes address space of pointer}}
- GLOB_HOST = static_cast<__attribute__((sycl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_host int *' is not allowed}}
- __attribute__((sycl_global_device)) int *GLOB_DEVICE;
+ GLOB_HOST = static_cast<__attribute__((opencl_global_host)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_host int *' is not allowed}}
+ __attribute__((opencl_global_device)) int *GLOB_DEVICE;
bar(*GLOB_DEVICE);
bar2(*GLOB_DEVICE);
GLOB = GLOB_DEVICE;
GLOB_DEVICE = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_device int *' changes address space of pointer}}
- GLOB_DEVICE = static_cast<__attribute__((sycl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_device int *' is not allowed}}
+ GLOB_DEVICE = static_cast<__attribute__((opencl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_device int *' is not allowed}}
}
diff --git a/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp b/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
index 7b4575e60c8ad..89c4fa4873086 100644
--- a/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
+++ b/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
@@ -10,9 +10,9 @@ void test_incompatible() {
// Address space attributes are resolved using mode of compilation and not the spelling itself. This results in the SYCL spelling
// being used in both instances of each diagnostic despite openCL spelling being used.
- opencl_global = sycl_local; // expected-error {{assigning 'sycl_local int *' to '__global int *' changes address space of pointer}}
- opencl_global = sycl_private; // expected-error {{assigning 'sycl_private int *' to '__global int *' changes address space of pointer}}
- sycl_local = opencl_global; // expected-error {{assigning '__global int *' to 'sycl_local int *' changes address space of pointer}}
+ opencl_global = sycl_local; // expected-error {{assigning 'sycl_local int *' to 'sycl_global int *' changes address space of pointer}}
+ opencl_global = sycl_private; // expected-error {{assigning 'sycl_private int *' to 'sycl_global int *' changes address space of pointer}}
+ sycl_local = opencl_global; // expected-error {{assigning 'sycl_global int *' to 'sycl_local int *' changes address space of pointer}}
}
void test_to_generic_mixed() {
@@ -23,13 +23,12 @@ void test_to_generic_mixed() {
int [[clang::sycl_local]] *sycl_local;
int [[clang::sycl_private]] *sycl_private;
- //FIXME: Why don't these throw an error?
opencl_gen = sycl_local;
opencl_gen = sycl_private;
- sycl_gen = opencl_global; // expected-error {{assigning '__global int *' to 'sycl_generic int *' changes address space of pointer}}
+ sycl_gen = opencl_global;
}
-void overload_test(__attribute__((opencl_global)) int *p) { (void)p; }
-void overload_test(__attribute__((sycl_global)) int *p) { (void)p; }
+void overload_test(__attribute__((opencl_global)) int *p) { (void)p; } // expected-note {{previous definition is here}}
+void overload_test(__attribute__((sycl_global)) int *p) { (void)p; } // expected-error {{redefinition of 'overload_test'}}
>From 7613ed9cff0d306de2428fcd4075d9ce0ec11dde Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews at intel.com>
Date: Wed, 10 Jun 2026 08:28:13 -0700
Subject: [PATCH 6/7] Apply review comments to separate out SYCL attributes.
---
clang/include/clang/Basic/AddressSpaces.h | 10 +-
clang/include/clang/Basic/Attr.td | 60 ++++--
clang/include/clang/Basic/AttrDocs.td | 98 +++++----
clang/include/clang/Sema/ParsedAttr.h | 31 +--
clang/lib/AST/TypePrinter.cpp | 15 +-
clang/lib/Headers/__clang_spirv_builtins.h | 192 ++++++++----------
clang/lib/Sema/ParsedAttr.cpp | 10 +-
clang/lib/Sema/SemaType.cpp | 21 +-
.../Builtins/generic_cast_to_ptr_explicit.c | 12 +-
.../CodeGenSYCL/address-space-conversions.cpp | 24 +--
.../CodeGenSYCL/address-space-mangling.cpp | 12 +-
.../amd-address-space-conversions.cpp | 24 +--
.../cuda-address-space-conversions.cpp | 24 +--
.../BuiltIns/generic_cast_to_ptr_explicit.c | 10 +-
.../SemaSYCL/address-space-conversions.cpp | 21 +-
.../address-space-opencl-sycl-compat.cpp | 34 ----
16 files changed, 301 insertions(+), 297 deletions(-)
delete mode 100644 clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h
index 58b04e50c2e3d..c3654ebbab82c 100644
--- a/clang/include/clang/Basic/AddressSpaces.h
+++ b/clang/include/clang/Basic/AddressSpaces.h
@@ -36,8 +36,9 @@ enum class LangAS : unsigned {
opencl_constant,
opencl_private,
opencl_generic,
- // TODO: Remove opencl_global_device and opencl_global_host after corresponding
- // attributes are deprecated for the required time.
+ // TODO: Remove opencl_global_device and opencl_global_host after
+ // corresponding attributes are deprecated for the required time.
+ // https://discourse.llvm.org/t/rfc-remove-opencl-global-device-and-opencl-global-host-address-space-attributes/90677
opencl_global_device,
opencl_global_host,
@@ -48,8 +49,9 @@ enum class LangAS : unsigned {
// SYCL specific address spaces.
sycl_global,
- // TODO: Remove sycl_global_device and sycl_global_host after corresponding attributes
- // are deprecated for the required time.
+ // TODO: Remove sycl_global_device and sycl_global_host after corresponding
+ // attributes are deprecated for the required time.
+ // https://discourse.llvm.org/t/rfc-remove-opencl-global-device-and-opencl-global-host-address-space-attributes/90677
sycl_global_device,
sycl_global_host,
sycl_local,
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index aa0fb8d913322..2f68fe1f91cc7 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1791,16 +1791,16 @@ def OpenCLAccess : Attr {
let Documentation = [OpenCLAccessDocs];
}
-def OffloadPrivateAddressSpace : TypeAttr {
+def OpenCLPrivateAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__private">, CustomKeyword<"private">,
- Clang<"opencl_private">, Clang<"sycl_private">];
- let Documentation = [OffloadAddressSpacePrivateDocs];
+ Clang<"opencl_private">];
+ let Documentation = [OpenCLAddressSpacePrivateDocs];
}
-def OffloadGlobalAddressSpace : TypeAttr {
+def OpenCLGlobalAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__global">, CustomKeyword<"global">,
- Clang<"opencl_global">, Clang<"sycl_global">];
- let Documentation = [OffloadAddressSpaceGlobalDocs];
+ Clang<"opencl_global">];
+ let Documentation = [OpenCLAddressSpaceGlobalDocs];
}
// TODO: Remove OpenCLGlobalDeviceAddressSpace after deprecation.
@@ -1815,22 +1815,52 @@ def OpenCLGlobalHostAddressSpace : TypeAttr {
let Documentation = [OpenCLAddressSpaceGlobalExtDocs];
}
-def OffloadLocalAddressSpace : TypeAttr {
+def OpenCLLocalAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__local">, CustomKeyword<"local">,
- Clang<"opencl_local">, Clang<"sycl_local">];
- let Documentation = [OffloadAddressSpaceLocalDocs];
+ Clang<"opencl_local">];
+ let Documentation = [OpenCLAddressSpaceLocalDocs];
}
-def OffloadConstantAddressSpace : TypeAttr {
+def OpenCLConstantAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__constant">, CustomKeyword<"constant">,
- Clang<"opencl_constant">, Clang<"sycl_constant">];
- let Documentation = [OffloadAddressSpaceConstantDocs];
+ Clang<"opencl_constant">];
+ let Documentation = [OpenCLAddressSpaceConstantDocs];
}
-def OffloadGenericAddressSpace : TypeAttr {
+def OpenCLGenericAddressSpace : TypeAttr {
let Spellings = [CustomKeyword<"__generic">, CustomKeyword<"generic">,
- Clang<"opencl_generic">, Clang<"sycl_generic">];
- let Documentation = [OffloadAddressSpaceGenericDocs];
+ Clang<"opencl_generic">];
+ let Documentation = [OpenCLAddressSpaceGenericDocs];
+}
+
+def SYCLPrivateAddressSpace : TypeAttr {
+ let Spellings = [CXX11<"clang", "sycl_private">];
+ let LangOpts = [SYCLHost, SYCLDevice];
+ let Documentation = [SYCLAddressSpaceDocs];
+}
+
+def SYCLGlobalAddressSpace : TypeAttr {
+ let Spellings = [CXX11<"clang", "sycl_global">];
+ let LangOpts = [SYCLHost, SYCLDevice];
+ let Documentation = [SYCLAddressSpaceDocs];
+}
+
+def SYCLLocalAddressSpace : TypeAttr {
+ let Spellings = [CXX11<"clang", "sycl_local">];
+ let LangOpts = [SYCLHost, SYCLDevice];
+ let Documentation = [SYCLAddressSpaceDocs];
+}
+
+def SYCLConstantAddressSpace : TypeAttr {
+ let Spellings = [CXX11<"clang", "sycl_constant">];
+ let LangOpts = [SYCLHost, SYCLDevice];
+ let Documentation = [SYCLAddressSpaceDocs];
+}
+
+def SYCLGenericAddressSpace : TypeAttr {
+ let Spellings = [CXX11<"clang", "sycl_generic">];
+ let LangOpts = [SYCLHost, SYCLDevice];
+ let Documentation = [SYCLAddressSpaceDocs];
}
def OpenCLNoSVM : Attr {
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 74a6d7edaa4a2..b058cc3d3fdbd 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -786,6 +786,27 @@ expression.
}];
}
+def SYCLAddressSpaceDocs : Documentation {
+ let Category = DocCatType;
+ let Heading = "Address space attributes for SYCL";
+ let Content = [{
+The memory model for SYCL devices is derived from the OpenCL memory model. Accordingly
+SYCL defines five address spaces: global, local, private, generic and constant. The
+following attributes correspond to these address spaces:
+
+[[clang::sycl_global]], [[clang::sycl_local]], [[clang::sycl_private]],
+[[clang::sycl_generic]] and [[clang::sycl_constant]] (deprecated)
+
+These attributes are intended for use in the implementation of SYCL run-time
+libraries. A direct declaration of pointers with address spaces is discouraged. Users
+should use the sycl::multi_ptr class to handle address space boundaries and
+interoperability.
+
+More details can be found in SYCL 2020 Specification, Section 3.8.2
+"SYCL device memory model" and Section 4.7.7, "Address space classes"
+ }];
+}
+
def SYCLSpecialClassDocs : Documentation {
let Category = DocCatStmt;
let Content = [{
@@ -5003,20 +5024,13 @@ More details can be found in the OpenCL C language Spec v2.0, Section 6.6.
}];
}
-def DocOffloadAddressSpaces : DocumentationCategory<"OpenCL and SYCL Address Spaces"> {
+def DocOpenCLAddressSpaces : DocumentationCategory<"OpenCL Address Spaces"> {
let Content = [{
The address space qualifier may be used to specify the region of memory that is
-used to allocate the object.
-
-OpenCL supports the following address spaces:
-
-__generic(generic), __global(global), __local(local), __private(private) and
+used to allocate the object. OpenCL supports the following address spaces:
+__generic(generic), __global(global), __local(local), __private(private),
__constant(constant).
-More details can be found in the OpenCL C language Spec v2.0, Section 6.5.
-
-Example:
-
.. code-block:: c
__constant int c = ...;
@@ -5028,32 +5042,14 @@ Example:
return l;
}
-The memory model for SYCL devices is derived from the OpenCL memory model. Accordingly
-SYCL defines five address spaces: global, local, private, generic and constant. The
-following attributes correspond to these address spaces:
-
-[[clang::sycl_global]], [[clang::sycl_local]], [[clang::sycl_private]],
-[[clang::sycl_generic]] and [[clang::sycl_constant]] (deprecated)
-
-These attributes are intended for use in the implementation of SYCL run-time
-libraries. A direct declaration of pointers with address spaces is discouraged. Users
-should use the sycl::multi_ptr class to handle address space boundaries and
-interoperability.
-
-More details can be found in SYCL 2020 Specification, Section 3.8.2
-"SYCL device memory model" and Section 4.7.7, "Address space classes"
+More details can be found in the OpenCL C language Spec v2.0, Section 6.5.
}];
-
}
-def OffloadAddressSpaceGenericDocs : Documentation {
- let Category = DocOffloadAddressSpaces;
- let Heading = "__generic, generic, [[clang::opencl_generic]], [[clang::sycl_generic]]";
+def OpenCLAddressSpaceGenericDocs : Documentation {
+ let Category = DocOpenCLAddressSpaces;
+ let Heading = "__generic, generic, [[clang::opencl_generic]]";
let Content = [{
-The generic address space is a virtual address space which overlaps the global, local
-and private address spaces.
-
-OpenCL:
The generic address space attribute is only available with OpenCL v2.0 and later.
It can be used with pointer types. Variables in global and local scope and
function parameters in non-kernel functions can have the generic address space
@@ -5063,35 +5059,33 @@ spaces.
}];
}
-def OffloadAddressSpaceConstantDocs : Documentation {
- let Category = DocOffloadAddressSpaces;
- let Heading = "__constant, constant, [[clang::opencl_constant]], [[clang::sycl_constant]]";
+def OpenCLAddressSpaceConstantDocs : Documentation {
+ let Category = DocOpenCLAddressSpaces;
+ let Heading = "__constant, constant, [[clang::opencl_constant]]";
let Content = [{
The constant address space attribute signals that an object is located in
a constant (non-modifiable) memory region. It is available to all work items.
Any type can be annotated with the constant address space attribute. Objects
with the constant address space qualifier can be declared in any scope and must
-have an initializer. The constant address space is deprecated in SYCL 2020
-specification.
+have an initializer.
}];
}
-def OffloadAddressSpaceGlobalDocs : Documentation {
- let Category = DocOffloadAddressSpaces;
- let Heading = "__global, global, [[clang::opencl_global]], [[clang::sycl_global]]";
+def OpenCLAddressSpaceGlobalDocs : Documentation {
+ let Category = DocOpenCLAddressSpaces;
+ let Heading = "__global, global, [[clang::opencl_global]]";
let Content = [{
The global address space attribute specifies that an object is allocated in
global memory, which is accessible by all work items. The content stored in this
-memory area persists between kernel executions.
-
-In OpenCL, pointer types to the global address space are allowed as function parameters or
-local variables. Starting with OpenCL v2.0, the global address space can be used with global
-(program scope) variables and static local variable as well.
+memory area persists between kernel executions. Pointer types to the global
+address space are allowed as function parameters or local variables. Starting
+with OpenCL v2.0, the global address space can be used with global (program
+scope) variables and static local variable as well.
}];
}
def OpenCLAddressSpaceGlobalExtDocs : Documentation {
- let Category = DocOffloadAddressSpaces;
+ let Category = DocOpenCLAddressSpaces;
let Heading = "[[clang::opencl_global_device]], [[clang::opencl_global_host]]";
let Content = [{
The ``global_device`` and ``global_host`` address space attributes specify that
@@ -5115,9 +5109,9 @@ As ``global_device`` and ``global_host`` are a subset of
}];
}
-def OffloadAddressSpaceLocalDocs : Documentation {
- let Category = DocOffloadAddressSpaces;
- let Heading = "__local, local, [[clang::opencl_local]], [[clang::sycl_local]]";
+def OpenCLAddressSpaceLocalDocs : Documentation {
+ let Category = DocOpenCLAddressSpaces;
+ let Heading = "__local, local, [[clang::opencl_local]]";
let Content = [{
The local address space specifies that an object is allocated in the local (work
group) memory area, which is accessible to all work items in the same work
@@ -5128,9 +5122,9 @@ space are allowed. Local address space variables cannot have an initializer.
}];
}
-def OffloadAddressSpacePrivateDocs : Documentation {
- let Category = DocOffloadAddressSpaces;
- let Heading = "__private, private, [[clang::opencl_private]], [[clang::sycl_private]]";
+def OpenCLAddressSpacePrivateDocs : Documentation {
+ let Category = DocOpenCLAddressSpaces;
+ let Heading = "__private, private, [[clang::opencl_private]]";
let Content = [{
The private address space specifies that an object is allocated in the private
(work item) memory. Other work items cannot access the same memory area and its
diff --git a/clang/include/clang/Sema/ParsedAttr.h b/clang/include/clang/Sema/ParsedAttr.h
index 9251c8aafdc71..7dd1290000f7c 100644
--- a/clang/include/clang/Sema/ParsedAttr.h
+++ b/clang/include/clang/Sema/ParsedAttr.h
@@ -553,44 +553,47 @@ class ParsedAttr final
/// a Spelling enumeration, the value UINT_MAX is returned.
unsigned getSemanticSpelling() const;
- /// If this is a named address space attribute for OpenCL compilation, returns its
- /// representation in LangAS, otherwise returns default address space.
+ /// If this is an OpenCL address space attribute, returns its representation
+ /// in LangAS, otherwise returns default address space.
LangAS asOpenCLLangAS() const {
switch (getParsedKind()) {
- case ParsedAttr::AT_OffloadConstantAddressSpace:
+ case ParsedAttr::AT_OpenCLConstantAddressSpace:
return LangAS::opencl_constant;
- case ParsedAttr::AT_OffloadGlobalAddressSpace:
+ case ParsedAttr::AT_OpenCLGlobalAddressSpace:
return LangAS::opencl_global;
case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
return LangAS::opencl_global_device;
case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
return LangAS::opencl_global_host;
- case ParsedAttr::AT_OffloadLocalAddressSpace:
+ case ParsedAttr::AT_OpenCLLocalAddressSpace:
return LangAS::opencl_local;
- case ParsedAttr::AT_OffloadPrivateAddressSpace:
+ case ParsedAttr::AT_OpenCLPrivateAddressSpace:
return LangAS::opencl_private;
- case ParsedAttr::AT_OffloadGenericAddressSpace:
+ case ParsedAttr::AT_OpenCLGenericAddressSpace:
return LangAS::opencl_generic;
default:
return LangAS::Default;
}
}
- /// If this is a named address space attribute for SYCL compilation, returns its
- /// representation in LangAS, otherwise returns default address space.
+ /// If this is a SYCL address space attribute, returns its SYCL
+ /// representation in LangAS.
LangAS asSYCLLangAS() const {
- switch (getKind()) {
- case ParsedAttr::AT_OffloadGlobalAddressSpace:
+ switch (getParsedKind()) {
+ case ParsedAttr::AT_SYCLGlobalAddressSpace:
return LangAS::sycl_global;
+ // TODO: OpenCLGlobalDeviceAddressSpace and OpenCLGlobalHostAddressSpace
+ // will be removed after deprecation.
+ // https://discourse.llvm.org/t/rfc-remove-opencl-global-device-and-opencl-global-host-address-space-attributes/90677
case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
return LangAS::sycl_global_device;
case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
return LangAS::sycl_global_host;
- case ParsedAttr::AT_OffloadLocalAddressSpace:
+ case ParsedAttr::AT_SYCLLocalAddressSpace:
return LangAS::sycl_local;
- case ParsedAttr::AT_OffloadPrivateAddressSpace:
+ case ParsedAttr::AT_SYCLPrivateAddressSpace:
return LangAS::sycl_private;
- case ParsedAttr::AT_OffloadGenericAddressSpace:
+ case ParsedAttr::AT_SYCLGenericAddressSpace:
return LangAS::sycl_generic;
default:
return LangAS::Default;
diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp
index ed5d95c360303..1ac5566c31904 100644
--- a/clang/lib/AST/TypePrinter.cpp
+++ b/clang/lib/AST/TypePrinter.cpp
@@ -1992,14 +1992,19 @@ void TypePrinter::printAttributedAfter(const AttributedType *T,
case attr::HLSLResourceDimension:
llvm_unreachable("HLSL resource type attributes handled separately");
- case attr::OffloadPrivateAddressSpace:
- case attr::OffloadGlobalAddressSpace:
+ case attr::OpenCLPrivateAddressSpace:
+ case attr::OpenCLGlobalAddressSpace:
case attr::OpenCLGlobalDeviceAddressSpace:
case attr::OpenCLGlobalHostAddressSpace:
- case attr::OffloadLocalAddressSpace:
- case attr::OffloadConstantAddressSpace:
- case attr::OffloadGenericAddressSpace:
+ case attr::OpenCLLocalAddressSpace:
+ case attr::OpenCLConstantAddressSpace:
+ case attr::OpenCLGenericAddressSpace:
case attr::HLSLGroupSharedAddressSpace:
+ case attr::SYCLPrivateAddressSpace:
+ case attr::SYCLGlobalAddressSpace:
+ case attr::SYCLLocalAddressSpace:
+ case attr::SYCLConstantAddressSpace:
+ case attr::SYCLGenericAddressSpace:
// FIXME: Update printAttributedBefore to print these once we generate
// AttributedType nodes for them.
break;
diff --git a/clang/lib/Headers/__clang_spirv_builtins.h b/clang/lib/Headers/__clang_spirv_builtins.h
index 9b3ac8c134ab1..4a5af96fdf1e0 100644
--- a/clang/lib/Headers/__clang_spirv_builtins.h
+++ b/clang/lib/Headers/__clang_spirv_builtins.h
@@ -27,17 +27,15 @@
#define __SPIRV_inline __attribute__((always_inline))
#ifdef __SYCL_DEVICE_ONLY__
-#define __global __attribute__((sycl_global))
-#define __local __attribute__((sycl_local))
-#define __private __attribute__((sycl_private))
-#define __constant __attribute__((sycl_constant))
-#define __generic __attribute__((sycl_generic))
+#define __GLOBALAS [[clang::sycl_global]]
+#define __LOCALAS [[clang::sycl_local]]
+#define __PRIVATEAS [[clang::sycl_private]]
+#define __GENERICAS [[clang::sycl_generic]]
#else
-#define __global __attribute__((opencl_global))
-#define __local __attribute__((opencl_local))
-#define __private __attribute__((opencl_private))
-#define __constant __attribute__((opencl_constant))
-#define __generic __attribute__((opencl_generic))
+#define __GLOBALAS __attribute__((opencl_global))
+#define __LOCALAS __attribute__((opencl_local))
+#define __PRIVATEAS __attribute__((opencl_private))
+#define __GENERICAS __attribute__((opencl_generic))
#endif
// Check if SPIR-V builtins are supported.
@@ -83,123 +81,113 @@ extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)
// OpGenericCastToPtrExplicit
+extern __SPIRV_overloadable __SPIRV_BUILTIN_ALIAS(
+ __builtin_spirv_generic_cast_to_ptr_explicit) void __GLOBALAS
+ *__spirv_GenericCastToPtrExplicit_ToGlobal(void __GENERICAS *,
+ int) __SPIRV_NOEXCEPT;
extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__global void *__spirv_GenericCastToPtrExplicit_ToGlobal(__generic void *,
- int) __SPIRV_NOEXCEPT;
+ __SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit) const
+ void __GLOBALAS *__spirv_GenericCastToPtrExplicit_ToGlobal(
+ const void __GENERICAS *, int) __SPIRV_NOEXCEPT;
+extern __SPIRV_overloadable __SPIRV_BUILTIN_ALIAS(
+ __builtin_spirv_generic_cast_to_ptr_explicit) volatile void __GLOBALAS
+ *__spirv_GenericCastToPtrExplicit_ToGlobal(volatile void __GENERICAS *,
+ int) __SPIRV_NOEXCEPT;
extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__global const void *
-__spirv_GenericCastToPtrExplicit_ToGlobal(__generic const void *,
- int) __SPIRV_NOEXCEPT;
+ __SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit) const
+ volatile void __GLOBALAS *__spirv_GenericCastToPtrExplicit_ToGlobal(
+ const volatile void __GENERICAS *, int) __SPIRV_NOEXCEPT;
+extern __SPIRV_overloadable __SPIRV_BUILTIN_ALIAS(
+ __builtin_spirv_generic_cast_to_ptr_explicit) void __LOCALAS
+ *__spirv_GenericCastToPtrExplicit_ToLocal(void __GENERICAS *,
+ int) __SPIRV_NOEXCEPT;
extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__global volatile void *
-__spirv_GenericCastToPtrExplicit_ToGlobal(__generic volatile void *,
- int) __SPIRV_NOEXCEPT;
+ __SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit) const
+ void __LOCALAS *__spirv_GenericCastToPtrExplicit_ToLocal(
+ const void __GENERICAS *, int) __SPIRV_NOEXCEPT;
+extern __SPIRV_overloadable __SPIRV_BUILTIN_ALIAS(
+ __builtin_spirv_generic_cast_to_ptr_explicit) volatile void __LOCALAS
+ *__spirv_GenericCastToPtrExplicit_ToLocal(volatile void __GENERICAS *,
+ int) __SPIRV_NOEXCEPT;
extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__global const volatile void *
-__spirv_GenericCastToPtrExplicit_ToGlobal(__generic const volatile void *,
- int) __SPIRV_NOEXCEPT;
+ __SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit) const
+ volatile void __LOCALAS *__spirv_GenericCastToPtrExplicit_ToLocal(
+ const volatile void __GENERICAS *, int) __SPIRV_NOEXCEPT;
+extern __SPIRV_overloadable __SPIRV_BUILTIN_ALIAS(
+ __builtin_spirv_generic_cast_to_ptr_explicit) void __PRIVATEAS
+ *__spirv_GenericCastToPtrExplicit_ToPrivate(void __GENERICAS *,
+ int) __SPIRV_NOEXCEPT;
extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__local void *__spirv_GenericCastToPtrExplicit_ToLocal(__generic void *,
- int) __SPIRV_NOEXCEPT;
+ __SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit) const
+ void __PRIVATEAS *__spirv_GenericCastToPtrExplicit_ToPrivate(
+ const void __GENERICAS *, int) __SPIRV_NOEXCEPT;
+extern __SPIRV_overloadable __SPIRV_BUILTIN_ALIAS(
+ __builtin_spirv_generic_cast_to_ptr_explicit) volatile void __PRIVATEAS
+ *__spirv_GenericCastToPtrExplicit_ToPrivate(volatile void __GENERICAS *,
+ int) __SPIRV_NOEXCEPT;
extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__local const void *
-__spirv_GenericCastToPtrExplicit_ToLocal(__generic const void *,
- int) __SPIRV_NOEXCEPT;
-extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__local volatile void *
-__spirv_GenericCastToPtrExplicit_ToLocal(__generic volatile void *,
- int) __SPIRV_NOEXCEPT;
-extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__local const volatile void *
-__spirv_GenericCastToPtrExplicit_ToLocal(__generic const volatile void *,
- int) __SPIRV_NOEXCEPT;
-extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__private void *
-__spirv_GenericCastToPtrExplicit_ToPrivate(__generic void *,
- int) __SPIRV_NOEXCEPT;
-extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__private const void *
-__spirv_GenericCastToPtrExplicit_ToPrivate(__generic const void *,
- int) __SPIRV_NOEXCEPT;
-extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__private volatile void *
-__spirv_GenericCastToPtrExplicit_ToPrivate(__generic volatile void *,
- int) __SPIRV_NOEXCEPT;
-extern __SPIRV_overloadable
-__SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit)
-__private const volatile void *
-__spirv_GenericCastToPtrExplicit_ToPrivate(__generic const volatile void *,
- int) __SPIRV_NOEXCEPT;
+ __SPIRV_BUILTIN_ALIAS(__builtin_spirv_generic_cast_to_ptr_explicit) const
+ volatile void __PRIVATEAS *__spirv_GenericCastToPtrExplicit_ToPrivate(
+ const volatile void __GENERICAS *, int) __SPIRV_NOEXCEPT;
// OpGenericCastToPtr
-static __SPIRV_overloadable __SPIRV_inline __global void *
-__spirv_GenericCastToPtr_ToGlobal(__generic void *p, int) __SPIRV_NOEXCEPT {
- return (__global void *)p;
+static __SPIRV_overloadable __SPIRV_inline void __GLOBALAS *
+__spirv_GenericCastToPtr_ToGlobal(void __GENERICAS *p, int) __SPIRV_NOEXCEPT {
+ return (void __GLOBALAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __global const void *
-__spirv_GenericCastToPtr_ToGlobal(__generic const void *p,
+static __SPIRV_overloadable __SPIRV_inline const void __GLOBALAS *
+__spirv_GenericCastToPtr_ToGlobal(const void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__global const void *)p;
+ return (const void __GLOBALAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __global volatile void *
-__spirv_GenericCastToPtr_ToGlobal(__generic volatile void *p,
+static __SPIRV_overloadable __SPIRV_inline volatile void __GLOBALAS *
+__spirv_GenericCastToPtr_ToGlobal(volatile void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__global volatile void *)p;
+ return (volatile void __GLOBALAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __global const volatile void *
-__spirv_GenericCastToPtr_ToGlobal(__generic const volatile void *p,
+static __SPIRV_overloadable __SPIRV_inline const volatile void __GLOBALAS *
+__spirv_GenericCastToPtr_ToGlobal(const volatile void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__global const volatile void *)p;
+ return (const volatile void __GLOBALAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __local void *
-__spirv_GenericCastToPtr_ToLocal(__generic void *p, int) __SPIRV_NOEXCEPT {
- return (__local void *)p;
+static __SPIRV_overloadable __SPIRV_inline void __LOCALAS *
+__spirv_GenericCastToPtr_ToLocal(void __GENERICAS *p, int) __SPIRV_NOEXCEPT {
+ return (void __LOCALAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __local const void *
-__spirv_GenericCastToPtr_ToLocal(__generic const void *p,
+static __SPIRV_overloadable __SPIRV_inline const void __LOCALAS *
+__spirv_GenericCastToPtr_ToLocal(const void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__local const void *)p;
+ return (const void __LOCALAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __local volatile void *
-__spirv_GenericCastToPtr_ToLocal(__generic volatile void *p,
+static __SPIRV_overloadable __SPIRV_inline volatile void __LOCALAS *
+__spirv_GenericCastToPtr_ToLocal(volatile void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__local volatile void *)p;
+ return (volatile void __LOCALAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __local const volatile void *
-__spirv_GenericCastToPtr_ToLocal(__generic const volatile void *p,
+static __SPIRV_overloadable __SPIRV_inline const volatile void __LOCALAS *
+__spirv_GenericCastToPtr_ToLocal(const volatile void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__local const volatile void *)p;
+ return (const volatile void __LOCALAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __private void *
-__spirv_GenericCastToPtr_ToPrivate(__generic void *p, int) __SPIRV_NOEXCEPT {
- return (__private void *)p;
+static __SPIRV_overloadable __SPIRV_inline void __PRIVATEAS *
+__spirv_GenericCastToPtr_ToPrivate(void __GENERICAS *p, int) __SPIRV_NOEXCEPT {
+ return (void __PRIVATEAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __private const void *
-__spirv_GenericCastToPtr_ToPrivate(__generic const void *p,
+static __SPIRV_overloadable __SPIRV_inline const void __PRIVATEAS *
+__spirv_GenericCastToPtr_ToPrivate(const void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__private const void *)p;
+ return (const void __PRIVATEAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __private volatile void *
-__spirv_GenericCastToPtr_ToPrivate(__generic volatile void *p,
+static __SPIRV_overloadable __SPIRV_inline volatile void __PRIVATEAS *
+__spirv_GenericCastToPtr_ToPrivate(volatile void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__private volatile void *)p;
+ return (volatile void __PRIVATEAS *)p;
}
-static __SPIRV_overloadable __SPIRV_inline __private const volatile void *
-__spirv_GenericCastToPtr_ToPrivate(__generic const volatile void *p,
+static __SPIRV_overloadable __SPIRV_inline const volatile void __PRIVATEAS *
+__spirv_GenericCastToPtr_ToPrivate(const volatile void __GENERICAS *p,
int) __SPIRV_NOEXCEPT {
- return (__private const volatile void *)p;
+ return (const volatile void __PRIVATEAS *)p;
}
#pragma pop_macro("__size_t")
@@ -210,10 +198,10 @@ __spirv_GenericCastToPtr_ToPrivate(__generic const volatile void *p,
#undef __SPIRV_convergent
#undef __SPIRV_inline
-#undef __global
-#undef __local
-#undef __constant
-#undef __generic
+#undef __GLOBALAS
+#undef __LOCALAS
+#undef __PRIVATEAS
+#undef __GENERICAS
#undef __SPIRV_BUILTIN_ALIAS
#undef __SPIRV_NOEXCEPT
diff --git a/clang/lib/Sema/ParsedAttr.cpp b/clang/lib/Sema/ParsedAttr.cpp
index 49dec6188d877..2b5ad33ad7b29 100644
--- a/clang/lib/Sema/ParsedAttr.cpp
+++ b/clang/lib/Sema/ParsedAttr.cpp
@@ -228,13 +228,13 @@ bool ParsedAttr::slidesFromDeclToDeclSpecLegacyBehavior() const {
// possible, we would like this list to go away entirely.
switch (getParsedKind()) {
case AT_AddressSpace:
- case AT_OffloadPrivateAddressSpace:
- case AT_OffloadGlobalAddressSpace:
+ case AT_OpenCLPrivateAddressSpace:
+ case AT_OpenCLGlobalAddressSpace:
case AT_OpenCLGlobalDeviceAddressSpace:
case AT_OpenCLGlobalHostAddressSpace:
- case AT_OffloadLocalAddressSpace:
- case AT_OffloadConstantAddressSpace:
- case AT_OffloadGenericAddressSpace:
+ case AT_OpenCLLocalAddressSpace:
+ case AT_OpenCLConstantAddressSpace:
+ case AT_OpenCLGenericAddressSpace:
case AT_NeonPolyVectorType:
case AT_NeonVectorType:
case AT_ArmMveStrictPolymorphism:
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 1d13e632c51d6..3c2fcffcde85b 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -6685,11 +6685,11 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
ASIdx = Attr.asHLSLLangAS();
if (ASIdx == LangAS::Default) {
- if (S.getLangOpts().SYCLIsDevice &&
- Attr.getKind() == ParsedAttr::AT_OffloadConstantAddressSpace)
+ if (Attr.getKind() == ParsedAttr::AT_SYCLConstantAddressSpace) {
S.Diag(Attr.getLoc(), diag::warn_deprecated_sycl_constant);
- else
+ } else {
llvm_unreachable("Invalid address space");
+ }
}
if (DiagnoseMultipleAddrSpaceAttributes(S, Type.getAddressSpace(), ASIdx,
@@ -9097,14 +9097,19 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type,
// it it breaks large amounts of Linux software.
attr.setUsedAsTypeAttr();
break;
- case ParsedAttr::AT_OffloadPrivateAddressSpace:
- case ParsedAttr::AT_OffloadGlobalAddressSpace:
+ case ParsedAttr::AT_OpenCLPrivateAddressSpace:
+ case ParsedAttr::AT_OpenCLGlobalAddressSpace:
case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
- case ParsedAttr::AT_OffloadLocalAddressSpace:
- case ParsedAttr::AT_OffloadConstantAddressSpace:
- case ParsedAttr::AT_OffloadGenericAddressSpace:
+ case ParsedAttr::AT_OpenCLLocalAddressSpace:
+ case ParsedAttr::AT_OpenCLConstantAddressSpace:
+ case ParsedAttr::AT_OpenCLGenericAddressSpace:
case ParsedAttr::AT_AddressSpace:
+ case ParsedAttr::AT_SYCLPrivateAddressSpace:
+ case ParsedAttr::AT_SYCLGlobalAddressSpace:
+ case ParsedAttr::AT_SYCLLocalAddressSpace:
+ case ParsedAttr::AT_SYCLConstantAddressSpace:
+ case ParsedAttr::AT_SYCLGenericAddressSpace:
HandleAddressSpaceTypeAttribute(type, attr, state);
attr.setUsedAsTypeAttr();
break;
diff --git a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
index 30f4ecb589a5c..640186a012444 100644
--- a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
+++ b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
@@ -4,8 +4,14 @@
#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_EXTERNAL [[clang::sycl_external]]
+#define GLOBALAS [[clang::sycl_global]]
+#define LOCALAS [[clang::sycl_local]]
+#define PRIVATEAS [[clang::sycl_private]]
#else
#define SYCL_EXTERNAL
+#define GLOBALAS __attribute__((opencl_global))
+#define LOCALAS __attribute__((opencl_local))
+#define PRIVATEAS __attribute__((opencl_private))
#endif
// CHECK: spir_func noundef ptr @{{.*}}test_cast_to_private{{.*}}(ptr addrspace(4) noundef readnone [[P:%.*]]
@@ -13,7 +19,7 @@
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr @llvm.spv.generic.cast.to.ptr.explicit.p0(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr [[SPV_CAST]]
//
-SYCL_EXTERNAL __attribute__((opencl_private)) int* test_cast_to_private(int* p) {
+SYCL_EXTERNAL int PRIVATEAS * test_cast_to_private(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 7);
}
@@ -22,7 +28,7 @@ SYCL_EXTERNAL __attribute__((opencl_private)) int* test_cast_to_private(int* p)
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(1) @llvm.spv.generic.cast.to.ptr.explicit.p1(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr addrspace(1) [[SPV_CAST]]
//
-SYCL_EXTERNAL __attribute__((opencl_global)) int* test_cast_to_global(int* p) {
+SYCL_EXTERNAL int GLOBALAS * test_cast_to_global(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 5);
}
@@ -31,6 +37,6 @@ SYCL_EXTERNAL __attribute__((opencl_global)) int* test_cast_to_global(int* p) {
// CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(3) @llvm.spv.generic.cast.to.ptr.explicit.p3(ptr addrspace(4) %p)
// CHECK-NEXT: ret ptr addrspace(3) [[SPV_CAST]]
//
-SYCL_EXTERNAL __attribute__((opencl_local)) int* test_cast_to_local(int* p) {
+SYCL_EXTERNAL int LOCALAS * test_cast_to_local(int* p) {
return __builtin_spirv_generic_cast_to_ptr_explicit(p, 4);
}
diff --git a/clang/test/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CodeGenSYCL/address-space-conversions.cpp
index 3eecdded06364..811f776cc845f 100644
--- a/clang/test/CodeGenSYCL/address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/address-space-conversions.cpp
@@ -3,13 +3,13 @@ void bar(int &Data) {}
// CHECK-DAG: define{{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) %
void bar2(int &Data) {}
// CHECK-DAG: define{{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) %
-void bar(__attribute__((sycl_local)) int &Data) {}
+void bar(int [[clang::sycl_local]] &Data) {}
// CHECK-DAG: define{{.*}} void [[LOC_REF:@[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) %
void foo(int *Data) {}
// CHECK-DAG: define{{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef %
void foo2(int *Data) {}
// CHECK-DAG: define{{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef %
-void foo(__attribute__((sycl_local)) int *Data) {}
+void foo(int [[clang::sycl_local]] *Data) {}
// CHECK-DAG: define{{.*}} void [[LOC_PTR:@[a-zA-Z0-9_]+]](ptr addrspace(3) noundef %
template <typename T>
@@ -19,15 +19,15 @@ void tmpl(T t) {}
[[clang::sycl_external]] void usages() {
int *NoAS;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr addrspace(4)
- __attribute__((sycl_global)) int *GLOB;
+ int [[clang::sycl_global]] *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1)
- __attribute__((sycl_local)) int *LOC;
+ int [[clang::sycl_local]] *LOC;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3)
- __attribute__((sycl_private)) int *PRIV;
+ int [[clang::sycl_private]] *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr
- __attribute__((opencl_global_device)) int *GLOBDEVICE;
+ int __attribute__((opencl_global_device)) *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5)
- __attribute__((opencl_global_host)) int *GLOBHOST;
+ int __attribute__((opencl_global_host)) *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(6)
// CHECK-DAG: [[NoAS]].ascast = addrspacecast ptr [[NoAS]] to ptr addrspace(4)
@@ -57,24 +57,24 @@ void tmpl(T t) {}
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[NoAS_CAST]], ptr addrspace(4) [[GLOB]].ascast
- GLOB = (__attribute__((sycl_global)) int *)NoAS;
+ GLOB = (int [[clang::sycl_global]] *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr addrspace(3)
// CHECK-DAG: store ptr addrspace(3) [[NoAS_CAST]], ptr addrspace(4) [[LOC]].ascast
- LOC = (__attribute__((sycl_local)) int *)NoAS;
+ LOC = (int [[clang::sycl_local]] *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr
// CHECK-DAG: store ptr [[NoAS_CAST]], ptr addrspace(4) [[PRIV]].ascast
- PRIV = (__attribute__((sycl_private)) int *)NoAS;
+ PRIV = (int [[clang::sycl_private]] *)NoAS;
// From opencl_global_[host/device] address spaces to sycl_global
// CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr addrspace(4) [[GLOB_DEVICE]].ascast
// CHECK-DAG: [[GLOBDEVICE_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[GLOBDEVICE_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_CAST]], ptr addrspace(4) [[GLOB]].ascast
- GLOB = (__attribute__((sycl_global)) int *)GLOBDEVICE;
+ GLOB = (int [[clang::sycl_global]] *)GLOBDEVICE;
// CHECK-DAG: [[GLOBHOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(6), ptr addrspace(4) [[GLOB_HOST]].ascast
// CHECK-DAG: [[GLOBHOST_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(6) [[GLOBHOST_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[GLOBHOST_CAST]], ptr addrspace(4) [[GLOB]].ascast
- GLOB = (__attribute__((sycl_global)) int *)GLOBHOST;
+ GLOB = (int [[clang::sycl_global]] *)GLOBHOST;
bar(*GLOB);
// CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast
diff --git a/clang/test/CodeGenSYCL/address-space-mangling.cpp b/clang/test/CodeGenSYCL/address-space-mangling.cpp
index 3e44be24ffc3e..c006905ab82cd 100644
--- a/clang/test/CodeGenSYCL/address-space-mangling.cpp
+++ b/clang/test/CodeGenSYCL/address-space-mangling.cpp
@@ -3,9 +3,9 @@
// REQUIRES: x86-registered-target
-void foo(__attribute__((sycl_global)) int *);
-void foo(__attribute__((sycl_local)) int *);
-void foo(__attribute__((sycl_private)) int *);
+void foo(int [[clang::sycl_global]] *);
+void foo(int [[clang::sycl_local]] *);
+void foo(int [[clang::sycl_private]] *);
void foo(int *);
// SPIR: declare spir_func void @_Z3fooPU3AS1i(ptr addrspace(1) noundef) #1
@@ -19,9 +19,9 @@ void foo(int *);
// X86: declare void @_Z3fooPi(ptr noundef) #1
[[clang::sycl_external]] void test() {
- __attribute__((sycl_global)) int *glob;
- __attribute__((sycl_local)) int *loc;
- __attribute__((sycl_private)) int *priv;
+ int [[clang::sycl_global]] *glob;
+ int [[clang::sycl_local]] *loc;
+ int [[clang::sycl_private]] *priv;
int *def;
foo(glob);
foo(loc);
diff --git a/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp b/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
index 17a98195318ad..3ead6a7533860 100644
--- a/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp
@@ -3,13 +3,13 @@ void bar(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
void bar2(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
-void bar(__attribute__((opencl_local)) int &Data) {}
+void bar(int [[clang::sycl_local]] &Data) {}
// CHECK-DAG: define {{.*}} void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) %
void foo(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef %
void foo2(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef %
-void foo(__attribute__((opencl_local)) int *Data) {}
+void foo(int [[clang::sycl_local]] *Data) {}
// CHECK-DAG: define {{.*}} void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef %
template <typename T>
@@ -19,15 +19,15 @@ void tmpl(T t) {}
[[clang::sycl_external]] void usages() {
int *NoAS;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8, addrspace(5)
- __attribute__((opencl_global)) int *GLOB;
+ int [[clang::sycl_global]] *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
- __attribute__((opencl_local)) int *LOC;
+ int [[clang::sycl_local]] *LOC;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 4, addrspace(5)
- __attribute__((opencl_private)) int *PRIV;
+ int [[clang::sycl_private]] *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5), align 4, addrspace(5)
- __attribute__((opencl_global_device)) int *GLOBDEVICE;
+ int __attribute__((opencl_global_device)) *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
- __attribute__((opencl_global_host)) int *GLOBHOST;
+ int __attribute__((opencl_global_host)) *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5)
LOC = nullptr;
// CHECK-DAG: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]].ascast, align 4
@@ -45,22 +45,22 @@ void tmpl(T t) {}
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[NoAS_LOAD]] to ptr
// CHECK-DAG: store ptr %5, ptr [[NoAS]].ascast, align 8
- GLOB = (__attribute__((opencl_global)) int *)NoAS;
+ GLOB = (int [[clang::sycl_global]] *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr %6 to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) %7, ptr [[GLOB]].ascast, align 8
- LOC = (__attribute__((opencl_local)) int *)NoAS;
+ LOC = (int [[clang::sycl_local]] *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3)
// CHECK-DAG: store ptr addrspace(3) %9, ptr [[LOC]].ascast, align 4
- PRIV = (__attribute__((opencl_private)) int *)NoAS;
+ PRIV = (int [[clang::sycl_private]] *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(5)
// CHECK-DAG: store ptr addrspace(5) [[NoAS_CAST]], ptr [[PRIV]].ascast, align 4
- GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE;
+ GLOB = (int [[clang::sycl_global]] *)GLOBDEVICE;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]DEVICE.ascast, align 8
// CHECK-DAG: store ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8
- GLOB = (__attribute__((opencl_global)) int *)GLOBHOST;
+ GLOB = (int [[clang::sycl_global]] *)GLOBHOST;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]HOST.ascast, align 8
// CHECK-DAG: tore ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8
bar(*GLOB);
diff --git a/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp b/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
index ffb601e62c118..305c86c0ef09b 100644
--- a/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
+++ b/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp
@@ -3,13 +3,13 @@ void bar(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
void bar2(int &Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) %
-void bar(__attribute__((opencl_local)) int &Data) {}
+void bar(int [[clang::sycl_local]] &Data) {}
// CHECK-DAG: define {{.*}} void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) %
void foo(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef %
void foo2(int *Data) {}
// CHECK-DAG: define {{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef %
-void foo(__attribute__((opencl_local)) int *Data) {}
+void foo(int [[clang::sycl_local]] *Data) {}
// CHECK-DAG: define {{.*}} void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef %
template <typename T>
@@ -19,15 +19,15 @@ void tmpl(T t);
[[clang::sycl_external]] void usages() {
int *NoAS;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8
- __attribute__((opencl_global)) int *GLOB;
+ int [[clang::sycl_global]] *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
- __attribute__((opencl_local)) int *LOC;
+ int [[clang::sycl_local]] *LOC;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 8
- __attribute__((opencl_private)) int *PRIV;
+ int [[clang::sycl_private]] *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr, align 8
- __attribute__((opencl_global_device)) int *GLOBDEVICE;
+ int __attribute__((opencl_global_device)) *GLOBDEVICE;
// CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
- __attribute__((opencl_global_host)) int *GLOBHOST;
+ int __attribute__((opencl_global_host)) *GLOBHOST;
// CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8
LOC = nullptr;
// CHECK-DAG: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]], align 8
@@ -44,21 +44,21 @@ void tmpl(T t);
NoAS = (int *)PRIV;
// CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[PRIV]], align 8
// CHECK-DAG: store ptr [[LOC_LOAD]], ptr [[NoAS]], align 8
- GLOB = (__attribute__((opencl_global)) int *)NoAS;
+ GLOB = (int [[clang::sycl_global]] *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(1)
// CHECK-DAG: store ptr addrspace(1) [[NoAS_CAST]], ptr [[GLOB]], align 8
- LOC = (__attribute__((opencl_local)) int *)NoAS;
+ LOC = (int [[clang::sycl_local]] *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3)
// CHECK-DAG: store ptr addrspace(3) [[NoAS_CAST]], ptr [[LOC]], align 8
- PRIV = (__attribute__((opencl_private)) int *)NoAS;
+ PRIV = (int [[clang::sycl_private]] *)NoAS;
// CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8
// CHECK-DAG: store ptr [[NoAS_LOAD]], ptr [[PRIV]], align 8
- GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE;
+ GLOB = (int [[clang::sycl_global]] *)GLOBDEVICE;
// CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_DEVICE]], align 8
// CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_LOAD]], ptr %GLOB, align 8
- GLOB = (__attribute__((opencl_global)) int *)GLOBHOST;
+ GLOB = (int [[clang::sycl_global]] *)GLOBHOST;
// CHECK-DAG: [[GLOB_HOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_HOST]], align 8
// CHECK-DAG: store ptr addrspace(1) [[GLOB_HOST_LOAD]], ptr [[GLOB]], align 8
bar(*GLOB);
diff --git a/clang/test/SemaSPIRV/BuiltIns/generic_cast_to_ptr_explicit.c b/clang/test/SemaSPIRV/BuiltIns/generic_cast_to_ptr_explicit.c
index 5a839961e20f5..13d595d87d114 100644
--- a/clang/test/SemaSPIRV/BuiltIns/generic_cast_to_ptr_explicit.c
+++ b/clang/test/SemaSPIRV/BuiltIns/generic_cast_to_ptr_explicit.c
@@ -1,7 +1,13 @@
-// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device -verify %s -o -
+// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device -x c++ -verify %s -o -
// RUN: %clang_cc1 -O1 -triple spirv64 -verify %s -cl-std=CL3.0 -x cl -o -
// RUN: %clang_cc1 -O1 -triple spirv32 -verify %s -cl-std=CL3.0 -x cl -o -
+#ifdef __SYCL_DEVICE_ONLY__
+#define LOCALAS [[clang::sycl_local]]
+#else
+#define LOCALAS __attribute__((opencl_local))
+#endif
+
void test_missing_arguments(int* p) {
__builtin_spirv_generic_cast_to_ptr_explicit(p);
// expected-error at -1 {{too few arguments to function call, expected 2, have 1}}
@@ -14,7 +20,7 @@ void test_wrong_flag_value(int* p) {
// expected-error at -1 {{invalid value for storage class argument}}
}
-void test_wrong_address_space(__attribute__((opencl_local)) int* p) {
+void test_wrong_address_space(int LOCALAS * p) {
__builtin_spirv_generic_cast_to_ptr_explicit(p, 14);
// expected-error at -1 {{expecting a pointer argument to the generic address space}}
}
diff --git a/clang/test/SemaSYCL/address-space-conversions.cpp b/clang/test/SemaSYCL/address-space-conversions.cpp
index 9d209dbe5f8d7..41dec2e926418 100644
--- a/clang/test/SemaSYCL/address-space-conversions.cpp
+++ b/clang/test/SemaSYCL/address-space-conversions.cpp
@@ -2,26 +2,25 @@
void bar(int &Data) {}
void bar2(int &Data) {}
-void bar(__attribute__((sycl_private)) int &Data) {}
+void bar(int [[clang::sycl_private]] &Data) {}
void foo(int *Data) {}
void foo2(int *Data) {}
-void foo(__attribute__((sycl_private)) int *Data) {}
-void baz(__attribute__((sycl_private)) int *Data) {} // expected-note {{candidate function not viable: cannot pass pointer to generic address space as a pointer to address space 'sycl_private' in 1st argument}}
+void foo(int [[clang::sycl_private]] *Data) {}
+void baz(int [[clang::sycl_private]] *Data) {} // expected-note {{candidate function not viable: cannot pass pointer to generic address space as a pointer to address space 'sycl_private' in 1st argument}}
template <typename T>
void tmpl(T *t) {}
void usages() {
- __attribute__((sycl_global)) int *GLOB;
- __attribute__((sycl_private)) int *PRIV;
- __attribute__((sycl_local)) int *LOC;
- __attribute__((sycl_constant)) int *ptr1; // expected-warning {{'sycl_constant' address space attribute is deprecated}}
+ int [[clang::sycl_global]] *GLOB;
+ int [[clang::sycl_private]] *PRIV;
+ int [[clang::sycl_local]] *LOC;
int *NoAS;
GLOB = PRIV; // expected-error {{assigning 'sycl_private int *' to 'sycl_global int *' changes address space of pointer}}
GLOB = LOC; // expected-error {{assigning 'sycl_local int *' to 'sycl_global int *' changes address space of pointer}}
- PRIV = static_cast<__attribute__((sycl_private)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to 'sycl_private int *' is not allowed}}
- PRIV = static_cast<__attribute__((sycl_private)) int *>(LOC); // expected-error {{static_cast from 'sycl_local int *' to 'sycl_private int *' is not allowed}}
+ PRIV = static_cast<int [[clang::sycl_private]] *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to 'sycl_private int *' is not allowed}}
+ PRIV = static_cast<int [[clang::sycl_private]] *>(LOC); // expected-error {{static_cast from 'sycl_local int *' to 'sycl_private int *' is not allowed}}
NoAS = GLOB + PRIV; // expected-error {{invalid operands to binary expression ('sycl_global int *' and 'sycl_private int *')}}
NoAS = GLOB + LOC; // expected-error {{invalid operands to binary expression ('sycl_global int *' and 'sycl_local int *')}}
NoAS += GLOB; // expected-error {{invalid operands to binary expression ('int *' and 'sycl_global int *')}}
@@ -54,10 +53,10 @@ void usages() {
// Implicit casts to named address space are disallowed
baz(NoAS); // expected-error {{no matching function for call to 'baz'}}
- __attribute__((sycl_local)) int *l = NoAS; // expected-error {{cannot initialize a variable of type 'sycl_local int *' with an lvalue of type 'int *'}}
+ int [[clang::sycl_local]] *l = NoAS; // expected-error {{cannot initialize a variable of type 'sycl_local int *' with an lvalue of type 'int *'}}
// Explicit casts between disjoint address spaces are disallowed
- GLOB = (__attribute__((sycl_global)) int *)PRIV; // expected-error {{C-style cast from 'sycl_private int *' to 'sycl_global int *' converts between mismatching address spaces}}
+ GLOB = (int [[clang::sycl_global]] *)PRIV; // expected-error {{C-style cast from 'sycl_private int *' to 'sycl_global int *' converts between mismatching address spaces}}
(void)static_cast<int *>(GLOB);
(void)static_cast<void *>(GLOB);
diff --git a/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp b/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
deleted file mode 100644
index 89c4fa4873086..0000000000000
--- a/clang/test/SemaSYCL/address-space-opencl-sycl-compat.cpp
+++ /dev/null
@@ -1,34 +0,0 @@
-// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
-
-// Test openCL and SYCL spelling conversions for address space
-// attributes.
-
-void test_incompatible() {
- __attribute__((opencl_global)) int *opencl_global;
- int [[clang::sycl_local]] *sycl_local;
- int [[clang::sycl_private]] *sycl_private;
-
- // Address space attributes are resolved using mode of compilation and not the spelling itself. This results in the SYCL spelling
- // being used in both instances of each diagnostic despite openCL spelling being used.
- opencl_global = sycl_local; // expected-error {{assigning 'sycl_local int *' to 'sycl_global int *' changes address space of pointer}}
- opencl_global = sycl_private; // expected-error {{assigning 'sycl_private int *' to 'sycl_global int *' changes address space of pointer}}
- sycl_local = opencl_global; // expected-error {{assigning 'sycl_global int *' to 'sycl_local int *' changes address space of pointer}}
-}
-
-void test_to_generic_mixed() {
- __attribute__((opencl_generic)) int *opencl_gen;
- int [[clang::sycl_generic]] *sycl_gen;
-
- __attribute__((opencl_global)) int *opencl_global;
- int [[clang::sycl_local]] *sycl_local;
- int [[clang::sycl_private]] *sycl_private;
-
- opencl_gen = sycl_local;
- opencl_gen = sycl_private;
- sycl_gen = opencl_global;
-
-}
-
-void overload_test(__attribute__((opencl_global)) int *p) { (void)p; } // expected-note {{previous definition is here}}
-void overload_test(__attribute__((sycl_global)) int *p) { (void)p; } // expected-error {{redefinition of 'overload_test'}}
-
>From d0e0728b79a56f2455fd95f0f5da21ac95aab1f4 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews at intel.com>
Date: Wed, 10 Jun 2026 11:35:01 -0700
Subject: [PATCH 7/7] Add sycl_constant LangAS based on review comments.
---
clang/include/clang/Basic/AddressSpaces.h | 1 +
clang/include/clang/Sema/ParsedAttr.h | 2 ++
clang/lib/AST/ItaniumMangle.cpp | 6 +++-
clang/lib/AST/TypePrinter.cpp | 2 ++
clang/lib/Basic/TargetInfo.cpp | 1 +
clang/lib/Basic/Targets/AArch64.h | 1 +
clang/lib/Basic/Targets/AMDGPU.cpp | 2 ++
clang/lib/Basic/Targets/DirectX.h | 1 +
clang/lib/Basic/Targets/NVPTX.h | 1 +
clang/lib/Basic/Targets/SPIR.h | 2 ++
clang/lib/Basic/Targets/SystemZ.h | 1 +
clang/lib/Basic/Targets/TCE.h | 1 +
clang/lib/Basic/Targets/WebAssembly.h | 1 +
clang/lib/Basic/Targets/X86.h | 1 +
clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 2 ++
clang/lib/Sema/SemaType.cpp | 12 ++++----
.../CodeGenSYCL/address-space-mangling.cpp | 14 ++++++++--
.../SemaSYCL/address-space-conversions.cpp | 28 +++++++++++++++++++
.../SemaTemplate/address_space-dependent.cpp | 4 +--
19 files changed, 71 insertions(+), 12 deletions(-)
diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h
index c3654ebbab82c..849770a0237f1 100644
--- a/clang/include/clang/Basic/AddressSpaces.h
+++ b/clang/include/clang/Basic/AddressSpaces.h
@@ -57,6 +57,7 @@ enum class LangAS : unsigned {
sycl_local,
sycl_private,
sycl_generic,
+ sycl_constant,
// Pointer size and extension address spaces.
ptr32_sptr,
diff --git a/clang/include/clang/Sema/ParsedAttr.h b/clang/include/clang/Sema/ParsedAttr.h
index 7dd1290000f7c..93e47622587a9 100644
--- a/clang/include/clang/Sema/ParsedAttr.h
+++ b/clang/include/clang/Sema/ParsedAttr.h
@@ -595,6 +595,8 @@ class ParsedAttr final
return LangAS::sycl_private;
case ParsedAttr::AT_SYCLGenericAddressSpace:
return LangAS::sycl_generic;
+ case ParsedAttr::AT_SYCLConstantAddressSpace:
+ return LangAS::sycl_constant;
default:
return LangAS::Default;
}
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index 6be573966781c..a11c8b5d0837e 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -2822,7 +2822,8 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals, const DependentAddressSp
ASString = "CLgeneric";
break;
// <SYCL-addrspace> ::= "SY" [ "global" | "local" | "private" |
- // "generic" | "device" | "host" ]
+ // "generic" | "constant" | "device" | "host"
+ // ]
case LangAS::sycl_global:
ASString = "SYglobal";
break;
@@ -2841,6 +2842,9 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals, const DependentAddressSp
case LangAS::sycl_generic:
ASString = "SYgeneric";
break;
+ case LangAS::sycl_constant:
+ ASString = "SYconstant";
+ break;
// <CUDA-addrspace> ::= "CU" [ "device" | "constant" | "shared" ]
case LangAS::cuda_device:
ASString = "CUdevice";
diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp
index 1ac5566c31904..9d0f12c23323b 100644
--- a/clang/lib/AST/TypePrinter.cpp
+++ b/clang/lib/AST/TypePrinter.cpp
@@ -2697,6 +2697,8 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
return "sycl_private";
case LangAS::sycl_generic:
return "sycl_generic";
+ case LangAS::sycl_constant:
+ return "sycl_constant";
case LangAS::cuda_device:
return "__device__";
case LangAS::cuda_constant:
diff --git a/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp
index 1d91bbaee21ef..cf589230812f0 100644
--- a/clang/lib/Basic/TargetInfo.cpp
+++ b/clang/lib/Basic/TargetInfo.cpp
@@ -45,6 +45,7 @@ static const LangASMap FakeAddrSpaceMap = {
3, // sycl_local
0, // sycl_private
4, // sycl_generic
+ 2, // sycl_constant
10, // ptr32_sptr
11, // ptr32_uptr
12, // ptr64
diff --git a/clang/lib/Basic/Targets/AArch64.h b/clang/lib/Basic/Targets/AArch64.h
index 90d8401149c37..8aad44a0964ad 100644
--- a/clang/lib/Basic/Targets/AArch64.h
+++ b/clang/lib/Basic/Targets/AArch64.h
@@ -42,6 +42,7 @@ static const unsigned ARM64AddrSpaceMap[] = {
0, // sycl_local
0, // sycl_private
0, // sycl_generic
+ 0, // sycl_constant
static_cast<unsigned>(AArch64AddrSpace::ptr32_sptr),
static_cast<unsigned>(AArch64AddrSpace::ptr32_uptr),
static_cast<unsigned>(AArch64AddrSpace::ptr64),
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index 5d102e015790f..5ca444353d34b 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -44,6 +44,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsGenMap = {
llvm::AMDGPUAS::LOCAL_ADDRESS, // sycl_local
llvm::AMDGPUAS::PRIVATE_ADDRESS, // sycl_private
llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_generic
+ llvm::AMDGPUAS::CONSTANT_ADDRESS, // sycl_constant
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_sptr
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_uptr
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr64
@@ -77,6 +78,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsPrivMap = {
llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_local
llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_private
llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_generic
+ llvm::AMDGPUAS::CONSTANT_ADDRESS, // sycl_constant
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_sptr
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_uptr
llvm::AMDGPUAS::FLAT_ADDRESS, // ptr64
diff --git a/clang/lib/Basic/Targets/DirectX.h b/clang/lib/Basic/Targets/DirectX.h
index 6eb770f4a960e..e72b2f20f12a6 100644
--- a/clang/lib/Basic/Targets/DirectX.h
+++ b/clang/lib/Basic/Targets/DirectX.h
@@ -39,6 +39,7 @@ static const unsigned DirectXAddrSpaceMap[] = {
0, // sycl_local
0, // sycl_private
0, // sycl_generic
+ 0, // sycl_constant
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h
index 62bf6a514d444..53a1a034f916e 100644
--- a/clang/lib/Basic/Targets/NVPTX.h
+++ b/clang/lib/Basic/Targets/NVPTX.h
@@ -43,6 +43,7 @@ static const unsigned NVPTXAddrSpaceMap[] = {
3, // sycl_local
0, // sycl_private
0, // sycl_generic
+ 0, // sycl_constant
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 88a9e5841b5a3..49df361832475 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -45,6 +45,7 @@ static const unsigned SPIRDefIsPrivMap[] = {
0, // sycl_local
0, // sycl_private
0, // sycl_generic
+ 0, // sycl_constant
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
@@ -84,6 +85,7 @@ static const unsigned SPIRDefIsGenMap[] = {
3, // sycl_local
0, // sycl_private
4, // sycl_generic
+ 2, // sycl_constant
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/SystemZ.h b/clang/lib/Basic/Targets/SystemZ.h
index bc597fe30165b..da996382ad15b 100644
--- a/clang/lib/Basic/Targets/SystemZ.h
+++ b/clang/lib/Basic/Targets/SystemZ.h
@@ -39,6 +39,7 @@ static const unsigned ZOSAddressMap[] = {
0, // sycl_local
0, // sycl_private
0, // sycl_generic
+ 0, // sycl_constant
0, // ptr32_sptr
1, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/TCE.h b/clang/lib/Basic/Targets/TCE.h
index d7086b23232c4..ab8a720517d83 100644
--- a/clang/lib/Basic/Targets/TCE.h
+++ b/clang/lib/Basic/Targets/TCE.h
@@ -48,6 +48,7 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = {
0, // sycl_local
0, // sycl_private
0, // sycl_generic
+ 0, // sycl_constant
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/WebAssembly.h b/clang/lib/Basic/Targets/WebAssembly.h
index b0ee4505625c2..aca9b9fa15f60 100644
--- a/clang/lib/Basic/Targets/WebAssembly.h
+++ b/clang/lib/Basic/Targets/WebAssembly.h
@@ -39,6 +39,7 @@ static const unsigned WebAssemblyAddrSpaceMap[] = {
0, // sycl_local
0, // sycl_private
0, // sycl_generic
+ 0, // sycl_constant
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
diff --git a/clang/lib/Basic/Targets/X86.h b/clang/lib/Basic/Targets/X86.h
index 31fb984caac51..dfd57f0aa1e60 100644
--- a/clang/lib/Basic/Targets/X86.h
+++ b/clang/lib/Basic/Targets/X86.h
@@ -43,6 +43,7 @@ static const unsigned X86AddrSpaceMap[] = {
0, // sycl_local
0, // sycl_private
0, // sycl_generic
+ 0, // sycl_constant
270, // ptr32_sptr
271, // ptr32_uptr
272, // ptr64
diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp
index 9f2342254a882..3c14f90e64ed0 100644
--- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp
@@ -1042,6 +1042,8 @@ cir::LangAddressSpace cir::toCIRLangAddressSpace(clang::LangAS langAS) {
case LangAS::sycl_global_host:
case LangAS::sycl_local:
case LangAS::sycl_private:
+ case LangAS::sycl_generic:
+ case LangAS::sycl_constant:
case LangAS::ptr32_sptr:
case LangAS::ptr32_uptr:
case LangAS::ptr64:
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 3c2fcffcde85b..70c6b9b895848 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -6684,13 +6684,11 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
if (S.getLangOpts().HLSL)
ASIdx = Attr.asHLSLLangAS();
- if (ASIdx == LangAS::Default) {
- if (Attr.getKind() == ParsedAttr::AT_SYCLConstantAddressSpace) {
- S.Diag(Attr.getLoc(), diag::warn_deprecated_sycl_constant);
- } else {
- llvm_unreachable("Invalid address space");
- }
- }
+ if (ASIdx == LangAS::sycl_constant)
+ S.Diag(Attr.getLoc(), diag::warn_deprecated_sycl_constant);
+
+ if (ASIdx == LangAS::Default)
+ llvm_unreachable("Invalid address space");
if (DiagnoseMultipleAddrSpaceAttributes(S, Type.getAddressSpace(), ASIdx,
Attr.getLoc())) {
diff --git a/clang/test/CodeGenSYCL/address-space-mangling.cpp b/clang/test/CodeGenSYCL/address-space-mangling.cpp
index c006905ab82cd..b92c93269026a 100644
--- a/clang/test/CodeGenSYCL/address-space-mangling.cpp
+++ b/clang/test/CodeGenSYCL/address-space-mangling.cpp
@@ -1,30 +1,40 @@
-// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s --check-prefix=SPIR
-// RUN: %clang_cc1 -triple x86_64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s --check-prefix=X86
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - -Wno-deprecated-attributes | FileCheck %s --check-prefix=SPIR
+// RUN: %clang_cc1 -triple x86_64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - -Wno-deprecated-attributes | FileCheck %s --check-prefix=X86
// REQUIRES: x86-registered-target
void foo(int [[clang::sycl_global]] *);
void foo(int [[clang::sycl_local]] *);
void foo(int [[clang::sycl_private]] *);
+void foo(int [[clang::sycl_generic]] *);
+void foo(int [[clang::sycl_constant]] *);
void foo(int *);
// SPIR: declare spir_func void @_Z3fooPU3AS1i(ptr addrspace(1) noundef) #1
// SPIR: declare spir_func void @_Z3fooPU3AS3i(ptr addrspace(3) noundef) #1
// SPIR: declare spir_func void @_Z3fooPU3AS0i(ptr noundef) #1
+// SPIR: declare spir_func void @_Z3fooPU3AS4i(ptr addrspace(4) noundef) #1
+// SPIR: declare spir_func void @_Z3fooPU3AS2i(ptr addrspace(2) noundef) #1
// SPIR: declare spir_func void @_Z3fooPi(ptr addrspace(4) noundef) #1
// X86: declare void @_Z3fooPU8SYglobali(ptr noundef) #1
// X86: declare void @_Z3fooPU7SYlocali(ptr noundef) #1
// X86: declare void @_Z3fooPU9SYprivatei(ptr noundef) #1
+// X86: declare void @_Z3fooPU9SYgenerici(ptr noundef) #1
+// X86: declare void @_Z3fooPU10SYconstanti(ptr noundef) #1
// X86: declare void @_Z3fooPi(ptr noundef) #1
[[clang::sycl_external]] void test() {
int [[clang::sycl_global]] *glob;
int [[clang::sycl_local]] *loc;
int [[clang::sycl_private]] *priv;
+ int [[clang::sycl_generic]] *gen;
+ int [[clang::sycl_constant]] *cnst;
int *def;
foo(glob);
foo(loc);
foo(priv);
+ foo(gen);
+ foo(cnst);
foo(def);
}
diff --git a/clang/test/SemaSYCL/address-space-conversions.cpp b/clang/test/SemaSYCL/address-space-conversions.cpp
index 41dec2e926418..5afb7748fa434 100644
--- a/clang/test/SemaSYCL/address-space-conversions.cpp
+++ b/clang/test/SemaSYCL/address-space-conversions.cpp
@@ -15,6 +15,7 @@ void usages() {
int [[clang::sycl_global]] *GLOB;
int [[clang::sycl_private]] *PRIV;
int [[clang::sycl_local]] *LOC;
+ int [[clang::sycl_constant]] *CONST; // expected-warning {{'sycl_constant' address space attribute is deprecated}}
int *NoAS;
GLOB = PRIV; // expected-error {{assigning 'sycl_private int *' to 'sycl_global int *' changes address space of pointer}}
@@ -77,4 +78,31 @@ void usages() {
GLOB = GLOB_DEVICE;
GLOB_DEVICE = GLOB; // expected-error {{assigning 'sycl_global int *' to '__global_device int *' changes address space of pointer}}
GLOB_DEVICE = static_cast<__attribute__((opencl_global_device)) int *>(GLOB); // expected-error {{static_cast from 'sycl_global int *' to '__global_device int *' is not allowed}}
+
+ // Test sycl_constant conversions
+ // constant -> constant: OK
+ int [[clang::sycl_constant]] *c2 = CONST; // expected-warning {{'sycl_constant' address space attribute is deprecated}}
+ (void)c2;
+
+ // constant -> other named: ERROR (disjoint address spaces)
+ GLOB = CONST; // expected-error {{assigning 'sycl_constant int *' to 'sycl_global int *' changes address space of pointer}}
+ PRIV = CONST; // expected-error {{assigning 'sycl_constant int *' to 'sycl_private int *' changes address space of pointer}}
+ LOC = CONST; // expected-error {{assigning 'sycl_constant int *' to 'sycl_local int *' changes address space of pointer}}
+
+ // constant -> generic: ERROR (constant not a subset of generic)
+ NoAS = CONST; // expected-error {{assigning 'sycl_constant int *' to 'int *' changes address space of pointer}}
+
+ // generic -> constant: ERROR (constant not a superset of generic)
+ CONST = NoAS; // expected-error {{assigning 'int *' to 'sycl_constant int *' changes address space of pointer}}
+
+ // other named -> constant: ERROR (disjoint address spaces)
+ CONST = GLOB; // expected-error {{assigning 'sycl_global int *' to 'sycl_constant int *' changes address space of pointer}}
+ CONST = PRIV; // expected-error {{assigning 'sycl_private int *' to 'sycl_constant int *' changes address space of pointer}}
+ CONST = LOC; // expected-error {{assigning 'sycl_local int *' to 'sycl_constant int *' changes address space of pointer}}
+
+ // Explicit casts between constant and other spaces are disallowed
+ GLOB = (int [[clang::sycl_global]] *)CONST; // expected-error {{C-style cast from 'sycl_constant int *' to 'sycl_global int *' converts between mismatching address spaces}}
+ CONST = (int [[clang::sycl_constant]] *)GLOB; // expected-warning {{'sycl_constant' address space attribute is deprecated}} expected-error {{C-style cast from 'sycl_global int *' to 'sycl_constant int *' converts between mismatching address spaces}}
+ PRIV = static_cast<int [[clang::sycl_private]] *>(CONST); // expected-error {{static_cast from 'sycl_constant int *' to 'sycl_private int *' is not allowed}}
+ CONST = static_cast<int [[clang::sycl_constant]] *>(PRIV); // expected-warning {{'sycl_constant' address space attribute is deprecated}} expected-error {{static_cast from 'sycl_private int *' to 'sycl_constant int *' is not allowed}}
}
diff --git a/clang/test/SemaTemplate/address_space-dependent.cpp b/clang/test/SemaTemplate/address_space-dependent.cpp
index d6f25923b69b5..f4d2af956a324 100644
--- a/clang/test/SemaTemplate/address_space-dependent.cpp
+++ b/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 (8388579)}}
+ __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388578)}}
}
template <long int I>
@@ -101,7 +101,7 @@ int main() {
car<1, 2, 3>(); // expected-note {{in instantiation of function template specialization 'car<1, 2, 3>' requested here}}
HasASTemplateFields<1> HASTF;
neg<-1>(); // expected-note {{in instantiation of function template specialization 'neg<-1>' requested here}}
- correct<0x7FFFE3>();
+ correct<0x7FFFE2>();
tooBig<8388650>(); // expected-note {{in instantiation of function template specialization 'tooBig<8388650L>' requested here}}
__attribute__((address_space(1))) char *x;
More information about the cfe-commits
mailing list