Author: Pierre van Houtryve Date: 2026-08-03T10:38:00+02:00 New Revision: b5d93cd7773b203e3523f36cf1b63c16f23c921b
URL: https://github.com/llvm/llvm-project/commit/b5d93cd7773b203e3523f36cf1b63c16f23c921b DIFF: https://github.com/llvm/llvm-project/commit/b5d93cd7773b203e3523f36cf1b63c16f23c921b.diff LOG: [clang][AMDGPU] Clean-up handling of named barrier type (#207687) - Allow the type in struct/classes in very limited circumstances. The goal is to enable creating trivial wrappers around the named barrier variable, but ensure we can't get into situations where things would get awkward. Currently this means we only allow the named barrier in RecordDecls with exactly 1 field, that have no base class, and are not inherited. - Use a `amdgpu_barrier` LangAS for this type that currently maps to the local AS. This allows easy switching to the barrier AS in a future patch. Added: clang/test/SemaHIP/amdgpu-barrier-spirv.hip Modified: clang/docs/AMDGPUSupport.md clang/include/clang/AST/TypeBase.h clang/include/clang/Basic/AddressSpaces.h clang/include/clang/Basic/Attr.td clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Basic/TargetInfo.h clang/include/clang/Sema/SemaAMDGPU.h clang/lib/AST/ASTContext.cpp clang/lib/AST/Type.cpp clang/lib/AST/TypePrinter.cpp clang/lib/Basic/TargetInfo.cpp clang/lib/Basic/Targets/AMDGPU.cpp clang/lib/Basic/Targets/SPIR.h clang/lib/CodeGen/CodeGenModule.cpp clang/lib/Sema/Sema.cpp clang/lib/Sema/SemaAMDGPU.cpp clang/lib/Sema/SemaDecl.cpp clang/test/CodeGenHIP/amdgpu-barrier-type.hip clang/test/SemaCXX/amdgpu-barrier.cpp clang/test/SemaHIP/amdgpu-barrier.hip clang/test/SemaOpenCL/amdgpu-barrier.cl clang/test/SemaTemplate/address_space-dependent.cpp Removed: ################################################################################ diff --git a/clang/docs/AMDGPUSupport.md b/clang/docs/AMDGPUSupport.md index 8731af371b23f..3306242d9d7e6 100644 --- a/clang/docs/AMDGPUSupport.md +++ b/clang/docs/AMDGPUSupport.md @@ -30,3 +30,46 @@ Please note that the specific architecture and feature names will vary depending Clang exposes AMDGPU hardware intrinsics as target-specific builtins with the `__builtin_amdgcn_` prefix. These are documented in {doc}`AMDGPUBuiltinReference`. + +## Target-Specific Types + +### Named Workgroup Barrier Type + +The `__amdgpu_named_workgroup_barrier_t` type is used to represent the GFX12.5 named barriers. + +This type is subject to certain restrictions when used to declare non-static fields of classes. +See the {ref}`named barrier wrappers<namedbarrierwrappers>` section below for more information. + +Example usage: + +```c +__amdgpu_named_workgroup_barrier_t x; +__amdgpu_named_workgroup_barrier_t arr[2]; // Arrays are also fine + +struct SimpleWrapper { + __amdgpu_named_workgroup_barrier_t foo; +}; + +void foo(int a) +{ + __builtin_amdgcn_s_barrier_init(&x, a); +} +``` + +(namedbarrierwrappers)= +### Named Barrier Wrappers + +A "named barrier wrapper" is a class that contains exactly one non-static field +of one of the following types: + +* `__amdgpu_named_workgroup_barrier_t`, or an array thereof. +* Another named barrier wrapper, or an array thereof. + +In C++, a class that inherits from a named barrier wrapper is also considered a +named barrier wrapper, and all named barrier wrappers must be standard-layout +types (see `std::is_standard_layout`). This means that named barrier wrappers: + +* May not have a virtual table: they cannot declare or inherit any virtual + functions, or inherit from a virtual base class. +* May not have any extra fields, either declared by the class or inherited + from a base class. diff --git a/clang/include/clang/AST/TypeBase.h b/clang/include/clang/AST/TypeBase.h index 530bfe72dac2b..4851c4e5185dd 100644 --- a/clang/include/clang/AST/TypeBase.h +++ b/clang/include/clang/AST/TypeBase.h @@ -2814,6 +2814,12 @@ class alignas(TypeAlignment) Type : public ExtQualsTypeCommonBase { /// Check if the type is the CUDA device builtin texture type. bool isCUDADeviceBuiltinTextureType() const; + /// Check if the type is the AMDGPU named barrier type, or an array thereof. + bool isAMDGPUNamedBarrierType() const; + /// Check if the type is the AMDGPU named barrier type/a RecordType of a named + /// barrier wrapper, or an array thereof. + bool isAMDGPUNamedBarrierTypeOrWrapper() const; + /// Return the implicit lifetime for this type, which must not be dependent. Qualifiers::ObjCLifetime getObjCARCImplicitLifetime() const; diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h index 2dfaa1c45ac55..ce38f60163b64 100644 --- a/clang/include/clang/Basic/AddressSpaces.h +++ b/clang/include/clang/Basic/AddressSpaces.h @@ -71,6 +71,9 @@ enum class LangAS : unsigned { // Wasm specific address spaces. wasm_funcref, + // AMDGPU address spaces + amdgpu_barrier, + // This denotes the count of language-specific address spaces and also // the offset added to the target-specific address spaces, which are usually // specified by address space attributes __attribute__(address_space(n))). diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 39c672322d515..ff8c8160ad6c1 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2532,6 +2532,13 @@ def AMDGPUMaxNumWorkGroups : InheritableAttr { let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; } +def AMDGPUNamedBarrierWrapper : InheritableAttr { + let Spellings = []; + let Args = []; + let Documentation = [InternalOnly]; + let SemaHandler = 0; +} + def BPFPreserveAccessIndex : InheritableAttr, TargetSpecificAttr<TargetBPF> { let Spellings = [Clang<"preserve_access_index">]; @@ -5333,7 +5340,7 @@ def HLSLVkLocation : HLSLAnnotationAttr { } // `row_major` / `column_major` are HLSL keywords that select the in-memory -// layout of a matrix-typed declaration. +// layout of a matrix-typed declaration. def HLSLRowMajor : TypeAttr { let Spellings = [CustomKeyword<"row_major">]; let LangOpts = [HLSL]; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 01527e87c903f..8d18f90f95aa9 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14282,6 +14282,22 @@ def note_acc_reduction_combiner_forming : Note<"while forming %select{|binary operator '%1'|conditional " "operator|final assignment operator}0">; +// AMDGCN type diagnostics +def err_amdgpu_target_ext_type_unsupported : Error< + "AMDGPU built-in type %0 is not supported on target '%1'">; + +def err_amdgpu_invalid_field_not_a_wrapper : Error< + "field with barrier type %0 seen in a structure that is not a named barrier wrapper">; +def note_amdgpu_not_a_named_barrier_wrapper_too_many_fields : Note< + "%0 is not a named barrier wrapper because it has more than one field">; + +def err_amdgpu_named_barrier_wrapper_non_standard_layout : Error< + "named barrier wrapper %0 must have a C++11 standard layout">; +def note_amdgpu_named_barrier_reason_field : Note< + "%0 is a named barrier wrapper because it has a named barrier or named barrier wrapper field %1 found here">; +def note_amdgpu_named_barrier_reason_inherited : Note< + "%0 is a named barrier wrapper because it inherits from named barrier wrapper %1">; + // AMDGCN builtins diagnostics def err_amdgcn_load_lds_size_invalid_value : Error<"invalid size value">; def note_amdgcn_load_lds_size_valid_value : Note<"size must be %select{1, 2, or 4|1, 2, 4, 12 or 16}0">; diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h index 968a0c1b129ef..6311b6b567a5e 100644 --- a/clang/include/clang/Basic/TargetInfo.h +++ b/clang/include/clang/Basic/TargetInfo.h @@ -287,6 +287,9 @@ class TargetInfo : public TransferrableTargetInfo, LLVM_PREFERRED_TYPE(bool) unsigned HasUnalignedAccess : 1; + LLVM_PREFERRED_TYPE(bool) + unsigned HasAMDGPUTypes : 1; + unsigned ARMCDECoprocMask : 8; unsigned MaxOpenCLWorkGroupSize; @@ -1073,6 +1076,10 @@ class TargetInfo : public TransferrableTargetInfo, /// available on this target. bool hasAArch64ACLETypes() const { return HasAArch64ACLETypes; } + /// Returns whether or not the AMDGPU built-in types are + /// available on this target. + bool hasAMDGPUTypes() const { return HasAMDGPUTypes; } + /// Returns whether or not the RISC-V V built-in types are /// available on this target. bool hasRISCVVTypes() const { return HasRISCVVTypes; } diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h index a6205534e0de3..0cf4f8adcf108 100644 --- a/clang/include/clang/Sema/SemaAMDGPU.h +++ b/clang/include/clang/Sema/SemaAMDGPU.h @@ -89,6 +89,13 @@ class SemaAMDGPU : public SemaBase { void AddPotentiallyUnguardedBuiltinUser(FunctionDecl *FD); bool HasPotentiallyUnguardedBuiltinUsage(FunctionDecl *FD) const; void DiagnoseUnguardedBuiltinUsage(FunctionDecl *FD); + + /// Check if \p Ty is supported on this AMDGPU target. + /// \returns false if \p Ty is unsupported and a diagnostic was emitted. + bool checkAMDGPUTypeSupport(QualType Ty, SourceLocation Loc); + + /// Called in `ActOnFields` - whenever a C/C++ Record is being finalized. + void checkNamedBarrierWrapper(RecordDecl *R); }; } // namespace clang diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 02a3f88431f58..5f1e5b30ee50c 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1478,13 +1478,7 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target, #include "clang/Basic/WebAssemblyReferenceTypes.def" } - if (Target.getTriple().isAMDGPU() || - (Target.getTriple().isSPIRV() && - Target.getTriple().getVendor() == llvm::Triple::AMD) || - (AuxTarget && - (AuxTarget->getTriple().isAMDGPU() || - ((AuxTarget->getTriple().isSPIRV() && - AuxTarget->getTriple().getVendor() == llvm::Triple::AMD))))) { + if (Target.hasAMDGPUTypes() || (AuxTarget && (AuxTarget->hasAMDGPUTypes()))) { #define AMDGPU_TYPE(Name, Id, SingletonId, Width, Align) \ InitBuiltinType(SingletonId, BuiltinType::Id); #include "clang/Basic/AMDGPUTypes.def" diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp index 42d148715bc40..e51e7de9f176a 100644 --- a/clang/lib/AST/Type.cpp +++ b/clang/lib/AST/Type.cpp @@ -93,7 +93,7 @@ bool Qualifiers::isTargetAddressSpaceSupersetOf(LangAS A, LangAS B, // to implicitly cast into the default address space. (A == LangAS::Default && (B == LangAS::cuda_constant || B == LangAS::cuda_device || - B == LangAS::cuda_shared)) || + B == LangAS::cuda_shared || B == LangAS::amdgpu_barrier)) || // In HLSL, the this pointer for member functions points to the default // address space. This causes a problem if the structure is in // a diff erent address space. We want to allow casting from these @@ -5492,6 +5492,31 @@ bool Type::isCUDADeviceBuiltinTextureType() const { return false; } +static bool isAMDGPUNamedBarrierTypeImpl(const Type *Ty, bool AllowWrappers) { + // This query does not care about qualifiers at all. + Ty = Ty->getUnqualifiedDesugaredType(); + + // Unwrap arrays. + while (isa<ArrayType>(Ty)) + Ty = Ty->getArrayElementTypeNoTypeQual()->getUnqualifiedDesugaredType(); + + if (const auto *BT = dyn_cast<BuiltinType>(Ty)) + return BT->getKind() == BuiltinType::AMDGPUNamedWorkgroupBarrier; + if (AllowWrappers) { + if (const auto *RT = dyn_cast<RecordType>(Ty)) + return RT->getDecl()->hasAttr<AMDGPUNamedBarrierWrapperAttr>(); + } + return false; +} + +bool Type::isAMDGPUNamedBarrierType() const { + return isAMDGPUNamedBarrierTypeImpl(this, /*AllowWrappers=*/false); +} + +bool Type::isAMDGPUNamedBarrierTypeOrWrapper() const { + return isAMDGPUNamedBarrierTypeImpl(this, /*AllowWrappers=*/true); +} + bool Type::hasSizedVLAType() const { if (!isVariablyModifiedType()) return false; diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp index 60964d2859790..75067f6ab3265 100644 --- a/clang/lib/AST/TypePrinter.cpp +++ b/clang/lib/AST/TypePrinter.cpp @@ -2766,6 +2766,8 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) { return "hlsl_push_constant"; case LangAS::wasm_funcref: return "__funcref"; + case LangAS::amdgpu_barrier: + return "amdgpu_barrier"; default: return std::to_string(toTargetAddressSpace(AS)); } diff --git a/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp index 103104ce4874b..1bb0026d8422e 100644 --- a/clang/lib/Basic/TargetInfo.cpp +++ b/clang/lib/Basic/TargetInfo.cpp @@ -21,6 +21,7 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/TargetParser/TargetParser.h" +#include "llvm/TargetParser/Triple.h" #include <cstdlib> using namespace clang; @@ -55,6 +56,7 @@ static constexpr LangASMap FakeAddrSpaceMap = { {LangAS::hlsl_output, 18}, {LangAS::hlsl_push_constant, 19}, {LangAS::wasm_funcref, 20}, + {LangAS::amdgpu_barrier, 21}, }; // TargetInfo Constructor. @@ -167,6 +169,7 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : Triple(T) { HasBuiltinZOSVaList = false; HasAArch64ACLETypes = false; HasRISCVVTypes = false; + HasAMDGPUTypes = false; AllowAMDGPUUnsafeFPAtomics = false; HasUnalignedAccess = false; ARMCDECoprocMask = 0; diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 1d74fcc3428b4..4109066ec910e 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -56,6 +56,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUAddrSpaceMap = { {LangAS::hlsl_input, llvm::AMDGPUAS::PRIVATE_ADDRESS}, {LangAS::hlsl_output, llvm::AMDGPUAS::PRIVATE_ADDRESS}, {LangAS::hlsl_push_constant, llvm::AMDGPUAS::GLOBAL_ADDRESS}, + {LangAS::amdgpu_barrier, llvm::AMDGPUAS::LOCAL_ADDRESS}, }; } // namespace targets @@ -202,6 +203,7 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple, AddrSpaceMap = &AMDGPUAddrSpaceMap; UseAddrSpaceMapMangling = true; + HasAMDGPUTypes = true; if (Triple.isAMDGCN()) { // __bf16 is always available as a load/store only type on AMDGCN. diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index a16c8ec79d2f2..a5d58dcf2c835 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -278,6 +278,7 @@ class LLVM_LIBRARY_VISIBILITY BaseSPIRVTargetInfo : public BaseSPIRTargetInfo { BaseSPIRVTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : BaseSPIRTargetInfo(Triple, Opts) { assert(Triple.isSPIRV() && "Invalid architecture for SPIR-V."); + HasAMDGPUTypes = (Triple.getVendor() == llvm::Triple::AMD); } llvm::SmallVector<Builtin::InfosShard> getTargetBuiltins() const override; diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index c284852c05790..9ce77917f40bc 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6266,6 +6266,9 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { if (LangOpts.CUDA && LangOpts.CUDAIsDevice) { if (D) { + if (D->getType()->isAMDGPUNamedBarrierTypeOrWrapper()) + return LangAS::amdgpu_barrier; + if (D->hasAttr<CUDAConstantAttr>()) return LangAS::cuda_constant; if (D->hasAttr<CUDASharedAttr>()) diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 66a88195722f8..2c229bb12cfc1 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -573,14 +573,9 @@ void Sema::Initialize() { #include "clang/Basic/WebAssemblyReferenceTypes.def" } - if (Context.getTargetInfo().getTriple().isAMDGPU() || - (Context.getTargetInfo().getTriple().isSPIRV() && - Context.getTargetInfo().getTriple().getVendor() == llvm::Triple::AMD) || + if (Context.getTargetInfo().hasAMDGPUTypes() || (Context.getAuxTargetInfo() && - (Context.getAuxTargetInfo()->getTriple().isAMDGPU() || - (Context.getAuxTargetInfo()->getTriple().isSPIRV() && - Context.getAuxTargetInfo()->getTriple().getVendor() == - llvm::Triple::AMD)))) { + (Context.getAuxTargetInfo()->hasAMDGPUTypes()))) { #define AMDGPU_TYPE(Name, Id, SingletonId, Width, Align) \ addImplicitTypedef(Name, Context.SingletonId); #include "clang/Basic/AMDGPUTypes.def" @@ -2411,6 +2406,9 @@ void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) { ARM().checkSVETypeSupport(Ty, Loc, FD, CallerFeatureMap); } + if (TI.hasAMDGPUTypes()) + AMDGPU().checkAMDGPUTypeSupport(Ty, Loc); + if (auto *VT = Ty->getAs<VectorType>(); VT && FD && (VT->getVectorKind() == VectorKind::SveFixedLengthData || diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 48230fa262d5c..11274458dd962 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -1079,4 +1079,107 @@ bool DiagnoseUnguardedBuiltins::VisitCallExpr(CallExpr *CE) { void SemaAMDGPU::DiagnoseUnguardedBuiltinUsage(FunctionDecl *FD) { DiagnoseUnguardedBuiltins(SemaRef).IssueDiagnostics(FD->getBody()); } + +bool SemaAMDGPU::checkAMDGPUTypeSupport(QualType Ty, SourceLocation Loc) { + ASTContext &Ctx = getASTContext(); + llvm::Triple TT = Ctx.getTargetInfo().getTriple(); + const Type *BaseTy = Ty->getPointeeOrArrayElementType(); + + if (Ctx.getTargetInfo().getTriple().isSPIRV()) { + // The AMDGPU named barrier type requires special handling in the back-end + // and is not supported for SPIR-V + if (BaseTy->isAMDGPUNamedBarrierType()) { + SemaRef.Diag(Loc, diag::err_amdgpu_target_ext_type_unsupported) + << Ty << TT.str(); + return false; + } + } + + return true; +} + +static FieldDecl *getNamedBarrierField(const RecordDecl *R) { + for (FieldDecl *FD : R->fields()) { + QualType FDTy = FD->getType(); + if (FDTy->isAMDGPUNamedBarrierTypeOrWrapper()) + return FD; + } + + return nullptr; +} + +void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { + ASTContext &Context = getASTContext(); + if (R->isInvalidDecl()) + return; + + if (!Context.getTargetInfo().hasAMDGPUTypes() && + (!Context.getAuxTargetInfo() || + !Context.getAuxTargetInfo()->hasAMDGPUTypes())) + return; + + bool IsWrapper = false; + std::function<void()> DiagWrapperNote; + + // First, check if this is a named barrier wrapper by virtue of the class + // declaring a named barrier field. This covers both C and C++. + if (FieldDecl *NamedBarrField = getNamedBarrierField(R)) { + // If this record contains a named barrier field, it must have only one + // field. + if (R->getNumFields() > 1) { + SemaRef.Diag(NamedBarrField->getLocation(), + diag::err_amdgpu_invalid_field_not_a_wrapper) + << NamedBarrField->getType(); + SemaRef.Diag( + R->getLocation(), + diag::note_amdgpu_not_a_named_barrier_wrapper_too_many_fields) + << R->getDeclName(); + return; + } + + IsWrapper = true; + DiagWrapperNote = [this, R, NamedBarrField]() { + SemaRef.Diag(NamedBarrField->getLocation(), + diag::note_amdgpu_named_barrier_reason_field) + << R->getDeclName() << NamedBarrField->getDeclName(); + }; + } + + // Then, for C++ classes, check if this is a named barrier wrapper by virtue + // of inheriting one. + const auto *CxxR = dyn_cast<CXXRecordDecl>(R); + if (CxxR && !IsWrapper) { + for (CXXBaseSpecifier BS : CxxR->bases()) { + const RecordDecl *Base = BS.getType()->getAsRecordDecl(); + if (!Base || !Base->hasAttr<AMDGPUNamedBarrierWrapperAttr>()) + continue; + + IsWrapper = true; + DiagWrapperNote = [this, BS, R]() { + // Print using the CXXBaseSpecifier type as it includes the template + // parameters. + SemaRef.Diag(BS.getBeginLoc(), + diag::note_amdgpu_named_barrier_reason_inherited) + << R->getDeclName() << BS.getType(); + }; + } + } + + if (!IsWrapper) + return; + + // Set the attribute even if the wrapper may be found to be invalid later. + R->addAttr( + AMDGPUNamedBarrierWrapperAttr::CreateImplicit(Context, SourceRange())); + + // This is a wrapper CXXRecordDecl, it must have a C++11 standard layout. + if (CxxR && !CxxR->isCXX11StandardLayout()) { + SemaRef.Diag(R->getLocation(), + diag::err_amdgpu_named_barrier_wrapper_non_standard_layout) + << R->getDeclName(); + assert(DiagWrapperNote && + "IsWrapper is set but no context diagnostic provided"); + DiagWrapperNote(); + } +} } // namespace clang diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index c46d2d780fad7..ae8c26808a164 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -9240,6 +9240,13 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) { CallerFeatureMap); } + if (Context.getTargetInfo().hasAMDGPUTypes()) { + if (!AMDGPU().checkAMDGPUTypeSupport(T, NewVD->getLocation())) { + NewVD->setInvalidDecl(); + return; + } + } + if (T.hasAddressSpace() && !CheckVarDeclSizeAddressSpace(NewVD, T.getAddressSpace())) { NewVD->setInvalidDecl(); @@ -19661,6 +19668,11 @@ FieldDecl *Sema::CheckFieldDecl(DeclarationName Name, QualType T, PPC().CheckPPCMMAType(T, NewFD->getLocation())) NewFD->setInvalidDecl(); + if (Context.getTargetInfo().hasAMDGPUTypes()) { + if (!AMDGPU().checkAMDGPUTypeSupport(T, NewFD->getLocation())) + NewFD->setInvalidDecl(); + } + NewFD->setAccess(AS); return NewFD; } @@ -20508,6 +20520,10 @@ void Sema::ActOnFields(Scope *S, SourceLocation RecLoc, Decl *EnclosingDecl, CDecl->setIvarRBraceLoc(RBrac); } } + + if (Record) + AMDGPU().checkNamedBarrierWrapper(Record); + if (Record && !isa<ClassTemplateSpecializationDecl>(Record)) ProcessAPINotes(Record); } diff --git a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip index b5f75dcf607ba..c5845a036aeb0 100644 --- a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip +++ b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip @@ -1,18 +1,28 @@ -// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature - // REQUIRES: amdgpu-registered-target - // RUN: %clang_cc1 -triple amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --global-value-regex "bar.*" +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -fcuda-is-device -triple amdgpu12.5-unknown-unknown -emit-llvm -o - %s | FileCheck %s #define __shared__ __attribute__((shared)) __shared__ __amdgpu_named_workgroup_barrier_t bar; -__shared__ __amdgpu_named_workgroup_barrier_t arr[2]; -__shared__ struct { +__shared__ __amdgpu_named_workgroup_barrier_t bar_arr[2]; + +//. +// CHECK: @bar = addrspace(3) global target("amdgcn.named.barrier", 0) undef, align 4 +// CHECK: @bar_arr = addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] undef, align 4 +// CHECK: @bar_wrapper_str = addrspace(3) global %struct.WrapperStruct undef, align 4 +// CHECK: @bar_wrapperwrapper_str = addrspace(3) global %struct.WrapperWrapperStruct undef, align 4 +//. +__shared__ struct WrapperStruct { __amdgpu_named_workgroup_barrier_t x; - __amdgpu_named_workgroup_barrier_t y; -} str; +} bar_wrapper_str; + +__shared__ struct WrapperWrapperStruct { + WrapperStruct x; +} bar_wrapperwrapper_str; -__amdgpu_named_workgroup_barrier_t *getBar(); -void useBar(__amdgpu_named_workgroup_barrier_t *); +__attribute__((device)) __amdgpu_named_workgroup_barrier_t *getBar(); +__attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); // CHECK-LABEL: define {{[^@]+}}@_Z7testSemPu34__amdgpu_named_workgroup_barrier_t // CHECK-SAME: (ptr noundef [[P:%.*]]) #[[ATTR0:[0-9]+]] { @@ -22,19 +32,21 @@ void useBar(__amdgpu_named_workgroup_barrier_t *); // CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8 // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[TMP0]]) #[[ATTR2:[0-9]+]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(1) @bar to ptr)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(1) @arr to ptr), i64 16)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(1) @str to ptr), i64 16)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @bar to ptr)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @bar_wrapper_str to ptr)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @bar_wrapperwrapper_str to ptr)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(3) @bar_arr to ptr), i64 16)) #[[ATTR2]] // CHECK-NEXT: [[CALL:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]] // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[CALL]]) #[[ATTR2]] // CHECK-NEXT: [[CALL1:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]] // CHECK-NEXT: ret ptr [[CALL1]] // -__amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_named_workgroup_barrier_t *p) { +__attribute__((device)) __amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_named_workgroup_barrier_t *p) { useBar(p); useBar(&bar); - useBar(&arr[1]); - useBar(&str.y); + useBar(&bar_wrapper_str.x); + useBar(&bar_wrapperwrapper_str.x.x); + useBar(&bar_arr[1]); useBar(getBar()); return getBar(); } diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index ffd8eb36aed68..3d43dcd1bba6b 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -13,5 +13,76 @@ void foo() { void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}} } +using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; + +struct TestSimple { + __amdgpu_named_workgroup_barrier_t x; +}; + +struct TestArray{ + __amdgpu_named_workgroup_barrier_t y[2]; +}; + +struct TestSugared { + SugaredArray z[2]; +}; + +struct TestSimpleWithStaticField { + __amdgpu_named_workgroup_barrier_t x; + static unsigned Harmless; +}; + +// Wrappers cannot have >1 field. +struct WrapperHasTooManyFields { // expected-note {{'WrapperHasTooManyFields' is not a named barrier wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{field with barrier type '__amdgpu_named_workgroup_barrier_t' seen in a structure that is not a named barrier wrapper}} + int other; +}; + +// Wrappers must have standard layout +struct WrapperBase { +__amdgpu_named_workgroup_barrier_t x; +}; + +struct WrapperWithBase : public WrapperBase { +}; + +// expected-error@+2 {{named barrier wrapper 'WrapperWithBaseNoStandardLayout' must have a C++11 standard layou}} +// expected-note@+1 {{'WrapperWithBaseNoStandardLayout' is a named barrier wrapper because it inherits from named barrier wrapper 'WrapperBase'}} +struct WrapperWithBaseNoStandardLayout : public WrapperBase { + unsigned K = 0; +}; + +// Wrappers of Wrappers have the same restrictions. +struct WrapperOfWrapperWithTooManyFields { // expected-note {{'WrapperOfWrapperWithTooManyFields' is not a named barrier wrapper because it has more than one field}} + TestSimple x; // expected-error {{field with barrier type 'TestSimple' seen in a structure that is not a named barrier wrapper}} + int other; +}; + +struct WrapperOfWrapper { + WrapperWithBase y; +}; + +// Check templated cases with a bit of complexity thrown in. +template<typename Derived> +class CRTPWrapperBase { + CRTPWrapperBase() { + static_cast<Derived*>(this)->sayHello(); + } + + WrapperOfWrapper wow; +}; + +class TemplatedWrapperImpl : public CRTPWrapperBase<TemplatedWrapperImpl> { + void sayHello() {} +}; + +// expected-error@+2 {{named barrier wrapper 'TemplatedWrapperImplNoStandardLayout' must have a C++11 standard layou}} +// expected-note@+1 {{'TemplatedWrapperImplNoStandardLayout' is a named barrier wrapper because it inherits from named barrier wrapper 'CRTPWrapperBase<TemplatedWrapperImpl>'}} +class TemplatedWrapperImplNoStandardLayout : public CRTPWrapperBase<TemplatedWrapperImpl> { + void sayHello() {} + + int k = 0; +}; + static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size"); static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 4, "wrong alignment"); diff --git a/clang/test/SemaHIP/amdgpu-barrier-spirv.hip b/clang/test/SemaHIP/amdgpu-barrier-spirv.hip new file mode 100644 index 0000000000000..9ea23248795f1 --- /dev/null +++ b/clang/test/SemaHIP/amdgpu-barrier-spirv.hip @@ -0,0 +1,20 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 %s -fsyntax-only -fcuda-is-device -std=c++17 -triple spirv64-amd-amdhsa -verify + +#define __device__ __attribute__((device)) + +__amdgpu_named_workgroup_barrier_t bar; // expected-error{{AMDGPU built-in type '__amdgpu_named_workgroup_barrier_t' is not supported on target 'spirv64-amd-amdhsa'}} + +__amdgpu_named_workgroup_barrier_t *getBar() { // expected-error{{AMDGPU built-in type '__amdgpu_named_workgroup_barrier_t *' is not supported on target 'spirv64-amd-amdhsa'}} + return &bar; +} + +void foo(__amdgpu_named_workgroup_barrier_t x) { // expected-error{{AMDGPU built-in type '__amdgpu_named_workgroup_barrier_t' is not supported on target 'spirv64-amd-amdhsa'}} + __amdgpu_named_workgroup_barrier_t v; // expected-error{{AMDGPU built-in type '__amdgpu_named_workgroup_barrier_t' is not supported on target 'spirv64-amd-amdhsa'}} +} + +struct Wrapper { + __amdgpu_named_workgroup_barrier_t x; // expected-error{{AMDGPU built-in type '__amdgpu_named_workgroup_barrier_t' is not supported on target 'spirv64-amd-amdhsa'}} +}; + +Wrapper W; // Decl is invalid so this isn't diagnosed. diff --git a/clang/test/SemaHIP/amdgpu-barrier.hip b/clang/test/SemaHIP/amdgpu-barrier.hip index b3edadcfa57d7..a14ca874a4098 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -16,5 +16,77 @@ __device__ void foo() { void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}} } +using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; + +struct TestSimple { + __amdgpu_named_workgroup_barrier_t x; +}; + +struct TestArray{ + __amdgpu_named_workgroup_barrier_t y[2]; +}; + +struct TestSugared { + SugaredArray z[2]; +}; + +struct TestSimpleWithStaticField { + __amdgpu_named_workgroup_barrier_t x; + static unsigned Harmless; +}; + +// Wrappers cannot have >1 field. +struct WrapperHasTooManyFields { // expected-note {{'WrapperHasTooManyFields' is not a named barrier wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{field with barrier type '__amdgpu_named_workgroup_barrier_t' seen in a structure that is not a named barrier wrapper}} + int other; +}; + +// Wrappers must have standard layout +struct WrapperBase { +__amdgpu_named_workgroup_barrier_t x; +}; + +struct WrapperWithBase : public WrapperBase { +}; + +// expected-error@+2 {{named barrier wrapper 'WrapperWithBaseNoStandardLayout' must have a C++11 standard layou}} +// expected-note@+1 {{'WrapperWithBaseNoStandardLayout' is a named barrier wrapper because it inherits from named barrier wrapper 'WrapperBase'}} +struct WrapperWithBaseNoStandardLayout : public WrapperBase { + unsigned K = 0; +}; + +// Wrappers of Wrappers have the same restrictions. +struct WrapperOfWrapperWithTooManyFields { // expected-note {{'WrapperOfWrapperWithTooManyFields' is not a named barrier wrapper because it has more than one field}} + TestSimple x; // expected-error {{field with barrier type 'TestSimple' seen in a structure that is not a named barrier wrapper}} + int other; +}; + +struct WrapperOfWrapper { + WrapperWithBase y; +}; + +// Check templated cases with a bit of complexity thrown in. +template<typename Derived> +class CRTPWrapperBase { + CRTPWrapperBase() { + static_cast<Derived*>(this)->sayHello(); + } + + WrapperOfWrapper wow; +}; + +class TemplatedWrapperImpl : public CRTPWrapperBase<TemplatedWrapperImpl> { + void sayHello() {} +}; + +// expected-error@+2 {{named barrier wrapper 'TemplatedWrapperImplNoStandardLayout' must have a C++11 standard layou}} +// expected-note@+1 {{'TemplatedWrapperImplNoStandardLayout' is a named barrier wrapper because it inherits from named barrier wrapper 'CRTPWrapperBase<TemplatedWrapperImpl>'}} +class TemplatedWrapperImplNoStandardLayout : public CRTPWrapperBase<TemplatedWrapperImpl> { + void sayHello() {} + + int k = 0; +}; + + static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size"); static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 4, "wrong alignment"); diff --git a/clang/test/SemaOpenCL/amdgpu-barrier.cl b/clang/test/SemaOpenCL/amdgpu-barrier.cl index 6d2c1570f11f0..4ed5b6e0a7094 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -3,6 +3,30 @@ // RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgpu-amd-amdhsa -Wno-unused-value %s void foo() { + typedef __amdgpu_named_workgroup_barrier_t SugaredArray[2]; + + struct TestSimple { + __amdgpu_named_workgroup_barrier_t x; + }; + + struct TestArray { + __amdgpu_named_workgroup_barrier_t y[2]; + }; + + struct TestSugared { + SugaredArray z[2]; + }; + + struct GoodWrapper { + __amdgpu_named_workgroup_barrier_t x; + }; + + // Wrappers cannot have >1 field. + struct WrapperHasTooManyFields { // expected-note {{'WrapperHasTooManyFields' is not a named barrier wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{field with barrier type '__amdgpu_named_workgroup_barrier_t' seen in a structure that is not a named barrier wrapper}} + int other; + }; + int n = 100; __amdgpu_named_workgroup_barrier_t v = 0; // expected-error {{initializing '__private __amdgpu_named_workgroup_barrier_t' with an expression of incompatible type 'int'}} int c = v; // expected-error {{initializing '__private int' with an expression of incompatible type '__private __amdgpu_named_workgroup_barrier_t'}} diff --git a/clang/test/SemaTemplate/address_space-dependent.cpp b/clang/test/SemaTemplate/address_space-dependent.cpp index 3fdccb2c71a76..d6f25923b69b5 100644 --- a/clang/test/SemaTemplate/address_space-dependent.cpp +++ b/clang/test/SemaTemplate/address_space-dependent.cpp @@ -43,7 +43,7 @@ void neg() { template <long int I> void tooBig() { - __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388580)}} + __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388579)}} } template <long int I> @@ -101,7 +101,7 @@ int main() { car<1, 2, 3>(); // expected-note {{in instantiation of function template specialization 'car<1, 2, 3>' requested here}} HasASTemplateFields<1> HASTF; neg<-1>(); // expected-note {{in instantiation of function template specialization 'neg<-1>' requested here}} - correct<0x7FFFE4>(); + correct<0x7FFFE3>(); tooBig<8388650>(); // expected-note {{in instantiation of function template specialization 'tooBig<8388650L>' requested here}} __attribute__((address_space(1))) char *x; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
