From 970aab0a930e38dfd266c01065112602bb274a5e Mon Sep 17 00:00:00 2001 From: Helena Kotas Date: Wed, 25 Sep 2024 15:48:18 -0700 Subject: [PATCH 1/5] [HLSL] Allow resource type attributes only on __hlsl_resource_t --- clang/include/clang/AST/Type.h | 4 +- .../clang/Basic/DiagnosticSemaKinds.td | 1 + clang/include/clang/Sema/SemaHLSL.h | 2 +- clang/lib/AST/ASTContext.cpp | 4 +- clang/lib/Sema/HLSLExternalSemaSource.cpp | 72 +++++++++---------- clang/lib/Sema/SemaHLSL.cpp | 10 ++- clang/lib/Sema/SemaType.cpp | 2 +- clang/test/AST/HLSL/RWBuffer-AST.hlsl | 22 ++---- clang/test/AST/HLSL/StructuredBuffer-AST.hlsl | 22 ++---- .../CodeGenHLSL/buffer-array-operator.hlsl | 3 + .../implicit-norecurse-attrib.hlsl | 2 +- .../hlsl_contained_type_attr_error.hlsl | 4 ++ .../ParserHLSL/hlsl_is_rov_attr_error.hlsl | 4 ++ .../hlsl_raw_buffer_attr_error.hlsl | 4 ++ .../hlsl_resource_class_attr_error.hlsl | 3 + .../hlsl_resource_handle_attrs.hlsl | 6 +- .../Types/Traits/IsIntangibleType.hlsl | 3 + 17 files changed, 85 insertions(+), 83 deletions(-) diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h index dc87b84153e74..67e75652a1664 100644 --- a/clang/include/clang/AST/Type.h +++ b/clang/include/clang/AST/Type.h @@ -6191,7 +6191,9 @@ class HLSLAttributedResourceType : public Type, public llvm::FoldingSetNode { HLSLAttributedResourceType(QualType Canon, QualType Wrapped, QualType Contained, const Attributes &Attrs) - : Type(HLSLAttributedResource, Canon, Wrapped->getDependence()), + : Type(HLSLAttributedResource, Canon, + Contained.isNull() ? TypeDependence::None + : Contained->getDependence()), WrappedType(Wrapped), ContainedType(Contained), Attrs(Attrs) {} public: diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index e4e04bff8b512..a9b2042fcff4e 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12394,6 +12394,7 @@ def err_hlsl_packoffset_alignment_mismatch : Error<"packoffset at 'y' not match def err_hlsl_pointers_unsupported : Error< "%select{pointers|references}0 are unsupported in HLSL">; def err_hlsl_missing_resource_class : Error<"HLSL resource needs to have [[hlsl::resource_class()]] attribute">; +def err_hlsl_attribute_needs_intangible_type: Error<"attribute %0 can be used only on HLSL intangible type %1">; def err_hlsl_operator_unsupported : Error< "the '%select{&|*|->}0' operator is unsupported in HLSL">; diff --git a/clang/include/clang/Sema/SemaHLSL.h b/clang/include/clang/Sema/SemaHLSL.h index e088254c566d3..311cd58bbcac2 100644 --- a/clang/include/clang/Sema/SemaHLSL.h +++ b/clang/include/clang/Sema/SemaHLSL.h @@ -70,7 +70,7 @@ class SemaHLSL : public SemaBase { void handleShaderAttr(Decl *D, const ParsedAttr &AL); void handleResourceBindingAttr(Decl *D, const ParsedAttr &AL); void handleParamModifierAttr(Decl *D, const ParsedAttr &AL); - bool handleResourceTypeAttr(const ParsedAttr &AL); + bool handleResourceTypeAttr(QualType T, const ParsedAttr &AL); bool CheckBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall); QualType ProcessResourceTypeAttributes(QualType Wrapped); diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index fd8aa8de79b49..ced18e1617e24 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -2272,8 +2272,8 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const { #include "clang/Basic/AMDGPUTypes.def" #define HLSL_INTANGIBLE_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/HLSLIntangibleTypes.def" - Width = 0; - Align = 8; + Width = Target->getPointerWidth(LangAS::Default); + Align = Target->getPointerAlign(LangAS::Default); break; } break; diff --git a/clang/lib/Sema/HLSLExternalSemaSource.cpp b/clang/lib/Sema/HLSLExternalSemaSource.cpp index d19f79b6ddefc..e72da9a071fb7 100644 --- a/clang/lib/Sema/HLSLExternalSemaSource.cpp +++ b/clang/lib/Sema/HLSLExternalSemaSource.cpp @@ -117,33 +117,30 @@ struct BuiltinTypeDeclBuilder { if (Record->isCompleteDefinition()) return *this; + ASTContext &Ctx = S.getASTContext(); TypeSourceInfo *ElementTypeInfo = nullptr; - QualType Ty = Record->getASTContext().VoidPtrTy; + QualType ElemTy = Ctx.Char8Ty; if (Template) { if (const auto *TTD = dyn_cast( Template->getTemplateParameters()->getParam(0))) { - Ty = Record->getASTContext().getPointerType( - QualType(TTD->getTypeForDecl(), 0)); - QualType ElemType = QualType(TTD->getTypeForDecl(), 0); - ElementTypeInfo = S.getASTContext().getTrivialTypeSourceInfo( - ElemType, SourceLocation()); + ElemTy = QualType(TTD->getTypeForDecl(), 0); } } + ElementTypeInfo = Ctx.getTrivialTypeSourceInfo(ElemTy, SourceLocation()); // add handle member with resource type attributes QualType AttributedResTy = QualType(); SmallVector Attrs = { - HLSLResourceClassAttr::CreateImplicit(Record->getASTContext(), RC), - IsROV ? HLSLROVAttr::CreateImplicit(Record->getASTContext()) : nullptr, - RawBuffer ? HLSLRawBufferAttr::CreateImplicit(Record->getASTContext()) - : nullptr, - ElementTypeInfo ? HLSLContainedTypeAttr::CreateImplicit( - Record->getASTContext(), ElementTypeInfo) - : nullptr}; - Attr *ResourceAttr = - HLSLResourceAttr::CreateImplicit(Record->getASTContext(), RK); - if (CreateHLSLAttributedResourceType(S, Ty, Attrs, AttributedResTy)) + HLSLResourceClassAttr::CreateImplicit(Ctx, RC), + IsROV ? HLSLROVAttr::CreateImplicit(Ctx) : nullptr, + RawBuffer ? HLSLRawBufferAttr::CreateImplicit(Ctx) : nullptr, + ElementTypeInfo + ? HLSLContainedTypeAttr::CreateImplicit(Ctx, ElementTypeInfo) + : nullptr}; + Attr *ResourceAttr = HLSLResourceAttr::CreateImplicit(Ctx, RK); + if (CreateHLSLAttributedResourceType(S, Ctx.HLSLResourceTy, Attrs, + AttributedResTy)) addMemberVariable("h", AttributedResTy, {ResourceAttr}, Access); return *this; } @@ -242,14 +239,14 @@ struct BuiltinTypeDeclBuilder { assert(Fields.count("h") > 0 && "Subscript operator must be added after the handle."); - FieldDecl *Handle = Fields["h"]; ASTContext &AST = Record->getASTContext(); - - assert(Handle->getType().getCanonicalType() != AST.VoidPtrTy && - "Not yet supported for void pointer handles."); - - QualType ElemTy = - QualType(Handle->getType()->getPointeeOrArrayElementType(), 0); + QualType ElemTy = AST.Char8Ty; + if (Template) { + if (const auto *TTD = dyn_cast( + Template->getTemplateParameters()->getParam(0))) { + ElemTy = QualType(TTD->getTypeForDecl(), 0); + } + } QualType ReturnTy = ElemTy; FunctionProtoType::ExtProtoInfo ExtInfo; @@ -285,22 +282,19 @@ struct BuiltinTypeDeclBuilder { auto FnProtoLoc = TSInfo->getTypeLoc().getAs(); FnProtoLoc.setParam(0, IdxParam); - auto *This = - CXXThisExpr::Create(AST, SourceLocation(), - MethodDecl->getFunctionObjectParameterType(), true); - auto *HandleAccess = MemberExpr::CreateImplicit( - AST, This, false, Handle, Handle->getType(), VK_LValue, OK_Ordinary); - - auto *IndexExpr = DeclRefExpr::Create( - AST, NestedNameSpecifierLoc(), SourceLocation(), IdxParam, false, - DeclarationNameInfo(IdxParam->getDeclName(), SourceLocation()), - AST.UnsignedIntTy, VK_PRValue); - - auto *Array = - new (AST) ArraySubscriptExpr(HandleAccess, IndexExpr, ElemTy, VK_LValue, - OK_Ordinary, SourceLocation()); - - auto *Return = ReturnStmt::Create(AST, SourceLocation(), Array, nullptr); + // FIXME: Placeholder to make sure we return the correct type - create + // field of element_type and return reference to it. This field will go + // away once indexing into resources is properly implemented in + // llvm/llvm-project#95956. + if (Fields.count("e") == 0) { + addMemberVariable("e", ElemTy, {}); + } + FieldDecl *ElemFieldDecl = Fields["e"]; + + auto *This = CXXThisExpr::Create(AST, SourceLocation(), MethodDecl->getFunctionObjectParameterType(),true); + Expr *ElemField = MemberExpr::CreateImplicit(AST, This, false, ElemFieldDecl, + ElemFieldDecl->getType(), VK_LValue,OK_Ordinary); + auto *Return = ReturnStmt::Create(AST, SourceLocation(), ElemField, nullptr); MethodDecl->setBody(CompoundStmt::Create(AST, {Return}, FPOptionsOverride(), SourceLocation(), diff --git a/clang/lib/Sema/SemaHLSL.cpp b/clang/lib/Sema/SemaHLSL.cpp index ebe76185cbb2d..1d8ccdda45573 100644 --- a/clang/lib/Sema/SemaHLSL.cpp +++ b/clang/lib/Sema/SemaHLSL.cpp @@ -693,13 +693,19 @@ bool clang::CreateHLSLAttributedResourceType( // HLSL resource. The attributes are collected in HLSLResourcesTypeAttrs and at // the end of the declaration they are applied to the declaration type by // wrapping it in HLSLAttributedResourceType. -bool SemaHLSL::handleResourceTypeAttr(const ParsedAttr &AL) { - Attr *A = nullptr; +bool SemaHLSL::handleResourceTypeAttr(QualType T, const ParsedAttr &AL) { + // only allow resource type attributes on intangible types + if (!T->isHLSLResourceType()) { + Diag(AL.getLoc(), diag::err_hlsl_attribute_needs_intangible_type) + << AL << getASTContext().HLSLResourceTy; + return false; + } // validate number of arguments if (!AL.checkExactlyNumArgs(SemaRef, AL.getMinArgs())) return false; + Attr *A = nullptr; switch (AL.getKind()) { case ParsedAttr::AT_HLSLResourceClass: { if (!AL.isArgIdent(0)) { diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 950bd6db0359d..a7beb9d222c3b 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -8860,7 +8860,7 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type, // decl-specifier-seq; do not collect attributes on declarations or those // that get to slide after declaration name. if (TAL == TAL_DeclSpec && - state.getSema().HLSL().handleResourceTypeAttr(attr)) + state.getSema().HLSL().handleResourceTypeAttr(type, attr)) attr.setUsedAsTypeAttr(); break; } diff --git a/clang/test/AST/HLSL/RWBuffer-AST.hlsl b/clang/test/AST/HLSL/RWBuffer-AST.hlsl index c3ba520e0f68e..55c0dfa2eaa53 100644 --- a/clang/test/AST/HLSL/RWBuffer-AST.hlsl +++ b/clang/test/AST/HLSL/RWBuffer-AST.hlsl @@ -29,36 +29,26 @@ RWBuffer Buffer; // CHECK-NEXT: CXXRecordDecl 0x{{[0-9A-Fa-f]+}} <> implicit class RWBuffer definition // CHECK: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: implicit h 'element_type * +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(element_type)]] -// CHECK-SAME:':'element_type *' +// CHECK-SAME: ':'__hlsl_resource_t' // CHECK-NEXT: HLSLResourceAttr 0x{{[0-9A-Fa-f]+}} <> Implicit TypedBuffer // CHECK: CXXMethodDecl 0x{{[0-9A-Fa-f]+}} <> operator[] 'element_type &const (unsigned int) const' // CHECK-NEXT: ParmVarDecl 0x{{[0-9A-Fa-f]+}} <> Idx 'unsigned int' // CHECK-NEXT: CompoundStmt 0x{{[0-9A-Fa-f]+}} <> // CHECK-NEXT: ReturnStmt 0x{{[0-9A-Fa-f]+}} <> -// CHECK-NEXT: ArraySubscriptExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type' lvalue -// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type * -// CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] -// CHECK-SAME{LITERAL}: [[hlsl::contained_type(element_type)]] -// CHECK-SAME: ':'element_type *' lvalue .h 0x{{[0-9A-Fa-f]+}} +// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type' lvalue .e 0x{{[0-9A-Fa-f]+}} // CHECK-NEXT: CXXThisExpr 0x{{[0-9A-Fa-f]+}} <> 'const RWBuffer' lvalue implicit this -// CHECK-NEXT: DeclRefExpr 0x{{[0-9A-Fa-f]+}} <> 'unsigned int' ParmVar 0x{{[0-9A-Fa-f]+}} 'Idx' 'unsigned int' // CHECK-NEXT: AlwaysInlineAttr 0x{{[0-9A-Fa-f]+}} <> Implicit always_inline // CHECK-NEXT: CXXMethodDecl 0x{{[0-9A-Fa-f]+}} <> operator[] 'element_type &(unsigned int)' // CHECK-NEXT: ParmVarDecl 0x{{[0-9A-Fa-f]+}} <> Idx 'unsigned int' // CHECK-NEXT: CompoundStmt 0x{{[0-9A-Fa-f]+}} <> // CHECK-NEXT: ReturnStmt 0x{{[0-9A-Fa-f]+}} <> -// CHECK-NEXT: ArraySubscriptExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type' lvalue -// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type * -// CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] -// CHECK-SAME{LITERAL}: [[hlsl::contained_type(element_type)]] -// CHECK-SAME: ':'element_type *' lvalue .h 0x{{[0-9A-Fa-f]+}} +// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type' lvalue .e 0x{{[0-9A-Fa-f]+}} // CHECK-NEXT: CXXThisExpr 0x{{[0-9A-Fa-f]+}} <> 'RWBuffer' lvalue implicit this -// CHECK-NEXT: DeclRefExpr 0x{{[0-9A-Fa-f]+}} <> 'unsigned int' ParmVar 0x{{[0-9A-Fa-f]+}} 'Idx' 'unsigned int' // CHECK-NEXT: AlwaysInlineAttr 0x{{[0-9A-Fa-f]+}} <> Implicit always_inline // CHECK: ClassTemplateSpecializationDecl 0x{{[0-9A-Fa-f]+}} <> class RWBuffer definition @@ -66,8 +56,8 @@ RWBuffer Buffer; // CHECK: TemplateArgument type 'float' // CHECK-NEXT: BuiltinType 0x{{[0-9A-Fa-f]+}} 'float' // CHECK-NEXT: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit referenced h 'float * +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]] -// CHECK-SAME: ':'float *' +// CHECK-SAME: ':'__hlsl_resource_t' // CHECK-NEXT: HLSLResourceAttr 0x{{[0-9A-Fa-f]+}} <> Implicit TypedBuffer diff --git a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl index 1a3deba5830fa..cdc7cc8de07c6 100644 --- a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl +++ b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl @@ -30,39 +30,27 @@ StructuredBuffer Buffer; // CHECK-NEXT: CXXRecordDecl 0x{{[0-9A-Fa-f]+}} <> implicit class StructuredBuffer definition // CHECK: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit h 'element_type * +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit referenced h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::raw_buffer]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(element_type)]] -// CHECK-SAME: ':'element_type *' +// CHECK-SAME: ':'__hlsl_resource_t' // CHECK-NEXT: HLSLResourceAttr 0x{{[0-9A-Fa-f]+}} <> Implicit TypedBuffer // CHECK: CXXMethodDecl 0x{{[0-9A-Fa-f]+}} <> operator[] 'element_type &const (unsigned int) const' // CHECK-NEXT: ParmVarDecl 0x{{[0-9A-Fa-f]+}} <> Idx 'unsigned int' // CHECK-NEXT: CompoundStmt 0x{{[0-9A-Fa-f]+}} <> // CHECK-NEXT: ReturnStmt 0x{{[0-9A-Fa-f]+}} <> -// CHECK-NEXT: ArraySubscriptExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type' lvalue -// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type * -// CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] -// CHECK-SAME{LITERAL}: [[hlsl::raw_buffer]] -// CHECK-SAME{LITERAL}: [[hlsl::contained_type(element_type)]] -// CHECK-SAME: ':'element_type *' lvalue .h 0x{{[0-9A-Fa-f]+}} +// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type' lvalue .e 0x{{[0-9A-Fa-f]+}} // CHECK-NEXT: CXXThisExpr 0x{{[0-9A-Fa-f]+}} <> 'const StructuredBuffer' lvalue implicit this -// CHECK-NEXT: DeclRefExpr 0x{{[0-9A-Fa-f]+}} <> 'unsigned int' ParmVar 0x{{[0-9A-Fa-f]+}} 'Idx' 'unsigned int' // CHECK-NEXT: AlwaysInlineAttr 0x{{[0-9A-Fa-f]+}} <> Implicit always_inline // CHECK-NEXT: CXXMethodDecl 0x{{[0-9A-Fa-f]+}} <> operator[] 'element_type &(unsigned int)' // CHECK-NEXT: ParmVarDecl 0x{{[0-9A-Fa-f]+}} <> Idx 'unsigned int' // CHECK-NEXT: CompoundStmt 0x{{[0-9A-Fa-f]+}} <> // CHECK-NEXT: ReturnStmt 0x{{[0-9A-Fa-f]+}} <> -// CHECK-NEXT: ArraySubscriptExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type' lvalue -// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type * -// CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] -// CHECK-SAME{LITERAL}: [[hlsl::raw_buffer]] -// CHECK-SAME{LITERAL}: [[hlsl::contained_type(element_type)]] -// CHECK-SAME: ':'element_type *' lvalue .h 0x{{[0-9A-Fa-f]+}} +// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <> 'element_type' lvalue .e 0x{{[0-9A-Fa-f]+}} // CHECK-NEXT: CXXThisExpr 0x{{[0-9A-Fa-f]+}} <> 'StructuredBuffer' lvalue implicit this -// CHECK-NEXT: DeclRefExpr 0x{{[0-9A-Fa-f]+}} <> 'unsigned int' ParmVar 0x{{[0-9A-Fa-f]+}} 'Idx' 'unsigned int' // CHECK-NEXT: AlwaysInlineAttr 0x{{[0-9A-Fa-f]+}} <> Implicit always_inline // CHECK: ClassTemplateSpecializationDecl 0x{{[0-9A-Fa-f]+}} <> class StructuredBuffer definition @@ -74,5 +62,5 @@ StructuredBuffer Buffer; // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::raw_buffer]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]] -// CHECK-SAME: ':'float *' +// CHECK-SAME: ':'__hlsl_resource_t' // CHECK-NEXT: HLSLResourceAttr 0x{{[0-9A-Fa-f]+}} <> Implicit TypedBuffer diff --git a/clang/test/CodeGenHLSL/buffer-array-operator.hlsl b/clang/test/CodeGenHLSL/buffer-array-operator.hlsl index 02e570ebdcb4f..f65cdbb43e27b 100644 --- a/clang/test/CodeGenHLSL/buffer-array-operator.hlsl +++ b/clang/test/CodeGenHLSL/buffer-array-operator.hlsl @@ -1,5 +1,8 @@ // RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s +// XFAIL: * +// Resource indexing will be properly implemented in llvm/llvm-project#95956 + const RWBuffer In; RWBuffer Out; diff --git a/clang/test/CodeGenHLSL/implicit-norecurse-attrib.hlsl b/clang/test/CodeGenHLSL/implicit-norecurse-attrib.hlsl index ae3a3b5f90199..f72fe059cb576 100644 --- a/clang/test/CodeGenHLSL/implicit-norecurse-attrib.hlsl +++ b/clang/test/CodeGenHLSL/implicit-norecurse-attrib.hlsl @@ -31,7 +31,7 @@ uint Find(Node SortedTree[MAX], uint key) { } // CHECK: Function Attrs:{{.*}}norecurse -// CHECK: define noundef i1 @"?InitTree@@YA_NY0GE@UNode@@V?$RWBuffer@T?$__vector@I$03@__clang@@@hlsl@@I@Z"(ptr noundef byval([100 x %struct.Node]) align 4 %tree, ptr noundef byval(%"class.hlsl::RWBuffer") align 4 %encodedTree, i32 noundef %maxDepth) [[ExtAttr:\#[0-9]+]] +// CHECK: define noundef i1 @"?InitTree@@YA_NY0GE@UNode@@V?$RWBuffer@T?$__vector@I$03@__clang@@@hlsl@@I@Z"(ptr noundef byval([100 x %struct.Node]) align 4 %tree, ptr noundef byval(%"class.hlsl::RWBuffer") align 16 %encodedTree, i32 noundef %maxDepth) [[ExtAttr:\#[0-9]+]] // CHECK: ret i1 // Initialize tree with given buffer // Imagine the inout works diff --git a/clang/test/ParserHLSL/hlsl_contained_type_attr_error.hlsl b/clang/test/ParserHLSL/hlsl_contained_type_attr_error.hlsl index 1c37d72de8614..b2d492d95945c 100644 --- a/clang/test/ParserHLSL/hlsl_contained_type_attr_error.hlsl +++ b/clang/test/ParserHLSL/hlsl_contained_type_attr_error.hlsl @@ -22,3 +22,7 @@ __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::contained_type(float)]] // expected-warning@+1{{attribute 'contained_type' is already applied with different arguments}} __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::contained_type(float)]] [[hlsl::contained_type(int)]] h8; + +// expected-error@+2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}} +// expected-error@+1{{attribute 'contained_type' can be used only on HLSL intangible type '__hlsl_resource_t'}} +float [[hlsl::resource_class(UAV)]] [[hlsl::contained_type(float)]] res5; diff --git a/clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl b/clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl index 15685bd1a3baa..3b2c12e7a96c5 100644 --- a/clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl +++ b/clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl @@ -14,3 +14,7 @@ __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::is_rov(gibberish)]] res3 // expected-warning@+1{{attribute 'is_rov' is already applied}} __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::is_rov]] [[hlsl::is_rov]] res4; + +// expected-error@+2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}} +// expected-error@+1{{attribute 'is_rov' can be used only on HLSL intangible type '__hlsl_resource_t'}} +float [[hlsl::resource_class(UAV)]] [[hlsl::is_rov]] res5; diff --git a/clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl b/clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl index 83273426017ed..77530cbf9e4d9 100644 --- a/clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl +++ b/clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl @@ -11,3 +11,7 @@ __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::raw_buffer(gibberish)]] // expected-warning@+1{{attribute 'raw_buffer' is already applied}} __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::raw_buffer]] [[hlsl::raw_buffer]] res4; + +// expected-error@+2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}} +// expected-error@+1{{attribute 'raw_buffer' can be used only on HLSL intangible type '__hlsl_resource_t'}} +float [[hlsl::resource_class(UAV)]] [[hlsl::raw_buffer]] res5; diff --git a/clang/test/ParserHLSL/hlsl_resource_class_attr_error.hlsl b/clang/test/ParserHLSL/hlsl_resource_class_attr_error.hlsl index 01ff1c007e2b5..63e39daff949b 100644 --- a/clang/test/ParserHLSL/hlsl_resource_class_attr_error.hlsl +++ b/clang/test/ParserHLSL/hlsl_resource_class_attr_error.hlsl @@ -17,3 +17,6 @@ __hlsl_resource_t [[hlsl::resource_class(SRV)]] [[hlsl::resource_class(SRV)]] e4 // expected-error@+1{{'resource_class' attribute takes one argument}} __hlsl_resource_t [[hlsl::resource_class(SRV, "aa")]] e5; + +// expected-error@+1{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}} +float [[hlsl::resource_class(UAV)]] e6; diff --git a/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl b/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl index 301d61c0e906e..b73c325fae071 100644 --- a/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl +++ b/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl @@ -3,10 +3,10 @@ // CHECK: -ClassTemplateSpecializationDecl 0x{{[0-9a-f]+}} <> class RWBuffer definition implicit_instantiation // CHECK: -TemplateArgument type 'float' // CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float' -// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit referenced h 'float * +// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit referenced h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]] -// CHECK-SAME: ':'float *' +// CHECK-SAME: ':'__hlsl_resource_t' // CHECK: -HLSLResourceAttr 0x{{[0-9a-f]+}} <> Implicit TypedBuffer RWBuffer Buffer1; @@ -18,6 +18,6 @@ RWBuffer Buffer1; // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)] // CHECK-SAME{LITERAL}: [[hlsl::is_rov]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(vector)]] -// CHECK-SAME: ':'vector' +// CHECK-SAME: ':'__hlsl_resource_t' // CHECK: -HLSLResourceAttr 0x{{[0-9a-f]+}} <> Implicit TypedBuffer RasterizerOrderedBuffer > BufferArray3[4]; diff --git a/clang/test/SemaHLSL/Types/Traits/IsIntangibleType.hlsl b/clang/test/SemaHLSL/Types/Traits/IsIntangibleType.hlsl index 92cba1dcd4bdf..8c0f8d6f271db 100644 --- a/clang/test/SemaHLSL/Types/Traits/IsIntangibleType.hlsl +++ b/clang/test/SemaHLSL/Types/Traits/IsIntangibleType.hlsl @@ -76,3 +76,6 @@ template struct SimpleTemplate { }; _Static_assert(__builtin_hlsl_is_intangible(SimpleTemplate<__hlsl_resource_t>), ""); _Static_assert(!__builtin_hlsl_is_intangible(SimpleTemplate), ""); + +_Static_assert(__builtin_hlsl_is_intangible(RWBuffer), ""); +_Static_assert(__builtin_hlsl_is_intangible(StructuredBuffer), ""); From 71c4a9361f07a606c95eb5232be50dcdff31c422 Mon Sep 17 00:00:00 2001 From: Helena Kotas Date: Tue, 24 Sep 2024 22:33:40 -0700 Subject: [PATCH 2/5] [HLSL] Remove `__builtin_hlsl_create_handle` The builtin `__builtin_hlsl_create_handle` was used in the constructor of resource buffer classes to initialize resource handles based on resource type and registry binding information. This is not possible because the registry binding information is not accessible from the constructor codegen. The resource handle initialization will be implemented later on as part 2/2 of llvm/llvm-project#105076 once the dependent tasks are completed. Part 1/2 of llvm/llvm-project#105076. (cherry picked from commit 9c0303e95a9e1db24c7b33a2f93b2759e54f5b31) --- clang/include/clang/Basic/Builtins.td | 6 ---- clang/lib/Sema/HLSLExternalSemaSource.cpp | 32 ++----------------- clang/test/AST/HLSL/StructuredBuffer-AST.hlsl | 4 +-- .../builtins/RWBuffer-constructor.hlsl | 8 ++++- .../StructuredBuffer-constructor.hlsl | 7 ++++ .../CodeGenHLSL/builtins/create_handle.hlsl | 7 ---- .../hlsl_resource_handle_attrs.hlsl | 4 +-- llvm/include/llvm/IR/IntrinsicsDirectX.td | 3 -- llvm/include/llvm/IR/IntrinsicsSPIRV.td | 2 -- llvm/unittests/IR/IntrinsicsTest.cpp | 1 - 10 files changed, 20 insertions(+), 54 deletions(-) delete mode 100644 clang/test/CodeGenHLSL/builtins/create_handle.hlsl diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td index 8c5d7ad763bf9..33791270800c9 100644 --- a/clang/include/clang/Basic/Builtins.td +++ b/clang/include/clang/Basic/Builtins.td @@ -4703,12 +4703,6 @@ def HLSLClamp : LangBuiltin<"HLSL_LANG"> { let Prototype = "void(...)"; } -def HLSLCreateHandle : LangBuiltin<"HLSL_LANG"> { - let Spellings = ["__builtin_hlsl_create_handle"]; - let Attributes = [NoThrow, Const]; - let Prototype = "void*(unsigned char)"; -} - def HLSLDotProduct : LangBuiltin<"HLSL_LANG"> { let Spellings = ["__builtin_hlsl_dot"]; let Attributes = [NoThrow, Const]; diff --git a/clang/lib/Sema/HLSLExternalSemaSource.cpp b/clang/lib/Sema/HLSLExternalSemaSource.cpp index e72da9a071fb7..b6bd8e778a617 100644 --- a/clang/lib/Sema/HLSLExternalSemaSource.cpp +++ b/clang/lib/Sema/HLSLExternalSemaSource.cpp @@ -190,36 +190,8 @@ struct BuiltinTypeDeclBuilder { ExplicitSpecifier(), false, true, false, ConstexprSpecKind::Unspecified); - DeclRefExpr *Fn = - lookupBuiltinFunction(AST, S, "__builtin_hlsl_create_handle"); - Expr *RCExpr = emitResourceClassExpr(AST, RC); - Expr *Call = CallExpr::Create(AST, Fn, {RCExpr}, AST.VoidPtrTy, VK_PRValue, - SourceLocation(), FPOptionsOverride()); - - CXXThisExpr *This = CXXThisExpr::Create( - AST, SourceLocation(), Constructor->getFunctionObjectParameterType(), - true); - Expr *Handle = MemberExpr::CreateImplicit(AST, This, false, Fields["h"], - Fields["h"]->getType(), VK_LValue, - OK_Ordinary); - - // If the handle isn't a void pointer, cast the builtin result to the - // correct type. - if (Handle->getType().getCanonicalType() != AST.VoidPtrTy) { - Call = CXXStaticCastExpr::Create( - AST, Handle->getType(), VK_PRValue, CK_Dependent, Call, nullptr, - AST.getTrivialTypeSourceInfo(Handle->getType(), SourceLocation()), - FPOptionsOverride(), SourceLocation(), SourceLocation(), - SourceRange()); - } - - BinaryOperator *Assign = BinaryOperator::Create( - AST, Handle, Call, BO_Assign, Handle->getType(), VK_LValue, OK_Ordinary, - SourceLocation(), FPOptionsOverride()); - - Constructor->setBody( - CompoundStmt::Create(AST, {Assign}, FPOptionsOverride(), - SourceLocation(), SourceLocation())); + Constructor->setBody(CompoundStmt::Create( + AST, {}, FPOptionsOverride(), SourceLocation(), SourceLocation())); Constructor->setAccess(AccessSpecifier::AS_public); Record->addDecl(Constructor); return *this; diff --git a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl index cdc7cc8de07c6..b31db8ce59f22 100644 --- a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl +++ b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl @@ -30,7 +30,7 @@ StructuredBuffer Buffer; // CHECK-NEXT: CXXRecordDecl 0x{{[0-9A-Fa-f]+}} <> implicit class StructuredBuffer definition // CHECK: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit referenced h '__hlsl_resource_t +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::raw_buffer]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(element_type)]] @@ -58,7 +58,7 @@ StructuredBuffer Buffer; // CHECK: TemplateArgument type 'float' // CHECK-NEXT: BuiltinType 0x{{[0-9A-Fa-f]+}} 'float' // CHECK-NEXT: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit referenced h 'float * +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::raw_buffer]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]] diff --git a/clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl b/clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl index 174f4c3eaaad2..19699dcf14d9f 100644 --- a/clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl +++ b/clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl @@ -1,6 +1,12 @@ // RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s // RUN: %clang_cc1 -triple spirv-vulkan-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s --check-prefix=CHECK-SPIRV +// XFAIL: * +// This expectedly fails because create.handle is no longer called +// from RWBuffer constructor and the replacement has not been +// implemented yet. This test should be updated to expect +// dx.create.handleFromBinding as part of issue #105076. + RWBuffer Buf; // CHECK: define linkonce_odr noundef ptr @"??0?$RWBuffer@M@hlsl@@QAA@XZ" @@ -10,4 +16,4 @@ RWBuffer Buf; // CHECK: store ptr %[[HandleRes]], ptr %h, align 4 // CHECK-SPIRV: %[[HandleRes:[0-9]+]] = call ptr @llvm.spv.create.handle(i8 1) -// CHECK-SPIRV: store ptr %[[HandleRes]], ptr %h, align 8 \ No newline at end of file +// CHECK-SPIRV: store ptr %[[HandleRes]], ptr %h, align 8 diff --git a/clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl b/clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl index 34019e5b18693..178332d03e640 100644 --- a/clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl +++ b/clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl @@ -1,5 +1,12 @@ +// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s // RUN: %clang_cc1 -triple spirv-vulkan-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s --check-prefix=CHECK-SPIRV +// XFAIL: * +// This expectedly fails because create.handle is no longer invoked +// from StructuredBuffer constructor and the replacement has not been +// implemented yet. This test should be updated to expect +// dx.create.handleFromBinding as part of issue #105076. + StructuredBuffer Buf; // CHECK: define linkonce_odr noundef ptr @"??0?$StructuredBuffer@M@hlsl@@QAA@XZ" diff --git a/clang/test/CodeGenHLSL/builtins/create_handle.hlsl b/clang/test/CodeGenHLSL/builtins/create_handle.hlsl deleted file mode 100644 index 61226c2b54e72..0000000000000 --- a/clang/test/CodeGenHLSL/builtins/create_handle.hlsl +++ /dev/null @@ -1,7 +0,0 @@ -// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s - -void fn() { - (void)__builtin_hlsl_create_handle(0); -} - -// CHECK: call ptr @llvm.dx.create.handle(i8 0) diff --git a/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl b/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl index b73c325fae071..e7d19c3da7216 100644 --- a/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl +++ b/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl @@ -3,7 +3,7 @@ // CHECK: -ClassTemplateSpecializationDecl 0x{{[0-9a-f]+}} <> class RWBuffer definition implicit_instantiation // CHECK: -TemplateArgument type 'float' // CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float' -// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit referenced h '__hlsl_resource_t +// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]] // CHECK-SAME: ':'__hlsl_resource_t' @@ -14,7 +14,7 @@ RWBuffer Buffer1; // CHECK: -TemplateArgument type 'vector' // CHECK: `-ExtVectorType 0x{{[0-9a-f]+}} 'vector' 4 // CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float' -// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit referenced h 'vector +// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)] // CHECK-SAME{LITERAL}: [[hlsl::is_rov]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(vector)]] diff --git a/llvm/include/llvm/IR/IntrinsicsDirectX.td b/llvm/include/llvm/IR/IntrinsicsDirectX.td index 3ce7b8b987ef8..555877e7aaf0e 100644 --- a/llvm/include/llvm/IR/IntrinsicsDirectX.td +++ b/llvm/include/llvm/IR/IntrinsicsDirectX.td @@ -17,9 +17,6 @@ def int_dx_group_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWi def int_dx_thread_id_in_group : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>; def int_dx_flattened_thread_id_in_group : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrWillReturn]>; -def int_dx_create_handle : ClangBuiltin<"__builtin_hlsl_create_handle">, - Intrinsic<[ llvm_ptr_ty ], [llvm_i8_ty], [IntrWillReturn]>; - // Create resource handle given binding information. Returns a `target("dx.")` // type appropriate for the kind of resource given a register space ID, lower // bound and range size of the binding, as well as an index and an indicator diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td index 7ac479f31386f..a9dbddb13adaa 100644 --- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td +++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td @@ -58,8 +58,6 @@ let TargetPrefix = "spv" in { // The following intrinsic(s) are mirrored from IntrinsicsDirectX.td for HLSL support. def int_spv_thread_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>; - def int_spv_create_handle : ClangBuiltin<"__builtin_hlsl_create_handle">, - Intrinsic<[ llvm_ptr_ty ], [llvm_i8_ty], [IntrWillReturn]>; def int_spv_all : DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_any_ty]>; def int_spv_any : DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_any_ty]>; def int_spv_frac : DefaultAttrsIntrinsic<[LLVMMatchType<0>], [llvm_anyfloat_ty]>; diff --git a/llvm/unittests/IR/IntrinsicsTest.cpp b/llvm/unittests/IR/IntrinsicsTest.cpp index a92ffe3cdeb7e..0c4af28a2ab57 100644 --- a/llvm/unittests/IR/IntrinsicsTest.cpp +++ b/llvm/unittests/IR/IntrinsicsTest.cpp @@ -94,7 +94,6 @@ TEST(IntrinsicNameLookup, ClangBuiltinLookup) { {"__builtin_amdgcn_workgroup_id_z", "amdgcn", amdgcn_workgroup_id_z}, {"__builtin_arm_cdp", "arm", arm_cdp}, {"__builtin_bpf_preserve_type_info", "bpf", bpf_preserve_type_info}, - {"__builtin_hlsl_create_handle", "dx", dx_create_handle}, {"__builtin_HEXAGON_A2_tfr", "hexagon", hexagon_A2_tfr}, {"__builtin_lasx_xbz_w", "loongarch", loongarch_lasx_xbz_w}, {"__builtin_mips_bitrev", "mips", mips_bitrev}, From aabec4b4eab91c49d709be04c55a3d83c4596c67 Mon Sep 17 00:00:00 2001 From: Helena Kotas Date: Wed, 25 Sep 2024 21:08:58 -0700 Subject: [PATCH 3/5] clang-format --- clang/lib/Sema/HLSLExternalSemaSource.cpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/HLSLExternalSemaSource.cpp b/clang/lib/Sema/HLSLExternalSemaSource.cpp index b6bd8e778a617..2913d16fca482 100644 --- a/clang/lib/Sema/HLSLExternalSemaSource.cpp +++ b/clang/lib/Sema/HLSLExternalSemaSource.cpp @@ -256,17 +256,21 @@ struct BuiltinTypeDeclBuilder { // FIXME: Placeholder to make sure we return the correct type - create // field of element_type and return reference to it. This field will go - // away once indexing into resources is properly implemented in + // away once indexing into resources is properly implemented in // llvm/llvm-project#95956. if (Fields.count("e") == 0) { addMemberVariable("e", ElemTy, {}); } FieldDecl *ElemFieldDecl = Fields["e"]; - - auto *This = CXXThisExpr::Create(AST, SourceLocation(), MethodDecl->getFunctionObjectParameterType(),true); - Expr *ElemField = MemberExpr::CreateImplicit(AST, This, false, ElemFieldDecl, - ElemFieldDecl->getType(), VK_LValue,OK_Ordinary); - auto *Return = ReturnStmt::Create(AST, SourceLocation(), ElemField, nullptr); + + auto *This = + CXXThisExpr::Create(AST, SourceLocation(), + MethodDecl->getFunctionObjectParameterType(), true); + Expr *ElemField = MemberExpr::CreateImplicit( + AST, This, false, ElemFieldDecl, ElemFieldDecl->getType(), VK_LValue, + OK_Ordinary); + auto *Return = + ReturnStmt::Create(AST, SourceLocation(), ElemField, nullptr); MethodDecl->setBody(CompoundStmt::Create(AST, {Return}, FPOptionsOverride(), SourceLocation(), From 0c17b5a2222a1b164f7eb37c78c69acc9d3ecc0a Mon Sep 17 00:00:00 2001 From: Helena Kotas Date: Thu, 26 Sep 2024 08:40:29 -0700 Subject: [PATCH 4/5] Revert "[HLSL] Remove `__builtin_hlsl_create_handle`" This reverts commit 71c4a9361f07a606c95eb5232be50dcdff31c422. --- clang/include/clang/Basic/Builtins.td | 6 ++++ clang/lib/Sema/HLSLExternalSemaSource.cpp | 32 +++++++++++++++++-- clang/test/AST/HLSL/StructuredBuffer-AST.hlsl | 4 +-- .../builtins/RWBuffer-constructor.hlsl | 8 +---- .../StructuredBuffer-constructor.hlsl | 7 ---- .../CodeGenHLSL/builtins/create_handle.hlsl | 7 ++++ .../hlsl_resource_handle_attrs.hlsl | 4 +-- llvm/include/llvm/IR/IntrinsicsDirectX.td | 3 ++ llvm/include/llvm/IR/IntrinsicsSPIRV.td | 2 ++ llvm/unittests/IR/IntrinsicsTest.cpp | 1 + 10 files changed, 54 insertions(+), 20 deletions(-) create mode 100644 clang/test/CodeGenHLSL/builtins/create_handle.hlsl diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td index 33791270800c9..8c5d7ad763bf9 100644 --- a/clang/include/clang/Basic/Builtins.td +++ b/clang/include/clang/Basic/Builtins.td @@ -4703,6 +4703,12 @@ def HLSLClamp : LangBuiltin<"HLSL_LANG"> { let Prototype = "void(...)"; } +def HLSLCreateHandle : LangBuiltin<"HLSL_LANG"> { + let Spellings = ["__builtin_hlsl_create_handle"]; + let Attributes = [NoThrow, Const]; + let Prototype = "void*(unsigned char)"; +} + def HLSLDotProduct : LangBuiltin<"HLSL_LANG"> { let Spellings = ["__builtin_hlsl_dot"]; let Attributes = [NoThrow, Const]; diff --git a/clang/lib/Sema/HLSLExternalSemaSource.cpp b/clang/lib/Sema/HLSLExternalSemaSource.cpp index 2913d16fca482..42d3358e4ddbe 100644 --- a/clang/lib/Sema/HLSLExternalSemaSource.cpp +++ b/clang/lib/Sema/HLSLExternalSemaSource.cpp @@ -190,8 +190,36 @@ struct BuiltinTypeDeclBuilder { ExplicitSpecifier(), false, true, false, ConstexprSpecKind::Unspecified); - Constructor->setBody(CompoundStmt::Create( - AST, {}, FPOptionsOverride(), SourceLocation(), SourceLocation())); + DeclRefExpr *Fn = + lookupBuiltinFunction(AST, S, "__builtin_hlsl_create_handle"); + Expr *RCExpr = emitResourceClassExpr(AST, RC); + Expr *Call = CallExpr::Create(AST, Fn, {RCExpr}, AST.VoidPtrTy, VK_PRValue, + SourceLocation(), FPOptionsOverride()); + + CXXThisExpr *This = CXXThisExpr::Create( + AST, SourceLocation(), Constructor->getFunctionObjectParameterType(), + true); + Expr *Handle = MemberExpr::CreateImplicit(AST, This, false, Fields["h"], + Fields["h"]->getType(), VK_LValue, + OK_Ordinary); + + // If the handle isn't a void pointer, cast the builtin result to the + // correct type. + if (Handle->getType().getCanonicalType() != AST.VoidPtrTy) { + Call = CXXStaticCastExpr::Create( + AST, Handle->getType(), VK_PRValue, CK_Dependent, Call, nullptr, + AST.getTrivialTypeSourceInfo(Handle->getType(), SourceLocation()), + FPOptionsOverride(), SourceLocation(), SourceLocation(), + SourceRange()); + } + + BinaryOperator *Assign = BinaryOperator::Create( + AST, Handle, Call, BO_Assign, Handle->getType(), VK_LValue, OK_Ordinary, + SourceLocation(), FPOptionsOverride()); + + Constructor->setBody( + CompoundStmt::Create(AST, {Assign}, FPOptionsOverride(), + SourceLocation(), SourceLocation())); Constructor->setAccess(AccessSpecifier::AS_public); Record->addDecl(Constructor); return *this; diff --git a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl index b31db8ce59f22..cdc7cc8de07c6 100644 --- a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl +++ b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl @@ -30,7 +30,7 @@ StructuredBuffer Buffer; // CHECK-NEXT: CXXRecordDecl 0x{{[0-9A-Fa-f]+}} <> implicit class StructuredBuffer definition // CHECK: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit h '__hlsl_resource_t +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit referenced h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::raw_buffer]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(element_type)]] @@ -58,7 +58,7 @@ StructuredBuffer Buffer; // CHECK: TemplateArgument type 'float' // CHECK-NEXT: BuiltinType 0x{{[0-9A-Fa-f]+}} 'float' // CHECK-NEXT: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit h '__hlsl_resource_t +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit referenced h 'float * // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::raw_buffer]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]] diff --git a/clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl b/clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl index 19699dcf14d9f..174f4c3eaaad2 100644 --- a/clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl +++ b/clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl @@ -1,12 +1,6 @@ // RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s // RUN: %clang_cc1 -triple spirv-vulkan-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s --check-prefix=CHECK-SPIRV -// XFAIL: * -// This expectedly fails because create.handle is no longer called -// from RWBuffer constructor and the replacement has not been -// implemented yet. This test should be updated to expect -// dx.create.handleFromBinding as part of issue #105076. - RWBuffer Buf; // CHECK: define linkonce_odr noundef ptr @"??0?$RWBuffer@M@hlsl@@QAA@XZ" @@ -16,4 +10,4 @@ RWBuffer Buf; // CHECK: store ptr %[[HandleRes]], ptr %h, align 4 // CHECK-SPIRV: %[[HandleRes:[0-9]+]] = call ptr @llvm.spv.create.handle(i8 1) -// CHECK-SPIRV: store ptr %[[HandleRes]], ptr %h, align 8 +// CHECK-SPIRV: store ptr %[[HandleRes]], ptr %h, align 8 \ No newline at end of file diff --git a/clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl b/clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl index 178332d03e640..34019e5b18693 100644 --- a/clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl +++ b/clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl @@ -1,12 +1,5 @@ -// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s // RUN: %clang_cc1 -triple spirv-vulkan-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s --check-prefix=CHECK-SPIRV -// XFAIL: * -// This expectedly fails because create.handle is no longer invoked -// from StructuredBuffer constructor and the replacement has not been -// implemented yet. This test should be updated to expect -// dx.create.handleFromBinding as part of issue #105076. - StructuredBuffer Buf; // CHECK: define linkonce_odr noundef ptr @"??0?$StructuredBuffer@M@hlsl@@QAA@XZ" diff --git a/clang/test/CodeGenHLSL/builtins/create_handle.hlsl b/clang/test/CodeGenHLSL/builtins/create_handle.hlsl new file mode 100644 index 0000000000000..61226c2b54e72 --- /dev/null +++ b/clang/test/CodeGenHLSL/builtins/create_handle.hlsl @@ -0,0 +1,7 @@ +// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s + +void fn() { + (void)__builtin_hlsl_create_handle(0); +} + +// CHECK: call ptr @llvm.dx.create.handle(i8 0) diff --git a/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl b/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl index e7d19c3da7216..b73c325fae071 100644 --- a/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl +++ b/clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl @@ -3,7 +3,7 @@ // CHECK: -ClassTemplateSpecializationDecl 0x{{[0-9a-f]+}} <> class RWBuffer definition implicit_instantiation // CHECK: -TemplateArgument type 'float' // CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float' -// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit h '__hlsl_resource_t +// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit referenced h '__hlsl_resource_t // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]] // CHECK-SAME: ':'__hlsl_resource_t' @@ -14,7 +14,7 @@ RWBuffer Buffer1; // CHECK: -TemplateArgument type 'vector' // CHECK: `-ExtVectorType 0x{{[0-9a-f]+}} 'vector' 4 // CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float' -// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit h '__hlsl_resource_t +// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit referenced h 'vector // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)] // CHECK-SAME{LITERAL}: [[hlsl::is_rov]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(vector)]] diff --git a/llvm/include/llvm/IR/IntrinsicsDirectX.td b/llvm/include/llvm/IR/IntrinsicsDirectX.td index 555877e7aaf0e..3ce7b8b987ef8 100644 --- a/llvm/include/llvm/IR/IntrinsicsDirectX.td +++ b/llvm/include/llvm/IR/IntrinsicsDirectX.td @@ -17,6 +17,9 @@ def int_dx_group_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWi def int_dx_thread_id_in_group : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>; def int_dx_flattened_thread_id_in_group : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrWillReturn]>; +def int_dx_create_handle : ClangBuiltin<"__builtin_hlsl_create_handle">, + Intrinsic<[ llvm_ptr_ty ], [llvm_i8_ty], [IntrWillReturn]>; + // Create resource handle given binding information. Returns a `target("dx.")` // type appropriate for the kind of resource given a register space ID, lower // bound and range size of the binding, as well as an index and an indicator diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td index a9dbddb13adaa..7ac479f31386f 100644 --- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td +++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td @@ -58,6 +58,8 @@ let TargetPrefix = "spv" in { // The following intrinsic(s) are mirrored from IntrinsicsDirectX.td for HLSL support. def int_spv_thread_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>; + def int_spv_create_handle : ClangBuiltin<"__builtin_hlsl_create_handle">, + Intrinsic<[ llvm_ptr_ty ], [llvm_i8_ty], [IntrWillReturn]>; def int_spv_all : DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_any_ty]>; def int_spv_any : DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_any_ty]>; def int_spv_frac : DefaultAttrsIntrinsic<[LLVMMatchType<0>], [llvm_anyfloat_ty]>; diff --git a/llvm/unittests/IR/IntrinsicsTest.cpp b/llvm/unittests/IR/IntrinsicsTest.cpp index 0c4af28a2ab57..a92ffe3cdeb7e 100644 --- a/llvm/unittests/IR/IntrinsicsTest.cpp +++ b/llvm/unittests/IR/IntrinsicsTest.cpp @@ -94,6 +94,7 @@ TEST(IntrinsicNameLookup, ClangBuiltinLookup) { {"__builtin_amdgcn_workgroup_id_z", "amdgcn", amdgcn_workgroup_id_z}, {"__builtin_arm_cdp", "arm", arm_cdp}, {"__builtin_bpf_preserve_type_info", "bpf", bpf_preserve_type_info}, + {"__builtin_hlsl_create_handle", "dx", dx_create_handle}, {"__builtin_HEXAGON_A2_tfr", "hexagon", hexagon_A2_tfr}, {"__builtin_lasx_xbz_w", "loongarch", loongarch_lasx_xbz_w}, {"__builtin_mips_bitrev", "mips", mips_bitrev}, From f0712844ef068a803e440893c9a88d72647bfbf4 Mon Sep 17 00:00:00 2001 From: Helena Kotas Date: Thu, 26 Sep 2024 09:24:46 -0700 Subject: [PATCH 5/5] Remove change that validates resource attributes are used only on __hlsl_resource_t It will go into a separate PR. --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 1 - clang/include/clang/Sema/SemaHLSL.h | 2 +- clang/lib/Sema/SemaHLSL.cpp | 10 ++-------- clang/lib/Sema/SemaType.cpp | 2 +- .../ParserHLSL/hlsl_contained_type_attr_error.hlsl | 4 ---- clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl | 4 ---- clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl | 4 ---- .../ParserHLSL/hlsl_resource_class_attr_error.hlsl | 3 --- 8 files changed, 4 insertions(+), 26 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a9b2042fcff4e..e4e04bff8b512 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12394,7 +12394,6 @@ def err_hlsl_packoffset_alignment_mismatch : Error<"packoffset at 'y' not match def err_hlsl_pointers_unsupported : Error< "%select{pointers|references}0 are unsupported in HLSL">; def err_hlsl_missing_resource_class : Error<"HLSL resource needs to have [[hlsl::resource_class()]] attribute">; -def err_hlsl_attribute_needs_intangible_type: Error<"attribute %0 can be used only on HLSL intangible type %1">; def err_hlsl_operator_unsupported : Error< "the '%select{&|*|->}0' operator is unsupported in HLSL">; diff --git a/clang/include/clang/Sema/SemaHLSL.h b/clang/include/clang/Sema/SemaHLSL.h index 311cd58bbcac2..e088254c566d3 100644 --- a/clang/include/clang/Sema/SemaHLSL.h +++ b/clang/include/clang/Sema/SemaHLSL.h @@ -70,7 +70,7 @@ class SemaHLSL : public SemaBase { void handleShaderAttr(Decl *D, const ParsedAttr &AL); void handleResourceBindingAttr(Decl *D, const ParsedAttr &AL); void handleParamModifierAttr(Decl *D, const ParsedAttr &AL); - bool handleResourceTypeAttr(QualType T, const ParsedAttr &AL); + bool handleResourceTypeAttr(const ParsedAttr &AL); bool CheckBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall); QualType ProcessResourceTypeAttributes(QualType Wrapped); diff --git a/clang/lib/Sema/SemaHLSL.cpp b/clang/lib/Sema/SemaHLSL.cpp index 1d8ccdda45573..ebe76185cbb2d 100644 --- a/clang/lib/Sema/SemaHLSL.cpp +++ b/clang/lib/Sema/SemaHLSL.cpp @@ -693,19 +693,13 @@ bool clang::CreateHLSLAttributedResourceType( // HLSL resource. The attributes are collected in HLSLResourcesTypeAttrs and at // the end of the declaration they are applied to the declaration type by // wrapping it in HLSLAttributedResourceType. -bool SemaHLSL::handleResourceTypeAttr(QualType T, const ParsedAttr &AL) { - // only allow resource type attributes on intangible types - if (!T->isHLSLResourceType()) { - Diag(AL.getLoc(), diag::err_hlsl_attribute_needs_intangible_type) - << AL << getASTContext().HLSLResourceTy; - return false; - } +bool SemaHLSL::handleResourceTypeAttr(const ParsedAttr &AL) { + Attr *A = nullptr; // validate number of arguments if (!AL.checkExactlyNumArgs(SemaRef, AL.getMinArgs())) return false; - Attr *A = nullptr; switch (AL.getKind()) { case ParsedAttr::AT_HLSLResourceClass: { if (!AL.isArgIdent(0)) { diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index a7beb9d222c3b..950bd6db0359d 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -8860,7 +8860,7 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type, // decl-specifier-seq; do not collect attributes on declarations or those // that get to slide after declaration name. if (TAL == TAL_DeclSpec && - state.getSema().HLSL().handleResourceTypeAttr(type, attr)) + state.getSema().HLSL().handleResourceTypeAttr(attr)) attr.setUsedAsTypeAttr(); break; } diff --git a/clang/test/ParserHLSL/hlsl_contained_type_attr_error.hlsl b/clang/test/ParserHLSL/hlsl_contained_type_attr_error.hlsl index b2d492d95945c..1c37d72de8614 100644 --- a/clang/test/ParserHLSL/hlsl_contained_type_attr_error.hlsl +++ b/clang/test/ParserHLSL/hlsl_contained_type_attr_error.hlsl @@ -22,7 +22,3 @@ __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::contained_type(float)]] // expected-warning@+1{{attribute 'contained_type' is already applied with different arguments}} __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::contained_type(float)]] [[hlsl::contained_type(int)]] h8; - -// expected-error@+2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}} -// expected-error@+1{{attribute 'contained_type' can be used only on HLSL intangible type '__hlsl_resource_t'}} -float [[hlsl::resource_class(UAV)]] [[hlsl::contained_type(float)]] res5; diff --git a/clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl b/clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl index 3b2c12e7a96c5..15685bd1a3baa 100644 --- a/clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl +++ b/clang/test/ParserHLSL/hlsl_is_rov_attr_error.hlsl @@ -14,7 +14,3 @@ __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::is_rov(gibberish)]] res3 // expected-warning@+1{{attribute 'is_rov' is already applied}} __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::is_rov]] [[hlsl::is_rov]] res4; - -// expected-error@+2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}} -// expected-error@+1{{attribute 'is_rov' can be used only on HLSL intangible type '__hlsl_resource_t'}} -float [[hlsl::resource_class(UAV)]] [[hlsl::is_rov]] res5; diff --git a/clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl b/clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl index 77530cbf9e4d9..83273426017ed 100644 --- a/clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl +++ b/clang/test/ParserHLSL/hlsl_raw_buffer_attr_error.hlsl @@ -11,7 +11,3 @@ __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::raw_buffer(gibberish)]] // expected-warning@+1{{attribute 'raw_buffer' is already applied}} __hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::raw_buffer]] [[hlsl::raw_buffer]] res4; - -// expected-error@+2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}} -// expected-error@+1{{attribute 'raw_buffer' can be used only on HLSL intangible type '__hlsl_resource_t'}} -float [[hlsl::resource_class(UAV)]] [[hlsl::raw_buffer]] res5; diff --git a/clang/test/ParserHLSL/hlsl_resource_class_attr_error.hlsl b/clang/test/ParserHLSL/hlsl_resource_class_attr_error.hlsl index 63e39daff949b..01ff1c007e2b5 100644 --- a/clang/test/ParserHLSL/hlsl_resource_class_attr_error.hlsl +++ b/clang/test/ParserHLSL/hlsl_resource_class_attr_error.hlsl @@ -17,6 +17,3 @@ __hlsl_resource_t [[hlsl::resource_class(SRV)]] [[hlsl::resource_class(SRV)]] e4 // expected-error@+1{{'resource_class' attribute takes one argument}} __hlsl_resource_t [[hlsl::resource_class(SRV, "aa")]] e5; - -// expected-error@+1{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}} -float [[hlsl::resource_class(UAV)]] e6;