https://github.com/elizabethandrews created https://github.com/llvm/llvm-project/pull/200849
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]], [[clang::sycl_generic]], [[clang::sycl_global_device]] and [[clang::sycl_global_host]]. LangAS::sycl_generic was added to support SYCL's generic address space. Commit 2 was implemented based on offline discussions to selected LangAS mapping 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 currently even though they technically refer to the same address space. >From 8df10139dac2986fe9df25745fe90e99f1663b33 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews <[email protected]> Date: Fri, 8 May 2026 04:10:12 -0700 Subject: [PATCH 1/3] [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 <[email protected]> Date: Fri, 29 May 2026 12:58:46 -0700 Subject: [PATCH 2/3] 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 <[email protected]> Date: Mon, 1 Jun 2026 08:34:55 -0700 Subject: [PATCH 3/3] 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}} } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
