[clang] [RFC][Clang] Add the __addrspaceof operator (PR #210242)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Tue Aug 11 19:49:35 PDT 2026
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/210242
>From 2f100f22dd72daea646f1d711fef35994bc54e6b Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Thu, 16 Jul 2026 13:53:21 -0400
Subject: [PATCH 1/7] [RFC][Clang] Add __builtin_pointee_address_space
Performance-sensitive code may need a compile-time address-space value to
choose an address-space-specific operation, overload, or template
specialization. This is easy to express when the address space is part of
the pointer type, as in OpenCL. It is harder in CUDA/HIP, where
source-level pointers are usually generic and device, shared, and
constant storage are represented by declaration attributes.
Add __builtin_pointee_address_space. It accepts a pointer or array
expression and returns a Clang address-space identifier. Function
designators are not accepted through function-to-pointer decay. Clang
predefines __CLANG_ADDRESS_SPACE_* macros for language-level address
spaces and emits them during preprocessing. CUDA and HIP use separate
predefined macro names for their declaration address spaces. Explicit
__attribute__((address_space(N))) address spaces are encoded as
__CLANG_ADDRESS_SPACE_TARGET_OFFSET + N.
For OpenCL and explicit address-space-qualified pointer types, the
builtin reports the AST pointee type address space. For CUDA/HIP,
declaration attributes are used only for direct array variable
expressions and direct address-of expressions naming non-array variables.
Other pointer expressions, such as pointer arithmetic, &array, casts, and
parameters, are treated by their expression type.
The result is a static frontend query. It is not runtime pointer
classification and does not use optimizer or backend analysis.
---
clang/docs/LanguageExtensions.md | 119 ++++++++
clang/docs/ReleaseNotes.md | 8 +
clang/include/clang/Basic/AddressSpaces.h | 106 +++++++
clang/include/clang/Basic/Builtins.td | 7 +
.../clang/Basic/DiagnosticSemaKinds.td | 3 +
clang/lib/AST/ExprConstant.cpp | 58 ++++
clang/lib/CodeGen/CGBuiltin.cpp | 4 +
clang/lib/Frontend/InitPreprocessor.cpp | 68 +++++
clang/lib/Sema/SemaChecking.cpp | 14 +
.../CodeGen/builtin-pointee-address-space.c | 42 +++
.../builtin-pointee-address-space.cu | 260 ++++++++++++++++++
.../builtin-pointee-address-space.cl | 18 ++
.../test/Preprocessor/address-space-macros.c | 33 +++
clang/test/Preprocessor/init-aarch64.c | 31 +++
clang/test/Preprocessor/init.c | 31 +++
.../SemaCUDA/builtin-pointee-address-space.cu | 43 +++
.../SemaCXX/builtin-pointee-address-space.cpp | 91 ++++++
17 files changed, 936 insertions(+)
create mode 100644 clang/test/CodeGen/builtin-pointee-address-space.c
create mode 100644 clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
create mode 100644 clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl
create mode 100644 clang/test/Preprocessor/address-space-macros.c
create mode 100644 clang/test/SemaCUDA/builtin-pointee-address-space.cu
create mode 100644 clang/test/SemaCXX/builtin-pointee-address-space.cpp
diff --git a/clang/docs/LanguageExtensions.md b/clang/docs/LanguageExtensions.md
index ee0b272c0a102..044ce29e8b38a 100644
--- a/clang/docs/LanguageExtensions.md
+++ b/clang/docs/LanguageExtensions.md
@@ -3983,6 +3983,125 @@ template<typename T> constexpr T *addressof(T &value) {
}
```
+### `__builtin_pointee_address_space`
+
+`__builtin_pointee_address_space` returns a Clang address-space identifier for
+the storage reached by a pointer or array expression.
+
+This is useful for performance-sensitive code that needs a compile-time value
+to choose an address-space-specific operation, overload, or template
+specialization. For example, CUDA/HIP code may want to select different helper
+code for a pointer to global, shared, or constant memory. The physical target
+address-space number can be target-specific and is not always defined by the
+source language, so this builtin returns Clang-defined values instead of
+backend address-space numbers for language-level address spaces.
+
+**Syntax**:
+
+```c++
+int __builtin_pointee_address_space(pointer-or-array)
+```
+
+The argument must have pointer or array type. Function designators are not
+accepted through function-to-pointer decay. The argument is not evaluated and
+is not converted to `void *` before its address space is queried.
+
+The result is an integer constant expression. It is based on the address space
+that Clang can determine from the expression in the frontend. It does not use
+optimizer or interprocedural analysis to infer a more precise address space.
+
+Clang predefines named macros for the values returned by this builtin. These
+macros are emitted by Clang during preprocessing. For language-level address
+spaces, such as OpenCL address spaces or CUDA/HIP
+`__device__`, `__shared__`, and `__constant__` variables, the result is one of
+the `__CLANG_ADDRESS_SPACE_*` values. For an address space written explicitly
+with `__attribute__((address_space(N)))`, the result is
+`__CLANG_ADDRESS_SPACE_TARGET_OFFSET + N`.
+CUDA and HIP use separate predefined macro names for their declaration address
+spaces. CUDA mode uses the `__CLANG_ADDRESS_SPACE_CUDA_*` values, and HIP mode
+uses the `__CLANG_ADDRESS_SPACE_HIP_*` values.
+
+For languages such as OpenCL, where address spaces are represented in AST
+types, the builtin returns the Clang address-space value of the pointee type of
+the expression as written. For an array expression, the pointee type is the
+array element type.
+
+CUDA/HIP variables are not represented as address-space-qualified pointer types
+in Clang's AST. For CUDA/HIP, the builtin uses variable declaration attributes
+only for two direct forms: an array variable expression, such as `arr`, and an
+address-of expression naming a non-array variable, such as `&var`. In those
+forms, the builtin can report the declaration address space from attributes
+such as `__device__`, `__shared__`, or `__constant__`. An explicit
+`__device__` `const` global or static data member that is promoted to constant
+memory is reported as the CUDA or HIP constant address space. This also works
+in host compilation because it does not depend on the auxiliary target's
+physical address-space map.
+
+Explicit user-written casts are respected. If a cast changes the pointer type
+seen by the builtin, the builtin reports the pointee address space of the cast
+type rather than looking through the cast to recover the original object.
+Other pointer expressions, such as pointer arithmetic or `&array`, are handled
+from the pointee type of the expression as written.
+
+The builtin is a static query. It does not classify an arbitrary runtime pointer
+value. For a parameter such as `int *p`, if the pointee type is generic/default,
+the result is `__CLANG_ADDRESS_SPACE_DEFAULT` even if the runtime value of `p`
+later points into a more specific memory region. Runtime pointer-value
+classification should use target-specific runtime interfaces instead.
+
+The CUDA/HIP declaration query is direct. It does not follow function
+parameters, even through a `consteval` or `constexpr` wrapper. Inside such a
+wrapper, the builtin reports the address space of the parameter type.
+
+**Example use**:
+
+```c++
+int *p;
+int __attribute__((address_space(3))) *p3;
+
+static_assert(__builtin_pointee_address_space(p) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(p3) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 3);
+```
+
+**OpenCL example use**:
+
+```c
+__global int *global_p;
+__local int *local_p;
+
+static_assert(__builtin_pointee_address_space(global_p) ==
+ __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL);
+static_assert(__builtin_pointee_address_space(local_p) ==
+ __CLANG_ADDRESS_SPACE_OPENCL_LOCAL);
+```
+
+**CUDA/HIP example use**:
+
+```c++
+__device__ int dev;
+__device__ int dev_arr[4];
+__constant__ int cst;
+__device__ const int dev_cst = 1;
+```
+
+In CUDA mode, `&dev` and `dev_arr` return
+`__CLANG_ADDRESS_SPACE_CUDA_DEVICE`, while `&cst` and `&dev_cst` return
+`__CLANG_ADDRESS_SPACE_CUDA_CONSTANT`. In HIP mode, they return the matching
+`__CLANG_ADDRESS_SPACE_HIP_*` values.
+
+The direct-query rule is not specific to shared variables. For a CUDA/HIP
+variable with a declaration address-space attribute, `&a` for a non-array
+variable and `b` for an array variable are direct declaration queries. For
+example, `&a` for `__shared__ int a` and `b` for `__shared__ int b[2]` report
+the CUDA or HIP shared address space.
+
+Forms such as `(void *)&a`, `&a + 1`, `&b[0]`, `&dev_arr`, `dev_arr + 1`,
+and calls through a function parameter, including `consteval` wrappers, are
+not direct declaration queries. They use the expression type and report
+`__CLANG_ADDRESS_SPACE_DEFAULT`.
+
### `__builtin_function_start`
`__builtin_function_start` returns the address of a function body.
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index b5aa71119e4b7..6b75f6914c9eb 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -182,6 +182,14 @@ features cannot lower the translation-unit ABI level;
- Clang now allows GNU computed `goto` extension in `constexpr` functions, matching the relaxed
`constexpr` function body rules introduced in C++23.
+- Added `__builtin_pointee_address_space`, which returns a Clang
+ address-space identifier for a pointer or array expression. It reports the
+ AST pointee type address space for OpenCL and explicit address-space-qualified
+ pointer types, and can report CUDA/HIP declaration address spaces for known
+ `__device__`, `__shared__`, and `__constant__` variables. Clang also now
+ emits predefined `__CLANG_ADDRESS_SPACE_*` macros for these values, including
+ separate CUDA and HIP macro names.
+
### New Compiler Flags
- New option `-fdefined-pointer-subtraction` added to preserve stable semantics
diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h
index ce38f60163b64..7c6bc7a9c2180 100644
--- a/clang/include/clang/Basic/AddressSpaces.h
+++ b/clang/include/clang/Basic/AddressSpaces.h
@@ -118,6 +118,112 @@ inline bool isPtrSizeAddressSpace(LangAS AS) {
AS == LangAS::ptr64);
}
+namespace PointeeAddressSpace {
+
+enum ID : unsigned {
+ Default = 0,
+ OpenCLGlobal = 1,
+ OpenCLLocal = 2,
+ OpenCLConstant = 3,
+ OpenCLPrivate = 4,
+ OpenCLGeneric = 5,
+ OpenCLGlobalDevice = 6,
+ OpenCLGlobalHost = 7,
+ CUDADevice = 8,
+ CUDAConstant = 9,
+ CUDAShared = 10,
+ SYCLGlobal = 11,
+ SYCLGlobalDevice = 12,
+ SYCLGlobalHost = 13,
+ SYCLLocal = 14,
+ SYCLPrivate = 15,
+ Ptr32Sptr = 16,
+ Ptr32Uptr = 17,
+ Ptr64 = 18,
+ HLSLGroupShared = 19,
+ HLSLConstant = 20,
+ HLSLPrivate = 21,
+ HLSLDevice = 22,
+ HLSLInput = 23,
+ HLSLOutput = 24,
+ HLSLPushConstant = 25,
+ WasmFuncRef = 26,
+ HIPDevice = 27,
+ HIPConstant = 28,
+ HIPShared = 29,
+
+ TargetOffset = 0x1000000
+};
+
+inline unsigned encode(LangAS AS, bool IsHIP = false) {
+ if (isTargetAddressSpace(AS))
+ return TargetOffset + toTargetAddressSpace(AS);
+
+ switch (AS) {
+ case LangAS::Default:
+ return Default;
+ case LangAS::opencl_global:
+ return OpenCLGlobal;
+ case LangAS::opencl_local:
+ return OpenCLLocal;
+ case LangAS::opencl_constant:
+ return OpenCLConstant;
+ case LangAS::opencl_private:
+ return OpenCLPrivate;
+ case LangAS::opencl_generic:
+ return OpenCLGeneric;
+ case LangAS::opencl_global_device:
+ return OpenCLGlobalDevice;
+ case LangAS::opencl_global_host:
+ return OpenCLGlobalHost;
+ case LangAS::cuda_device:
+ return IsHIP ? HIPDevice : CUDADevice;
+ case LangAS::cuda_constant:
+ return IsHIP ? HIPConstant : CUDAConstant;
+ case LangAS::cuda_shared:
+ return IsHIP ? HIPShared : CUDAShared;
+ case LangAS::sycl_global:
+ return SYCLGlobal;
+ case LangAS::sycl_global_device:
+ return SYCLGlobalDevice;
+ case LangAS::sycl_global_host:
+ return SYCLGlobalHost;
+ case LangAS::sycl_local:
+ return SYCLLocal;
+ case LangAS::sycl_private:
+ return SYCLPrivate;
+ case LangAS::ptr32_sptr:
+ return Ptr32Sptr;
+ case LangAS::ptr32_uptr:
+ return Ptr32Uptr;
+ case LangAS::ptr64:
+ return Ptr64;
+ case LangAS::hlsl_groupshared:
+ return HLSLGroupShared;
+ case LangAS::hlsl_constant:
+ return HLSLConstant;
+ case LangAS::hlsl_private:
+ return HLSLPrivate;
+ case LangAS::hlsl_device:
+ return HLSLDevice;
+ case LangAS::hlsl_input:
+ return HLSLInput;
+ case LangAS::hlsl_output:
+ return HLSLOutput;
+ case LangAS::hlsl_push_constant:
+ return HLSLPushConstant;
+ case LangAS::wasm_funcref:
+ return WasmFuncRef;
+ case LangAS::FirstTargetAddressSpace:
+ break;
+ }
+
+ assert(false && "unknown language address space");
+ return Default;
+}
+
+} // namespace PointeeAddressSpace
+
} // namespace clang
#endif // LLVM_CLANG_BASIC_ADDRESSSPACES_H
diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td
index a54f91069acd0..67dee1d024431 100644
--- a/clang/include/clang/Basic/Builtins.td
+++ b/clang/include/clang/Basic/Builtins.td
@@ -1046,6 +1046,13 @@ def BuiltinClassifyType : Builtin {
let Prototype = "int(...)";
}
+def BuiltinPointeeAddressSpace : Builtin {
+ let Spellings = ["__builtin_pointee_address_space"];
+ let Attributes = [NoThrow, Const, CustomTypeChecking, UnevaluatedArguments,
+ Constexpr];
+ let Prototype = "int(...)";
+}
+
def BuiltinCFStringMakeConstantString : Builtin {
let Spellings = ["__builtin___CFStringMakeConstantString"];
let Attributes = [NoThrow, Const, Constexpr];
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index b314c17ad27bd..0c5e099cea882 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -864,6 +864,9 @@ def warn_redecl_library_builtin : Warning<
def err_builtin_definition : Error<"definition of builtin function %0">;
def err_builtin_redeclare : Error<"cannot redeclare builtin function %0">;
def err_invalid_builtin_argument : Error<"invalid argument '%0' to %1">;
+def err_builtin_pointee_address_space_arg_not_pointer : Error<
+ "argument to '__builtin_pointee_address_space' must be a pointer or array "
+ "expression">;
def err_arm_invalid_specialreg : Error<"invalid special register for builtin">;
def err_arm_invalid_coproc : Error<"coprocessor %0 must be configured as "
diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp
index 480d5119a5363..fed6b645b1e58 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/clang/lib/AST/ExprConstant.cpp
@@ -15937,6 +15937,54 @@ bool ArrayExprEvaluator::VisitDesignatedInitUpdateExpr(
//===----------------------------------------------------------------------===//
namespace {
+
+static QualType getPointeeAddressSpaceType(ASTContext &Ctx, QualType ArgTy) {
+ return ArgTy->isArrayType() ? Ctx.getAsArrayType(ArgTy)->getElementType()
+ : ArgTy->getPointeeType();
+}
+
+static std::optional<LangAS> getCUDAPointeeDeclAddressSpace(ASTContext &Ctx,
+ const Expr *Arg) {
+ if (!Ctx.getLangOpts().CUDA)
+ return std::nullopt;
+
+ Arg = Arg->IgnoreParens();
+ if (isa<ExplicitCastExpr>(Arg))
+ return std::nullopt;
+
+ const VarDecl *VD = nullptr;
+ const Expr *NoImpCasts = Arg->IgnoreImpCasts()->IgnoreParens();
+ if (const auto *UO = dyn_cast<UnaryOperator>(NoImpCasts);
+ UO && UO->getOpcode() == UO_AddrOf) {
+ if (const auto *DRE =
+ dyn_cast<DeclRefExpr>(UO->getSubExpr()->IgnoreParenImpCasts())) {
+ const auto *CandidateVD = dyn_cast<VarDecl>(DRE->getDecl());
+ if (CandidateVD && !CandidateVD->getType()->isArrayType())
+ VD = CandidateVD;
+ }
+ } else if (const auto *DRE = dyn_cast<DeclRefExpr>(NoImpCasts)) {
+ if (NoImpCasts->getType()->isArrayType())
+ VD = dyn_cast<VarDecl>(DRE->getDecl());
+ }
+
+ if (!VD)
+ return std::nullopt;
+ if (VD->hasAttr<CUDAConstantAttr>())
+ return LangAS::cuda_constant;
+ if (VD->hasAttr<CUDASharedAttr>())
+ return LangAS::cuda_shared;
+ // Host compilation does not attach the implicit CUDAConstantAttr that
+ // SemaCUDA adds in device compilation, but the device-side storage is still
+ // constant memory.
+ if (!Ctx.getLangOpts().CUDAIsDevice && VD->hasAttr<CUDADeviceAttr>() &&
+ (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
+ (VD->isConstexpr() || VD->getType().isConstQualified()))
+ return LangAS::cuda_constant;
+ if (VD->hasAttr<CUDADeviceAttr>())
+ return LangAS::cuda_device;
+ return std::nullopt;
+}
+
class IntExprEvaluator
: public ExprEvaluatorBase<IntExprEvaluator> {
APValue &Result;
@@ -17003,6 +17051,16 @@ bool IntExprEvaluator::VisitBuiltinCallExpr(const CallExpr *E,
llvm_unreachable("unexpected EvalMode");
}
+ case Builtin::BI__builtin_pointee_address_space: {
+ QualType ArgTy = E->getArg(0)->getType();
+ LangAS AS =
+ getCUDAPointeeDeclAddressSpace(Info.Ctx, E->getArg(0))
+ .value_or(
+ getPointeeAddressSpaceType(Info.Ctx, ArgTy).getAddressSpace());
+ return Success(PointeeAddressSpace::encode(AS, Info.Ctx.getLangOpts().HIP),
+ E);
+ }
+
case Builtin::BI__builtin_os_log_format_buffer_size: {
analyze_os_log::OSLogBufferLayout Layout;
analyze_os_log::computeOSLogBufferLayout(Info.Ctx, E, Layout);
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 3521fc10f1387..ad91a9c5f7bc8 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -4261,6 +4261,10 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
Result = Builder.CreateIntCast(Result, ResultType, /*isSigned*/false);
return RValue::get(Result);
}
+ case Builtin::BI__builtin_pointee_address_space: {
+ unsigned AS = E->EvaluateKnownConstInt(getContext()).getZExtValue();
+ return RValue::get(ConstantInt::get(ConvertType(E->getType()), AS));
+ }
case Builtin::BI__builtin_dynamic_object_size:
case Builtin::BI__builtin_object_size: {
unsigned Type =
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 349628c3439bf..604a04e597db8 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -10,6 +10,7 @@
//
//===----------------------------------------------------------------------===//
+#include "clang/Basic/AddressSpaces.h"
#include "clang/Basic/DiagnosticFrontend.h"
#include "clang/Basic/DiagnosticLex.h"
#include "clang/Basic/HLSLRuntime.h"
@@ -914,6 +915,73 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
Builder.defineMacro("__OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES", "3");
Builder.defineMacro("__OPENCL_MEMORY_SCOPE_SUB_GROUP", "4");
+ auto DefinePointeeASMacro = [&](StringRef Name, PointeeAddressSpace::ID AS) {
+ Builder.defineMacro(Name, Twine(static_cast<unsigned>(AS)));
+ };
+
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_DEFAULT",
+ PointeeAddressSpace::Default);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL",
+ PointeeAddressSpace::OpenCLGlobal);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_LOCAL",
+ PointeeAddressSpace::OpenCLLocal);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_CONSTANT",
+ PointeeAddressSpace::OpenCLConstant);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_PRIVATE",
+ PointeeAddressSpace::OpenCLPrivate);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GENERIC",
+ PointeeAddressSpace::OpenCLGeneric);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_DEVICE",
+ PointeeAddressSpace::OpenCLGlobalDevice);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_HOST",
+ PointeeAddressSpace::OpenCLGlobalHost);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_CUDA_DEVICE",
+ PointeeAddressSpace::CUDADevice);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_CUDA_CONSTANT",
+ PointeeAddressSpace::CUDAConstant);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_CUDA_SHARED",
+ PointeeAddressSpace::CUDAShared);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL",
+ PointeeAddressSpace::SYCLGlobal);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL_DEVICE",
+ PointeeAddressSpace::SYCLGlobalDevice);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL_HOST",
+ PointeeAddressSpace::SYCLGlobalHost);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_LOCAL",
+ PointeeAddressSpace::SYCLLocal);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_PRIVATE",
+ PointeeAddressSpace::SYCLPrivate);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_PTR32_SPTR",
+ PointeeAddressSpace::Ptr32Sptr);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_PTR32_UPTR",
+ PointeeAddressSpace::Ptr32Uptr);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_PTR64",
+ PointeeAddressSpace::Ptr64);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED",
+ PointeeAddressSpace::HLSLGroupShared);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_CONSTANT",
+ PointeeAddressSpace::HLSLConstant);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_PRIVATE",
+ PointeeAddressSpace::HLSLPrivate);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_DEVICE",
+ PointeeAddressSpace::HLSLDevice);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_INPUT",
+ PointeeAddressSpace::HLSLInput);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_OUTPUT",
+ PointeeAddressSpace::HLSLOutput);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT",
+ PointeeAddressSpace::HLSLPushConstant);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_WASM_FUNCREF",
+ PointeeAddressSpace::WasmFuncRef);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HIP_DEVICE",
+ PointeeAddressSpace::HIPDevice);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HIP_CONSTANT",
+ PointeeAddressSpace::HIPConstant);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HIP_SHARED",
+ PointeeAddressSpace::HIPShared);
+ DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_TARGET_OFFSET",
+ PointeeAddressSpace::TargetOffset);
+
// Define macros for floating-point data classes, used in __builtin_isfpclass.
Builder.defineMacro("__FPCLASS_SNAN", "0x0001");
Builder.defineMacro("__FPCLASS_QNAN", "0x0002");
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index f2f38c84dc5f8..537e952443d0d 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -3284,6 +3284,20 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID,
if (BuiltinComplex(TheCall))
return ExprError();
break;
+ case Builtin::BI__builtin_pointee_address_space: {
+ if (checkArgCount(TheCall, 1))
+ return true;
+ Expr *Arg = TheCall->getArg(0);
+ if (!Arg->isTypeDependent() && !Arg->getType()->isPointerType() &&
+ !Arg->getType()->isArrayType()) {
+ Diag(Arg->getBeginLoc(),
+ diag::err_builtin_pointee_address_space_arg_not_pointer)
+ << Arg->getSourceRange();
+ return ExprError();
+ }
+ TheCall->setType(Context.IntTy);
+ break;
+ }
case Builtin::BI__builtin_classify_type:
case Builtin::BI__builtin_constant_p: {
if (checkArgCount(TheCall, 1))
diff --git a/clang/test/CodeGen/builtin-pointee-address-space.c b/clang/test/CodeGen/builtin-pointee-address-space.c
new file mode 100644
index 0000000000000..7fe5d48c74b9c
--- /dev/null
+++ b/clang/test/CodeGen/builtin-pointee-address-space.c
@@ -0,0 +1,42 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | FileCheck %s
+
+int test_default(int *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_default(
+// CHECK: ret i32 0
+
+int test_address_space(int __attribute__((address_space(4))) *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_address_space(
+// CHECK: ret i32 16777220
+
+int __attribute__((address_space(7))) *side_effect(void);
+
+int test_unevaluated(void) {
+ return __builtin_pointee_address_space(side_effect());
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_unevaluated(
+// CHECK-NOT: call
+// CHECK: ret i32 16777223
+
+int array_default[4];
+int __attribute__((address_space(5))) array_address_space[4];
+
+int test_array_default(void) {
+ return __builtin_pointee_address_space(array_default);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_array_default(
+// CHECK: ret i32 0
+
+int test_array_address_space(void) {
+ return __builtin_pointee_address_space(array_address_space);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_array_address_space(
+// CHECK: ret i32 16777221
diff --git a/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu b/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
new file mode 100644
index 0000000000000..e51548b59bbe6
--- /dev/null
+++ b/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
@@ -0,0 +1,260 @@
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++20 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,CUDA %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip -std=c++20 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,HIP %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -x hip -std=c++20 -DHOST_TEST -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s
+
+#include "Inputs/cuda.h"
+
+__device__ int device_var;
+__device__ int device_array[4];
+__device__ int *device_ptr;
+__constant__ int constant_var;
+__device__ const int const_device_var = 1;
+
+#if defined(__HIP__)
+#define EXPECTED_DEVICE_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_HIP_DEVICE
+#define EXPECTED_CONSTANT_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_HIP_CONSTANT
+#define EXPECTED_SHARED_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_HIP_SHARED
+#else
+#define EXPECTED_DEVICE_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_CUDA_DEVICE
+#define EXPECTED_CONSTANT_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_CUDA_CONSTANT
+#define EXPECTED_SHARED_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_CUDA_SHARED
+#endif
+
+#ifdef HOST_TEST
+
+template <class T> consteval int host_consteval_address_space(T *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+template <class T> constexpr int host_constexpr_address_space(T *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+static_assert(__builtin_pointee_address_space(&device_var) ==
+ EXPECTED_DEVICE_ADDRESS_SPACE);
+static_assert(__builtin_pointee_address_space(&constant_var) ==
+ EXPECTED_CONSTANT_ADDRESS_SPACE);
+static_assert(__builtin_pointee_address_space(&const_device_var) ==
+ EXPECTED_CONSTANT_ADDRESS_SPACE);
+static_assert(__builtin_pointee_address_space(device_array) ==
+ EXPECTED_DEVICE_ADDRESS_SPACE);
+static_assert(__builtin_pointee_address_space(&device_array) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(device_array + 1) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(host_consteval_address_space(&device_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(host_consteval_address_space(&constant_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(host_consteval_address_space(&const_device_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(host_consteval_address_space((int *)&constant_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(host_constexpr_address_space(&constant_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+
+extern "C" int test_host_device_var() {
+ return __builtin_pointee_address_space(&device_var);
+}
+
+// HOST-LABEL: define{{.*}} i32 @test_host_device_var(
+// HOST: ret i32 27
+
+extern "C" int test_host_constant_var() {
+ return __builtin_pointee_address_space(&constant_var);
+}
+
+// HOST-LABEL: define{{.*}} i32 @test_host_constant_var(
+// HOST: ret i32 28
+
+extern "C" int test_host_const_device_var() {
+ return __builtin_pointee_address_space(&const_device_var);
+}
+
+// HOST-LABEL: define{{.*}} i32 @test_host_const_device_var(
+// HOST: ret i32 28
+
+extern "C" int test_host_device_array_address() {
+ return __builtin_pointee_address_space(&device_array);
+}
+
+// HOST-LABEL: define{{.*}} i32 @test_host_device_array_address(
+// HOST: ret i32 0
+
+#else
+
+template <class T> consteval int consteval_address_space(T *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+template <class T> constexpr int constexpr_address_space(T *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+template <int AS> struct AddressSpaceSpecialization;
+template <>
+struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_DEFAULT> {
+ static constexpr int value = __CLANG_ADDRESS_SPACE_DEFAULT;
+};
+template <> struct AddressSpaceSpecialization<EXPECTED_DEVICE_ADDRESS_SPACE> {
+ static constexpr int value = EXPECTED_DEVICE_ADDRESS_SPACE;
+};
+template <> struct AddressSpaceSpecialization<EXPECTED_SHARED_ADDRESS_SPACE> {
+ static constexpr int value = EXPECTED_SHARED_ADDRESS_SPACE;
+};
+template <>
+struct AddressSpaceSpecialization<EXPECTED_CONSTANT_ADDRESS_SPACE> {
+ static constexpr int value = EXPECTED_CONSTANT_ADDRESS_SPACE;
+};
+
+static_assert(__builtin_pointee_address_space((int *)&constant_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(consteval_address_space(&device_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(consteval_address_space(&constant_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(consteval_address_space(&const_device_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(consteval_address_space((int *)&constant_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(constexpr_address_space(&constant_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(device_array) ==
+ EXPECTED_DEVICE_ADDRESS_SPACE);
+static_assert(__builtin_pointee_address_space(&device_array) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(device_array + 1) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(device_ptr) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(
+ AddressSpaceSpecialization<
+ __builtin_pointee_address_space(&device_var)>::value ==
+ EXPECTED_DEVICE_ADDRESS_SPACE);
+static_assert(
+ AddressSpaceSpecialization<
+ __builtin_pointee_address_space(&constant_var)>::value ==
+ EXPECTED_CONSTANT_ADDRESS_SPACE);
+static_assert(
+ AddressSpaceSpecialization<
+ consteval_address_space(&constant_var)>::value ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+
+extern "C" __device__ int test_generic_pointer(int *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_generic_pointer(
+// CHECK: ret i32 0
+
+extern "C" __device__ int test_shared_local() {
+ __shared__ int shared_var;
+ __shared__ int shared_array[4];
+ static_assert(__builtin_pointee_address_space(&shared_var) ==
+ EXPECTED_SHARED_ADDRESS_SPACE);
+ static_assert(__builtin_pointee_address_space((void *)&shared_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+ static_assert(__builtin_pointee_address_space(&shared_var + 1) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+ static_assert(__builtin_pointee_address_space(shared_array) ==
+ EXPECTED_SHARED_ADDRESS_SPACE);
+ static_assert(__builtin_pointee_address_space(&shared_array) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+ static_assert(__builtin_pointee_address_space(&shared_array[0]) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+ static_assert(__builtin_pointee_address_space(shared_array + 1) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+ static_assert(consteval_address_space(&shared_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+ static_assert(consteval_address_space(shared_array) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+ static_assert(constexpr_address_space(&shared_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+ static_assert(
+ AddressSpaceSpecialization<
+ __builtin_pointee_address_space(&shared_var)>::value ==
+ EXPECTED_SHARED_ADDRESS_SPACE);
+ return __builtin_pointee_address_space(&shared_var);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_shared_local(
+// CUDA: ret i32 10
+// HIP: ret i32 29
+
+extern "C" __device__ int test_device_var() {
+ return __builtin_pointee_address_space(&device_var);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_device_var(
+// CUDA: ret i32 8
+// HIP: ret i32 27
+
+extern "C" __device__ int test_device_array() {
+ return __builtin_pointee_address_space(device_array);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_device_array(
+// CUDA: ret i32 8
+// HIP: ret i32 27
+
+extern "C" __device__ int test_device_array_address() {
+ return __builtin_pointee_address_space(&device_array);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_device_array_address(
+// CHECK: ret i32 0
+
+extern "C" __device__ int test_device_array_arithmetic() {
+ return __builtin_pointee_address_space(device_array + 1);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_device_array_arithmetic(
+// CHECK: ret i32 0
+
+extern "C" __device__ int test_device_pointer_value() {
+ return __builtin_pointee_address_space(device_ptr);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_device_pointer_value(
+// CHECK: ret i32 0
+
+extern "C" __device__ int test_constant_var() {
+ return __builtin_pointee_address_space(&constant_var);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_constant_var(
+// CUDA: ret i32 9
+// HIP: ret i32 28
+
+extern "C" __device__ int test_const_device_var() {
+ return __builtin_pointee_address_space(&const_device_var);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_const_device_var(
+// CUDA: ret i32 9
+// HIP: ret i32 28
+
+extern "C" __device__ int test_explicit_cast() {
+ return __builtin_pointee_address_space((int *)&constant_var);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_explicit_cast(
+// CHECK: ret i32 0
+
+extern "C" __device__ int
+test_target_address_space_3(int __attribute__((address_space(3))) *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_target_address_space_3(
+// CHECK: ret i32 16777219
+
+extern "C" __device__ int
+test_target_address_space_4(int __attribute__((address_space(4))) *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @test_target_address_space_4(
+// CHECK: ret i32 16777220
+
+#endif
diff --git a/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl b/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl
new file mode 100644
index 0000000000000..3f002801c15d7
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl
@@ -0,0 +1,18 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -emit-llvm -o - %s | FileCheck %s
+
+int global_as(__global int *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @global_as(
+// CHECK-SAME: ptr addrspace(1)
+// CHECK: ret i32 1
+
+int local_as(__local int *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @local_as(
+// CHECK-SAME: ptr addrspace(3)
+// CHECK: ret i32 2
diff --git a/clang/test/Preprocessor/address-space-macros.c b/clang/test/Preprocessor/address-space-macros.c
new file mode 100644
index 0000000000000..8e2baefe5333c
--- /dev/null
+++ b/clang/test/Preprocessor/address-space-macros.c
@@ -0,0 +1,33 @@
+// RUN: %clang_cc1 -E -dM %s | FileCheck %s
+
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_DEFAULT 0
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL 1
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_LOCAL 2
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_DEVICE 6
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_HOST 7
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 8
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 9
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_SHARED 10
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 11
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_DEVICE 12
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_HOST 13
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 14
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 15
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR32_SPTR 16
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR32_UPTR 17
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR64 18
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 19
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 20
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 21
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 22
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_INPUT 23
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 24
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 25
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 26
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_DEVICE 27
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 28
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_SHARED 29
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
diff --git a/clang/test/Preprocessor/init-aarch64.c b/clang/test/Preprocessor/init-aarch64.c
index 44b92ea2331b0..c596f79512208 100644
--- a/clang/test/Preprocessor/init-aarch64.c
+++ b/clang/test/Preprocessor/init-aarch64.c
@@ -52,6 +52,37 @@
// AARCH64-NEXT: #define __CHAR16_TYPE__ unsigned short
// AARCH64-NEXT: #define __CHAR32_TYPE__ unsigned int
// AARCH64-NEXT: #define __CHAR_BIT__ 8
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 9
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 8
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_SHARED 10
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_DEFAULT 0
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 28
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_DEVICE 27
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_SHARED 29
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 20
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 22
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 19
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_INPUT 23
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 24
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 21
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 25
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL 1
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_DEVICE 6
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_HOST 7
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_LOCAL 2
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR32_SPTR 16
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR32_UPTR 17
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR64 18
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 11
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_DEVICE 12
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_HOST 13
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 14
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 15
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 26
// AARCH64-NEXT: #define __CLANG_ATOMIC_BOOL_LOCK_FREE 2
// AARCH64-NEXT: #define __CLANG_ATOMIC_CHAR16_T_LOCK_FREE 2
// AARCH64-NEXT: #define __CLANG_ATOMIC_CHAR32_T_LOCK_FREE 2
diff --git a/clang/test/Preprocessor/init.c b/clang/test/Preprocessor/init.c
index 21d61de8c0aaf..ae8f6c3980a09 100644
--- a/clang/test/Preprocessor/init.c
+++ b/clang/test/Preprocessor/init.c
@@ -1716,6 +1716,37 @@
// WEBASSEMBLY-NEXT:#define __CHAR32_TYPE__ unsigned int
// WEBASSEMBLY-NEXT:#define __CHAR_BIT__ 8
// WEBASSEMBLY-NOT:#define __CHAR_UNSIGNED__
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 9
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 8
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_SHARED 10
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_DEFAULT 0
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 28
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_DEVICE 27
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_SHARED 29
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 20
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 22
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 19
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_INPUT 23
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 24
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 21
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 25
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL 1
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_DEVICE 6
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_HOST 7
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_LOCAL 2
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR32_SPTR 16
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR32_UPTR 17
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR64 18
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 11
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_DEVICE 12
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_HOST 13
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 14
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 15
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 26
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_BOOL_LOCK_FREE 2
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_CHAR16_T_LOCK_FREE 2
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_CHAR32_T_LOCK_FREE 2
diff --git a/clang/test/SemaCUDA/builtin-pointee-address-space.cu b/clang/test/SemaCUDA/builtin-pointee-address-space.cu
new file mode 100644
index 0000000000000..cd3ee9aa3a56b
--- /dev/null
+++ b/clang/test/SemaCUDA/builtin-pointee-address-space.cu
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fsyntax-only -verify=noaux %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -x hip -fsyntax-only -verify=aux %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip -fsyntax-only -verify=device %s
+
+// noaux-no-diagnostics
+// aux-no-diagnostics
+// device-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+__device__ int device_var;
+__device__ int device_array[4];
+__device__ int *device_ptr;
+__constant__ int constant_var;
+__device__ const int const_device_var = 1;
+
+static_assert(__builtin_pointee_address_space(&device_var) ==
+ __CLANG_ADDRESS_SPACE_HIP_DEVICE);
+static_assert(__builtin_pointee_address_space(device_array) ==
+ __CLANG_ADDRESS_SPACE_HIP_DEVICE);
+static_assert(__builtin_pointee_address_space(&device_array) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(device_array + 1) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(&constant_var) ==
+ __CLANG_ADDRESS_SPACE_HIP_CONSTANT);
+static_assert(__builtin_pointee_address_space(&const_device_var) ==
+ __CLANG_ADDRESS_SPACE_HIP_CONSTANT);
+static_assert(__builtin_pointee_address_space((int *)&constant_var) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(device_ptr) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+
+void host_queries() {
+ (void)__builtin_pointee_address_space(&device_var);
+ (void)__builtin_pointee_address_space(device_array);
+ (void)__builtin_pointee_address_space(&device_array);
+ (void)__builtin_pointee_address_space(device_array + 1);
+ (void)__builtin_pointee_address_space(&constant_var);
+ (void)__builtin_pointee_address_space(&const_device_var);
+ (void)__builtin_pointee_address_space((int *)&constant_var);
+ (void)__builtin_pointee_address_space(device_ptr);
+}
diff --git a/clang/test/SemaCXX/builtin-pointee-address-space.cpp b/clang/test/SemaCXX/builtin-pointee-address-space.cpp
new file mode 100644
index 0000000000000..209fc315339be
--- /dev/null
+++ b/clang/test/SemaCXX/builtin-pointee-address-space.cpp
@@ -0,0 +1,91 @@
+// RUN: %clang_cc1 -fsyntax-only -verify -std=c++20 %s
+
+#if !__has_builtin(__builtin_pointee_address_space)
+#error "missing __builtin_pointee_address_space"
+#endif
+
+int *p0;
+int __attribute__((address_space(1))) *p1;
+const int __attribute__((address_space(7))) *p7;
+
+static_assert(__builtin_pointee_address_space(p0) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(p1) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+static_assert(__builtin_pointee_address_space(p7) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 7);
+static_assert(__builtin_pointee_address_space(
+ (int __attribute__((address_space(3))) *)0) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 3);
+
+int global;
+int __attribute__((address_space(4))) global_as4;
+int arr[4];
+int __attribute__((address_space(5))) arr_as5[4];
+
+static_assert(__builtin_pointee_address_space(&global) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(&global_as4) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 4);
+static_assert(__builtin_pointee_address_space(arr) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(arr_as5) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 5);
+
+template <class T> constexpr int get_as(T *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+static_assert(get_as((int *)0) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(get_as((int __attribute__((address_space(1))) *)0) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+
+template <class T> struct PointeeAddressSpace {
+ static constexpr int value =
+ __builtin_pointee_address_space((T *)0);
+};
+
+static_assert(PointeeAddressSpace<int>::value ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(PointeeAddressSpace<
+ int __attribute__((address_space(2)))>::value ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 2);
+
+template <int AS> struct AddressSpaceSpecialization;
+template <>
+struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_DEFAULT> {
+ static constexpr int value = __CLANG_ADDRESS_SPACE_DEFAULT;
+};
+template <>
+struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1> {
+ static constexpr int value = __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1;
+};
+template <>
+struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_TARGET_OFFSET + 5> {
+ static constexpr int value = __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 5;
+};
+
+static_assert(
+ AddressSpaceSpecialization<
+ __builtin_pointee_address_space(p1)>::value ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+static_assert(
+ AddressSpaceSpecialization<
+ __builtin_pointee_address_space(arr_as5)>::value ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 5);
+
+void errors() {
+ int i;
+ void f();
+
+ (void)__builtin_pointee_address_space();
+ // expected-error at -1 {{too few arguments}}
+ (void)__builtin_pointee_address_space(p0, p1);
+ // expected-error at -1 {{too many arguments}}
+ (void)__builtin_pointee_address_space(i);
+ // expected-error at -1 {{argument to '__builtin_pointee_address_space' must be a pointer or array expression}}
+ (void)__builtin_pointee_address_space(0);
+ // expected-error at -1 {{argument to '__builtin_pointee_address_space' must be a pointer or array expression}}
+ (void)__builtin_pointee_address_space(f);
+ // expected-error at -1 {{argument to '__builtin_pointee_address_space' must be a pointer or array expression}}
+}
>From a4a1c5b3ba7da1c10c12b7a9995fd4e7b7541713 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Thu, 23 Jul 2026 23:39:52 -0400
Subject: [PATCH 2/7] [Clang] Drop deprecated pointee address-space values
OpenCL and SYCL global_device and global_host are deprecated and will be removed. Do not expose separate public values for them.
Report them as their corresponding global address space instead.
---
clang/docs/LanguageExtensions.md | 4 +-
clang/include/clang/Basic/AddressSpaces.h | 51 ++++++++-----------
clang/lib/Frontend/InitPreprocessor.cpp | 8 ---
.../builtin-pointee-address-space.cu | 26 +++++-----
.../builtin-pointee-address-space.cl | 18 ++++++-
.../builtin-pointee-address-space.cpp | 23 +++++++++
.../test/Preprocessor/address-space-macros.c | 44 ++++++++--------
clang/test/Preprocessor/init-aarch64.c | 44 ++++++++--------
clang/test/Preprocessor/init.c | 44 ++++++++--------
9 files changed, 138 insertions(+), 124 deletions(-)
create mode 100644 clang/test/CodeGenSYCL/builtin-pointee-address-space.cpp
diff --git a/clang/docs/LanguageExtensions.md b/clang/docs/LanguageExtensions.md
index 044ce29e8b38a..563fc4bca61cc 100644
--- a/clang/docs/LanguageExtensions.md
+++ b/clang/docs/LanguageExtensions.md
@@ -4024,7 +4024,9 @@ uses the `__CLANG_ADDRESS_SPACE_HIP_*` values.
For languages such as OpenCL, where address spaces are represented in AST
types, the builtin returns the Clang address-space value of the pointee type of
the expression as written. For an array expression, the pointee type is the
-array element type.
+array element type. The deprecated OpenCL and SYCL `global_device` and
+`global_host` address spaces are reported as their corresponding global address
+space.
CUDA/HIP variables are not represented as address-space-qualified pointer types
in Clang's AST. For CUDA/HIP, the builtin uses variable declaration attributes
diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h
index 7c6bc7a9c2180..df490e60ed493 100644
--- a/clang/include/clang/Basic/AddressSpaces.h
+++ b/clang/include/clang/Basic/AddressSpaces.h
@@ -127,30 +127,26 @@ enum ID : unsigned {
OpenCLConstant = 3,
OpenCLPrivate = 4,
OpenCLGeneric = 5,
- OpenCLGlobalDevice = 6,
- OpenCLGlobalHost = 7,
- CUDADevice = 8,
- CUDAConstant = 9,
- CUDAShared = 10,
- SYCLGlobal = 11,
- SYCLGlobalDevice = 12,
- SYCLGlobalHost = 13,
- SYCLLocal = 14,
- SYCLPrivate = 15,
- Ptr32Sptr = 16,
- Ptr32Uptr = 17,
- Ptr64 = 18,
- HLSLGroupShared = 19,
- HLSLConstant = 20,
- HLSLPrivate = 21,
- HLSLDevice = 22,
- HLSLInput = 23,
- HLSLOutput = 24,
- HLSLPushConstant = 25,
- WasmFuncRef = 26,
- HIPDevice = 27,
- HIPConstant = 28,
- HIPShared = 29,
+ CUDADevice = 6,
+ CUDAConstant = 7,
+ CUDAShared = 8,
+ SYCLGlobal = 9,
+ SYCLLocal = 10,
+ SYCLPrivate = 11,
+ Ptr32Sptr = 12,
+ Ptr32Uptr = 13,
+ Ptr64 = 14,
+ HLSLGroupShared = 15,
+ HLSLConstant = 16,
+ HLSLPrivate = 17,
+ HLSLDevice = 18,
+ HLSLInput = 19,
+ HLSLOutput = 20,
+ HLSLPushConstant = 21,
+ WasmFuncRef = 22,
+ HIPDevice = 23,
+ HIPConstant = 24,
+ HIPShared = 25,
TargetOffset = 0x1000000
};
@@ -173,9 +169,8 @@ inline unsigned encode(LangAS AS, bool IsHIP = false) {
case LangAS::opencl_generic:
return OpenCLGeneric;
case LangAS::opencl_global_device:
- return OpenCLGlobalDevice;
case LangAS::opencl_global_host:
- return OpenCLGlobalHost;
+ return OpenCLGlobal;
case LangAS::cuda_device:
return IsHIP ? HIPDevice : CUDADevice;
case LangAS::cuda_constant:
@@ -183,11 +178,9 @@ inline unsigned encode(LangAS AS, bool IsHIP = false) {
case LangAS::cuda_shared:
return IsHIP ? HIPShared : CUDAShared;
case LangAS::sycl_global:
- return SYCLGlobal;
case LangAS::sycl_global_device:
- return SYCLGlobalDevice;
case LangAS::sycl_global_host:
- return SYCLGlobalHost;
+ return SYCLGlobal;
case LangAS::sycl_local:
return SYCLLocal;
case LangAS::sycl_private:
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 604a04e597db8..ee0598ee85bb8 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -931,10 +931,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
PointeeAddressSpace::OpenCLPrivate);
DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GENERIC",
PointeeAddressSpace::OpenCLGeneric);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_DEVICE",
- PointeeAddressSpace::OpenCLGlobalDevice);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_HOST",
- PointeeAddressSpace::OpenCLGlobalHost);
DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_CUDA_DEVICE",
PointeeAddressSpace::CUDADevice);
DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_CUDA_CONSTANT",
@@ -943,10 +939,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
PointeeAddressSpace::CUDAShared);
DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL",
PointeeAddressSpace::SYCLGlobal);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL_DEVICE",
- PointeeAddressSpace::SYCLGlobalDevice);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL_HOST",
- PointeeAddressSpace::SYCLGlobalHost);
DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_LOCAL",
PointeeAddressSpace::SYCLLocal);
DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_PRIVATE",
diff --git a/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu b/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
index e51548b59bbe6..a07b1c649193c 100644
--- a/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
+++ b/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
@@ -58,21 +58,21 @@ extern "C" int test_host_device_var() {
}
// HOST-LABEL: define{{.*}} i32 @test_host_device_var(
-// HOST: ret i32 27
+// HOST: ret i32 23
extern "C" int test_host_constant_var() {
return __builtin_pointee_address_space(&constant_var);
}
// HOST-LABEL: define{{.*}} i32 @test_host_constant_var(
-// HOST: ret i32 28
+// HOST: ret i32 24
extern "C" int test_host_const_device_var() {
return __builtin_pointee_address_space(&const_device_var);
}
// HOST-LABEL: define{{.*}} i32 @test_host_const_device_var(
-// HOST: ret i32 28
+// HOST: ret i32 24
extern "C" int test_host_device_array_address() {
return __builtin_pointee_address_space(&device_array);
@@ -178,24 +178,24 @@ extern "C" __device__ int test_shared_local() {
}
// CHECK-LABEL: define{{.*}} i32 @test_shared_local(
-// CUDA: ret i32 10
-// HIP: ret i32 29
+// CUDA: ret i32 8
+// HIP: ret i32 25
extern "C" __device__ int test_device_var() {
return __builtin_pointee_address_space(&device_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_device_var(
-// CUDA: ret i32 8
-// HIP: ret i32 27
+// CUDA: ret i32 6
+// HIP: ret i32 23
extern "C" __device__ int test_device_array() {
return __builtin_pointee_address_space(device_array);
}
// CHECK-LABEL: define{{.*}} i32 @test_device_array(
-// CUDA: ret i32 8
-// HIP: ret i32 27
+// CUDA: ret i32 6
+// HIP: ret i32 23
extern "C" __device__ int test_device_array_address() {
return __builtin_pointee_address_space(&device_array);
@@ -223,16 +223,16 @@ extern "C" __device__ int test_constant_var() {
}
// CHECK-LABEL: define{{.*}} i32 @test_constant_var(
-// CUDA: ret i32 9
-// HIP: ret i32 28
+// CUDA: ret i32 7
+// HIP: ret i32 24
extern "C" __device__ int test_const_device_var() {
return __builtin_pointee_address_space(&const_device_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_const_device_var(
-// CUDA: ret i32 9
-// HIP: ret i32 28
+// CUDA: ret i32 7
+// HIP: ret i32 24
extern "C" __device__ int test_explicit_cast() {
return __builtin_pointee_address_space((int *)&constant_var);
diff --git a/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl b/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl
index 3f002801c15d7..0fd1569daebad 100644
--- a/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl
+++ b/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl
@@ -1,5 +1,5 @@
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -Wno-deprecated-attributes -emit-llvm -o - %s | FileCheck %s
int global_as(__global int *p) {
return __builtin_pointee_address_space(p);
@@ -16,3 +16,19 @@ int local_as(__local int *p) {
// CHECK-LABEL: define{{.*}} i32 @local_as(
// CHECK-SAME: ptr addrspace(3)
// CHECK: ret i32 2
+
+int global_device_as(__attribute__((opencl_global_device)) int *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @global_device_as(
+// CHECK-SAME: ptr addrspace(1)
+// CHECK: ret i32 1
+
+int global_host_as(__attribute__((opencl_global_host)) int *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @global_host_as(
+// CHECK-SAME: ptr addrspace(1)
+// CHECK: ret i32 1
diff --git a/clang/test/CodeGenSYCL/builtin-pointee-address-space.cpp b/clang/test/CodeGenSYCL/builtin-pointee-address-space.cpp
new file mode 100644
index 0000000000000..5ffc70b72cf28
--- /dev/null
+++ b/clang/test/CodeGenSYCL/builtin-pointee-address-space.cpp
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -Wno-deprecated-attributes -emit-llvm -o - %s | FileCheck %s
+
+[[clang::sycl_external]] int
+global_device_as(__attribute__((opencl_global_device)) int *p) {
+ static_assert(__builtin_pointee_address_space(p) ==
+ __CLANG_ADDRESS_SPACE_SYCL_GLOBAL);
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @_Z16global_device_asPU3AS5i(
+// CHECK-SAME: ptr addrspace(5)
+// CHECK: ret i32 9
+
+[[clang::sycl_external]] int
+global_host_as(__attribute__((opencl_global_host)) int *p) {
+ static_assert(__builtin_pointee_address_space(p) ==
+ __CLANG_ADDRESS_SPACE_SYCL_GLOBAL);
+ return __builtin_pointee_address_space(p);
+}
+
+// CHECK-LABEL: define{{.*}} i32 @_Z14global_host_asPU3AS6i(
+// CHECK-SAME: ptr addrspace(6)
+// CHECK: ret i32 9
diff --git a/clang/test/Preprocessor/address-space-macros.c b/clang/test/Preprocessor/address-space-macros.c
index 8e2baefe5333c..5772f778b0e1a 100644
--- a/clang/test/Preprocessor/address-space-macros.c
+++ b/clang/test/Preprocessor/address-space-macros.c
@@ -6,28 +6,24 @@
// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_DEVICE 6
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_HOST 7
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 8
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 9
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_SHARED 10
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 11
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_DEVICE 12
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_HOST 13
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 14
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 15
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR32_SPTR 16
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR32_UPTR 17
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR64 18
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 19
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 20
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 21
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 22
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_INPUT 23
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 24
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 25
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 26
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_DEVICE 27
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 28
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_SHARED 29
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 6
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 7
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_SHARED 8
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 9
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 10
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 11
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR32_SPTR 12
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR32_UPTR 13
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR64 14
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 15
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 16
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 17
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 18
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_INPUT 19
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 20
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 21
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 22
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_DEVICE 23
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 24
+// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_SHARED 25
// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
diff --git a/clang/test/Preprocessor/init-aarch64.c b/clang/test/Preprocessor/init-aarch64.c
index c596f79512208..97cc695cc4ef4 100644
--- a/clang/test/Preprocessor/init-aarch64.c
+++ b/clang/test/Preprocessor/init-aarch64.c
@@ -52,37 +52,33 @@
// AARCH64-NEXT: #define __CHAR16_TYPE__ unsigned short
// AARCH64-NEXT: #define __CHAR32_TYPE__ unsigned int
// AARCH64-NEXT: #define __CHAR_BIT__ 8
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 9
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 8
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_SHARED 10
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 7
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 6
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_SHARED 8
// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_DEFAULT 0
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 28
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_DEVICE 27
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_SHARED 29
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 20
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 22
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 19
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_INPUT 23
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 24
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 21
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 25
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 24
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_DEVICE 23
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_SHARED 25
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 16
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 18
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 15
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_INPUT 19
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 20
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 17
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 21
// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL 1
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_DEVICE 6
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_HOST 7
// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_LOCAL 2
// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR32_SPTR 16
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR32_UPTR 17
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR64 18
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 11
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_DEVICE 12
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_HOST 13
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 14
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 15
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR32_SPTR 12
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR32_UPTR 13
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR64 14
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 9
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 10
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 11
// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 26
+// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 22
// AARCH64-NEXT: #define __CLANG_ATOMIC_BOOL_LOCK_FREE 2
// AARCH64-NEXT: #define __CLANG_ATOMIC_CHAR16_T_LOCK_FREE 2
// AARCH64-NEXT: #define __CLANG_ATOMIC_CHAR32_T_LOCK_FREE 2
diff --git a/clang/test/Preprocessor/init.c b/clang/test/Preprocessor/init.c
index ae8f6c3980a09..6f8c0409b6abb 100644
--- a/clang/test/Preprocessor/init.c
+++ b/clang/test/Preprocessor/init.c
@@ -1716,37 +1716,33 @@
// WEBASSEMBLY-NEXT:#define __CHAR32_TYPE__ unsigned int
// WEBASSEMBLY-NEXT:#define __CHAR_BIT__ 8
// WEBASSEMBLY-NOT:#define __CHAR_UNSIGNED__
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 9
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 8
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_SHARED 10
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 7
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 6
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_SHARED 8
// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_DEFAULT 0
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 28
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_DEVICE 27
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_SHARED 29
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 20
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 22
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 19
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_INPUT 23
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 24
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 21
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 25
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 24
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_DEVICE 23
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_SHARED 25
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 16
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 18
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 15
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_INPUT 19
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 20
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 17
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 21
// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL 1
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_DEVICE 6
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL_HOST 7
// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_LOCAL 2
// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR32_SPTR 16
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR32_UPTR 17
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR64 18
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 11
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_DEVICE 12
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL_HOST 13
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 14
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 15
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR32_SPTR 12
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR32_UPTR 13
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR64 14
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 9
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 10
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 11
// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 26
+// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 22
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_BOOL_LOCK_FREE 2
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_CHAR16_T_LOCK_FREE 2
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_CHAR32_T_LOCK_FREE 2
>From 4b8397093b33cf74735fe52742ebc80982751ef4 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 24 Jul 2026 11:42:05 -0400
Subject: [PATCH 3/7] [Clang] Replace pointee address-space builtin with
__addrspaceof
The builtin only accepts pointer and array expressions. This makes it difficult to query storage address spaces directly and does not provide decltype-like entity semantics.
Replace it with __addrspaceof. The operator accepts a type, variable or data member, or lvalue expression and returns a stable Clang address-space value. Parentheses distinguish direct entity queries from expression queries.
---
clang/docs/LanguageExtensions.md | 156 ++++++++----------
clang/docs/ReleaseNotes.md | 12 +-
clang/include/clang/AST/Expr.h | 2 +
clang/include/clang/Basic/AddressSpaces.h | 4 +-
clang/include/clang/Basic/BuiltinTraits.td | 5 +
clang/include/clang/Basic/Builtins.td | 7 -
.../clang/Basic/DiagnosticSemaKinds.td | 7 +-
clang/include/clang/Basic/Features.def | 1 +
clang/include/clang/Sema/Sema.h | 12 +-
clang/lib/AST/ByteCode/Compiler.cpp | 3 +
clang/lib/AST/Expr.cpp | 56 +++++++
clang/lib/AST/ExprConstant.cpp | 60 +------
clang/lib/AST/ItaniumMangle.cpp | 3 +
clang/lib/AST/StmtPrinter.cpp | 4 +
clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 7 +-
clang/lib/CodeGen/CGBuiltin.cpp | 4 -
clang/lib/Frontend/InitPreprocessor.cpp | 110 ++++++------
clang/lib/Parse/ParseExpr.cpp | 32 +++-
clang/lib/Sema/SemaChecking.cpp | 14 --
clang/lib/Sema/SemaExpr.cpp | 57 +++++--
clang/lib/Sema/TreeTransform.h | 6 +-
...-pointee-address-space.c => addrspaceof.c} | 10 +-
...ointee-address-space.cu => addrspaceof.cu} | 90 +++++-----
clang/test/CodeGenCXX/mangle-addrspaceof.cpp | 19 +++
...ointee-address-space.cl => addrspaceof.cl} | 8 +-
...ntee-address-space.cpp => addrspaceof.cpp} | 8 +-
clang/test/PCH/exprs.c | 1 +
clang/test/PCH/exprs.h | 1 +
...ointee-address-space.cu => addrspaceof.cu} | 32 ++--
clang/test/SemaCXX/addrspaceof.cpp | 87 ++++++++++
.../SemaCXX/builtin-pointee-address-space.cpp | 91 ----------
31 files changed, 483 insertions(+), 426 deletions(-)
rename clang/test/CodeGen/{builtin-pointee-address-space.c => addrspaceof.c} (76%)
rename clang/test/CodeGenCUDA/{builtin-pointee-address-space.cu => addrspaceof.cu} (74%)
create mode 100644 clang/test/CodeGenCXX/mangle-addrspaceof.cpp
rename clang/test/CodeGenOpenCL/{builtin-pointee-address-space.cl => addrspaceof.cl} (81%)
rename clang/test/CodeGenSYCL/{builtin-pointee-address-space.cpp => addrspaceof.cpp} (76%)
rename clang/test/SemaCUDA/{builtin-pointee-address-space.cu => addrspaceof.cu} (50%)
create mode 100644 clang/test/SemaCXX/addrspaceof.cpp
delete mode 100644 clang/test/SemaCXX/builtin-pointee-address-space.cpp
diff --git a/clang/docs/LanguageExtensions.md b/clang/docs/LanguageExtensions.md
index 563fc4bca61cc..566be2bd47dbb 100644
--- a/clang/docs/LanguageExtensions.md
+++ b/clang/docs/LanguageExtensions.md
@@ -3983,87 +3983,76 @@ template<typename T> constexpr T *addressof(T &value) {
}
```
-### `__builtin_pointee_address_space`
+### `__addrspaceof`
-`__builtin_pointee_address_space` returns a Clang address-space identifier for
-the storage reached by a pointer or array expression.
+`__addrspaceof` returns a Clang address-space identifier for a type, object, or
+lvalue expression.
This is useful for performance-sensitive code that needs a compile-time value
to choose an address-space-specific operation, overload, or template
-specialization. For example, CUDA/HIP code may want to select different helper
-code for a pointer to global, shared, or constant memory. The physical target
-address-space number can be target-specific and is not always defined by the
-source language, so this builtin returns Clang-defined values instead of
-backend address-space numbers for language-level address spaces.
+specialization. For example, GPU code may select different loads for global,
+shared/local, or constant memory.
**Syntax**:
```c++
-int __builtin_pointee_address_space(pointer-or-array)
-```
-
-The argument must have pointer or array type. Function designators are not
-accepted through function-to-pointer decay. The argument is not evaluated and
-is not converted to `void *` before its address space is queried.
-
-The result is an integer constant expression. It is based on the address space
-that Clang can determine from the expression in the frontend. It does not use
-optimizer or interprocedural analysis to infer a more precise address space.
-
-Clang predefines named macros for the values returned by this builtin. These
-macros are emitted by Clang during preprocessing. For language-level address
-spaces, such as OpenCL address spaces or CUDA/HIP
-`__device__`, `__shared__`, and `__constant__` variables, the result is one of
-the `__CLANG_ADDRESS_SPACE_*` values. For an address space written explicitly
-with `__attribute__((address_space(N)))`, the result is
-`__CLANG_ADDRESS_SPACE_TARGET_OFFSET + N`.
-CUDA and HIP use separate predefined macro names for their declaration address
-spaces. CUDA mode uses the `__CLANG_ADDRESS_SPACE_CUDA_*` values, and HIP mode
-uses the `__CLANG_ADDRESS_SPACE_HIP_*` values.
-
-For languages such as OpenCL, where address spaces are represented in AST
-types, the builtin returns the Clang address-space value of the pointee type of
-the expression as written. For an array expression, the pointee type is the
-array element type. The deprecated OpenCL and SYCL `global_device` and
-`global_host` address spaces are reported as their corresponding global address
-space.
-
-CUDA/HIP variables are not represented as address-space-qualified pointer types
-in Clang's AST. For CUDA/HIP, the builtin uses variable declaration attributes
-only for two direct forms: an array variable expression, such as `arr`, and an
-address-of expression naming a non-array variable, such as `&var`. In those
-forms, the builtin can report the declaration address space from attributes
-such as `__device__`, `__shared__`, or `__constant__`. An explicit
-`__device__` `const` global or static data member that is promoted to constant
-memory is reported as the CUDA or HIP constant address space. This also works
-in host compilation because it does not depend on the auxiliary target's
-physical address-space map.
-
-Explicit user-written casts are respected. If a cast changes the pointer type
-seen by the builtin, the builtin reports the pointee address space of the cast
-type rather than looking through the cast to recover the original object.
-Other pointer expressions, such as pointer arithmetic or `&array`, are handled
-from the pointee type of the expression as written.
-
-The builtin is a static query. It does not classify an arbitrary runtime pointer
-value. For a parameter such as `int *p`, if the pointee type is generic/default,
-the result is `__CLANG_ADDRESS_SPACE_DEFAULT` even if the runtime value of `p`
-later points into a more specific memory region. Runtime pointer-value
-classification should use target-specific runtime interfaces instead.
-
-The CUDA/HIP declaration query is direct. It does not follow function
-parameters, even through a `consteval` or `constexpr` wrapper. Inside such a
-wrapper, the builtin reports the address space of the parameter type.
+int __addrspaceof(type-or-expression)
+```
+
+The parentheses are required. The operand is not evaluated. The result is an
+integer constant expression.
+Reference types are adjusted to their referenced type before the address space
+is queried.
+
+For a type operand, the operator returns its top-level explicit address space,
+or the default address space if it has none.
+
+For an unparenthesized id-expression or unparenthesized member access,
+the operator queries the directly named variable or data member. It first uses
+the top-level explicit address space of the entity's type. For a CUDA/HIP
+entity without a top-level explicit address space, it uses the storage address
+space specified by `__device__`, `__shared__`, or `__constant__`. Otherwise, it
+returns the default address space. An explicit `__device__` `const` global or
+static data member that is promoted to constant memory is reported as the CUDA
+or HIP constant address space.
+
+Other expression operands must be lvalues. The operator returns the top-level
+explicit address space of the expression's type, or the default address space
+if it has none. Prvalue operands, including address-of expressions, are
+ill-formed. As with `decltype`, extra parentheses can change a direct entity
+query into an expression query.
+
+```c++
+__device__ int *p;
+
+static_assert(__addrspaceof(p) == __CLANG_ADDRESS_SPACE_CUDA_DEVICE);
+static_assert(__addrspaceof(*p) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof((p)) == __CLANG_ADDRESS_SPACE_DEFAULT);
+```
+
+The first query reports where the pointer variable is stored. The second
+reports the address space of the pointed-to object type. The third treats `p`
+as an lvalue expression because of the extra parentheses.
+
+Clang predefines named macros for the values returned by this operator. For
+language-level address spaces, the result is one of the
+`__CLANG_ADDRESS_SPACE_*` values. For an address space written explicitly with
+`__attribute__((address_space(N)))`, the result is
+`__CLANG_ADDRESS_SPACE_TARGET_OFFSET + N`. CUDA mode uses the
+`__CLANG_ADDRESS_SPACE_CUDA_*` values, and HIP mode uses the corresponding
+`__CLANG_ADDRESS_SPACE_HIP_*` values. Deprecated OpenCL and SYCL
+`global_device` and `global_host` address spaces are reported as their
+corresponding global address space.
**Example use**:
```c++
-int *p;
-int __attribute__((address_space(3))) *p3;
+using local_int = int __attribute__((address_space(3)));
+local_int *p;
-static_assert(__builtin_pointee_address_space(p) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(p3) ==
+static_assert(__addrspaceof(local_int) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 3);
+static_assert(__addrspaceof(*p) ==
__CLANG_ADDRESS_SPACE_TARGET_OFFSET + 3);
```
@@ -4073,9 +4062,9 @@ static_assert(__builtin_pointee_address_space(p3) ==
__global int *global_p;
__local int *local_p;
-static_assert(__builtin_pointee_address_space(global_p) ==
+static_assert(__addrspaceof(*global_p) ==
__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL);
-static_assert(__builtin_pointee_address_space(local_p) ==
+static_assert(__addrspaceof(*local_p) ==
__CLANG_ADDRESS_SPACE_OPENCL_LOCAL);
```
@@ -4086,23 +4075,22 @@ __device__ int dev;
__device__ int dev_arr[4];
__constant__ int cst;
__device__ const int dev_cst = 1;
-```
-In CUDA mode, `&dev` and `dev_arr` return
-`__CLANG_ADDRESS_SPACE_CUDA_DEVICE`, while `&cst` and `&dev_cst` return
-`__CLANG_ADDRESS_SPACE_CUDA_CONSTANT`. In HIP mode, they return the matching
-`__CLANG_ADDRESS_SPACE_HIP_*` values.
+static_assert(__addrspaceof(dev) ==
+ __CLANG_ADDRESS_SPACE_CUDA_DEVICE);
+static_assert(__addrspaceof(dev_arr) ==
+ __CLANG_ADDRESS_SPACE_CUDA_DEVICE);
+static_assert(__addrspaceof(cst) ==
+ __CLANG_ADDRESS_SPACE_CUDA_CONSTANT);
+static_assert(__addrspaceof(dev_cst) ==
+ __CLANG_ADDRESS_SPACE_CUDA_CONSTANT);
+```
-The direct-query rule is not specific to shared variables. For a CUDA/HIP
-variable with a declaration address-space attribute, `&a` for a non-array
-variable and `b` for an array variable are direct declaration queries. For
-example, `&a` for `__shared__ int a` and `b` for `__shared__ int b[2]` report
-the CUDA or HIP shared address space.
+HIP compilation returns the corresponding `__CLANG_ADDRESS_SPACE_HIP_*`
+values. The operator is a static frontend query. It does not classify an
+arbitrary runtime pointer value and does not use optimizer or backend analysis.
-Forms such as `(void *)&a`, `&a + 1`, `&b[0]`, `&dev_arr`, `dev_arr + 1`,
-and calls through a function parameter, including `consteval` wrappers, are
-not direct declaration queries. They use the expression type and report
-`__CLANG_ADDRESS_SPACE_DEFAULT`.
+Query for this extension with `__has_extension(addrspaceof)`.
### `__builtin_function_start`
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index 6b75f6914c9eb..c9bb3e995c6c8 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -182,13 +182,11 @@ features cannot lower the translation-unit ABI level;
- Clang now allows GNU computed `goto` extension in `constexpr` functions, matching the relaxed
`constexpr` function body rules introduced in C++23.
-- Added `__builtin_pointee_address_space`, which returns a Clang
- address-space identifier for a pointer or array expression. It reports the
- AST pointee type address space for OpenCL and explicit address-space-qualified
- pointer types, and can report CUDA/HIP declaration address spaces for known
- `__device__`, `__shared__`, and `__constant__` variables. Clang also now
- emits predefined `__CLANG_ADDRESS_SPACE_*` macros for these values, including
- separate CUDA and HIP macro names.
+- Added the `__addrspaceof` operator, which returns a Clang address-space
+ identifier for a type, directly named entity, or lvalue expression. It can
+ report CUDA/HIP storage address spaces for known `__device__`, `__shared__`,
+ and `__constant__` variables. Clang also now emits predefined
+ `__CLANG_ADDRESS_SPACE_*` macros for these values.
### New Compiler Flags
diff --git a/clang/include/clang/AST/Expr.h b/clang/include/clang/AST/Expr.h
index f95f87cc4e8e0..f48629c91abbb 100644
--- a/clang/include/clang/AST/Expr.h
+++ b/clang/include/clang/AST/Expr.h
@@ -2701,6 +2701,8 @@ class UnaryExprOrTypeTraitExpr : public Expr {
return isArgumentType() ? getArgumentType() : getArgumentExpr()->getType();
}
+ unsigned getAddressSpaceQueryResult(const ASTContext &Ctx) const;
+
SourceLocation getOperatorLoc() const { return OpLoc; }
void setOperatorLoc(SourceLocation L) { OpLoc = L; }
diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h
index df490e60ed493..ecb3430392095 100644
--- a/clang/include/clang/Basic/AddressSpaces.h
+++ b/clang/include/clang/Basic/AddressSpaces.h
@@ -118,7 +118,7 @@ inline bool isPtrSizeAddressSpace(LangAS AS) {
AS == LangAS::ptr64);
}
-namespace PointeeAddressSpace {
+namespace AddressSpaceQuery {
enum ID : unsigned {
Default = 0,
@@ -215,7 +215,7 @@ inline unsigned encode(LangAS AS, bool IsHIP = false) {
return Default;
}
-} // namespace PointeeAddressSpace
+} // namespace AddressSpaceQuery
} // namespace clang
diff --git a/clang/include/clang/Basic/BuiltinTraits.td b/clang/include/clang/Basic/BuiltinTraits.td
index f5546acdb17c4..2318f4363ef28 100644
--- a/clang/include/clang/Basic/BuiltinTraits.td
+++ b/clang/include/clang/Basic/BuiltinTraits.td
@@ -82,6 +82,11 @@ def DataSizeOf : UnaryExprOrTypeTrait {
let Spelling = "__datasizeof";
}
+def AddrSpaceOf : UnaryExprOrTypeTrait {
+ let Spelling = "__addrspaceof";
+ let KeyFlags = [KEYALL];
+}
+
// C2y
def CountOf : UnaryExprOrTypeTrait {
let Spelling = "_Countof";
diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td
index 67dee1d024431..a54f91069acd0 100644
--- a/clang/include/clang/Basic/Builtins.td
+++ b/clang/include/clang/Basic/Builtins.td
@@ -1046,13 +1046,6 @@ def BuiltinClassifyType : Builtin {
let Prototype = "int(...)";
}
-def BuiltinPointeeAddressSpace : Builtin {
- let Spellings = ["__builtin_pointee_address_space"];
- let Attributes = [NoThrow, Const, CustomTypeChecking, UnevaluatedArguments,
- Constexpr];
- let Prototype = "int(...)";
-}
-
def BuiltinCFStringMakeConstantString : Builtin {
let Spellings = ["__builtin___CFStringMakeConstantString"];
let Attributes = [NoThrow, Const, Constexpr];
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 0c5e099cea882..afb28a5f31ae9 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -864,9 +864,10 @@ def warn_redecl_library_builtin : Warning<
def err_builtin_definition : Error<"definition of builtin function %0">;
def err_builtin_redeclare : Error<"cannot redeclare builtin function %0">;
def err_invalid_builtin_argument : Error<"invalid argument '%0' to %1">;
-def err_builtin_pointee_address_space_arg_not_pointer : Error<
- "argument to '__builtin_pointee_address_space' must be a pointer or array "
- "expression">;
+def err_addrspaceof_invalid_expression : Error<
+ "expression operand of '__addrspaceof' must be an lvalue">;
+def err_addrspaceof_invalid_entity : Error<
+ "entity operand of '__addrspaceof' must name a variable or data member">;
def err_arm_invalid_specialreg : Error<"invalid special register for builtin">;
def err_arm_invalid_coproc : Error<"coprocessor %0 must be configured as "
diff --git a/clang/include/clang/Basic/Features.def b/clang/include/clang/Basic/Features.def
index cfd5bf69f938c..fc51471f7c244 100644
--- a/clang/include/clang/Basic/Features.def
+++ b/clang/include/clang/Basic/Features.def
@@ -377,6 +377,7 @@ EXTENSION(matrix_types, LangOpts.MatrixTypes)
EXTENSION(matrix_types_scalar_division, true)
EXTENSION(cxx_attributes_on_using_declarations, LangOpts.CPlusPlus11)
EXTENSION(datasizeof, LangOpts.CPlusPlus)
+EXTENSION(addrspaceof, true)
FEATURE(cxx_abi_relative_vtable, LangOpts.CPlusPlus && LangOpts.RelativeCXXABIVTables)
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index b1d2488d2163b..dd16fa829b375 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -7464,20 +7464,18 @@ class Sema final : public SemaBase {
bool CheckAlignasTypeArgument(StringRef KWName, TypeSourceInfo *TInfo,
SourceLocation OpLoc, SourceRange R);
- /// Build a sizeof or alignof expression given a type operand.
+ /// Build a unary expression or type trait given a type operand.
ExprResult CreateUnaryExprOrTypeTraitExpr(TypeSourceInfo *TInfo,
SourceLocation OpLoc,
UnaryExprOrTypeTrait ExprKind,
SourceRange R);
- /// Build a sizeof or alignof expression given an expression
- /// operand.
+ /// Build a unary expression or type trait given an expression operand.
ExprResult CreateUnaryExprOrTypeTraitExpr(Expr *E, SourceLocation OpLoc,
- UnaryExprOrTypeTrait ExprKind);
+ UnaryExprOrTypeTrait ExprKind,
+ SourceLocation RParenLoc = {});
- /// ActOnUnaryExprOrTypeTraitExpr - Handle @c sizeof(type) and @c sizeof @c
- /// expr and the same for @c alignof and @c __alignof
- /// Note that the ArgRange is invalid if isType is false.
+ /// Handle a unary expression or type trait.
ExprResult ActOnUnaryExprOrTypeTraitExpr(SourceLocation OpLoc,
UnaryExprOrTypeTrait ExprKind,
bool IsType, void *TyOrEx,
diff --git a/clang/lib/AST/ByteCode/Compiler.cpp b/clang/lib/AST/ByteCode/Compiler.cpp
index 9b2c9191f729a..a32ba05dca541 100644
--- a/clang/lib/AST/ByteCode/Compiler.cpp
+++ b/clang/lib/AST/ByteCode/Compiler.cpp
@@ -2745,6 +2745,9 @@ bool Compiler<Emitter>::VisitUnaryExprOrTypeTraitExpr(
UnaryExprOrTypeTrait Kind = E->getKind();
const ASTContext &ASTCtx = Ctx.getASTContext();
+ if (Kind == UETT_AddrSpaceOf)
+ return this->emitConst(E->getAddressSpaceQueryResult(ASTCtx), E);
+
if (Kind == UETT_SizeOf || Kind == UETT_DataSizeOf) {
QualType ArgType = E->getTypeOfArgument();
diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index 5d7ee4710481c..898ec6e84cb55 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/lib/AST/Expr.cpp
@@ -27,6 +27,7 @@
#include "clang/AST/RecordLayout.h"
#include "clang/AST/StmtVisitor.h"
#include "clang/AST/TypeBase.h"
+#include "clang/Basic/AddressSpaces.h"
#include "clang/Basic/Builtins.h"
#include "clang/Basic/CharInfo.h"
#include "clang/Basic/SourceManager.h"
@@ -1715,6 +1716,61 @@ UnaryExprOrTypeTraitExpr::UnaryExprOrTypeTraitExpr(
setDependence(computeDependence(this));
}
+static const ValueDecl *getAddrSpaceOfEntity(const Expr *E) {
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(E))
+ return dyn_cast<ValueDecl>(DRE->getDecl());
+ if (const auto *ME = dyn_cast<MemberExpr>(E))
+ return ME->getMemberDecl();
+ return nullptr;
+}
+
+static std::optional<LangAS> getCUDADeclAddressSpace(const ASTContext &Ctx,
+ const ValueDecl *D) {
+ if (!Ctx.getLangOpts().CUDA)
+ return std::nullopt;
+
+ if (D->hasAttr<CUDAConstantAttr>())
+ return LangAS::cuda_constant;
+ if (D->hasAttr<CUDASharedAttr>())
+ return LangAS::cuda_shared;
+
+ const auto *VD = dyn_cast<VarDecl>(D);
+ if (!VD)
+ return std::nullopt;
+
+ // Host compilation does not attach the implicit CUDAConstantAttr that
+ // SemaCUDA adds in device compilation.
+ if (!Ctx.getLangOpts().CUDAIsDevice && VD->hasAttr<CUDADeviceAttr>() &&
+ (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
+ (VD->isConstexpr() || VD->getType().isConstQualified()))
+ return LangAS::cuda_constant;
+ if (VD->hasAttr<CUDADeviceAttr>())
+ return LangAS::cuda_device;
+ return std::nullopt;
+}
+
+unsigned UnaryExprOrTypeTraitExpr::getAddressSpaceQueryResult(
+ const ASTContext &Ctx) const {
+ assert(getKind() == UETT_AddrSpaceOf && "not an address-space query");
+
+ QualType T;
+ if (isArgumentType()) {
+ T = getArgumentType();
+ } else if (const ValueDecl *D = getAddrSpaceOfEntity(getArgumentExpr())) {
+ T = D->getType();
+ LangAS AS = T.getNonReferenceType().getAddressSpace();
+ if (AS != LangAS::Default)
+ return AddressSpaceQuery::encode(AS, Ctx.getLangOpts().HIP);
+ if (std::optional<LangAS> AS = getCUDADeclAddressSpace(Ctx, D))
+ return AddressSpaceQuery::encode(*AS, Ctx.getLangOpts().HIP);
+ } else {
+ T = getArgumentExpr()->getType();
+ }
+
+ return AddressSpaceQuery::encode(T.getNonReferenceType().getAddressSpace(),
+ Ctx.getLangOpts().HIP);
+}
+
MemberExpr::MemberExpr(Expr *Base, bool IsArrow, SourceLocation OperatorLoc,
NestedNameSpecifierLoc QualifierLoc,
SourceLocation TemplateKWLoc, ValueDecl *MemberDecl,
diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp
index fed6b645b1e58..100421876d281 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/clang/lib/AST/ExprConstant.cpp
@@ -15937,54 +15937,6 @@ bool ArrayExprEvaluator::VisitDesignatedInitUpdateExpr(
//===----------------------------------------------------------------------===//
namespace {
-
-static QualType getPointeeAddressSpaceType(ASTContext &Ctx, QualType ArgTy) {
- return ArgTy->isArrayType() ? Ctx.getAsArrayType(ArgTy)->getElementType()
- : ArgTy->getPointeeType();
-}
-
-static std::optional<LangAS> getCUDAPointeeDeclAddressSpace(ASTContext &Ctx,
- const Expr *Arg) {
- if (!Ctx.getLangOpts().CUDA)
- return std::nullopt;
-
- Arg = Arg->IgnoreParens();
- if (isa<ExplicitCastExpr>(Arg))
- return std::nullopt;
-
- const VarDecl *VD = nullptr;
- const Expr *NoImpCasts = Arg->IgnoreImpCasts()->IgnoreParens();
- if (const auto *UO = dyn_cast<UnaryOperator>(NoImpCasts);
- UO && UO->getOpcode() == UO_AddrOf) {
- if (const auto *DRE =
- dyn_cast<DeclRefExpr>(UO->getSubExpr()->IgnoreParenImpCasts())) {
- const auto *CandidateVD = dyn_cast<VarDecl>(DRE->getDecl());
- if (CandidateVD && !CandidateVD->getType()->isArrayType())
- VD = CandidateVD;
- }
- } else if (const auto *DRE = dyn_cast<DeclRefExpr>(NoImpCasts)) {
- if (NoImpCasts->getType()->isArrayType())
- VD = dyn_cast<VarDecl>(DRE->getDecl());
- }
-
- if (!VD)
- return std::nullopt;
- if (VD->hasAttr<CUDAConstantAttr>())
- return LangAS::cuda_constant;
- if (VD->hasAttr<CUDASharedAttr>())
- return LangAS::cuda_shared;
- // Host compilation does not attach the implicit CUDAConstantAttr that
- // SemaCUDA adds in device compilation, but the device-side storage is still
- // constant memory.
- if (!Ctx.getLangOpts().CUDAIsDevice && VD->hasAttr<CUDADeviceAttr>() &&
- (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
- (VD->isConstexpr() || VD->getType().isConstQualified()))
- return LangAS::cuda_constant;
- if (VD->hasAttr<CUDADeviceAttr>())
- return LangAS::cuda_device;
- return std::nullopt;
-}
-
class IntExprEvaluator
: public ExprEvaluatorBase<IntExprEvaluator> {
APValue &Result;
@@ -17051,16 +17003,6 @@ bool IntExprEvaluator::VisitBuiltinCallExpr(const CallExpr *E,
llvm_unreachable("unexpected EvalMode");
}
- case Builtin::BI__builtin_pointee_address_space: {
- QualType ArgTy = E->getArg(0)->getType();
- LangAS AS =
- getCUDAPointeeDeclAddressSpace(Info.Ctx, E->getArg(0))
- .value_or(
- getPointeeAddressSpaceType(Info.Ctx, ArgTy).getAddressSpace());
- return Success(PointeeAddressSpace::encode(AS, Info.Ctx.getLangOpts().HIP),
- E);
- }
-
case Builtin::BI__builtin_os_log_format_buffer_size: {
analyze_os_log::OSLogBufferLayout Layout;
analyze_os_log::computeOSLogBufferLayout(Info.Ctx, E, Layout);
@@ -19704,6 +19646,8 @@ bool IntExprEvaluator::VisitBinaryOperator(const BinaryOperator *E) {
bool IntExprEvaluator::VisitUnaryExprOrTypeTraitExpr(
const UnaryExprOrTypeTraitExpr *E) {
switch(E->getKind()) {
+ case UETT_AddrSpaceOf:
+ return Success(E->getAddressSpaceQueryResult(Info.Ctx), E);
case UETT_PreferredAlignOf:
case UETT_AlignOf: {
if (E->isArgumentType())
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index f6c4ca1ae6ba8..2e90f15ab2329 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -5471,6 +5471,9 @@ void CXXNameMangler::mangleExpression(const Expr *E, unsigned Arity,
MangleAlignofSizeofArg();
break;
+ case UETT_AddrSpaceOf:
+ MangleExtensionBuiltin(SAE);
+ break;
case UETT_CountOf:
case UETT_VectorElements:
case UETT_OpenMPRequiredSimdAlign:
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index eeb377c794e05..04bb9bbd3a18b 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1714,6 +1714,10 @@ void StmtPrinter::VisitUnaryExprOrTypeTraitExpr(
OS << '(';
Node->getArgumentType().print(OS, Policy);
OS << ')';
+ } else if (Node->getKind() == UETT_AddrSpaceOf) {
+ OS << '(';
+ PrintExpr(Node->getArgumentExpr());
+ OS << ')';
} else {
OS << " ";
PrintExpr(Node->getArgumentExpr());
diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
index 8d660a0a2c721..ad0955647c82a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
@@ -2772,10 +2772,11 @@ mlir::Value ScalarExprEmitter::VisitUnaryExprOrTypeTraitExpr(
loc, cir::IntAttr::get(cgf.cgm.sizeTy, vecTy.getSize()));
}
- // The result type is size_t (target-dependent width); use it so the IntAttr
- // width matches the APInt from EvaluateKnownConstInt.
+ mlir::Type resultType = e->getKind() == UETT_AddrSpaceOf
+ ? convertType(e->getType())
+ : cgf.cgm.sizeTy;
return builder.getConstant(
- loc, cir::IntAttr::get(cgf.cgm.sizeTy,
+ loc, cir::IntAttr::get(resultType,
e->EvaluateKnownConstInt(cgf.getContext())));
}
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index ad91a9c5f7bc8..3521fc10f1387 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -4261,10 +4261,6 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
Result = Builder.CreateIntCast(Result, ResultType, /*isSigned*/false);
return RValue::get(Result);
}
- case Builtin::BI__builtin_pointee_address_space: {
- unsigned AS = E->EvaluateKnownConstInt(getContext()).getZExtValue();
- return RValue::get(ConstantInt::get(ConvertType(E->getType()), AS));
- }
case Builtin::BI__builtin_dynamic_object_size:
case Builtin::BI__builtin_object_size: {
unsigned Type =
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index ee0598ee85bb8..068448b580155 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -915,64 +915,64 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
Builder.defineMacro("__OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES", "3");
Builder.defineMacro("__OPENCL_MEMORY_SCOPE_SUB_GROUP", "4");
- auto DefinePointeeASMacro = [&](StringRef Name, PointeeAddressSpace::ID AS) {
+ auto DefineAddressSpaceMacro = [&](StringRef Name, AddressSpaceQuery::ID AS) {
Builder.defineMacro(Name, Twine(static_cast<unsigned>(AS)));
};
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_DEFAULT",
- PointeeAddressSpace::Default);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL",
- PointeeAddressSpace::OpenCLGlobal);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_LOCAL",
- PointeeAddressSpace::OpenCLLocal);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_CONSTANT",
- PointeeAddressSpace::OpenCLConstant);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_PRIVATE",
- PointeeAddressSpace::OpenCLPrivate);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_OPENCL_GENERIC",
- PointeeAddressSpace::OpenCLGeneric);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_CUDA_DEVICE",
- PointeeAddressSpace::CUDADevice);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_CUDA_CONSTANT",
- PointeeAddressSpace::CUDAConstant);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_CUDA_SHARED",
- PointeeAddressSpace::CUDAShared);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL",
- PointeeAddressSpace::SYCLGlobal);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_LOCAL",
- PointeeAddressSpace::SYCLLocal);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_SYCL_PRIVATE",
- PointeeAddressSpace::SYCLPrivate);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_PTR32_SPTR",
- PointeeAddressSpace::Ptr32Sptr);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_PTR32_UPTR",
- PointeeAddressSpace::Ptr32Uptr);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_PTR64",
- PointeeAddressSpace::Ptr64);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED",
- PointeeAddressSpace::HLSLGroupShared);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_CONSTANT",
- PointeeAddressSpace::HLSLConstant);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_PRIVATE",
- PointeeAddressSpace::HLSLPrivate);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_DEVICE",
- PointeeAddressSpace::HLSLDevice);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_INPUT",
- PointeeAddressSpace::HLSLInput);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_OUTPUT",
- PointeeAddressSpace::HLSLOutput);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT",
- PointeeAddressSpace::HLSLPushConstant);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_WASM_FUNCREF",
- PointeeAddressSpace::WasmFuncRef);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HIP_DEVICE",
- PointeeAddressSpace::HIPDevice);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HIP_CONSTANT",
- PointeeAddressSpace::HIPConstant);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_HIP_SHARED",
- PointeeAddressSpace::HIPShared);
- DefinePointeeASMacro("__CLANG_ADDRESS_SPACE_TARGET_OFFSET",
- PointeeAddressSpace::TargetOffset);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_DEFAULT",
+ AddressSpaceQuery::Default);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL",
+ AddressSpaceQuery::OpenCLGlobal);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_LOCAL",
+ AddressSpaceQuery::OpenCLLocal);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_CONSTANT",
+ AddressSpaceQuery::OpenCLConstant);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_PRIVATE",
+ AddressSpaceQuery::OpenCLPrivate);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_GENERIC",
+ AddressSpaceQuery::OpenCLGeneric);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_CUDA_DEVICE",
+ AddressSpaceQuery::CUDADevice);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_CUDA_CONSTANT",
+ AddressSpaceQuery::CUDAConstant);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_CUDA_SHARED",
+ AddressSpaceQuery::CUDAShared);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL",
+ AddressSpaceQuery::SYCLGlobal);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_SYCL_LOCAL",
+ AddressSpaceQuery::SYCLLocal);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_SYCL_PRIVATE",
+ AddressSpaceQuery::SYCLPrivate);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_PTR32_SPTR",
+ AddressSpaceQuery::Ptr32Sptr);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_PTR32_UPTR",
+ AddressSpaceQuery::Ptr32Uptr);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_PTR64",
+ AddressSpaceQuery::Ptr64);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED",
+ AddressSpaceQuery::HLSLGroupShared);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_CONSTANT",
+ AddressSpaceQuery::HLSLConstant);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_PRIVATE",
+ AddressSpaceQuery::HLSLPrivate);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_DEVICE",
+ AddressSpaceQuery::HLSLDevice);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_INPUT",
+ AddressSpaceQuery::HLSLInput);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_OUTPUT",
+ AddressSpaceQuery::HLSLOutput);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT",
+ AddressSpaceQuery::HLSLPushConstant);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_WASM_FUNCREF",
+ AddressSpaceQuery::WasmFuncRef);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HIP_DEVICE",
+ AddressSpaceQuery::HIPDevice);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HIP_CONSTANT",
+ AddressSpaceQuery::HIPConstant);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HIP_SHARED",
+ AddressSpaceQuery::HIPShared);
+ DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_TARGET_OFFSET",
+ AddressSpaceQuery::TargetOffset);
// Define macros for floating-point data classes, used in __builtin_isfpclass.
Builder.defineMacro("__FPCLASS_SNAN", "0x0001");
diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp
index 87cd7a01451cf..d7e4499ad8173 100644
--- a/clang/lib/Parse/ParseExpr.cpp
+++ b/clang/lib/Parse/ParseExpr.cpp
@@ -1216,6 +1216,7 @@ Parser::ParseCastExpression(CastParseKind ParseKind, bool isAddressOfOperand,
// unary-expression: '__datasizeof' unary-expression
// unary-expression: '__datasizeof' '(' type-name ')'
case tok::kw___datasizeof:
+ case tok::kw___addrspaceof:
case tok::kw_vec_step: // unary-expression: OpenCL 'vec_step' expression
// unary-expression: '__builtin_omp_required_simd_align' '(' type-name ')'
case tok::kw___builtin_omp_required_simd_align:
@@ -2083,8 +2084,9 @@ Parser::ParseExprAfterUnaryExprOrTypeTrait(const Token &OpTok,
SourceRange &CastRange) {
assert(OpTok.isOneOf(tok::kw_typeof, tok::kw_typeof_unqual, tok::kw_sizeof,
- tok::kw___datasizeof, tok::kw___alignof, tok::kw_alignof,
- tok::kw__Alignof, tok::kw_vec_step,
+ tok::kw___datasizeof, tok::kw___addrspaceof,
+ tok::kw___alignof, tok::kw_alignof, tok::kw__Alignof,
+ tok::kw_vec_step,
tok::kw___builtin_omp_required_simd_align,
tok::kw___builtin_vectorelements, tok::kw__Countof) &&
"Not a typeof/sizeof/alignof/vec_step expression!");
@@ -2093,6 +2095,10 @@ Parser::ParseExprAfterUnaryExprOrTypeTrait(const Token &OpTok,
// If the operand doesn't start with an '(', it must be an expression.
if (Tok.isNot(tok::l_paren)) {
+ if (OpTok.is(tok::kw___addrspaceof))
+ return ExprError(Diag(Tok, diag::err_expected_lparen_after)
+ << OpTok.getName());
+
// If construct allows a form without parenthesis, user may forget to put
// pathenthesis around type name.
if (OpTok.isOneOf(tok::kw_sizeof, tok::kw___datasizeof, tok::kw___alignof,
@@ -2156,8 +2162,8 @@ Parser::ParseExprAfterUnaryExprOrTypeTrait(const Token &OpTok,
// them looking like a compound literal, as in sizeof (int){}; where the
// parens could be part of a parenthesized type name or for a cast
// expression of some kind.
- bool ParenKnownToBeNonCast =
- OpTok.isOneOf(tok::kw_typeof, tok::kw_typeof_unqual);
+ bool ParenKnownToBeNonCast = OpTok.isOneOf(
+ tok::kw_typeof, tok::kw_typeof_unqual, tok::kw___addrspaceof);
ParenParseOption ExprType = ParenParseOption::CastExpr;
SourceLocation LParenLoc = Tok.getLocation(), RParenLoc;
@@ -2175,8 +2181,9 @@ Parser::ParseExprAfterUnaryExprOrTypeTrait(const Token &OpTok,
return ExprEmpty();
}
- if (getLangOpts().CPlusPlus ||
- !OpTok.isOneOf(tok::kw_typeof, tok::kw_typeof_unqual)) {
+ if (OpTok.isNot(tok::kw___addrspaceof) &&
+ (getLangOpts().CPlusPlus ||
+ !OpTok.isOneOf(tok::kw_typeof, tok::kw_typeof_unqual))) {
// GNU typeof in C requires the expression to be parenthesized. Not so for
// sizeof/alignof or in C++. Therefore, the parenthesized expression is
// the start of a unary-expression, but doesn't include any postfix
@@ -2184,6 +2191,13 @@ Parser::ParseExprAfterUnaryExprOrTypeTrait(const Token &OpTok,
if (!Operand.isInvalid())
Operand = ParsePostfixExpressionSuffix(Operand.get());
}
+
+ // The first parentheses belong to the operator and do not change an entity
+ // operand into an expression operand.
+ if (OpTok.is(tok::kw___addrspaceof) && !Operand.isInvalid()) {
+ if (auto *PE = dyn_cast<ParenExpr>(Operand.get()))
+ Operand = PE->getSubExpr();
+ }
}
// If we get here, the operand to the typeof/sizeof/alignof was an expression.
@@ -2219,7 +2233,8 @@ ExprResult Parser::ParseSYCLUniqueStableNameExpression() {
ExprResult Parser::ParseUnaryExprOrTypeTraitExpression() {
assert(Tok.isOneOf(tok::kw_sizeof, tok::kw___datasizeof, tok::kw___alignof,
- tok::kw_alignof, tok::kw__Alignof, tok::kw_vec_step,
+ tok::kw___addrspaceof, tok::kw_alignof, tok::kw__Alignof,
+ tok::kw_vec_step,
tok::kw___builtin_omp_required_simd_align,
tok::kw___builtin_vectorelements, tok::kw__Countof) &&
"Not a sizeof/alignof/vec_step expression!");
@@ -2311,6 +2326,9 @@ ExprResult Parser::ParseUnaryExprOrTypeTraitExpression() {
case tok::kw___datasizeof:
ExprKind = UETT_DataSizeOf;
break;
+ case tok::kw___addrspaceof:
+ ExprKind = UETT_AddrSpaceOf;
+ break;
case tok::kw___builtin_vectorelements:
ExprKind = UETT_VectorElements;
break;
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index 537e952443d0d..f2f38c84dc5f8 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -3284,20 +3284,6 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID,
if (BuiltinComplex(TheCall))
return ExprError();
break;
- case Builtin::BI__builtin_pointee_address_space: {
- if (checkArgCount(TheCall, 1))
- return true;
- Expr *Arg = TheCall->getArg(0);
- if (!Arg->isTypeDependent() && !Arg->getType()->isPointerType() &&
- !Arg->getType()->isArrayType()) {
- Diag(Arg->getBeginLoc(),
- diag::err_builtin_pointee_address_space_arg_not_pointer)
- << Arg->getSourceRange();
- return ExprError();
- }
- TheCall->setType(Context.IntTy);
- break;
- }
case Builtin::BI__builtin_classify_type:
case Builtin::BI__builtin_constant_p: {
if (checkArgCount(TheCall, 1))
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 90e360e2b05fc..66e69557005b0 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -4537,6 +4537,32 @@ bool Sema::CheckUnaryExprOrTypeTraitOperand(Expr *E,
return false;
}
+static bool CheckAddrSpaceOfExpr(Sema &S, Expr *E) {
+ if (E->isTypeDependent())
+ return false;
+
+ const ValueDecl *D = nullptr;
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(E))
+ D = dyn_cast<ValueDecl>(DRE->getDecl());
+ else if (const auto *ME = dyn_cast<MemberExpr>(E))
+ D = ME->getMemberDecl();
+
+ if (D) {
+ if (isa<VarDecl, FieldDecl>(D))
+ return false;
+ S.Diag(E->getExprLoc(), diag::err_addrspaceof_invalid_entity)
+ << E->getSourceRange();
+ return true;
+ }
+
+ if (!E->isLValue()) {
+ S.Diag(E->getExprLoc(), diag::err_addrspaceof_invalid_expression)
+ << E->getSourceRange();
+ return true;
+ }
+ return false;
+}
+
static bool CheckAlignOfExpr(Sema &S, Expr *E, UnaryExprOrTypeTrait ExprKind) {
// Cannot know anything else if the expression is dependent.
if (E->isTypeDependent())
@@ -4839,7 +4865,7 @@ ExprResult Sema::CreateUnaryExprOrTypeTraitExpr(TypeSourceInfo *TInfo,
QualType T = TInfo->getType();
- if (!T->isDependentType() &&
+ if (ExprKind != UETT_AddrSpaceOf && !T->isDependentType() &&
CheckUnaryExprOrTypeTraitOperand(T, OpLoc, R, ExprKind,
getTraitSpelling(ExprKind)))
return ExprError();
@@ -4856,14 +4882,15 @@ ExprResult Sema::CreateUnaryExprOrTypeTraitExpr(TypeSourceInfo *TInfo,
if (!TInfo)
return ExprError();
- // C99 6.5.3.4p4: the type (an unsigned integer type) is size_t.
- return new (Context) UnaryExprOrTypeTraitExpr(
- ExprKind, TInfo, Context.getSizeType(), OpLoc, R.getEnd());
+ QualType ResultType =
+ ExprKind == UETT_AddrSpaceOf ? Context.IntTy : Context.getSizeType();
+ return new (Context)
+ UnaryExprOrTypeTraitExpr(ExprKind, TInfo, ResultType, OpLoc, R.getEnd());
}
-ExprResult
-Sema::CreateUnaryExprOrTypeTraitExpr(Expr *E, SourceLocation OpLoc,
- UnaryExprOrTypeTrait ExprKind) {
+ExprResult Sema::CreateUnaryExprOrTypeTraitExpr(Expr *E, SourceLocation OpLoc,
+ UnaryExprOrTypeTrait ExprKind,
+ SourceLocation RParenLoc) {
ExprResult PE = CheckPlaceholderExpr(E);
if (PE.isInvalid())
return ExprError();
@@ -4874,6 +4901,8 @@ Sema::CreateUnaryExprOrTypeTraitExpr(Expr *E, SourceLocation OpLoc,
bool isInvalid = false;
if (E->isTypeDependent()) {
// Delay type-checking for type-dependent expressions.
+ } else if (ExprKind == UETT_AddrSpaceOf) {
+ isInvalid = CheckAddrSpaceOfExpr(*this, E);
} else if (ExprKind == UETT_AlignOf || ExprKind == UETT_PreferredAlignOf) {
isInvalid = CheckAlignOfExpr(*this, E, ExprKind);
} else if (ExprKind == UETT_VecStep) {
@@ -4899,9 +4928,12 @@ Sema::CreateUnaryExprOrTypeTraitExpr(Expr *E, SourceLocation OpLoc,
E = PE.get();
}
- // C99 6.5.3.4p4: the type (an unsigned integer type) is size_t.
- return new (Context) UnaryExprOrTypeTraitExpr(
- ExprKind, E, Context.getSizeType(), OpLoc, E->getSourceRange().getEnd());
+ QualType ResultType =
+ ExprKind == UETT_AddrSpaceOf ? Context.IntTy : Context.getSizeType();
+ if (RParenLoc.isInvalid())
+ RParenLoc = E->getSourceRange().getEnd();
+ return new (Context)
+ UnaryExprOrTypeTraitExpr(ExprKind, E, ResultType, OpLoc, RParenLoc);
}
ExprResult
@@ -4918,8 +4950,9 @@ Sema::ActOnUnaryExprOrTypeTraitExpr(SourceLocation OpLoc,
}
Expr *ArgEx = (Expr *)TyOrEx;
- ExprResult Result = CreateUnaryExprOrTypeTraitExpr(ArgEx, OpLoc, ExprKind);
- return Result;
+ SourceLocation RParenLoc =
+ ExprKind == UETT_AddrSpaceOf ? ArgRange.getEnd() : SourceLocation();
+ return CreateUnaryExprOrTypeTraitExpr(ArgEx, OpLoc, ExprKind, RParenLoc);
}
bool Sema::CheckAlignasTypeArgument(StringRef KWName, TypeSourceInfo *TInfo,
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index e6abb3ad577c6..20b593c4e9792 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -2857,8 +2857,10 @@ class TreeTransform {
ExprResult RebuildUnaryExprOrTypeTrait(Expr *SubExpr, SourceLocation OpLoc,
UnaryExprOrTypeTrait ExprKind,
SourceRange R) {
- ExprResult Result
- = getSema().CreateUnaryExprOrTypeTraitExpr(SubExpr, OpLoc, ExprKind);
+ SourceLocation RParenLoc =
+ ExprKind == UETT_AddrSpaceOf ? R.getEnd() : SourceLocation();
+ ExprResult Result = getSema().CreateUnaryExprOrTypeTraitExpr(
+ SubExpr, OpLoc, ExprKind, RParenLoc);
if (Result.isInvalid())
return ExprError();
diff --git a/clang/test/CodeGen/builtin-pointee-address-space.c b/clang/test/CodeGen/addrspaceof.c
similarity index 76%
rename from clang/test/CodeGen/builtin-pointee-address-space.c
rename to clang/test/CodeGen/addrspaceof.c
index 7fe5d48c74b9c..c1ef505f3b143 100644
--- a/clang/test/CodeGen/builtin-pointee-address-space.c
+++ b/clang/test/CodeGen/addrspaceof.c
@@ -1,14 +1,14 @@
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | FileCheck %s
int test_default(int *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @test_default(
// CHECK: ret i32 0
int test_address_space(int __attribute__((address_space(4))) *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @test_address_space(
@@ -17,7 +17,7 @@ int test_address_space(int __attribute__((address_space(4))) *p) {
int __attribute__((address_space(7))) *side_effect(void);
int test_unevaluated(void) {
- return __builtin_pointee_address_space(side_effect());
+ return __addrspaceof(*side_effect());
}
// CHECK-LABEL: define{{.*}} i32 @test_unevaluated(
@@ -28,14 +28,14 @@ int array_default[4];
int __attribute__((address_space(5))) array_address_space[4];
int test_array_default(void) {
- return __builtin_pointee_address_space(array_default);
+ return __addrspaceof(array_default);
}
// CHECK-LABEL: define{{.*}} i32 @test_array_default(
// CHECK: ret i32 0
int test_array_address_space(void) {
- return __builtin_pointee_address_space(array_address_space);
+ return __addrspaceof(array_address_space);
}
// CHECK-LABEL: define{{.*}} i32 @test_array_address_space(
diff --git a/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu b/clang/test/CodeGenCUDA/addrspaceof.cu
similarity index 74%
rename from clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
rename to clang/test/CodeGenCUDA/addrspaceof.cu
index a07b1c649193c..1010cbe085948 100644
--- a/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
+++ b/clang/test/CodeGenCUDA/addrspaceof.cu
@@ -20,27 +20,35 @@ __device__ const int const_device_var = 1;
#define EXPECTED_SHARED_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_CUDA_SHARED
#endif
+static_assert(__addrspaceof(device_ptr) == EXPECTED_DEVICE_ADDRESS_SPACE);
+static_assert(__addrspaceof(*device_ptr) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof((device_ptr)) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof((device_var)) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+
#ifdef HOST_TEST
template <class T> consteval int host_consteval_address_space(T *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
template <class T> constexpr int host_constexpr_address_space(T *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
-static_assert(__builtin_pointee_address_space(&device_var) ==
+static_assert(__addrspaceof(device_var) ==
EXPECTED_DEVICE_ADDRESS_SPACE);
-static_assert(__builtin_pointee_address_space(&constant_var) ==
+static_assert(__addrspaceof(constant_var) ==
EXPECTED_CONSTANT_ADDRESS_SPACE);
-static_assert(__builtin_pointee_address_space(&const_device_var) ==
+static_assert(__addrspaceof(const_device_var) ==
EXPECTED_CONSTANT_ADDRESS_SPACE);
-static_assert(__builtin_pointee_address_space(device_array) ==
+static_assert(__addrspaceof(device_array) ==
EXPECTED_DEVICE_ADDRESS_SPACE);
-static_assert(__builtin_pointee_address_space(&device_array) ==
+static_assert(__addrspaceof(*&device_array) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(device_array + 1) ==
+static_assert(__addrspaceof(*(device_array + 1)) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
static_assert(host_consteval_address_space(&device_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
@@ -54,28 +62,28 @@ static_assert(host_constexpr_address_space(&constant_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
extern "C" int test_host_device_var() {
- return __builtin_pointee_address_space(&device_var);
+ return __addrspaceof(device_var);
}
// HOST-LABEL: define{{.*}} i32 @test_host_device_var(
// HOST: ret i32 23
extern "C" int test_host_constant_var() {
- return __builtin_pointee_address_space(&constant_var);
+ return __addrspaceof(constant_var);
}
// HOST-LABEL: define{{.*}} i32 @test_host_constant_var(
// HOST: ret i32 24
extern "C" int test_host_const_device_var() {
- return __builtin_pointee_address_space(&const_device_var);
+ return __addrspaceof(const_device_var);
}
// HOST-LABEL: define{{.*}} i32 @test_host_const_device_var(
// HOST: ret i32 24
extern "C" int test_host_device_array_address() {
- return __builtin_pointee_address_space(&device_array);
+ return __addrspaceof(*&device_array);
}
// HOST-LABEL: define{{.*}} i32 @test_host_device_array_address(
@@ -84,11 +92,11 @@ extern "C" int test_host_device_array_address() {
#else
template <class T> consteval int consteval_address_space(T *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
template <class T> constexpr int constexpr_address_space(T *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
template <int AS> struct AddressSpaceSpecialization;
@@ -107,7 +115,7 @@ struct AddressSpaceSpecialization<EXPECTED_CONSTANT_ADDRESS_SPACE> {
static constexpr int value = EXPECTED_CONSTANT_ADDRESS_SPACE;
};
-static_assert(__builtin_pointee_address_space((int *)&constant_var) ==
+static_assert(__addrspaceof(*(int *)&constant_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
static_assert(consteval_address_space(&device_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
@@ -119,21 +127,21 @@ static_assert(consteval_address_space((int *)&constant_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
static_assert(constexpr_address_space(&constant_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(device_array) ==
+static_assert(__addrspaceof(device_array) ==
EXPECTED_DEVICE_ADDRESS_SPACE);
-static_assert(__builtin_pointee_address_space(&device_array) ==
+static_assert(__addrspaceof(*&device_array) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(device_array + 1) ==
+static_assert(__addrspaceof(*(device_array + 1)) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(device_ptr) ==
+static_assert(__addrspaceof(*device_ptr) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
static_assert(
AddressSpaceSpecialization<
- __builtin_pointee_address_space(&device_var)>::value ==
+ __addrspaceof(device_var)>::value ==
EXPECTED_DEVICE_ADDRESS_SPACE);
static_assert(
AddressSpaceSpecialization<
- __builtin_pointee_address_space(&constant_var)>::value ==
+ __addrspaceof(constant_var)>::value ==
EXPECTED_CONSTANT_ADDRESS_SPACE);
static_assert(
AddressSpaceSpecialization<
@@ -141,7 +149,7 @@ static_assert(
__CLANG_ADDRESS_SPACE_DEFAULT);
extern "C" __device__ int test_generic_pointer(int *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @test_generic_pointer(
@@ -150,19 +158,19 @@ extern "C" __device__ int test_generic_pointer(int *p) {
extern "C" __device__ int test_shared_local() {
__shared__ int shared_var;
__shared__ int shared_array[4];
- static_assert(__builtin_pointee_address_space(&shared_var) ==
+ static_assert(__addrspaceof(shared_var) ==
EXPECTED_SHARED_ADDRESS_SPACE);
- static_assert(__builtin_pointee_address_space((void *)&shared_var) ==
+ static_assert(__addrspaceof(*(char *)&shared_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
- static_assert(__builtin_pointee_address_space(&shared_var + 1) ==
+ static_assert(__addrspaceof(*(&shared_var + 1)) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
- static_assert(__builtin_pointee_address_space(shared_array) ==
+ static_assert(__addrspaceof(shared_array) ==
EXPECTED_SHARED_ADDRESS_SPACE);
- static_assert(__builtin_pointee_address_space(&shared_array) ==
+ static_assert(__addrspaceof(*&shared_array) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
- static_assert(__builtin_pointee_address_space(&shared_array[0]) ==
+ static_assert(__addrspaceof(shared_array[0]) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
- static_assert(__builtin_pointee_address_space(shared_array + 1) ==
+ static_assert(__addrspaceof(*(shared_array + 1)) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
static_assert(consteval_address_space(&shared_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
@@ -172,9 +180,9 @@ extern "C" __device__ int test_shared_local() {
__CLANG_ADDRESS_SPACE_DEFAULT);
static_assert(
AddressSpaceSpecialization<
- __builtin_pointee_address_space(&shared_var)>::value ==
+ __addrspaceof(shared_var)>::value ==
EXPECTED_SHARED_ADDRESS_SPACE);
- return __builtin_pointee_address_space(&shared_var);
+ return __addrspaceof(shared_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_shared_local(
@@ -182,7 +190,7 @@ extern "C" __device__ int test_shared_local() {
// HIP: ret i32 25
extern "C" __device__ int test_device_var() {
- return __builtin_pointee_address_space(&device_var);
+ return __addrspaceof(device_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_device_var(
@@ -190,7 +198,7 @@ extern "C" __device__ int test_device_var() {
// HIP: ret i32 23
extern "C" __device__ int test_device_array() {
- return __builtin_pointee_address_space(device_array);
+ return __addrspaceof(device_array);
}
// CHECK-LABEL: define{{.*}} i32 @test_device_array(
@@ -198,28 +206,28 @@ extern "C" __device__ int test_device_array() {
// HIP: ret i32 23
extern "C" __device__ int test_device_array_address() {
- return __builtin_pointee_address_space(&device_array);
+ return __addrspaceof(*&device_array);
}
// CHECK-LABEL: define{{.*}} i32 @test_device_array_address(
// CHECK: ret i32 0
extern "C" __device__ int test_device_array_arithmetic() {
- return __builtin_pointee_address_space(device_array + 1);
+ return __addrspaceof(*(device_array + 1));
}
// CHECK-LABEL: define{{.*}} i32 @test_device_array_arithmetic(
// CHECK: ret i32 0
extern "C" __device__ int test_device_pointer_value() {
- return __builtin_pointee_address_space(device_ptr);
+ return __addrspaceof(*device_ptr);
}
// CHECK-LABEL: define{{.*}} i32 @test_device_pointer_value(
// CHECK: ret i32 0
extern "C" __device__ int test_constant_var() {
- return __builtin_pointee_address_space(&constant_var);
+ return __addrspaceof(constant_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_constant_var(
@@ -227,7 +235,7 @@ extern "C" __device__ int test_constant_var() {
// HIP: ret i32 24
extern "C" __device__ int test_const_device_var() {
- return __builtin_pointee_address_space(&const_device_var);
+ return __addrspaceof(const_device_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_const_device_var(
@@ -235,7 +243,7 @@ extern "C" __device__ int test_const_device_var() {
// HIP: ret i32 24
extern "C" __device__ int test_explicit_cast() {
- return __builtin_pointee_address_space((int *)&constant_var);
+ return __addrspaceof(*(int *)&constant_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_explicit_cast(
@@ -243,7 +251,7 @@ extern "C" __device__ int test_explicit_cast() {
extern "C" __device__ int
test_target_address_space_3(int __attribute__((address_space(3))) *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @test_target_address_space_3(
@@ -251,7 +259,7 @@ test_target_address_space_3(int __attribute__((address_space(3))) *p) {
extern "C" __device__ int
test_target_address_space_4(int __attribute__((address_space(4))) *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @test_target_address_space_4(
diff --git a/clang/test/CodeGenCXX/mangle-addrspaceof.cpp b/clang/test/CodeGenCXX/mangle-addrspaceof.cpp
new file mode 100644
index 0000000000000..a8e521e033ee7
--- /dev/null
+++ b/clang/test/CodeGenCXX/mangle-addrspaceof.cpp
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -std=c++20 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -std=c++20 -ast-print %s | FileCheck %s --check-prefix=PRINT
+
+int global;
+constexpr int direct_entity = __addrspaceof(global);
+constexpr int parenthesized_expression = __addrspaceof((global));
+
+// PRINT: constexpr int direct_entity = __addrspaceof(global);
+// PRINT: constexpr int parenthesized_expression = __addrspaceof((global));
+
+template <class T> void type_operand(decltype(__addrspaceof(T))) {}
+template void type_operand<int>(int);
+
+template <class T>
+void expression_operand(T &value, decltype(__addrspaceof(value))) {}
+template void expression_operand<int>(int &, int);
+
+// CHECK-DAG: define weak_odr void @_Z12type_operandIiEvDTu13__addrspaceofT_EE(
+// CHECK-DAG: define weak_odr void @_Z18expression_operandIiEvRT_DTu13__addrspaceofXfL0p_EEE(
diff --git a/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl b/clang/test/CodeGenOpenCL/addrspaceof.cl
similarity index 81%
rename from clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl
rename to clang/test/CodeGenOpenCL/addrspaceof.cl
index 0fd1569daebad..05ed9bf72be53 100644
--- a/clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl
+++ b/clang/test/CodeGenOpenCL/addrspaceof.cl
@@ -2,7 +2,7 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -Wno-deprecated-attributes -emit-llvm -o - %s | FileCheck %s
int global_as(__global int *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @global_as(
@@ -10,7 +10,7 @@ int global_as(__global int *p) {
// CHECK: ret i32 1
int local_as(__local int *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @local_as(
@@ -18,7 +18,7 @@ int local_as(__local int *p) {
// CHECK: ret i32 2
int global_device_as(__attribute__((opencl_global_device)) int *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @global_device_as(
@@ -26,7 +26,7 @@ int global_device_as(__attribute__((opencl_global_device)) int *p) {
// CHECK: ret i32 1
int global_host_as(__attribute__((opencl_global_host)) int *p) {
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @global_host_as(
diff --git a/clang/test/CodeGenSYCL/builtin-pointee-address-space.cpp b/clang/test/CodeGenSYCL/addrspaceof.cpp
similarity index 76%
rename from clang/test/CodeGenSYCL/builtin-pointee-address-space.cpp
rename to clang/test/CodeGenSYCL/addrspaceof.cpp
index 5ffc70b72cf28..427992536b14c 100644
--- a/clang/test/CodeGenSYCL/builtin-pointee-address-space.cpp
+++ b/clang/test/CodeGenSYCL/addrspaceof.cpp
@@ -2,9 +2,9 @@
[[clang::sycl_external]] int
global_device_as(__attribute__((opencl_global_device)) int *p) {
- static_assert(__builtin_pointee_address_space(p) ==
+ static_assert(__addrspaceof(*p) ==
__CLANG_ADDRESS_SPACE_SYCL_GLOBAL);
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @_Z16global_device_asPU3AS5i(
@@ -13,9 +13,9 @@ global_device_as(__attribute__((opencl_global_device)) int *p) {
[[clang::sycl_external]] int
global_host_as(__attribute__((opencl_global_host)) int *p) {
- static_assert(__builtin_pointee_address_space(p) ==
+ static_assert(__addrspaceof(*p) ==
__CLANG_ADDRESS_SPACE_SYCL_GLOBAL);
- return __builtin_pointee_address_space(p);
+ return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @_Z14global_host_asPU3AS6i(
diff --git a/clang/test/PCH/exprs.c b/clang/test/PCH/exprs.c
index 1244b2faaf7f2..95b4838376f0d 100644
--- a/clang/test/PCH/exprs.c
+++ b/clang/test/PCH/exprs.c
@@ -47,6 +47,7 @@ offsetof_type *offsetof_ptr = &size_type_value;
typeof(sizeof(float)) size_t_value;
typeof_sizeof *size_t_ptr = &size_t_value;
typeof_sizeof2 *size_t_ptr2 = &size_t_value;
+_Static_assert(AddressSpaceOfI == __CLANG_ADDRESS_SPACE_DEFAULT, "");
// ArraySubscriptExpr
array_subscript *double_ptr1_5 = &floating;
diff --git a/clang/test/PCH/exprs.h b/clang/test/PCH/exprs.h
index d6735a727a538..2f163d1e26f9f 100644
--- a/clang/test/PCH/exprs.h
+++ b/clang/test/PCH/exprs.h
@@ -41,6 +41,7 @@ typedef typeof(__builtin_offsetof(struct Z, y.array[1 + 2].member))
// UnaryExprOrTypeTraitExpr
typedef typeof(sizeof(int)) typeof_sizeof;
typedef typeof(sizeof(Enumerator)) typeof_sizeof2;
+enum { AddressSpaceOfI = __addrspaceof(i) };
// ArraySubscriptExpr
extern double values[];
diff --git a/clang/test/SemaCUDA/builtin-pointee-address-space.cu b/clang/test/SemaCUDA/addrspaceof.cu
similarity index 50%
rename from clang/test/SemaCUDA/builtin-pointee-address-space.cu
rename to clang/test/SemaCUDA/addrspaceof.cu
index cd3ee9aa3a56b..01b3389ab7ad4 100644
--- a/clang/test/SemaCUDA/builtin-pointee-address-space.cu
+++ b/clang/test/SemaCUDA/addrspaceof.cu
@@ -14,30 +14,30 @@ __device__ int *device_ptr;
__constant__ int constant_var;
__device__ const int const_device_var = 1;
-static_assert(__builtin_pointee_address_space(&device_var) ==
+static_assert(__addrspaceof(device_var) ==
__CLANG_ADDRESS_SPACE_HIP_DEVICE);
-static_assert(__builtin_pointee_address_space(device_array) ==
+static_assert(__addrspaceof(device_array) ==
__CLANG_ADDRESS_SPACE_HIP_DEVICE);
-static_assert(__builtin_pointee_address_space(&device_array) ==
+static_assert(__addrspaceof(*&device_array) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(device_array + 1) ==
+static_assert(__addrspaceof(*(device_array + 1)) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(&constant_var) ==
+static_assert(__addrspaceof(constant_var) ==
__CLANG_ADDRESS_SPACE_HIP_CONSTANT);
-static_assert(__builtin_pointee_address_space(&const_device_var) ==
+static_assert(__addrspaceof(const_device_var) ==
__CLANG_ADDRESS_SPACE_HIP_CONSTANT);
-static_assert(__builtin_pointee_address_space((int *)&constant_var) ==
+static_assert(__addrspaceof(*(int *)&constant_var) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(device_ptr) ==
+static_assert(__addrspaceof(*device_ptr) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
void host_queries() {
- (void)__builtin_pointee_address_space(&device_var);
- (void)__builtin_pointee_address_space(device_array);
- (void)__builtin_pointee_address_space(&device_array);
- (void)__builtin_pointee_address_space(device_array + 1);
- (void)__builtin_pointee_address_space(&constant_var);
- (void)__builtin_pointee_address_space(&const_device_var);
- (void)__builtin_pointee_address_space((int *)&constant_var);
- (void)__builtin_pointee_address_space(device_ptr);
+ (void)__addrspaceof(device_var);
+ (void)__addrspaceof(device_array);
+ (void)__addrspaceof(*&device_array);
+ (void)__addrspaceof(*(device_array + 1));
+ (void)__addrspaceof(constant_var);
+ (void)__addrspaceof(const_device_var);
+ (void)__addrspaceof(*(int *)&constant_var);
+ (void)__addrspaceof(*device_ptr);
}
diff --git a/clang/test/SemaCXX/addrspaceof.cpp b/clang/test/SemaCXX/addrspaceof.cpp
new file mode 100644
index 0000000000000..58d195ffcb92c
--- /dev/null
+++ b/clang/test/SemaCXX/addrspaceof.cpp
@@ -0,0 +1,87 @@
+// RUN: %clang_cc1 -fsyntax-only -verify -std=c++20 %s
+// RUN: %clang_cc1 -fsyntax-only -verify -std=c++20 -fexperimental-new-constant-interpreter %s
+
+#if !__has_extension(addrspaceof)
+#error "missing addrspaceof extension"
+#endif
+
+using AS1 = int __attribute__((address_space(1)));
+using AS2 = int __attribute__((address_space(2)));
+
+static_assert(__addrspaceof(int) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(AS1) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+static_assert(__addrspaceof(AS2 &) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 2);
+
+int *p0;
+AS1 *p1;
+
+static_assert(__addrspaceof(p0) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(p1) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(*p0) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(*p1) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+
+int global;
+AS1 global_as1;
+int array[4];
+AS2 array_as2[4];
+
+static_assert(__addrspaceof(global) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(global_as1) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+static_assert(__addrspaceof(array) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(array_as2) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 2);
+static_assert(__addrspaceof(array_as2[0]) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 2);
+
+struct S {
+ int member;
+ static AS1 member_as1;
+};
+AS1 S::member_as1;
+
+S object;
+static_assert(__addrspaceof(object.member) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(object.member_as1) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+
+template <class T> constexpr int type_address_space() {
+ return __addrspaceof(T);
+}
+
+template <class T> constexpr int expression_address_space(T &value) {
+ return __addrspaceof(value);
+}
+
+static_assert(type_address_space<AS1>() ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+static_assert(expression_address_space(global_as1) ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+
+template <int AS> struct AddressSpaceSpecialization;
+template <>
+struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1> {
+ static constexpr int value = __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1;
+};
+
+static_assert(AddressSpaceSpecialization<__addrspaceof(AS1)>::value ==
+ __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+
+void function();
+
+void errors() {
+ (void)__addrspaceof global;
+ // expected-error at -1 {{expected '(' after '__addrspaceof'}}
+ (void)__addrspaceof(0);
+ // expected-error at -1 {{expression operand of '__addrspaceof' must be an lvalue}}
+ (void)__addrspaceof(&global);
+ // expected-error at -1 {{expression operand of '__addrspaceof' must be an lvalue}}
+ (void)__addrspaceof(p0 + 1);
+ // expected-error at -1 {{expression operand of '__addrspaceof' must be an lvalue}}
+ (void)__addrspaceof(function);
+ // expected-error at -1 {{entity operand of '__addrspaceof' must name a variable or data member}}
+}
diff --git a/clang/test/SemaCXX/builtin-pointee-address-space.cpp b/clang/test/SemaCXX/builtin-pointee-address-space.cpp
deleted file mode 100644
index 209fc315339be..0000000000000
--- a/clang/test/SemaCXX/builtin-pointee-address-space.cpp
+++ /dev/null
@@ -1,91 +0,0 @@
-// RUN: %clang_cc1 -fsyntax-only -verify -std=c++20 %s
-
-#if !__has_builtin(__builtin_pointee_address_space)
-#error "missing __builtin_pointee_address_space"
-#endif
-
-int *p0;
-int __attribute__((address_space(1))) *p1;
-const int __attribute__((address_space(7))) *p7;
-
-static_assert(__builtin_pointee_address_space(p0) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(p1) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
-static_assert(__builtin_pointee_address_space(p7) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 7);
-static_assert(__builtin_pointee_address_space(
- (int __attribute__((address_space(3))) *)0) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 3);
-
-int global;
-int __attribute__((address_space(4))) global_as4;
-int arr[4];
-int __attribute__((address_space(5))) arr_as5[4];
-
-static_assert(__builtin_pointee_address_space(&global) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(&global_as4) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 4);
-static_assert(__builtin_pointee_address_space(arr) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__builtin_pointee_address_space(arr_as5) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 5);
-
-template <class T> constexpr int get_as(T *p) {
- return __builtin_pointee_address_space(p);
-}
-
-static_assert(get_as((int *)0) == __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(get_as((int __attribute__((address_space(1))) *)0) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
-
-template <class T> struct PointeeAddressSpace {
- static constexpr int value =
- __builtin_pointee_address_space((T *)0);
-};
-
-static_assert(PointeeAddressSpace<int>::value ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(PointeeAddressSpace<
- int __attribute__((address_space(2)))>::value ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 2);
-
-template <int AS> struct AddressSpaceSpecialization;
-template <>
-struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_DEFAULT> {
- static constexpr int value = __CLANG_ADDRESS_SPACE_DEFAULT;
-};
-template <>
-struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1> {
- static constexpr int value = __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1;
-};
-template <>
-struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_TARGET_OFFSET + 5> {
- static constexpr int value = __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 5;
-};
-
-static_assert(
- AddressSpaceSpecialization<
- __builtin_pointee_address_space(p1)>::value ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
-static_assert(
- AddressSpaceSpecialization<
- __builtin_pointee_address_space(arr_as5)>::value ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 5);
-
-void errors() {
- int i;
- void f();
-
- (void)__builtin_pointee_address_space();
- // expected-error at -1 {{too few arguments}}
- (void)__builtin_pointee_address_space(p0, p1);
- // expected-error at -1 {{too many arguments}}
- (void)__builtin_pointee_address_space(i);
- // expected-error at -1 {{argument to '__builtin_pointee_address_space' must be a pointer or array expression}}
- (void)__builtin_pointee_address_space(0);
- // expected-error at -1 {{argument to '__builtin_pointee_address_space' must be a pointer or array expression}}
- (void)__builtin_pointee_address_space(f);
- // expected-error at -1 {{argument to '__builtin_pointee_address_space' must be a pointer or array expression}}
-}
>From 1cdf7e7ab834782b0162988f8cf352fd9d50593f Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 24 Jul 2026 13:29:55 -0400
Subject: [PATCH 4/7] [Clang] Distinguish addrspaceof forms in mangling
Parentheses change an entity query into an expression query, but normal expression mangling ignores them. Record the query form so dependent function templates do not get the same symbol name.
---
clang/lib/AST/ItaniumMangle.cpp | 38 +++++++++++++-------
clang/test/CodeGenCUDA/mangle-addrspaceof.cu | 26 ++++++++++++++
clang/test/CodeGenCXX/mangle-addrspaceof.cpp | 4 ++-
3 files changed, 55 insertions(+), 13 deletions(-)
create mode 100644 clang/test/CodeGenCUDA/mangle-addrspaceof.cu
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index 2e90f15ab2329..d1da36cc21fbd 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -4552,6 +4552,15 @@ void CXXNameMangler::mangleType(const TypeOfExprType *T) {
Out << "u6typeof";
}
+static bool isUnparenthesizedIdOrMemberExpr(const Expr *E) {
+ // This is the exhaustive list used for decltype's entity form. Do not
+ // ignore parentheses because they change the meaning of decltype and
+ // __addrspaceof.
+ return isa<DeclRefExpr, MemberExpr, UnresolvedLookupExpr,
+ DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr,
+ UnresolvedMemberExpr>(E);
+}
+
void CXXNameMangler::mangleType(const DecltypeType *T) {
Expr *E = T->getUnderlyingExpr();
@@ -4559,16 +4568,7 @@ void CXXNameMangler::mangleType(const DecltypeType *T) {
// # or class member access
// ::= DT <expression> E # decltype of an expression
- // This purports to be an exhaustive list of id-expressions and
- // class member accesses. Note that we do not ignore parentheses;
- // parentheses change the semantics of decltype for these
- // expressions (and cause the mangler to use the other form).
- if (isa<DeclRefExpr>(E) ||
- isa<MemberExpr>(E) ||
- isa<UnresolvedLookupExpr>(E) ||
- isa<DependentScopeDeclRefExpr>(E) ||
- isa<CXXDependentScopeMemberExpr>(E) ||
- isa<UnresolvedMemberExpr>(E))
+ if (isUnparenthesizedIdOrMemberExpr(E))
Out << "Dt";
else
Out << "DT";
@@ -5471,9 +5471,23 @@ void CXXNameMangler::mangleExpression(const Expr *E, unsigned Arity,
MangleAlignofSizeofArg();
break;
- case UETT_AddrSpaceOf:
- MangleExtensionBuiltin(SAE);
+ case UETT_AddrSpaceOf: {
+ if (SAE->isArgumentType()) {
+ MangleExtensionBuiltin(SAE);
+ break;
+ }
+
+ // Normal expression mangling ignores parentheses, but __addrspaceof(x)
+ // and __addrspaceof((x)) can have different values. Record whether the
+ // operand uses the entity form before mangling the operand itself.
+ mangleVendorType(getTraitSpelling(SAE->getKind()));
+ bool IsEntity =
+ isUnparenthesizedIdOrMemberExpr(SAE->getArgumentExpr());
+ Out << (IsEntity ? "Lb1E" : "Lb0E");
+ mangleTemplateArgExpr(SAE->getArgumentExpr());
+ Out << 'E';
break;
+ }
case UETT_CountOf:
case UETT_VectorElements:
case UETT_OpenMPRequiredSimdAlign:
diff --git a/clang/test/CodeGenCUDA/mangle-addrspaceof.cu b/clang/test/CodeGenCUDA/mangle-addrspaceof.cu
new file mode 100644
index 0000000000000..7866e3dd6b052
--- /dev/null
+++ b/clang/test/CodeGenCUDA/mangle-addrspaceof.cu
@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++20 -emit-llvm -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+template <class T> __device__ T value;
+template <int> struct Result {};
+
+// Parentheses change an entity query into an expression query. Both templates
+// must therefore have different mangled names even though normal expression
+// mangling ignores parentheses.
+template <class T>
+__device__ Result<__addrspaceof(value<T>)> query(T) {
+ return {};
+}
+
+template <class T>
+__device__ Result<__addrspaceof((value<T>))> query(T) {
+ return {};
+}
+
+template __device__ Result<__CLANG_ADDRESS_SPACE_CUDA_DEVICE> query<int>(int);
+template __device__ Result<__CLANG_ADDRESS_SPACE_DEFAULT> query<int>(int);
+
+// Lb1E records the entity form; Lb0E records the expression form.
+// CHECK-DAG: define {{.*}}@_Z5queryIiE6ResultIXu13__addrspaceofLb1EX5valueIT_EEEEES1_(
+// CHECK-DAG: define {{.*}}@_Z5queryIiE6ResultIXu13__addrspaceofLb0EX5valueIT_EEEEES1_(
diff --git a/clang/test/CodeGenCXX/mangle-addrspaceof.cpp b/clang/test/CodeGenCXX/mangle-addrspaceof.cpp
index a8e521e033ee7..112a9a165347c 100644
--- a/clang/test/CodeGenCXX/mangle-addrspaceof.cpp
+++ b/clang/test/CodeGenCXX/mangle-addrspaceof.cpp
@@ -16,4 +16,6 @@ void expression_operand(T &value, decltype(__addrspaceof(value))) {}
template void expression_operand<int>(int &, int);
// CHECK-DAG: define weak_odr void @_Z12type_operandIiEvDTu13__addrspaceofT_EE(
-// CHECK-DAG: define weak_odr void @_Z18expression_operandIiEvRT_DTu13__addrspaceofXfL0p_EEE(
+// The boolean template argument records the entity form because ordinary
+// expression mangling does not preserve parentheses.
+// CHECK-DAG: define weak_odr void @_Z18expression_operandIiEvRT_DTu13__addrspaceofLb1EXfL0p_EEE(
>From fb0abdc15378ba5df8512ce1f9c08f4ccdf06df7 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 24 Jul 2026 14:15:13 -0400
Subject: [PATCH 5/7] [Clang] Handle external device constants in addrspaceof
An external const device declaration is not promoted to constant memory. Report its device address space consistently in host and device compilations.
---
clang/lib/AST/Expr.cpp | 1 +
clang/test/CodeGenCUDA/addrspaceof.cu | 3 +++
2 files changed, 4 insertions(+)
diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index 898ec6e84cb55..26722a9727b97 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/lib/AST/Expr.cpp
@@ -1741,6 +1741,7 @@ static std::optional<LangAS> getCUDADeclAddressSpace(const ASTContext &Ctx,
// Host compilation does not attach the implicit CUDAConstantAttr that
// SemaCUDA adds in device compilation.
if (!Ctx.getLangOpts().CUDAIsDevice && VD->hasAttr<CUDADeviceAttr>() &&
+ VD->isThisDeclarationADefinition() == VarDecl::Definition &&
(VD->isFileVarDecl() || VD->isStaticDataMember()) &&
(VD->isConstexpr() || VD->getType().isConstQualified()))
return LangAS::cuda_constant;
diff --git a/clang/test/CodeGenCUDA/addrspaceof.cu b/clang/test/CodeGenCUDA/addrspaceof.cu
index 1010cbe085948..b9fd475f8c7b0 100644
--- a/clang/test/CodeGenCUDA/addrspaceof.cu
+++ b/clang/test/CodeGenCUDA/addrspaceof.cu
@@ -9,6 +9,7 @@ __device__ int device_array[4];
__device__ int *device_ptr;
__constant__ int constant_var;
__device__ const int const_device_var = 1;
+extern __device__ const int extern_const_device_var;
#if defined(__HIP__)
#define EXPECTED_DEVICE_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_HIP_DEVICE
@@ -27,6 +28,8 @@ static_assert(__addrspaceof((device_ptr)) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
static_assert(__addrspaceof((device_var)) ==
__CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(extern_const_device_var) ==
+ EXPECTED_DEVICE_ADDRESS_SPACE);
#ifdef HOST_TEST
>From 03c6ab22b45443d97937c37dd4c542fa887a656d Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 24 Jul 2026 21:11:58 -0400
Subject: [PATCH 6/7] [Clang] Fix addrspaceof CI failures
The mangling test assumed the host uses the Itanium ABI and failed on Windows. Use an explicit target triple and format the recent code.
---
clang/lib/AST/ItaniumMangle.cpp | 3 +--
clang/lib/Sema/SemaExpr.cpp | 6 +++---
clang/test/CodeGenCXX/mangle-addrspaceof.cpp | 2 +-
3 files changed, 5 insertions(+), 6 deletions(-)
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index d1da36cc21fbd..5f0e2db170cfc 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -5481,8 +5481,7 @@ void CXXNameMangler::mangleExpression(const Expr *E, unsigned Arity,
// and __addrspaceof((x)) can have different values. Record whether the
// operand uses the entity form before mangling the operand itself.
mangleVendorType(getTraitSpelling(SAE->getKind()));
- bool IsEntity =
- isUnparenthesizedIdOrMemberExpr(SAE->getArgumentExpr());
+ bool IsEntity = isUnparenthesizedIdOrMemberExpr(SAE->getArgumentExpr());
Out << (IsEntity ? "Lb1E" : "Lb0E");
mangleTemplateArgExpr(SAE->getArgumentExpr());
Out << 'E';
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 66e69557005b0..99d938cbcfb6a 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -4908,9 +4908,9 @@ ExprResult Sema::CreateUnaryExprOrTypeTraitExpr(Expr *E, SourceLocation OpLoc,
} else if (ExprKind == UETT_VecStep) {
isInvalid = CheckVecStepExpr(E);
} else if (ExprKind == UETT_OpenMPRequiredSimdAlign) {
- Diag(E->getExprLoc(), diag::err_openmp_default_simd_align_expr);
- isInvalid = true;
- } else if (E->refersToBitField()) { // C99 6.5.3.4p1.
+ Diag(E->getExprLoc(), diag::err_openmp_default_simd_align_expr);
+ isInvalid = true;
+ } else if (E->refersToBitField()) { // C99 6.5.3.4p1.
Diag(E->getExprLoc(), diag::err_sizeof_alignof_typeof_bitfield) << 0;
isInvalid = true;
} else if (ExprKind == UETT_VectorElements || ExprKind == UETT_SizeOf ||
diff --git a/clang/test/CodeGenCXX/mangle-addrspaceof.cpp b/clang/test/CodeGenCXX/mangle-addrspaceof.cpp
index 112a9a165347c..8e47a19d88a67 100644
--- a/clang/test/CodeGenCXX/mangle-addrspaceof.cpp
+++ b/clang/test/CodeGenCXX/mangle-addrspaceof.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -std=c++20 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -std=c++20 -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -std=c++20 -ast-print %s | FileCheck %s --check-prefix=PRINT
int global;
>From f399d4bb2a3fed056a575157c230b896f119be0e Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Tue, 11 Aug 2026 21:14:54 -0400
Subject: [PATCH 7/7] [Clang] Use common addrspaceof macro values
The predefined values for __addrspaceof should describe common memory regions instead of exposing every Clang LangAS value.
Collapse equivalent OpenCL, CUDA, HIP, SYCL, and HLSL regions to shared values, rename the public macros to __addrspace_*, and avoid publishing names for internal address spaces.
---
clang/docs/LanguageExtensions.md | 44 ++++-----
clang/docs/ReleaseNotes.md | 2 +-
clang/include/clang/Basic/AddressSpaces.h | 74 ++++++--------
clang/lib/AST/Expr.cpp | 7 +-
clang/lib/Frontend/InitPreprocessor.cpp | 60 ++----------
clang/test/CodeGenCUDA/addrspaceof.cu | 97 +++++++++----------
clang/test/CodeGenCUDA/mangle-addrspaceof.cu | 4 +-
clang/test/CodeGenSYCL/addrspaceof.cpp | 8 +-
clang/test/PCH/exprs.c | 2 +-
.../test/Preprocessor/address-space-macros.c | 37 ++-----
clang/test/Preprocessor/init-aarch64.c | 37 ++-----
clang/test/Preprocessor/init.c | 37 ++-----
clang/test/SemaCUDA/addrspaceof.cu | 16 +--
clang/test/SemaCXX/addrspaceof.cpp | 38 ++++----
14 files changed, 178 insertions(+), 285 deletions(-)
diff --git a/clang/docs/LanguageExtensions.md b/clang/docs/LanguageExtensions.md
index 566be2bd47dbb..b83815e5059c0 100644
--- a/clang/docs/LanguageExtensions.md
+++ b/clang/docs/LanguageExtensions.md
@@ -4025,24 +4025,24 @@ query into an expression query.
```c++
__device__ int *p;
-static_assert(__addrspaceof(p) == __CLANG_ADDRESS_SPACE_CUDA_DEVICE);
-static_assert(__addrspaceof(*p) == __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__addrspaceof((p)) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(p) == __ADDRSPACE_GLOBAL);
+static_assert(__addrspaceof(*p) == __ADDRSPACE_DEFAULT);
+static_assert(__addrspaceof((p)) == __ADDRSPACE_DEFAULT);
```
The first query reports where the pointer variable is stored. The second
reports the address space of the pointed-to object type. The third treats `p`
as an lvalue expression because of the extra parentheses.
-Clang predefines named macros for the values returned by this operator. For
-language-level address spaces, the result is one of the
-`__CLANG_ADDRESS_SPACE_*` values. For an address space written explicitly with
+Clang predefines macros for the values returned by this operator. Most language
+address spaces are reported using common memory regions:
+`__ADDRSPACE_DEFAULT`, `__ADDRSPACE_GLOBAL`, `__ADDRSPACE_LOCAL`,
+`__ADDRSPACE_CONSTANT`, `__ADDRSPACE_PRIVATE`, and `__ADDRSPACE_GENERIC`.
+OpenCL, CUDA, HIP, SYCL, and HLSL address spaces that describe the same memory
+region produce the same value. For an address space written explicitly with
`__attribute__((address_space(N)))`, the result is
-`__CLANG_ADDRESS_SPACE_TARGET_OFFSET + N`. CUDA mode uses the
-`__CLANG_ADDRESS_SPACE_CUDA_*` values, and HIP mode uses the corresponding
-`__CLANG_ADDRESS_SPACE_HIP_*` values. Deprecated OpenCL and SYCL
-`global_device` and `global_host` address spaces are reported as their
-corresponding global address space.
+`__ADDRSPACE_TARGET_OFFSET + N`. Deprecated OpenCL and SYCL `global_device` and
+`global_host` address spaces are reported as the global address space.
**Example use**:
@@ -4051,9 +4051,9 @@ using local_int = int __attribute__((address_space(3)));
local_int *p;
static_assert(__addrspaceof(local_int) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 3);
+ __ADDRSPACE_TARGET_OFFSET + 3);
static_assert(__addrspaceof(*p) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 3);
+ __ADDRSPACE_TARGET_OFFSET + 3);
```
**OpenCL example use**:
@@ -4063,9 +4063,9 @@ __global int *global_p;
__local int *local_p;
static_assert(__addrspaceof(*global_p) ==
- __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL);
+ __ADDRSPACE_GLOBAL);
static_assert(__addrspaceof(*local_p) ==
- __CLANG_ADDRESS_SPACE_OPENCL_LOCAL);
+ __ADDRSPACE_LOCAL);
```
**CUDA/HIP example use**:
@@ -4077,18 +4077,18 @@ __constant__ int cst;
__device__ const int dev_cst = 1;
static_assert(__addrspaceof(dev) ==
- __CLANG_ADDRESS_SPACE_CUDA_DEVICE);
+ __ADDRSPACE_GLOBAL);
static_assert(__addrspaceof(dev_arr) ==
- __CLANG_ADDRESS_SPACE_CUDA_DEVICE);
+ __ADDRSPACE_GLOBAL);
static_assert(__addrspaceof(cst) ==
- __CLANG_ADDRESS_SPACE_CUDA_CONSTANT);
+ __ADDRSPACE_CONSTANT);
static_assert(__addrspaceof(dev_cst) ==
- __CLANG_ADDRESS_SPACE_CUDA_CONSTANT);
+ __ADDRSPACE_CONSTANT);
```
-HIP compilation returns the corresponding `__CLANG_ADDRESS_SPACE_HIP_*`
-values. The operator is a static frontend query. It does not classify an
-arbitrary runtime pointer value and does not use optimizer or backend analysis.
+CUDA and HIP compilation return the same common memory-region values. The
+operator is a static frontend query. It does not classify an arbitrary runtime
+pointer value and does not use optimizer or backend analysis.
Query for this extension with `__has_extension(addrspaceof)`.
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index c9bb3e995c6c8..8e958d786cd62 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -186,7 +186,7 @@ features cannot lower the translation-unit ABI level;
identifier for a type, directly named entity, or lvalue expression. It can
report CUDA/HIP storage address spaces for known `__device__`, `__shared__`,
and `__constant__` variables. Clang also now emits predefined
- `__CLANG_ADDRESS_SPACE_*` macros for these values.
+ `__ADDRSPACE_*` macros for these values.
### New Compiler Flags
diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h
index ecb3430392095..da75887cefbcb 100644
--- a/clang/include/clang/Basic/AddressSpaces.h
+++ b/clang/include/clang/Basic/AddressSpaces.h
@@ -122,36 +122,24 @@ namespace AddressSpaceQuery {
enum ID : unsigned {
Default = 0,
- OpenCLGlobal = 1,
- OpenCLLocal = 2,
- OpenCLConstant = 3,
- OpenCLPrivate = 4,
- OpenCLGeneric = 5,
- CUDADevice = 6,
- CUDAConstant = 7,
- CUDAShared = 8,
- SYCLGlobal = 9,
- SYCLLocal = 10,
- SYCLPrivate = 11,
- Ptr32Sptr = 12,
- Ptr32Uptr = 13,
- Ptr64 = 14,
- HLSLGroupShared = 15,
- HLSLConstant = 16,
- HLSLPrivate = 17,
- HLSLDevice = 18,
- HLSLInput = 19,
- HLSLOutput = 20,
- HLSLPushConstant = 21,
- WasmFuncRef = 22,
- HIPDevice = 23,
- HIPConstant = 24,
- HIPShared = 25,
+ Global = 1,
+ Local = 2,
+ Constant = 3,
+ Private = 4,
+ Generic = 5,
+ HLSLInput = 6,
+ HLSLOutput = 7,
+ HLSLPushConstant = 8,
+ Ptr32Sptr = 9,
+ Ptr32Uptr = 10,
+ Ptr64 = 11,
+ WasmFuncRef = 12,
+ AMDGPUBarrier = 13,
TargetOffset = 0x1000000
};
-inline unsigned encode(LangAS AS, bool IsHIP = false) {
+inline unsigned encode(LangAS AS) {
if (isTargetAddressSpace(AS))
return TargetOffset + toTargetAddressSpace(AS);
@@ -159,32 +147,32 @@ inline unsigned encode(LangAS AS, bool IsHIP = false) {
case LangAS::Default:
return Default;
case LangAS::opencl_global:
- return OpenCLGlobal;
+ return Global;
case LangAS::opencl_local:
- return OpenCLLocal;
+ return Local;
case LangAS::opencl_constant:
- return OpenCLConstant;
+ return Constant;
case LangAS::opencl_private:
- return OpenCLPrivate;
+ return Private;
case LangAS::opencl_generic:
- return OpenCLGeneric;
+ return Generic;
case LangAS::opencl_global_device:
case LangAS::opencl_global_host:
- return OpenCLGlobal;
+ return Global;
case LangAS::cuda_device:
- return IsHIP ? HIPDevice : CUDADevice;
+ return Global;
case LangAS::cuda_constant:
- return IsHIP ? HIPConstant : CUDAConstant;
+ return Constant;
case LangAS::cuda_shared:
- return IsHIP ? HIPShared : CUDAShared;
+ return Local;
case LangAS::sycl_global:
case LangAS::sycl_global_device:
case LangAS::sycl_global_host:
- return SYCLGlobal;
+ return Global;
case LangAS::sycl_local:
- return SYCLLocal;
+ return Local;
case LangAS::sycl_private:
- return SYCLPrivate;
+ return Private;
case LangAS::ptr32_sptr:
return Ptr32Sptr;
case LangAS::ptr32_uptr:
@@ -192,13 +180,13 @@ inline unsigned encode(LangAS AS, bool IsHIP = false) {
case LangAS::ptr64:
return Ptr64;
case LangAS::hlsl_groupshared:
- return HLSLGroupShared;
+ return Local;
case LangAS::hlsl_constant:
- return HLSLConstant;
+ return Constant;
case LangAS::hlsl_private:
- return HLSLPrivate;
+ return Private;
case LangAS::hlsl_device:
- return HLSLDevice;
+ return Global;
case LangAS::hlsl_input:
return HLSLInput;
case LangAS::hlsl_output:
@@ -207,6 +195,8 @@ inline unsigned encode(LangAS AS, bool IsHIP = false) {
return HLSLPushConstant;
case LangAS::wasm_funcref:
return WasmFuncRef;
+ case LangAS::amdgpu_barrier:
+ return AMDGPUBarrier;
case LangAS::FirstTargetAddressSpace:
break;
}
diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index 26722a9727b97..7077f7816faee 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/lib/AST/Expr.cpp
@@ -1761,15 +1761,14 @@ unsigned UnaryExprOrTypeTraitExpr::getAddressSpaceQueryResult(
T = D->getType();
LangAS AS = T.getNonReferenceType().getAddressSpace();
if (AS != LangAS::Default)
- return AddressSpaceQuery::encode(AS, Ctx.getLangOpts().HIP);
+ return AddressSpaceQuery::encode(AS);
if (std::optional<LangAS> AS = getCUDADeclAddressSpace(Ctx, D))
- return AddressSpaceQuery::encode(*AS, Ctx.getLangOpts().HIP);
+ return AddressSpaceQuery::encode(*AS);
} else {
T = getArgumentExpr()->getType();
}
- return AddressSpaceQuery::encode(T.getNonReferenceType().getAddressSpace(),
- Ctx.getLangOpts().HIP);
+ return AddressSpaceQuery::encode(T.getNonReferenceType().getAddressSpace());
}
MemberExpr::MemberExpr(Expr *Base, bool IsArrow, SourceLocation OperatorLoc,
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 068448b580155..959506444b159 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -919,59 +919,19 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
Builder.defineMacro(Name, Twine(static_cast<unsigned>(AS)));
};
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_DEFAULT",
- AddressSpaceQuery::Default);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_GLOBAL",
- AddressSpaceQuery::OpenCLGlobal);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_LOCAL",
- AddressSpaceQuery::OpenCLLocal);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_CONSTANT",
- AddressSpaceQuery::OpenCLConstant);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_PRIVATE",
- AddressSpaceQuery::OpenCLPrivate);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_OPENCL_GENERIC",
- AddressSpaceQuery::OpenCLGeneric);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_CUDA_DEVICE",
- AddressSpaceQuery::CUDADevice);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_CUDA_CONSTANT",
- AddressSpaceQuery::CUDAConstant);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_CUDA_SHARED",
- AddressSpaceQuery::CUDAShared);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_SYCL_GLOBAL",
- AddressSpaceQuery::SYCLGlobal);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_SYCL_LOCAL",
- AddressSpaceQuery::SYCLLocal);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_SYCL_PRIVATE",
- AddressSpaceQuery::SYCLPrivate);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_PTR32_SPTR",
- AddressSpaceQuery::Ptr32Sptr);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_PTR32_UPTR",
- AddressSpaceQuery::Ptr32Uptr);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_PTR64",
- AddressSpaceQuery::Ptr64);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED",
- AddressSpaceQuery::HLSLGroupShared);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_CONSTANT",
- AddressSpaceQuery::HLSLConstant);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_PRIVATE",
- AddressSpaceQuery::HLSLPrivate);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_DEVICE",
- AddressSpaceQuery::HLSLDevice);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_INPUT",
+ DefineAddressSpaceMacro("__ADDRSPACE_DEFAULT", AddressSpaceQuery::Default);
+ DefineAddressSpaceMacro("__ADDRSPACE_GLOBAL", AddressSpaceQuery::Global);
+ DefineAddressSpaceMacro("__ADDRSPACE_LOCAL", AddressSpaceQuery::Local);
+ DefineAddressSpaceMacro("__ADDRSPACE_CONSTANT", AddressSpaceQuery::Constant);
+ DefineAddressSpaceMacro("__ADDRSPACE_PRIVATE", AddressSpaceQuery::Private);
+ DefineAddressSpaceMacro("__ADDRSPACE_GENERIC", AddressSpaceQuery::Generic);
+ DefineAddressSpaceMacro("__ADDRSPACE_HLSL_INPUT",
AddressSpaceQuery::HLSLInput);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_OUTPUT",
+ DefineAddressSpaceMacro("__ADDRSPACE_HLSL_OUTPUT",
AddressSpaceQuery::HLSLOutput);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT",
+ DefineAddressSpaceMacro("__ADDRSPACE_HLSL_PUSH_CONSTANT",
AddressSpaceQuery::HLSLPushConstant);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_WASM_FUNCREF",
- AddressSpaceQuery::WasmFuncRef);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HIP_DEVICE",
- AddressSpaceQuery::HIPDevice);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HIP_CONSTANT",
- AddressSpaceQuery::HIPConstant);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_HIP_SHARED",
- AddressSpaceQuery::HIPShared);
- DefineAddressSpaceMacro("__CLANG_ADDRESS_SPACE_TARGET_OFFSET",
+ DefineAddressSpaceMacro("__ADDRSPACE_TARGET_OFFSET",
AddressSpaceQuery::TargetOffset);
// Define macros for floating-point data classes, used in __builtin_isfpclass.
diff --git a/clang/test/CodeGenCUDA/addrspaceof.cu b/clang/test/CodeGenCUDA/addrspaceof.cu
index b9fd475f8c7b0..d0ffc63ce7cc0 100644
--- a/clang/test/CodeGenCUDA/addrspaceof.cu
+++ b/clang/test/CodeGenCUDA/addrspaceof.cu
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++20 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,CUDA %s
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip -std=c++20 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,HIP %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++20 -emit-llvm -o - %s | FileCheck --check-prefix=CHECK %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip -std=c++20 -emit-llvm -o - %s | FileCheck --check-prefix=CHECK %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -x hip -std=c++20 -DHOST_TEST -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s
#include "Inputs/cuda.h"
@@ -12,22 +12,22 @@ __device__ const int const_device_var = 1;
extern __device__ const int extern_const_device_var;
#if defined(__HIP__)
-#define EXPECTED_DEVICE_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_HIP_DEVICE
-#define EXPECTED_CONSTANT_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_HIP_CONSTANT
-#define EXPECTED_SHARED_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_HIP_SHARED
+#define EXPECTED_DEVICE_ADDRESS_SPACE __ADDRSPACE_GLOBAL
+#define EXPECTED_CONSTANT_ADDRESS_SPACE __ADDRSPACE_CONSTANT
+#define EXPECTED_SHARED_ADDRESS_SPACE __ADDRSPACE_LOCAL
#else
-#define EXPECTED_DEVICE_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_CUDA_DEVICE
-#define EXPECTED_CONSTANT_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_CUDA_CONSTANT
-#define EXPECTED_SHARED_ADDRESS_SPACE __CLANG_ADDRESS_SPACE_CUDA_SHARED
+#define EXPECTED_DEVICE_ADDRESS_SPACE __ADDRSPACE_GLOBAL
+#define EXPECTED_CONSTANT_ADDRESS_SPACE __ADDRSPACE_CONSTANT
+#define EXPECTED_SHARED_ADDRESS_SPACE __ADDRSPACE_LOCAL
#endif
static_assert(__addrspaceof(device_ptr) == EXPECTED_DEVICE_ADDRESS_SPACE);
static_assert(__addrspaceof(*device_ptr) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof((device_ptr)) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof((device_var)) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(extern_const_device_var) ==
EXPECTED_DEVICE_ADDRESS_SPACE);
@@ -50,40 +50,40 @@ static_assert(__addrspaceof(const_device_var) ==
static_assert(__addrspaceof(device_array) ==
EXPECTED_DEVICE_ADDRESS_SPACE);
static_assert(__addrspaceof(*&device_array) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(*(device_array + 1)) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(host_consteval_address_space(&device_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(host_consteval_address_space(&constant_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(host_consteval_address_space(&const_device_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(host_consteval_address_space((int *)&constant_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(host_constexpr_address_space(&constant_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
extern "C" int test_host_device_var() {
return __addrspaceof(device_var);
}
// HOST-LABEL: define{{.*}} i32 @test_host_device_var(
-// HOST: ret i32 23
+// HOST: ret i32 1
extern "C" int test_host_constant_var() {
return __addrspaceof(constant_var);
}
// HOST-LABEL: define{{.*}} i32 @test_host_constant_var(
-// HOST: ret i32 24
+// HOST: ret i32 3
extern "C" int test_host_const_device_var() {
return __addrspaceof(const_device_var);
}
// HOST-LABEL: define{{.*}} i32 @test_host_const_device_var(
-// HOST: ret i32 24
+// HOST: ret i32 3
extern "C" int test_host_device_array_address() {
return __addrspaceof(*&device_array);
@@ -104,8 +104,8 @@ template <class T> constexpr int constexpr_address_space(T *p) {
template <int AS> struct AddressSpaceSpecialization;
template <>
-struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_DEFAULT> {
- static constexpr int value = __CLANG_ADDRESS_SPACE_DEFAULT;
+struct AddressSpaceSpecialization<__ADDRSPACE_DEFAULT> {
+ static constexpr int value = __ADDRSPACE_DEFAULT;
};
template <> struct AddressSpaceSpecialization<EXPECTED_DEVICE_ADDRESS_SPACE> {
static constexpr int value = EXPECTED_DEVICE_ADDRESS_SPACE;
@@ -119,25 +119,25 @@ struct AddressSpaceSpecialization<EXPECTED_CONSTANT_ADDRESS_SPACE> {
};
static_assert(__addrspaceof(*(int *)&constant_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(consteval_address_space(&device_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(consteval_address_space(&constant_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(consteval_address_space(&const_device_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(consteval_address_space((int *)&constant_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(constexpr_address_space(&constant_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(device_array) ==
EXPECTED_DEVICE_ADDRESS_SPACE);
static_assert(__addrspaceof(*&device_array) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(*(device_array + 1)) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(*device_ptr) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(
AddressSpaceSpecialization<
__addrspaceof(device_var)>::value ==
@@ -149,7 +149,7 @@ static_assert(
static_assert(
AddressSpaceSpecialization<
consteval_address_space(&constant_var)>::value ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
extern "C" __device__ int test_generic_pointer(int *p) {
return __addrspaceof(*p);
@@ -164,23 +164,23 @@ extern "C" __device__ int test_shared_local() {
static_assert(__addrspaceof(shared_var) ==
EXPECTED_SHARED_ADDRESS_SPACE);
static_assert(__addrspaceof(*(char *)&shared_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(*(&shared_var + 1)) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(shared_array) ==
EXPECTED_SHARED_ADDRESS_SPACE);
static_assert(__addrspaceof(*&shared_array) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(shared_array[0]) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(*(shared_array + 1)) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(consteval_address_space(&shared_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(consteval_address_space(shared_array) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(constexpr_address_space(&shared_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(
AddressSpaceSpecialization<
__addrspaceof(shared_var)>::value ==
@@ -189,24 +189,21 @@ extern "C" __device__ int test_shared_local() {
}
// CHECK-LABEL: define{{.*}} i32 @test_shared_local(
-// CUDA: ret i32 8
-// HIP: ret i32 25
+// CHECK: ret i32 2
extern "C" __device__ int test_device_var() {
return __addrspaceof(device_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_device_var(
-// CUDA: ret i32 6
-// HIP: ret i32 23
+// CHECK: ret i32 1
extern "C" __device__ int test_device_array() {
return __addrspaceof(device_array);
}
// CHECK-LABEL: define{{.*}} i32 @test_device_array(
-// CUDA: ret i32 6
-// HIP: ret i32 23
+// CHECK: ret i32 1
extern "C" __device__ int test_device_array_address() {
return __addrspaceof(*&device_array);
@@ -234,16 +231,14 @@ extern "C" __device__ int test_constant_var() {
}
// CHECK-LABEL: define{{.*}} i32 @test_constant_var(
-// CUDA: ret i32 7
-// HIP: ret i32 24
+// CHECK: ret i32 3
extern "C" __device__ int test_const_device_var() {
return __addrspaceof(const_device_var);
}
// CHECK-LABEL: define{{.*}} i32 @test_const_device_var(
-// CUDA: ret i32 7
-// HIP: ret i32 24
+// CHECK: ret i32 3
extern "C" __device__ int test_explicit_cast() {
return __addrspaceof(*(int *)&constant_var);
diff --git a/clang/test/CodeGenCUDA/mangle-addrspaceof.cu b/clang/test/CodeGenCUDA/mangle-addrspaceof.cu
index 7866e3dd6b052..df42711551c01 100644
--- a/clang/test/CodeGenCUDA/mangle-addrspaceof.cu
+++ b/clang/test/CodeGenCUDA/mangle-addrspaceof.cu
@@ -18,8 +18,8 @@ __device__ Result<__addrspaceof((value<T>))> query(T) {
return {};
}
-template __device__ Result<__CLANG_ADDRESS_SPACE_CUDA_DEVICE> query<int>(int);
-template __device__ Result<__CLANG_ADDRESS_SPACE_DEFAULT> query<int>(int);
+template __device__ Result<__ADDRSPACE_GLOBAL> query<int>(int);
+template __device__ Result<__ADDRSPACE_DEFAULT> query<int>(int);
// Lb1E records the entity form; Lb0E records the expression form.
// CHECK-DAG: define {{.*}}@_Z5queryIiE6ResultIXu13__addrspaceofLb1EX5valueIT_EEEEES1_(
diff --git a/clang/test/CodeGenSYCL/addrspaceof.cpp b/clang/test/CodeGenSYCL/addrspaceof.cpp
index 427992536b14c..3bde4705356f2 100644
--- a/clang/test/CodeGenSYCL/addrspaceof.cpp
+++ b/clang/test/CodeGenSYCL/addrspaceof.cpp
@@ -3,21 +3,21 @@
[[clang::sycl_external]] int
global_device_as(__attribute__((opencl_global_device)) int *p) {
static_assert(__addrspaceof(*p) ==
- __CLANG_ADDRESS_SPACE_SYCL_GLOBAL);
+ __ADDRSPACE_GLOBAL);
return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @_Z16global_device_asPU3AS5i(
// CHECK-SAME: ptr addrspace(5)
-// CHECK: ret i32 9
+// CHECK: ret i32 1
[[clang::sycl_external]] int
global_host_as(__attribute__((opencl_global_host)) int *p) {
static_assert(__addrspaceof(*p) ==
- __CLANG_ADDRESS_SPACE_SYCL_GLOBAL);
+ __ADDRSPACE_GLOBAL);
return __addrspaceof(*p);
}
// CHECK-LABEL: define{{.*}} i32 @_Z14global_host_asPU3AS6i(
// CHECK-SAME: ptr addrspace(6)
-// CHECK: ret i32 9
+// CHECK: ret i32 1
diff --git a/clang/test/PCH/exprs.c b/clang/test/PCH/exprs.c
index 95b4838376f0d..1a45779714e7b 100644
--- a/clang/test/PCH/exprs.c
+++ b/clang/test/PCH/exprs.c
@@ -47,7 +47,7 @@ offsetof_type *offsetof_ptr = &size_type_value;
typeof(sizeof(float)) size_t_value;
typeof_sizeof *size_t_ptr = &size_t_value;
typeof_sizeof2 *size_t_ptr2 = &size_t_value;
-_Static_assert(AddressSpaceOfI == __CLANG_ADDRESS_SPACE_DEFAULT, "");
+_Static_assert(AddressSpaceOfI == __ADDRSPACE_DEFAULT, "");
// ArraySubscriptExpr
array_subscript *double_ptr1_5 = &floating;
diff --git a/clang/test/Preprocessor/address-space-macros.c b/clang/test/Preprocessor/address-space-macros.c
index 5772f778b0e1a..0d89c4efbc6ca 100644
--- a/clang/test/Preprocessor/address-space-macros.c
+++ b/clang/test/Preprocessor/address-space-macros.c
@@ -1,29 +1,12 @@
// RUN: %clang_cc1 -E -dM %s | FileCheck %s
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_DEFAULT 0
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL 1
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_LOCAL 2
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 6
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 7
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_CUDA_SHARED 8
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 9
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 10
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 11
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR32_SPTR 12
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR32_UPTR 13
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_PTR64 14
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 15
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 16
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 17
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 18
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_INPUT 19
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 20
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 21
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 22
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_DEVICE 23
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 24
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_HIP_SHARED 25
-// CHECK-DAG: #define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
+// CHECK-DAG: #define __ADDRSPACE_DEFAULT 0
+// CHECK-DAG: #define __ADDRSPACE_GLOBAL 1
+// CHECK-DAG: #define __ADDRSPACE_LOCAL 2
+// CHECK-DAG: #define __ADDRSPACE_CONSTANT 3
+// CHECK-DAG: #define __ADDRSPACE_PRIVATE 4
+// CHECK-DAG: #define __ADDRSPACE_GENERIC 5
+// CHECK-DAG: #define __ADDRSPACE_HLSL_INPUT 6
+// CHECK-DAG: #define __ADDRSPACE_HLSL_OUTPUT 7
+// CHECK-DAG: #define __ADDRSPACE_HLSL_PUSH_CONSTANT 8
+// CHECK-DAG: #define __ADDRSPACE_TARGET_OFFSET 16777216
diff --git a/clang/test/Preprocessor/init-aarch64.c b/clang/test/Preprocessor/init-aarch64.c
index 97cc695cc4ef4..41b57763c215c 100644
--- a/clang/test/Preprocessor/init-aarch64.c
+++ b/clang/test/Preprocessor/init-aarch64.c
@@ -11,6 +11,16 @@
// AARCH64_BE-NEXT: #define __AARCH_BIG_ENDIAN 1
// AARCH64_LE-NEXT: #define __AARCH64EL__ 1
// AARCH64_LE-NEXT: #define __AARCH64_CMODEL_SMALL__ 1
+// AARCH64-NEXT: #define __ADDRSPACE_CONSTANT 3
+// AARCH64-NEXT: #define __ADDRSPACE_DEFAULT 0
+// AARCH64-NEXT: #define __ADDRSPACE_GENERIC 5
+// AARCH64-NEXT: #define __ADDRSPACE_GLOBAL 1
+// AARCH64-NEXT: #define __ADDRSPACE_HLSL_INPUT 6
+// AARCH64-NEXT: #define __ADDRSPACE_HLSL_OUTPUT 7
+// AARCH64-NEXT: #define __ADDRSPACE_HLSL_PUSH_CONSTANT 8
+// AARCH64-NEXT: #define __ADDRSPACE_LOCAL 2
+// AARCH64-NEXT: #define __ADDRSPACE_PRIVATE 4
+// AARCH64-NEXT: #define __ADDRSPACE_TARGET_OFFSET 16777216
// AARCH64-NEXT: #define __ARM_64BIT_STATE 1
// AARCH64-NEXT: #define __ARM_ACLE 202420
// AARCH64-NEXT: #define __ARM_ACLE_VERSION(year,quarter,patch) (100 * (year) + 10 * (quarter) + (patch))
@@ -52,33 +62,6 @@
// AARCH64-NEXT: #define __CHAR16_TYPE__ unsigned short
// AARCH64-NEXT: #define __CHAR32_TYPE__ unsigned int
// AARCH64-NEXT: #define __CHAR_BIT__ 8
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 7
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 6
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_CUDA_SHARED 8
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_DEFAULT 0
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 24
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_DEVICE 23
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HIP_SHARED 25
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 16
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 18
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 15
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_INPUT 19
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 20
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 17
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 21
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL 1
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_LOCAL 2
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR32_SPTR 12
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR32_UPTR 13
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_PTR64 14
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 9
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 10
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 11
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
-// AARCH64-NEXT: #define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 22
// AARCH64-NEXT: #define __CLANG_ATOMIC_BOOL_LOCK_FREE 2
// AARCH64-NEXT: #define __CLANG_ATOMIC_CHAR16_T_LOCK_FREE 2
// AARCH64-NEXT: #define __CLANG_ATOMIC_CHAR32_T_LOCK_FREE 2
diff --git a/clang/test/Preprocessor/init.c b/clang/test/Preprocessor/init.c
index 6f8c0409b6abb..601b701d18a97 100644
--- a/clang/test/Preprocessor/init.c
+++ b/clang/test/Preprocessor/init.c
@@ -1702,6 +1702,16 @@
// WEBASSEMBLY64-NOT:#define _ILP32
// WEBASSEMBLY64:#define _LP64 1
// EMSCRIPTEN-THREADS:#define _REENTRANT 1
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_CONSTANT 3
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_DEFAULT 0
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_GENERIC 5
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_GLOBAL 1
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_HLSL_INPUT 6
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_HLSL_OUTPUT 7
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_HLSL_PUSH_CONSTANT 8
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_LOCAL 2
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_PRIVATE 4
+// WEBASSEMBLY-NEXT:#define __ADDRSPACE_TARGET_OFFSET 16777216
// WEBASSEMBLY-NEXT:#define __ATOMIC_ACQUIRE 2
// WEBASSEMBLY-NEXT:#define __ATOMIC_ACQ_REL 4
// WEBASSEMBLY-NEXT:#define __ATOMIC_CONSUME 1
@@ -1716,33 +1726,6 @@
// WEBASSEMBLY-NEXT:#define __CHAR32_TYPE__ unsigned int
// WEBASSEMBLY-NEXT:#define __CHAR_BIT__ 8
// WEBASSEMBLY-NOT:#define __CHAR_UNSIGNED__
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_CONSTANT 7
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_DEVICE 6
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_CUDA_SHARED 8
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_DEFAULT 0
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_CONSTANT 24
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_DEVICE 23
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HIP_SHARED 25
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_CONSTANT 16
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_DEVICE 18
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_GROUPSHARED 15
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_INPUT 19
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_OUTPUT 20
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_PRIVATE 17
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_HLSL_PUSH_CONSTANT 21
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_CONSTANT 3
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GENERIC 5
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_GLOBAL 1
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_LOCAL 2
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_OPENCL_PRIVATE 4
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR32_SPTR 12
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR32_UPTR 13
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_PTR64 14
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_GLOBAL 9
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_LOCAL 10
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_SYCL_PRIVATE 11
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_TARGET_OFFSET 16777216
-// WEBASSEMBLY-NEXT:#define __CLANG_ADDRESS_SPACE_WASM_FUNCREF 22
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_BOOL_LOCK_FREE 2
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_CHAR16_T_LOCK_FREE 2
// WEBASSEMBLY-NEXT:#define __CLANG_ATOMIC_CHAR32_T_LOCK_FREE 2
diff --git a/clang/test/SemaCUDA/addrspaceof.cu b/clang/test/SemaCUDA/addrspaceof.cu
index 01b3389ab7ad4..7fb5bf83dce33 100644
--- a/clang/test/SemaCUDA/addrspaceof.cu
+++ b/clang/test/SemaCUDA/addrspaceof.cu
@@ -15,21 +15,21 @@ __constant__ int constant_var;
__device__ const int const_device_var = 1;
static_assert(__addrspaceof(device_var) ==
- __CLANG_ADDRESS_SPACE_HIP_DEVICE);
+ __ADDRSPACE_GLOBAL);
static_assert(__addrspaceof(device_array) ==
- __CLANG_ADDRESS_SPACE_HIP_DEVICE);
+ __ADDRSPACE_GLOBAL);
static_assert(__addrspaceof(*&device_array) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(*(device_array + 1)) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(constant_var) ==
- __CLANG_ADDRESS_SPACE_HIP_CONSTANT);
+ __ADDRSPACE_CONSTANT);
static_assert(__addrspaceof(const_device_var) ==
- __CLANG_ADDRESS_SPACE_HIP_CONSTANT);
+ __ADDRSPACE_CONSTANT);
static_assert(__addrspaceof(*(int *)&constant_var) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(*device_ptr) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
void host_queries() {
(void)__addrspaceof(device_var);
diff --git a/clang/test/SemaCXX/addrspaceof.cpp b/clang/test/SemaCXX/addrspaceof.cpp
index 58d195ffcb92c..d71e1e8c73367 100644
--- a/clang/test/SemaCXX/addrspaceof.cpp
+++ b/clang/test/SemaCXX/addrspaceof.cpp
@@ -8,34 +8,34 @@
using AS1 = int __attribute__((address_space(1)));
using AS2 = int __attribute__((address_space(2)));
-static_assert(__addrspaceof(int) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(int) == __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(AS1) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+ __ADDRSPACE_TARGET_OFFSET + 1);
static_assert(__addrspaceof(AS2 &) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 2);
+ __ADDRSPACE_TARGET_OFFSET + 2);
int *p0;
AS1 *p1;
-static_assert(__addrspaceof(p0) == __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__addrspaceof(p1) == __CLANG_ADDRESS_SPACE_DEFAULT);
-static_assert(__addrspaceof(*p0) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(p0) == __ADDRSPACE_DEFAULT);
+static_assert(__addrspaceof(p1) == __ADDRSPACE_DEFAULT);
+static_assert(__addrspaceof(*p0) == __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(*p1) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+ __ADDRSPACE_TARGET_OFFSET + 1);
int global;
AS1 global_as1;
int array[4];
AS2 array_as2[4];
-static_assert(__addrspaceof(global) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__addrspaceof(global) == __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(global_as1) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
-static_assert(__addrspaceof(array) == __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_TARGET_OFFSET + 1);
+static_assert(__addrspaceof(array) == __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(array_as2) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 2);
+ __ADDRSPACE_TARGET_OFFSET + 2);
static_assert(__addrspaceof(array_as2[0]) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 2);
+ __ADDRSPACE_TARGET_OFFSET + 2);
struct S {
int member;
@@ -45,9 +45,9 @@ AS1 S::member_as1;
S object;
static_assert(__addrspaceof(object.member) ==
- __CLANG_ADDRESS_SPACE_DEFAULT);
+ __ADDRSPACE_DEFAULT);
static_assert(__addrspaceof(object.member_as1) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+ __ADDRSPACE_TARGET_OFFSET + 1);
template <class T> constexpr int type_address_space() {
return __addrspaceof(T);
@@ -58,18 +58,18 @@ template <class T> constexpr int expression_address_space(T &value) {
}
static_assert(type_address_space<AS1>() ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+ __ADDRSPACE_TARGET_OFFSET + 1);
static_assert(expression_address_space(global_as1) ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+ __ADDRSPACE_TARGET_OFFSET + 1);
template <int AS> struct AddressSpaceSpecialization;
template <>
-struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1> {
- static constexpr int value = __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1;
+struct AddressSpaceSpecialization<__ADDRSPACE_TARGET_OFFSET + 1> {
+ static constexpr int value = __ADDRSPACE_TARGET_OFFSET + 1;
};
static_assert(AddressSpaceSpecialization<__addrspaceof(AS1)>::value ==
- __CLANG_ADDRESS_SPACE_TARGET_OFFSET + 1);
+ __ADDRSPACE_TARGET_OFFSET + 1);
void function();
More information about the cfe-commits
mailing list