llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Yaxun (Sam) Liu (yxsamliu) <details> <summary>Changes</summary> Some code needs to choose a code path from the address space of a pointee. This is easier in OpenCL because address spaces are part of the language type system. It is harder in CUDA/HIP: source-level pointers are generic, and the compiler may cast pointers to the generic address space before use. CUDA predicates such as `__isShared(p)` test a runtime pointer value, are not constant expressions, and do not expose target-specific address space numbers. Add `__builtin_pointee_address_space`. It accepts a pointer or array expression and returns the target address space number for the pointee type, array element type, or known CUDA/HIP device object. The argument is not evaluated and is not converted to `void *` before its type is queried. User-written casts are respected. During constant evaluation, simple constexpr forwarding can preserve the address space of a known object. In CUDA/HIP host compilation, known device objects are mapped with auxiliary target information. If no auxiliary target is available, Clang warns and returns the default address space. Document constexpr behavior, invalid arguments, unevaluated operands, runtime pointer limitations, and target address-space values. --- Patch is 24.74 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/210242.diff 12 Files Affected: - (modified) clang/docs/LanguageExtensions.md (+76) - (modified) clang/docs/ReleaseNotes.md (+4) - (modified) clang/include/clang/Basic/Builtins.td (+7) - (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+7) - (modified) clang/lib/AST/ExprConstant.cpp (+83) - (modified) clang/lib/CodeGen/CGBuiltin.cpp (+4) - (modified) clang/lib/Sema/SemaChecking.cpp (+37) - (added) clang/test/CodeGen/builtin-pointee-address-space.c (+42) - (added) clang/test/CodeGenCUDA/builtin-pointee-address-space.cu (+165) - (added) clang/test/CodeGenOpenCL/builtin-pointee-address-space.cl (+18) - (added) clang/test/SemaCUDA/builtin-pointee-address-space.cu (+28) - (added) clang/test/SemaCXX/builtin-pointee-address-space.cpp (+72) ``````````diff diff --git a/clang/docs/LanguageExtensions.md b/clang/docs/LanguageExtensions.md index f5313e0378ca0..d7de5d1c7966a 100644 --- a/clang/docs/LanguageExtensions.md +++ b/clang/docs/LanguageExtensions.md @@ -3983,6 +3983,82 @@ template<typename T> constexpr T *addressof(T &value) { } ``` +### `__builtin_pointee_address_space` + +`__builtin_pointee_address_space` returns the target address space number +of the pointee type of a pointer expression, or the element type of an +array expression. + +**Syntax**: + +```c++ +int __builtin_pointee_address_space(pointer-or-array) +``` + +The argument must have pointer or array type. For an array expression, the +result is computed from the array element type. The argument is not evaluated +and is not converted to `void *` before its type is inspected. + +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. + +For a pointer expression, the builtin normally returns the address space of the +pointee type of the expression as written. This makes the result stable for +ordinary pointers and for explicitly address-spaced pointer types. In CUDA/HIP +compilation, Clang can also determine the address space of known `__device__`, +`__shared__`, and `__constant__` objects. In host compilation, this requires +auxiliary target information for the device target. If auxiliary target +information is not available, Clang returns the default address space and emits +a warning. This is useful because CUDA/HIP pointer types are source-level +generic, so this information is not otherwise visible from a plain `T *` type. +An explicit `__device__` `const` global or static data member that is promoted +to constant memory is reported as the constant address space. + +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. + +When the builtin is evaluated in a constant expression, Clang may look through a +constexpr function parameter to the argument passed by the caller, as long as +doing so does not look through an explicit user-written cast. This allows simple +constexpr wrappers to preserve the address space of a known 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 the target address space for the generic/default pointee type 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. + +**Example use**: + +```c++ +int *p; +int __attribute__((address_space(3))) *p3; + +static_assert(__builtin_pointee_address_space(p) == 0); +static_assert(__builtin_pointee_address_space(p3) == 3); +``` + +**CUDA/HIP example use**: + +```c++ +__device__ int dev; +__constant__ int cst; +__device__ const int dev_cst = 1; + +constexpr int pointee_as(int *p) { + return __builtin_pointee_address_space(p); +} + +static_assert(__builtin_pointee_address_space(&dev) == 1); +static_assert(__builtin_pointee_address_space(&cst) == 4); +static_assert(__builtin_pointee_address_space(&dev_cst) == 4); +static_assert(pointee_as(&cst) == 4); +static_assert(pointee_as((int *)&cst) == 0); +``` + ### `__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..4d14889f042a2 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -111,6 +111,10 @@ 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 the target address + space number for the pointee type of a pointer expression, the element type + of an array expression, or a known CUDA/HIP device object. + ### New Compiler Flags ### Deprecated Compiler Flags 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..46a5cff919058 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -864,6 +864,13 @@ 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 warn_builtin_pointee_address_space_no_aux_target : Warning< + "cannot determine CUDA/HIP device address space without auxiliary target " + "information; returning default address space">, + InGroup<DiagGroup<"builtin-pointee-address-space">>; 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..f0a228aadb4e3 100644 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -15926,6 +15926,79 @@ bool ArrayExprEvaluator::VisitDesignatedInitUpdateExpr( //===----------------------------------------------------------------------===// namespace { + +static QualType getPointeeAddressSpaceType(ASTContext &Ctx, QualType ArgTy) { + return ArgTy->isArrayType() ? Ctx.getAsArrayType(ArgTy)->getElementType() + : ArgTy->getPointeeType(); +} + +static bool isCUDADeviceAddressSpace(LangAS AS) { + return AS == LangAS::cuda_device || AS == LangAS::cuda_constant || + AS == LangAS::cuda_shared; +} + +static unsigned getTargetAddressSpaceForPointeeAddressSpace(ASTContext &Ctx, + LangAS AS) { + if (Ctx.getLangOpts().CUDA && !Ctx.getLangOpts().CUDAIsDevice && + isCUDADeviceAddressSpace(AS)) { + if (const TargetInfo *AuxTarget = Ctx.getAuxTargetInfo()) + return AuxTarget->getTargetAddressSpace(AS); + return Ctx.getTargetAddressSpace(LangAS::Default); + } + return Ctx.getTargetAddressSpace(AS); +} + +static std::optional<LangAS> +getCUDAPointeeDeclAddressSpace(ASTContext &Ctx, const Expr *Arg, + const CallStackFrame *Frame) { + 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 (const auto *PVD = dyn_cast<ParmVarDecl>(DRE->getDecl())) { + if (Frame && Frame->CallExpr && Frame->Caller) { + if (const auto *CE = dyn_cast<CallExpr>(Frame->CallExpr)) { + unsigned Index = PVD->getFunctionScopeIndex(); + if (Index < CE->getNumArgs()) + return getCUDAPointeeDeclAddressSpace(Ctx, CE->getArg(Index), + Frame->Caller); + } + } + } else { + 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 +17067,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), Info.CurrentCall) + .value_or( + getPointeeAddressSpaceType(Info.Ctx, ArgTy).getAddressSpace()); + return Success(getTargetAddressSpaceForPointeeAddressSpace(Info.Ctx, 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/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp index 2a8395e26d0bd..1eada6250bf99 100644 --- a/clang/lib/Sema/SemaChecking.cpp +++ b/clang/lib/Sema/SemaChecking.cpp @@ -3202,6 +3202,43 @@ 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(); + } + if (Context.getLangOpts().CUDA && !Context.getLangOpts().CUDAIsDevice && + !Context.getAuxTargetInfo()) { + const Expr *PointeeArg = Arg->IgnoreParens(); + if (!isa<ExplicitCastExpr>(PointeeArg)) { + const VarDecl *VD = nullptr; + const Expr *NoImpCasts = PointeeArg->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 && + (VD->hasAttr<CUDAConstantAttr>() || VD->hasAttr<CUDASharedAttr>() || + VD->hasAttr<CUDADeviceAttr>())) + Diag(Arg->getBeginLoc(), + diag::warn_builtin_pointee_address_space_no_aux_target) + << Arg->getSourceRange(); + } + } + 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..f76b9131c0299 --- /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 4 + +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 7 + +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 5 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..0f1021fef8152 --- /dev/null +++ b/clang/test/CodeGenCUDA/builtin-pointee-address-space.cu @@ -0,0 +1,165 @@ +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -x hip -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> constexpr int get_host_as(T *p) { + return __builtin_pointee_address_space(p); +} + +static_assert(__builtin_pointee_address_space(&device_var) == 1); +static_assert(__builtin_pointee_address_space(&constant_var) == 4); +static_assert(__builtin_pointee_address_space(&const_device_var) == 4); +static_assert(get_host_as(&device_var) == 1); +static_assert(get_host_as(&constant_var) == 4); +static_assert(get_host_as(&const_device_var) == 4); + +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 1 + +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 4 + +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 4 + +#else + +constexpr int get_as(int *p) { + return __builtin_pointee_address_space(p); +} + +template <class T> constexpr int get_template_as(T *p) { + return __builtin_pointee_address_space(p); +} + +template <int AS> struct AddressSpaceSpecialization; +template <> struct AddressSpaceSpecialization<0> { + static constexpr int value = 0; +}; +template <> struct AddressSpaceSpecialization<1> { + static constexpr int value = 1; +}; +template <> struct AddressSpaceSpecialization<3> { + static constexpr int value = 3; +}; +template <> struct AddressSpaceSpecialization<4> { + static constexpr int value = 4; +}; + +static_assert(get_as(&device_var) == 1); +static_assert(get_as(&constant_var) == 4); +static_assert(get_as((int *)&constant_var) == 0); +static_assert(get_template_as(&device_var) == 1); +static_assert(get_template_as(&constant_var) == 4); +static_assert(get_template_as(&const_device_var) == 4); +static_assert(__builtin_pointee_address_space(device_array) == 1); +static_assert(__builtin_pointee_address_space(device_ptr) == 0); +static_assert( + AddressSpaceSpecialization< + __builtin_pointee_address_space(&device_var)>::value == 1); +static_assert( + AddressSpaceSpecialization< + __builtin_pointee_address_space(&constant_var)>::value == 4); +static_assert( + AddressSpaceSpecialization<get_template_as(&constant_var)>::value == 4); + +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; + static_assert(__builtin_pointee_address_space(&shared_var) == 3); + static_assert( + AddressSpaceSpecialization< + __builtin_pointee_address_space(&shared_var)>::value == 3); + return __builtin_pointee_address_space(&shared_var); +} + +// CHECK-LABEL: define{{.*}} i32 @test_shared_local( +// CHECK: ret i32 3 + +extern "C" __device__ int test_device_var() { + return __builtin_pointee_address_space(&device_var); +} + +// CHECK-LABEL: define{{.*}} i32 @test_device_var( +// CHECK: ret i32 1 + +extern "C" __device__ int test_device_array() { + return __builtin_pointee_address_space(device_array); +} + +// CHECK-LABEL: define{{.*}} i32 @test_device_array( +// CHECK: ret i32 1 + +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 4 + +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 4 + +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 3 + +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 4 + +#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..54c9e9728cd84 --- /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{{.*... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/210242 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
