https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/209679
>From 38c02e9357e4e209e4eea09afecacd75b0dd4de8 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> Date: Tue, 14 Jul 2026 21:59:44 -0400 Subject: [PATCH] [AMDGPU] Add builtin to test for constant memory libhipc++ needs this to implement a CUDA-compatible API that checks whether a device pointer points into constant memory. Add `__builtin_amdgcn_is_constant` and lower it to a range check against linker-provided start and stop symbols for the AMDGPU constant section. Place target constant-address-space globals in that section so the check covers both explicit constant variables and variables that Clang places in constant address space. Diagnose explicit section attributes that would break this invariant. --- clang/include/clang/AST/ASTContext.h | 10 ++ clang/include/clang/Basic/BuiltinsAMDGPU.td | 1 + .../clang/Basic/DiagnosticSemaKinds.td | 4 + clang/include/clang/Basic/TargetInfo.h | 6 ++ clang/include/clang/Sema/SemaAMDGPU.h | 2 + clang/lib/AST/ASTContext.cpp | 16 +++ clang/lib/Basic/Targets/AMDGPU.h | 4 + clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 94 ++++++++++++++++ clang/lib/CodeGen/Targets/AMDGPU.cpp | 17 +++ clang/lib/Sema/SemaAMDGPU.cpp | 57 ++++++++++ clang/lib/Sema/SemaDecl.cpp | 2 + .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 1 + clang/test/CodeGen/amdgpu-address-spaces.cpp | 4 +- .../template-class-static-member.cu | 4 +- clang/test/CodeGenHIP/amdgpu-is-constant.hip | 102 ++++++++++++++++++ .../test/CodeGenOpenCL/amdgpu-is-constant.cl | 81 ++++++++++++++ 16 files changed, 401 insertions(+), 4 deletions(-) create mode 100644 clang/test/CodeGenHIP/amdgpu-is-constant.hip create mode 100644 clang/test/CodeGenOpenCL/amdgpu-is-constant.cl diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index 7ed6509c3c16c..2cad8be3d3cf3 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -139,6 +139,7 @@ class TemplateTypeParmDecl; class TypeConstraint; class UnresolvedSetIterator; class UsingShadowDecl; +class VarDecl; class VarTemplateDecl; class VTableContextBase; class XRayFunctionFilter; @@ -3289,6 +3290,15 @@ class ASTContext : public RefCountedBase<ASTContext> { LangAS getLangASForBuiltinAddressSpace(unsigned AS) const; + /// Return whether the variable is in a distinct target constant address + /// space. + /// + /// This checks the address-space placement chosen by the frontend rather + /// than source spelling. LangAS::Default is not considered a target constant + /// address space, even if TargetInfo::getConstantAddressSpace() returns it + /// as a fallback placement address space. + bool isTargetConstantAddressSpaceObject(const VarDecl *D) const; + /// Get target-dependent integer value for null pointer which is used for /// constant folding. uint64_t getTargetNullPointerValue(QualType QT) const; diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 8c64116066971..aa4f60baa7f44 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -356,6 +356,7 @@ def __builtin_amdgcn_fcmpf : AMDGPUBuiltin<"uint64_t(float, float, _Constant int //===----------------------------------------------------------------------===// def __builtin_amdgcn_is_shared : AMDGPUBuiltin<"bool(void const address_space<0> *)", [Const]>; def __builtin_amdgcn_is_private : AMDGPUBuiltin<"bool(void const address_space<0> *)", [Const]>; +def __builtin_amdgcn_is_constant : AMDGPUBuiltin<"bool(void const address_space<0> *)", [Const]>; //===----------------------------------------------------------------------===// // GWS builtins. diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3ff7e30d9f1f7..aafad2df406d8 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3758,6 +3758,10 @@ def err_attribute_section_invalid_for_target : Error< "argument to %select{'code_seg'|'section'}1 attribute is not valid for this target: %0">; def err_pragma_section_invalid_for_target : Error< "argument to #pragma section is not valid for this target: %0">; +def err_amdgpu_constant_address_space_section : Error< + "AMDGPU constant address space variables must use section '%0'">; +def err_amdgpu_reserved_constant_section : Error< + "section '%0' is reserved for AMDGPU constant address space variables">; def warn_attribute_section_drectve : Warning< "#pragma %0(\".drectve\") has undefined behavior, " "use #pragma comment(linker, ...) instead">, InGroup<MicrosoftDrectveSection>; diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h index 3aba4d261a651..2c7269ca133f1 100644 --- a/clang/include/clang/Basic/TargetInfo.h +++ b/clang/include/clang/Basic/TargetInfo.h @@ -1729,6 +1729,12 @@ class TargetInfo : public TransferrableTargetInfo, return LangAS::Default; } + /// Return the section name used for objects in the target's constant address + /// space, if the target uses a dedicated section. + virtual std::optional<StringRef> getConstantAddressSpaceSectionName() const { + return std::nullopt; + } + // access target-specific GPU grid values that must be consistent between // host RTL (plugin), deviceRTL and clang. virtual const llvm::omp::GV &getGridValue() const { diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h index a6205534e0de3..0a46e608e8eff 100644 --- a/clang/include/clang/Sema/SemaAMDGPU.h +++ b/clang/include/clang/Sema/SemaAMDGPU.h @@ -25,6 +25,7 @@ class ParsedAttr; class SemaAMDGPU : public SemaBase { llvm::SmallPtrSet<Expr *, 32> ExpandedPredicates; llvm::SmallPtrSet<FunctionDecl *, 32> PotentiallyUnguardedBuiltinUsers; + llvm::SmallPtrSet<VarDecl *, 32> InvalidConstantSectionVars; public: SemaAMDGPU(Sema &S); @@ -43,6 +44,7 @@ class SemaAMDGPU : public SemaBase { bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs, unsigned NumDataArgs); + void checkConstantAddressSpaceSection(VarDecl *VD); /// Create an AMDGPUWavesPerEUAttr attribute. AMDGPUFlatWorkGroupSizeAttr * diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 2228811546c0f..b142ed222a280 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -13990,6 +13990,22 @@ unsigned ASTContext::getTargetAddressSpace(LangAS AS) const { return getTargetInfo().getTargetAddressSpace(AS); } +bool ASTContext::isTargetConstantAddressSpaceObject(const VarDecl *D) const { + std::optional<LangAS> ConstantAS = getTargetInfo().getConstantAddressSpace(); + // TargetInfo::getConstantAddressSpace() may return LangAS::Default as a + // placement fallback. That is not a distinct target constant address space. + if (!ConstantAS || *ConstantAS == LangAS::Default) + return false; + + unsigned TargetConstantAS = getTargetAddressSpace(*ConstantAS); + + if (LangOpts.CUDA && LangOpts.CUDAIsDevice && D->hasAttr<CUDAConstantAttr>()) + return getTargetAddressSpace(LangAS::cuda_constant) == TargetConstantAS; + + LangAS AS = D->getType().getAddressSpace(); + return AS != LangAS::Default && getTargetAddressSpace(AS) == TargetConstantAS; +} + bool ASTContext::hasSameExpr(const Expr *X, const Expr *Y) const { if (X == Y) return true; diff --git a/clang/lib/Basic/Targets/AMDGPU.h b/clang/lib/Basic/Targets/AMDGPU.h index 89ba561ef302d..7893bf1f37211 100644 --- a/clang/lib/Basic/Targets/AMDGPU.h +++ b/clang/lib/Basic/Targets/AMDGPU.h @@ -398,6 +398,10 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo { return getLangASFromTargetAS(llvm::AMDGPUAS::CONSTANT_ADDRESS); } + std::optional<StringRef> getConstantAddressSpaceSectionName() const override { + return "__amdgpu_constant"; + } + const llvm::omp::GV &getGridValue() const override { switch (WavefrontSize) { case 32: diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 66e4e688e33ee..d3e1f368cfdc7 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -22,8 +22,11 @@ #include "llvm/IR/IntrinsicsR600.h" #include "llvm/IR/IntrinsicsSPIRV.h" #include "llvm/IR/MemoryModelRelaxationAnnotations.h" +#include "llvm/IR/Module.h" #include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/AtomicOrdering.h" +#include <optional> +#include <string> using namespace clang; using namespace CodeGen; @@ -47,6 +50,95 @@ static Value *emitAMDGPUSBufferLoadBuiltin(CodeGenFunction &CGF, CGF.EmitScalarExpr(E->getArg(2))}); } +constexpr char AMDGPUConstantSectionMarker[] = + "__clang_amdgpu_constant_section_marker"; + +static llvm::GlobalVariable * +getAMDGPUConstantSectionBoundary(CodeGenModule &CGM, StringRef Name, + SourceLocation Loc) { + llvm::Module &M = CGM.getModule(); + auto *Int8Ty = llvm::Type::getInt8Ty(CGM.getLLVMContext()); + + if (auto *V = M.getNamedValue(Name)) { + auto *GV = dyn_cast<llvm::GlobalVariable>(V); + if (GV && GV->isDeclaration() && GV->hasExternalLinkage() && + GV->getAddressSpace() == llvm::AMDGPUAS::CONSTANT_ADDRESS && + GV->getValueType() == Int8Ty) + return GV; + + std::string Message = "AMDGPU constant section boundary symbol '"; + Message += Name; + Message += "' conflicts with existing symbol"; + CGM.Error(Loc, Message); + return nullptr; + } + + return new llvm::GlobalVariable( + M, Int8Ty, /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, + /*Initializer=*/nullptr, Name, /*InsertBefore=*/nullptr, + llvm::GlobalVariable::NotThreadLocal, llvm::AMDGPUAS::CONSTANT_ADDRESS); +} + +static void ensureAMDGPUConstantSection(CodeGenModule &CGM, + StringRef ConstantSection) { + if (auto *GV = CGM.getModule().getNamedGlobal(AMDGPUConstantSectionMarker)) + if (GV->getSection() == ConstantSection) + return; + + auto *Int8Ty = llvm::Type::getInt8Ty(CGM.getLLVMContext()); + auto *MarkerTy = llvm::ArrayType::get(Int8Ty, 0); + auto *Marker = new llvm::GlobalVariable( + CGM.getModule(), MarkerTy, /*isConstant=*/true, + llvm::GlobalValue::InternalLinkage, + llvm::ConstantAggregateZero::get(MarkerTy), AMDGPUConstantSectionMarker, + /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal, + llvm::AMDGPUAS::CONSTANT_ADDRESS); + Marker->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global); + Marker->setSection(ConstantSection); + CGM.addCompilerUsedGlobal(Marker); +} + +static std::string getAMDGPUConstantSectionStartSymbol(StringRef Section) { + // AMDGPU uses an ELF section name chosen to support GNU-style linker-provided + // __start_<section> and __stop_<section> range boundary symbols. + return ("__start_" + Section).str(); +} + +static std::string getAMDGPUConstantSectionStopSymbol(StringRef Section) { + return ("__stop_" + Section).str(); +} + +static Value *emitAMDGPUIsConstant(CodeGenFunction &CGF, const CallExpr *E) { + Value *Ptr = CGF.EmitScalarExpr(E->getArg(0)); + + if (!CGF.CGM.getTriple().isAMDGCN()) + return CGF.Builder.getFalse(); + std::optional<StringRef> ConstantSection = + CGF.CGM.getTarget().getConstantAddressSpaceSectionName(); + if (!ConstantSection) + return CGF.Builder.getFalse(); + + ensureAMDGPUConstantSection(CGF.CGM, *ConstantSection); + + Value *PtrInt = CGF.Builder.CreatePtrToInt(Ptr, CGF.IntPtrTy); + + std::string StartName = getAMDGPUConstantSectionStartSymbol(*ConstantSection); + std::string StopName = getAMDGPUConstantSectionStopSymbol(*ConstantSection); + auto *Start = + getAMDGPUConstantSectionBoundary(CGF.CGM, StartName, E->getExprLoc()); + auto *Stop = + getAMDGPUConstantSectionBoundary(CGF.CGM, StopName, E->getExprLoc()); + if (!Start || !Stop) + return CGF.Builder.getFalse(); + + Value *StartInt = CGF.Builder.CreatePtrToInt(Start, CGF.IntPtrTy); + Value *StopInt = CGF.Builder.CreatePtrToInt(Stop, CGF.IntPtrTy); + + Value *Offset = CGF.Builder.CreateSub(PtrInt, StartInt); + Value *Size = CGF.Builder.CreateSub(StopInt, StartInt); + return CGF.Builder.CreateICmpULT(Offset, Size); +} + // Has second type mangled argument. static Value * emitBinaryExpMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, const CallExpr *E, @@ -537,6 +629,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent; llvm::SyncScope::ID SSID; switch (BuiltinID) { + case AMDGPU::BI__builtin_amdgcn_is_constant: + return emitAMDGPUIsConstant(*this, E); case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32: case AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f32: case AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f64: diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 7b37f3f7f9b6e..18fcaa8812a9c 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -11,6 +11,7 @@ #include "clang/AST/DeclCXX.h" #include "llvm/ADT/StringExtras.h" #include "llvm/Support/AMDGPUAddrSpace.h" +#include <optional> using namespace clang; using namespace clang::CodeGen; @@ -337,6 +338,19 @@ static bool requiresAMDGPUProtectedVisibility(const Decl *D, cast<VarDecl>(D)->getType()->isCUDADeviceBuiltinTextureType()))); } +static void setAMDGPUConstantSection(CodeGenModule &M, const VarDecl *D, + llvm::GlobalValue *GV) { + if (GV->isDeclaration()) + return; + std::optional<StringRef> ConstantSection = + M.getTarget().getConstantAddressSpaceSectionName(); + if (!ConstantSection) + return; + + if (M.getContext().isTargetConstantAddressSpaceObject(D)) + cast<llvm::GlobalVariable>(GV)->setSection(*ConstantSection); +} + void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( const FunctionDecl *FD, llvm::Function *F, CodeGenModule &M) const { const auto *ReqdWGS = @@ -428,6 +442,9 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes( GV->setDSOLocal(true); } + if (const auto *VD = dyn_cast_or_null<VarDecl>(D)) + setAMDGPUConstantSection(M, VD, GV); + if (GV->isDeclaration()) return; diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index f1de44e3d2ed7..03a638dd242c5 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -28,12 +28,69 @@ #include "llvm/Support/AtomicOrdering.h" #include "llvm/TargetParser/AMDGPUTargetParser.h" #include <cstdint> +#include <optional> #include <utility> namespace clang { SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {} +namespace { + +static std::pair<StringRef, SourceLocation> +getExplicitAMDGPUSectionName(const VarDecl *D) { + if (auto *SA = D->getAttr<SectionAttr>()) + return {SA->getName(), SA->getLocation()}; + return {StringRef(), SourceLocation()}; +} + +static bool isDependentVarDecl(const VarDecl *D) { + if (D->getType()->isDependentType()) + return true; + if (const auto *Init = D->getInit()) + return Init->isValueDependent(); + return false; +} + +} // namespace + +void SemaAMDGPU::checkConstantAddressSpaceSection(VarDecl *VD) { + if (!SemaRef.Context.getTargetInfo().getTriple().isAMDGCN()) + return; + std::optional<StringRef> ConstantSection = + SemaRef.Context.getTargetInfo().getConstantAddressSpaceSectionName(); + if (!ConstantSection) + return; + if (VD->isInvalidDecl() || !VD->hasGlobalStorage() || + !VD->isThisDeclarationADefinition() || isDependentVarDecl(VD)) + return; + VarDecl *CanonicalVD = VD->getCanonicalDecl(); + if (InvalidConstantSectionVars.contains(CanonicalVD)) + return; + + bool IsConstantAS = SemaRef.Context.isTargetConstantAddressSpaceObject(VD); + + if (IsConstantAS) { + auto [ExplicitSection, ExplicitDiagLoc] = getExplicitAMDGPUSectionName(VD); + if (!ExplicitSection.empty() && ExplicitSection != *ConstantSection) { + SemaRef.Diag(ExplicitDiagLoc, + diag::err_amdgpu_constant_address_space_section) + << *ConstantSection; + InvalidConstantSectionVars.insert(CanonicalVD); + VD->setInvalidDecl(); + } + return; + } + + auto [Section, DiagLoc] = getExplicitAMDGPUSectionName(VD); + if (Section == *ConstantSection) { + SemaRef.Diag(DiagLoc, diag::err_amdgpu_reserved_constant_section) + << *ConstantSection; + InvalidConstantSectionVars.insert(CanonicalVD); + VD->setInvalidDecl(); + } +} + bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall) { // position of memory order and scope arguments in the builtin diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 7e1b23c971a9c..fbd0c34a48811 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -15355,6 +15355,8 @@ void Sema::FinalizeDeclaration(Decl *ThisDecl) { } } + AMDGPU().checkConstantAddressSpaceSection(VD); + CheckInvalidBuiltinCountedByRef(VD->getInit(), BuiltinCountedByRefKind::Initializer); diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index c2e90624b8a6e..973d71d3a2e9c 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -6556,6 +6556,7 @@ void Sema::InstantiateVariableDefinition(SourceLocation PointOfInstantiation, OldVar->getPointOfInstantiation()); // Emit any deferred warnings for the variable's initializer AnalysisWarnings.issueWarningsForRegisteredVarDecl(Var); + AMDGPU().checkConstantAddressSpaceSection(Var); } // This variable may have local implicit instantiations that need to be diff --git a/clang/test/CodeGen/amdgpu-address-spaces.cpp b/clang/test/CodeGen/amdgpu-address-spaces.cpp index 71e2ebd2c8a87..330810f20eca8 100644 --- a/clang/test/CodeGen/amdgpu-address-spaces.cpp +++ b/clang/test/CodeGen/amdgpu-address-spaces.cpp @@ -19,12 +19,12 @@ int [[clang::address_space(999)]] bbb = 1234; //. // CHECK: @a = addrspace(1) global i32 100, align 4 // CHECK: @b = global i32 42, align 4 -// CHECK: @c = addrspace(4) constant i32 999, align 4 +// CHECK: @c = addrspace(4) constant i32 999, section "__amdgpu_constant", align 4 // CHECK: @d = addrspace(3) global i32 undef, align 4 // CHECK: @e = addrspace(5) global i32 undef, align 4 // CHECK: @x = addrspace(1) global i32 100, align 4 // CHECK: @y = global i32 42, align 4 -// CHECK: @z = addrspace(4) global i32 999, align 4 +// CHECK: @z = addrspace(4) global i32 999, section "__amdgpu_constant", align 4 // CHECK: @w = addrspace(3) global i32 undef, align 4 // CHECK: @u = addrspace(5) global i32 undef, align 4 // CHECK: @aaa = addrspace(6) global i32 1000, align 4 diff --git a/clang/test/CodeGenCUDA/template-class-static-member.cu b/clang/test/CodeGenCUDA/template-class-static-member.cu index b614cd9dcbb14..866d548200caf 100644 --- a/clang/test/CodeGenCUDA/template-class-static-member.cu +++ b/clang/test/CodeGenCUDA/template-class-static-member.cu @@ -38,9 +38,9 @@ const int A<T>::const_member; template class A<int>; //DEV-DAG: @_ZN1AIiE8d_memberE = internal addrspace(1) global i32 0, comdat, align 4 -//DEV-DAG: @_ZN1AIiE8c_memberE = internal addrspace(4) constant i32 0, comdat, align 4 +//DEV-DAG: @_ZN1AIiE8c_memberE = internal addrspace(4) constant i32 0, section "__amdgpu_constant", comdat, align 4 //DEV-DAG: @_ZN1AIiE8m_memberE = internal addrspace(1) externally_initialized global ptr addrspace(1) null -//DEV-DAG: @_ZN1AIiE12const_memberE = internal addrspace(4) constant i32 0, comdat, align 4 +//DEV-DAG: @_ZN1AIiE12const_memberE = internal addrspace(4) constant i32 0, section "__amdgpu_constant", comdat, align 4 //DEV-NEG-NOT: @_ZN1AIiE8h_memberE //HOST-DAG: @_ZN1AIiE8h_memberE = weak_odr global i32 0, comdat, align 4 diff --git a/clang/test/CodeGenHIP/amdgpu-is-constant.hip b/clang/test/CodeGenHIP/amdgpu-is-constant.hip new file mode 100644 index 0000000000000..01bc547bec1fc --- /dev/null +++ b/clang/test/CodeGenHIP/amdgpu-is-constant.hip @@ -0,0 +1,102 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -emit-llvm -o - %s | FileCheck --check-prefixes=COMMON,WITH %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -emit-llvm -o - -DNO_CONSTANTS %s | FileCheck --check-prefixes=COMMON,EMPTY %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DCONSTANT_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DRESERVED_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DDEVICE_CONST_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DPRAGMA_CONSTANT_NO_DIAGNOSTICS -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DPRAGMA_RESERVED_NO_DIAGNOSTICS -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DPRAGMA_DEVICE_CONST_NO_DIAGNOSTICS -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DPRECEDENCE_NO_DIAGNOSTICS -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DPRECEDENCE_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DTEMPLATE_DEPENDENT_NO_INSTANTIATION -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fsyntax-only -DTEMPLATE_DEPENDENT_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -emit-llvm -o /dev/null -DBOUNDARY_SYMBOL_COLLISION_ERROR -verify %s + +#define __constant__ __attribute__((constant)) +#define __device__ __attribute__((device)) + +#ifdef CONSTANT_SECTION_ERROR +__constant__ int constant_wrong_section __attribute__((section("user_constant"))) = 1; +// expected-error@-1 {{AMDGPU constant address space variables must use section '__amdgpu_constant'}} +#elif defined(RESERVED_SECTION_ERROR) +__device__ int device_wrong_section __attribute__((section("__amdgpu_constant"))) = 1; +// expected-error@-1 {{section '__amdgpu_constant' is reserved for AMDGPU constant address space variables}} +#elif defined(DEVICE_CONST_SECTION_ERROR) +__device__ const int device_const_wrong_section __attribute__((section("user_constant"))) = 1; +// expected-error@-1 {{AMDGPU constant address space variables must use section '__amdgpu_constant'}} +#elif defined(PRAGMA_CONSTANT_NO_DIAGNOSTICS) +// expected-no-diagnostics +#pragma clang section data = "user_constant" +__constant__ int constant_wrong_pragma = 1; +#pragma clang section data = "" +#elif defined(PRAGMA_RESERVED_NO_DIAGNOSTICS) +// expected-no-diagnostics +#pragma clang section data = "__amdgpu_constant" +__device__ int device_wrong_pragma = 1; +#pragma clang section data = "" +#elif defined(PRAGMA_DEVICE_CONST_NO_DIAGNOSTICS) +// expected-no-diagnostics +#pragma clang section data = "user_constant" +__device__ const int device_const_wrong_pragma = 1; +#pragma clang section data = "" +#elif defined(PRECEDENCE_NO_DIAGNOSTICS) +// expected-no-diagnostics +#pragma clang section data = "user_constant" +__constant__ int explicit_constant_section_precedence + __attribute__((section("__amdgpu_constant"))) = 1; +#pragma clang section data = "" +#pragma clang section data = "__amdgpu_constant" +__device__ int explicit_device_section_precedence + __attribute__((section("user_data"))) = 1; +#pragma clang section data = "" +#elif defined(PRECEDENCE_SECTION_ERROR) +#pragma clang section data = "__amdgpu_constant" +__constant__ int explicit_wrong_constant_section_precedence + __attribute__((section("user_constant"))) = 1; +// expected-error@-1 {{AMDGPU constant address space variables must use section '__amdgpu_constant'}} +#pragma clang section data = "" +#elif defined(TEMPLATE_DEPENDENT_NO_INSTANTIATION) +// expected-no-diagnostics +template <class T> +__device__ const T dependent_wrong_section __attribute__((section("user_constant"))) = T(1); +#elif defined(TEMPLATE_DEPENDENT_SECTION_ERROR) +template <class T> +__device__ const T dependent_wrong_section __attribute__((section("user_constant"))) = T(1); +// expected-error@-1 {{AMDGPU constant address space variables must use section '__amdgpu_constant'}} +__device__ int use_dependent_wrong_section() { + return dependent_wrong_section<int>; + // expected-note@-1 {{in instantiation of variable template specialization 'dependent_wrong_section<int>' requested here}} +} +#elif defined(BOUNDARY_SYMBOL_COLLISION_ERROR) +extern "C" __device__ int __start___amdgpu_constant; +__device__ bool boundary_symbol_collision(const void *ptr) { + (void)&__start___amdgpu_constant; + return __builtin_amdgcn_is_constant(ptr); + // expected-error@-1 {{AMDGPU constant section boundary symbol '__start___amdgpu_constant' conflicts with existing symbol}} +} +#elif !defined(NO_CONSTANTS) +// WITH-DAG: @constant_var = addrspace(4) externally_initialized constant i32 1, section "__amdgpu_constant" +__constant__ int constant_var = 1; +#endif + +// COMMON-DAG: @device_var = addrspace(1) externally_initialized global i32 0 +__device__ int device_var; + +// COMMON-DAG: @device_const_var = addrspace(4) constant i32 2, section "__amdgpu_constant" +extern "C" __device__ const int device_const_var = 2; + +// COMMON-DAG: @__clang_amdgpu_constant_section_marker = internal unnamed_addr addrspace(4) constant [0 x i8] zeroinitializer, section "__amdgpu_constant" +// COMMON-DAG: @__start___amdgpu_constant = external addrspace(4) constant i8 +// COMMON-DAG: @__stop___amdgpu_constant = external addrspace(4) constant i8 +// COMMON-DAG: @llvm.compiler.used = {{.*}}@__clang_amdgpu_constant_section_marker{{.*}}@__hip_cuid_{{.*}}, section "llvm.metadata" +// EMPTY-NOT: @constant_var = + +// COMMON-LABEL: define{{.*}} noundef zeroext i1 @_Z11is_constantPKv( +// COMMON: [[PTR:%.*]] = load ptr, ptr {{.*}}, align 8 +// COMMON: [[PTR_INT:%.*]] = ptrtoint ptr [[PTR]] to i64 +// COMMON: [[OFFSET:%.*]] = sub i64 [[PTR_INT]], ptrtoint (ptr addrspace(4) @__start___amdgpu_constant to i64) +// COMMON: [[CMP:%.*]] = icmp ult i64 [[OFFSET]], sub (i64 ptrtoint (ptr addrspace(4) @__stop___amdgpu_constant to i64), i64 ptrtoint (ptr addrspace(4) @__start___amdgpu_constant to i64)) +// COMMON: ret i1 [[CMP]] +__device__ bool is_constant(const void *ptr) { + return __builtin_amdgcn_is_constant(ptr); +} diff --git a/clang/test/CodeGenOpenCL/amdgpu-is-constant.cl b/clang/test/CodeGenOpenCL/amdgpu-is-constant.cl new file mode 100644 index 0000000000000..6a7872a7db108 --- /dev/null +++ b/clang/test/CodeGenOpenCL/amdgpu-is-constant.cl @@ -0,0 +1,81 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -DCONSTANT_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -DPRAGMA_CONSTANT_NO_DIAGNOSTICS -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -DRESERVED_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -DPRAGMA_RESERVED_NO_DIAGNOSTICS -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -DGLOBAL_CONST_RESERVED_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -DPRECEDENCE_CONSTANT_NO_DIAGNOSTICS -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -DPRECEDENCE_GLOBAL_NO_DIAGNOSTICS -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -DPRECEDENCE_SECTION_ERROR -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -emit-llvm -o /dev/null -DBOUNDARY_SYMBOL_COLLISION_ERROR -verify %s + +#ifdef CONSTANT_SECTION_ERROR +__constant int constant_wrong_section __attribute__((section("user_constant"))) = 1; +// expected-error@-1 {{AMDGPU constant address space variables must use section '__amdgpu_constant'}} +#elif defined(PRAGMA_CONSTANT_NO_DIAGNOSTICS) +// expected-no-diagnostics +#pragma clang section data = "user_constant" +__constant int constant_wrong_pragma = 1; +#pragma clang section data = "" +#elif defined(RESERVED_SECTION_ERROR) +__global int global_wrong_section __attribute__((section("__amdgpu_constant"))) = 1; +// expected-error@-1 {{section '__amdgpu_constant' is reserved for AMDGPU constant address space variables}} +#elif defined(PRAGMA_RESERVED_NO_DIAGNOSTICS) +// expected-no-diagnostics +#pragma clang section data = "__amdgpu_constant" +__global int global_wrong_pragma = 1; +#pragma clang section data = "" +#elif defined(GLOBAL_CONST_RESERVED_SECTION_ERROR) +__global const int global_const_wrong_section __attribute__((section("__amdgpu_constant"))) = 1; +// expected-error@-1 {{section '__amdgpu_constant' is reserved for AMDGPU constant address space variables}} +#elif defined(PRECEDENCE_CONSTANT_NO_DIAGNOSTICS) +// expected-no-diagnostics +#pragma clang section data = "user_constant" +__constant int explicit_constant_section_precedence + __attribute__((section("__amdgpu_constant"))) = 1; +#pragma clang section data = "" +#elif defined(PRECEDENCE_GLOBAL_NO_DIAGNOSTICS) +// expected-no-diagnostics +#pragma clang section data = "__amdgpu_constant" +__global int explicit_global_section_precedence + __attribute__((section("user_global"))) = 1; +#pragma clang section data = "" +#elif defined(PRECEDENCE_SECTION_ERROR) +#pragma clang section data = "__amdgpu_constant" +__constant int explicit_wrong_constant_section_precedence + __attribute__((section("user_constant"))) = 1; +// expected-error@-1 {{AMDGPU constant address space variables must use section '__amdgpu_constant'}} +#pragma clang section data = "" +#elif defined(BOUNDARY_SYMBOL_COLLISION_ERROR) +__global int __start___amdgpu_constant; +bool boundary_symbol_collision(const void *ptr) { + (void)&__start___amdgpu_constant; + return __builtin_amdgcn_is_constant(ptr); + // expected-error@-1 {{AMDGPU constant section boundary symbol '__start___amdgpu_constant' conflicts with existing symbol}} +} +#else +// CHECK-DAG: @constant_var = {{.*}}addrspace(4) constant i32 1, section "__amdgpu_constant" +__constant int constant_var = 1; + +// CHECK-DAG: @global_var = {{.*}}addrspace(1) global i32 0 +__global int global_var; + +// CHECK-DAG: @global_const_var = {{.*}}addrspace(1) constant i32 2, align 4{{$}} +__global const int global_const_var = 2; + +// CHECK-DAG: @__clang_amdgpu_constant_section_marker = internal unnamed_addr addrspace(4) constant [0 x i8] zeroinitializer, section "__amdgpu_constant" +// CHECK-DAG: @__start___amdgpu_constant = external addrspace(4) constant i8 +// CHECK-DAG: @__stop___amdgpu_constant = external addrspace(4) constant i8 +// CHECK-DAG: @llvm.compiler.used = {{.*}}@__clang_amdgpu_constant_section_marker{{.*}}, section "llvm.metadata" +#endif + +// CHECK-LABEL: define{{.*}} zeroext i1 @is_constant( +// CHECK-SAME: ptr noundef [[PTR:%[[:alnum:]_.]+]] +// CHECK: [[PTR_INT:%.*]] = ptrtoint ptr [[PTR]] to i64 +// CHECK: [[OFFSET:%.*]] = sub i64 [[PTR_INT]], ptrtoint (ptr addrspace(4) @__start___amdgpu_constant to i64) +// CHECK: [[CMP:%.*]] = icmp ult i64 [[OFFSET]], sub (i64 ptrtoint (ptr addrspace(4) @__stop___amdgpu_constant to i64), i64 ptrtoint (ptr addrspace(4) @__start___amdgpu_constant to i64)) +// CHECK: ret i1 [[CMP]] +bool is_constant(const void *ptr) { + return __builtin_amdgcn_is_constant(ptr); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
