[clang] [RFC][Clang] Add __builtin_pointee_address_space (PR #210242)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 22 21:56:36 PDT 2026
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/210242
>From b4c750329ab4844c9b4453a15e578e86f4c17dfb 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] [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 | 131 +++++++++
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 | 254 ++++++++++++++++++
.../builtin-pointee-address-space.cl | 18 ++
.../test/Preprocessor/address-space-macros.c | 33 +++
clang/test/Preprocessor/init.c | 31 +++
.../SemaCUDA/builtin-pointee-address-space.cu | 43 +++
.../SemaCXX/builtin-pointee-address-space.cpp | 91 +++++++
16 files changed, 911 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 f5313e0378ca0..9170d70d2e78d 100644
--- a/clang/docs/LanguageExtensions.md
+++ b/clang/docs/LanguageExtensions.md
@@ -3983,6 +3983,137 @@ 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;
+
+#if defined(__HIP__)
+#define DEVICE_AS __CLANG_ADDRESS_SPACE_HIP_DEVICE
+#define CONSTANT_AS __CLANG_ADDRESS_SPACE_HIP_CONSTANT
+#else
+#define DEVICE_AS __CLANG_ADDRESS_SPACE_CUDA_DEVICE
+#define CONSTANT_AS __CLANG_ADDRESS_SPACE_CUDA_CONSTANT
+#endif
+
+template <class T> consteval int pointee_as(T *p) {
+ return __builtin_pointee_address_space(p);
+}
+
+static_assert(__builtin_pointee_address_space(&dev) ==
+ DEVICE_AS);
+static_assert(__builtin_pointee_address_space(&cst) ==
+ CONSTANT_AS);
+static_assert(__builtin_pointee_address_space(&dev_cst) ==
+ CONSTANT_AS);
+static_assert(__builtin_pointee_address_space(dev_arr) ==
+ DEVICE_AS);
+static_assert(__builtin_pointee_address_space(&dev_arr) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space(dev_arr + 1) ==
+ __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(pointee_as(&cst) == __CLANG_ADDRESS_SPACE_DEFAULT);
+static_assert(__builtin_pointee_address_space((int *)&cst) ==
+ __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 052838ef0af0d..efb8d7d8f6e4d 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -111,6 +111,14 @@ latest release, please see the [Clang Web Site](https://clang.llvm.org) or the
### Non-comprehensive list of changes in this release
+- 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
### Deprecated Compiler Flags
diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h
index a941805423bca..ebefc83f0bde6 100644
--- a/clang/include/clang/Basic/AddressSpaces.h
+++ b/clang/include/clang/Basic/AddressSpaces.h
@@ -99,6 +99,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 344a712ddc585..396111d6991cb 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 f14288dd2967d..db5f35a3e17d4 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 574dd8b04e779..5106ca7cd1c61 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/clang/lib/AST/ExprConstant.cpp
@@ -15926,6 +15926,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;
@@ -16994,6 +17042,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 37846bbb0b5ec..8851636f63c52 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -4255,6 +4255,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 8b6ff844d0daa..de6dc2d652fb2 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"
@@ -915,6 +916,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 2a8395e26d0bd..00eedfe406a7a 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -3202,6 +3202,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..b1705a2e52fc1
--- /dev/null
+++ b/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu
@@ -0,0 +1,254 @@
+// 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(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 + 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.c b/clang/test/Preprocessor/init.c
index 1d49af5c8d62f..c3f5a5b919462 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}}
+}
More information about the cfe-commits
mailing list