https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/210242
>From a0943097099a33025c502384a1c9d433bb7ec266 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> 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. Clang predefines __CLANG_ADDRESS_SPACE_* macros for language-level address spaces and emits them during preprocessing. 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, it can report __device__, __shared__, and __constant__ declaration address spaces when the argument directly names or directly takes the address of a known variable declaration. Explicit casts are respected. Function parameters are treated by their type; the builtin does not recover caller declaration metadata through constexpr or consteval wrappers. 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 | 109 ++++++++++ clang/docs/ReleaseNotes.md | 7 + clang/include/clang/Basic/AddressSpaces.h | 103 +++++++++ clang/include/clang/Basic/Builtins.td | 7 + .../clang/Basic/DiagnosticSemaKinds.td | 3 + clang/lib/AST/ExprConstant.cpp | 54 +++++ clang/lib/CodeGen/CGBuiltin.cpp | 4 + clang/lib/Frontend/InitPreprocessor.cpp | 62 ++++++ clang/lib/Sema/SemaChecking.cpp | 14 ++ .../CodeGen/builtin-pointee-address-space.c | 42 ++++ .../builtin-pointee-address-space.cu | 202 ++++++++++++++++++ .../builtin-pointee-address-space.cl | 18 ++ .../test/Preprocessor/address-space-macros.c | 9 + .../SemaCUDA/builtin-pointee-address-space.cu | 37 ++++ .../SemaCXX/builtin-pointee-address-space.cpp | 88 ++++++++ 15 files changed, 759 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..5cedf1299dacb 100644 --- a/clang/docs/LanguageExtensions.md +++ b/clang/docs/LanguageExtensions.md @@ -3983,6 +3983,115 @@ 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. For an array expression, Clang +uses the array element type. 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`. + +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. + +CUDA/HIP variables are not represented as address-space-qualified pointer types +in Clang's AST. If the pointer or array expression directly names, or directly +takes the address of, a known CUDA/HIP variable declaration, the builtin can +still report the declaration address space from the variable 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 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. + +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; +__constant__ int cst; +__device__ const int dev_cst = 1; + +template <class T> consteval int pointee_as(T *p) { + return __builtin_pointee_address_space(p); +} + +static_assert(__builtin_pointee_address_space(&dev) == + __CLANG_ADDRESS_SPACE_CUDA_DEVICE); +static_assert(__builtin_pointee_address_space(&cst) == + __CLANG_ADDRESS_SPACE_CUDA_CONSTANT); +static_assert(__builtin_pointee_address_space(&dev_cst) == + __CLANG_ADDRESS_SPACE_CUDA_CONSTANT); +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..9874494930f45 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -111,6 +111,13 @@ 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. + ### New Compiler Flags ### Deprecated Compiler Flags diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h index a941805423bca..0b31b26280474 100644 --- a/clang/include/clang/Basic/AddressSpaces.h +++ b/clang/include/clang/Basic/AddressSpaces.h @@ -99,6 +99,109 @@ 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, + + TargetOffset = 0x1000000 +}; + +inline unsigned encode(LangAS AS) { + 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 CUDADevice; + case LangAS::cuda_constant: + return CUDAConstant; + case LangAS::cuda_shared: + return 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..a40736fefb027 100644 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -15926,6 +15926,51 @@ 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())) + VD = dyn_cast<VarDecl>(DRE->getDecl()); + } 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 +17039,15 @@ 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), 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..cf94a927286b7 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,67 @@ 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_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..dc51afb8567e4 --- /dev/null +++ b/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu @@ -0,0 +1,202 @@ +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++20 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip -std=c++20 -emit-llvm -o - %s | FileCheck %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; + +#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) == + __CLANG_ADDRESS_SPACE_CUDA_DEVICE); +static_assert(__builtin_pointee_address_space(&constant_var) == + __CLANG_ADDRESS_SPACE_CUDA_CONSTANT); +static_assert(__builtin_pointee_address_space(&const_device_var) == + __CLANG_ADDRESS_SPACE_CUDA_CONSTANT); +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 8 + +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 9 + +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 9 + +#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<__CLANG_ADDRESS_SPACE_CUDA_DEVICE> { + static constexpr int value = __CLANG_ADDRESS_SPACE_CUDA_DEVICE; +}; +template <> struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_CUDA_SHARED> { + static constexpr int value = __CLANG_ADDRESS_SPACE_CUDA_SHARED; +}; +template <> +struct AddressSpaceSpecialization<__CLANG_ADDRESS_SPACE_CUDA_CONSTANT> { + static constexpr int value = __CLANG_ADDRESS_SPACE_CUDA_CONSTANT; +}; + +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) == + __CLANG_ADDRESS_SPACE_CUDA_DEVICE); +static_assert(__builtin_pointee_address_space(device_ptr) == + __CLANG_ADDRESS_SPACE_DEFAULT); +static_assert( + AddressSpaceSpecialization< + __builtin_pointee_address_space(&device_var)>::value == + __CLANG_ADDRESS_SPACE_CUDA_DEVICE); +static_assert( + AddressSpaceSpecialization< + __builtin_pointee_address_space(&constant_var)>::value == + __CLANG_ADDRESS_SPACE_CUDA_CONSTANT); +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) == + __CLANG_ADDRESS_SPACE_CUDA_SHARED); + 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 == + __CLANG_ADDRESS_SPACE_CUDA_SHARED); + return __builtin_pointee_address_space(&shared_var); +} + +// CHECK-LABEL: define{{.*}} i32 @test_shared_local( +// CHECK: ret i32 10 + +extern "C" __device__ int test_device_var() { + return __builtin_pointee_address_space(&device_var); +} + +// CHECK-LABEL: define{{.*}} i32 @test_device_var( +// CHECK: ret i32 8 + +extern "C" __device__ int test_device_array() { + return __builtin_pointee_address_space(device_array); +} + +// CHECK-LABEL: define{{.*}} i32 @test_device_array( +// CHECK: ret i32 8 + +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( +// CHECK: ret i32 9 + +extern "C" __device__ int test_const_device_var() { + return __builtin_pointee_address_space(&const_device_var); +} + +// CHECK-LABEL: define{{.*}} i32 @test_const_device_var( +// CHECK: ret i32 9 + +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..7728e305196b3 --- /dev/null +++ b/clang/test/Preprocessor/address-space-macros.c @@ -0,0 +1,9 @@ +// 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_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_TARGET_OFFSET 16777216 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..4f58b4d189823 --- /dev/null +++ b/clang/test/SemaCUDA/builtin-pointee-address-space.cu @@ -0,0 +1,37 @@ +// 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_CUDA_DEVICE); +static_assert(__builtin_pointee_address_space(device_array) == + __CLANG_ADDRESS_SPACE_CUDA_DEVICE); +static_assert(__builtin_pointee_address_space(&constant_var) == + __CLANG_ADDRESS_SPACE_CUDA_CONSTANT); +static_assert(__builtin_pointee_address_space(&const_device_var) == + __CLANG_ADDRESS_SPACE_CUDA_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(&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..da387cafcee50 --- /dev/null +++ b/clang/test/SemaCXX/builtin-pointee-address-space.cpp @@ -0,0 +1,88 @@ +// 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)__builtin_pointee_address_space(); + // expected-error@-1 {{too few arguments}} + (void)__builtin_pointee_address_space(p0, p1); + // expected-error@-1 {{too many arguments}} + (void)__builtin_pointee_address_space(i); + // expected-error@-1 {{argument to '__builtin_pointee_address_space' must be a pointer or array expression}} + (void)__builtin_pointee_address_space(0); + // expected-error@-1 {{argument to '__builtin_pointee_address_space' must be a pointer or array expression}} +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
