https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/210242
>From 9465bfa05d08983c51cadab1046b663d79ba348d 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. 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 f5313e0378ca0..4d527db8851d0 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 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..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 6796cb3bdd912..b8559a6d37be4 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 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@-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}} + (void)__builtin_pointee_address_space(f); + // 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
