[clang] [llvm] [HLSL] Replace `element_type*` handles in HLSLExternalSemaSource with `__hlsl_resource_t` builtin type (PR #110079)
Helena Kotas via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 27 08:22:49 PDT 2024
https://github.com/hekota updated https://github.com/llvm/llvm-project/pull/110079
>From 970aab0a930e38dfd266c01065112602bb274a5e Mon Sep 17 00:00:00 2001
From: Helena Kotas <hekotas at microsoft.com>
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 dc87b84153e74a..67e75652a16649 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 e4e04bff8b5120..a9b2042fcff4e3 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 e088254c566d3e..311cd58bbcac2c 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 fd8aa8de79b49f..ced18e1617e248 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 d19f79b6ddefcd..e72da9a071fb71 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<TemplateTypeParmDecl>(
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<const Attr *> 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<TemplateTypeParmDecl>(
+ 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<FunctionProtoTypeLoc>();
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 ebe76185cbb2d5..1d8ccdda45573f 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 950bd6db0359d1..a7beb9d222c3b5 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 c3ba520e0f68e4..55c0dfa2eaa533 100644
--- a/clang/test/AST/HLSL/RWBuffer-AST.hlsl
+++ b/clang/test/AST/HLSL/RWBuffer-AST.hlsl
@@ -29,36 +29,26 @@ RWBuffer<float> Buffer;
// CHECK-NEXT: CXXRecordDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit class RWBuffer definition
// CHECK: FinalAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit final
-// CHECK-NEXT: implicit h 'element_type *
+// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> 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]+}} <<invalid sloc>> Implicit TypedBuffer
// CHECK: CXXMethodDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> operator[] 'element_type &const (unsigned int) const'
// CHECK-NEXT: ParmVarDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> Idx 'unsigned int'
// CHECK-NEXT: CompoundStmt 0x{{[0-9A-Fa-f]+}} <<invalid sloc>>
// CHECK-NEXT: ReturnStmt 0x{{[0-9A-Fa-f]+}} <<invalid sloc>>
-// CHECK-NEXT: ArraySubscriptExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'element_type' lvalue
-// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> '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]+}} <<invalid sloc>> 'element_type' lvalue .e 0x{{[0-9A-Fa-f]+}}
// CHECK-NEXT: CXXThisExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'const RWBuffer<element_type>' lvalue implicit this
-// CHECK-NEXT: DeclRefExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'unsigned int' ParmVar 0x{{[0-9A-Fa-f]+}} 'Idx' 'unsigned int'
// CHECK-NEXT: AlwaysInlineAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit always_inline
// CHECK-NEXT: CXXMethodDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> operator[] 'element_type &(unsigned int)'
// CHECK-NEXT: ParmVarDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> Idx 'unsigned int'
// CHECK-NEXT: CompoundStmt 0x{{[0-9A-Fa-f]+}} <<invalid sloc>>
// CHECK-NEXT: ReturnStmt 0x{{[0-9A-Fa-f]+}} <<invalid sloc>>
-// CHECK-NEXT: ArraySubscriptExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'element_type' lvalue
-// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> '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]+}} <<invalid sloc>> 'element_type' lvalue .e 0x{{[0-9A-Fa-f]+}}
// CHECK-NEXT: CXXThisExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'RWBuffer<element_type>' lvalue implicit this
-// CHECK-NEXT: DeclRefExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'unsigned int' ParmVar 0x{{[0-9A-Fa-f]+}} 'Idx' 'unsigned int'
// CHECK-NEXT: AlwaysInlineAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit always_inline
// CHECK: ClassTemplateSpecializationDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> class RWBuffer definition
@@ -66,8 +56,8 @@ RWBuffer<float> Buffer;
// CHECK: TemplateArgument type 'float'
// CHECK-NEXT: BuiltinType 0x{{[0-9A-Fa-f]+}} 'float'
// CHECK-NEXT: FinalAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit final
-// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit referenced h 'float *
+// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> 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]+}} <<invalid sloc>> Implicit TypedBuffer
diff --git a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
index 1a3deba5830fa7..cdc7cc8de07c63 100644
--- a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
+++ b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
@@ -30,39 +30,27 @@ StructuredBuffer<float> Buffer;
// CHECK-NEXT: CXXRecordDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit class StructuredBuffer definition
// CHECK: FinalAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit final
-// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit h 'element_type *
+// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> 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]+}} <<invalid sloc>> Implicit TypedBuffer
// CHECK: CXXMethodDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> operator[] 'element_type &const (unsigned int) const'
// CHECK-NEXT: ParmVarDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> Idx 'unsigned int'
// CHECK-NEXT: CompoundStmt 0x{{[0-9A-Fa-f]+}} <<invalid sloc>>
// CHECK-NEXT: ReturnStmt 0x{{[0-9A-Fa-f]+}} <<invalid sloc>>
-// CHECK-NEXT: ArraySubscriptExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'element_type' lvalue
-// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> '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]+}} <<invalid sloc>> 'element_type' lvalue .e 0x{{[0-9A-Fa-f]+}}
// CHECK-NEXT: CXXThisExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'const StructuredBuffer<element_type>' lvalue implicit this
-// CHECK-NEXT: DeclRefExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'unsigned int' ParmVar 0x{{[0-9A-Fa-f]+}} 'Idx' 'unsigned int'
// CHECK-NEXT: AlwaysInlineAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit always_inline
// CHECK-NEXT: CXXMethodDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> operator[] 'element_type &(unsigned int)'
// CHECK-NEXT: ParmVarDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> Idx 'unsigned int'
// CHECK-NEXT: CompoundStmt 0x{{[0-9A-Fa-f]+}} <<invalid sloc>>
// CHECK-NEXT: ReturnStmt 0x{{[0-9A-Fa-f]+}} <<invalid sloc>>
-// CHECK-NEXT: ArraySubscriptExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'element_type' lvalue
-// CHECK-NEXT: MemberExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> '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]+}} <<invalid sloc>> 'element_type' lvalue .e 0x{{[0-9A-Fa-f]+}}
// CHECK-NEXT: CXXThisExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'StructuredBuffer<element_type>' lvalue implicit this
-// CHECK-NEXT: DeclRefExpr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> 'unsigned int' ParmVar 0x{{[0-9A-Fa-f]+}} 'Idx' 'unsigned int'
// CHECK-NEXT: AlwaysInlineAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit always_inline
// CHECK: ClassTemplateSpecializationDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> class StructuredBuffer definition
@@ -74,5 +62,5 @@ StructuredBuffer<float> 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]+}} <<invalid sloc>> Implicit TypedBuffer
diff --git a/clang/test/CodeGenHLSL/buffer-array-operator.hlsl b/clang/test/CodeGenHLSL/buffer-array-operator.hlsl
index 02e570ebdcb4fc..f65cdbb43e27b5 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<float> In;
RWBuffer<float> Out;
diff --git a/clang/test/CodeGenHLSL/implicit-norecurse-attrib.hlsl b/clang/test/CodeGenHLSL/implicit-norecurse-attrib.hlsl
index ae3a3b5f90199f..f72fe059cb5763 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 at UNode@@V?$RWBuffer at T?$__vector at I$03 at __clang@@@hlsl@@I at 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 at UNode@@V?$RWBuffer at T?$__vector at I$03 at __clang@@@hlsl@@I at 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 1c37d72de86148..b2d492d95945c1 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 at +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 at +2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}}
+// expected-error at +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 15685bd1a3baa7..3b2c12e7a96c5c 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 at +1{{attribute 'is_rov' is already applied}}
__hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::is_rov]] [[hlsl::is_rov]] res4;
+
+// expected-error at +2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}}
+// expected-error at +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 83273426017ed0..77530cbf9e4d92 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 at +1{{attribute 'raw_buffer' is already applied}}
__hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::raw_buffer]] [[hlsl::raw_buffer]] res4;
+
+// expected-error at +2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}}
+// expected-error at +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 01ff1c007e2b57..63e39daff949b4 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 at +1{{'resource_class' attribute takes one argument}}
__hlsl_resource_t [[hlsl::resource_class(SRV, "aa")]] e5;
+
+// expected-error at +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 301d61c0e906ed..b73c325fae0715 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]+}} <<invalid sloc>> <invalid sloc> class RWBuffer definition implicit_instantiation
// CHECK: -TemplateArgument type 'float'
// CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float'
-// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> implicit referenced h 'float *
+// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> 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]+}} <<invalid sloc>> Implicit TypedBuffer
RWBuffer<float> Buffer1;
@@ -18,6 +18,6 @@ RWBuffer<float> Buffer1;
// CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]
// CHECK-SAME{LITERAL}: [[hlsl::is_rov]]
// CHECK-SAME{LITERAL}: [[hlsl::contained_type(vector<float, 4>)]]
-// CHECK-SAME: ':'vector<float *, 4>'
+// CHECK-SAME: ':'__hlsl_resource_t'
// CHECK: -HLSLResourceAttr 0x{{[0-9a-f]+}} <<invalid sloc>> Implicit TypedBuffer
RasterizerOrderedBuffer<vector<float, 4> > BufferArray3[4];
diff --git a/clang/test/SemaHLSL/Types/Traits/IsIntangibleType.hlsl b/clang/test/SemaHLSL/Types/Traits/IsIntangibleType.hlsl
index 92cba1dcd4bdfe..8c0f8d6f271dbd 100644
--- a/clang/test/SemaHLSL/Types/Traits/IsIntangibleType.hlsl
+++ b/clang/test/SemaHLSL/Types/Traits/IsIntangibleType.hlsl
@@ -76,3 +76,6 @@ template<typename T> struct SimpleTemplate {
};
_Static_assert(__builtin_hlsl_is_intangible(SimpleTemplate<__hlsl_resource_t>), "");
_Static_assert(!__builtin_hlsl_is_intangible(SimpleTemplate<float>), "");
+
+_Static_assert(__builtin_hlsl_is_intangible(RWBuffer<float>), "");
+_Static_assert(__builtin_hlsl_is_intangible(StructuredBuffer<Simple>), "");
>From 71c4a9361f07a606c95eb5232be50dcdff31c422 Mon Sep 17 00:00:00 2001
From: Helena Kotas <hekotas at microsoft.com>
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 8c5d7ad763bf97..33791270800c9d 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 e72da9a071fb71..b6bd8e778a6179 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 cdc7cc8de07c63..b31db8ce59f224 100644
--- a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
+++ b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
@@ -30,7 +30,7 @@ StructuredBuffer<float> Buffer;
// CHECK-NEXT: CXXRecordDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit class StructuredBuffer definition
// CHECK: FinalAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit final
-// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit referenced h '__hlsl_resource_t
+// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> 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<float> Buffer;
// CHECK: TemplateArgument type 'float'
// CHECK-NEXT: BuiltinType 0x{{[0-9A-Fa-f]+}} 'float'
// CHECK-NEXT: FinalAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit final
-// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit referenced h 'float *
+// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> 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 174f4c3eaaad26..19699dcf14d9f4 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<float> Buf;
// CHECK: define linkonce_odr noundef ptr @"??0?$RWBuffer at M@hlsl@@QAA at XZ"
@@ -10,4 +16,4 @@ RWBuffer<float> 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 34019e5b186931..178332d03e6404 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<float> Buf;
// CHECK: define linkonce_odr noundef ptr @"??0?$StructuredBuffer at M@hlsl@@QAA at XZ"
diff --git a/clang/test/CodeGenHLSL/builtins/create_handle.hlsl b/clang/test/CodeGenHLSL/builtins/create_handle.hlsl
deleted file mode 100644
index 61226c2b54e726..00000000000000
--- 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 b73c325fae0715..e7d19c3da7216d 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]+}} <<invalid sloc>> <invalid sloc> class RWBuffer definition implicit_instantiation
// CHECK: -TemplateArgument type 'float'
// CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float'
-// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> implicit referenced h '__hlsl_resource_t
+// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> 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<float> Buffer1;
// CHECK: -TemplateArgument type 'vector<float, 4>'
// CHECK: `-ExtVectorType 0x{{[0-9a-f]+}} 'vector<float, 4>' 4
// CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float'
-// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> implicit referenced h 'vector<float *, 4>
+// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> 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<float, 4>)]]
diff --git a/llvm/include/llvm/IR/IntrinsicsDirectX.td b/llvm/include/llvm/IR/IntrinsicsDirectX.td
index 3ce7b8b987ef86..555877e7aaf0e5 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 7ac479f31386f9..a9dbddb13adaa7 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 a92ffe3cdeb7ec..0c4af28a2ab57b 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 <hekotas at microsoft.com>
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 b6bd8e778a6179..2913d16fca4823 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 <hekotas at microsoft.com>
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 33791270800c9d..8c5d7ad763bf97 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 2913d16fca4823..42d3358e4ddbe6 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 b31db8ce59f224..cdc7cc8de07c63 100644
--- a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
+++ b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
@@ -30,7 +30,7 @@ StructuredBuffer<float> Buffer;
// CHECK-NEXT: CXXRecordDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit class StructuredBuffer definition
// CHECK: FinalAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit final
-// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit h '__hlsl_resource_t
+// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> 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<float> Buffer;
// CHECK: TemplateArgument type 'float'
// CHECK-NEXT: BuiltinType 0x{{[0-9A-Fa-f]+}} 'float'
// CHECK-NEXT: FinalAttr 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> Implicit final
-// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> implicit h '__hlsl_resource_t
+// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <<invalid sloc>> <invalid sloc> 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 19699dcf14d9f4..174f4c3eaaad26 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<float> Buf;
// CHECK: define linkonce_odr noundef ptr @"??0?$RWBuffer at M@hlsl@@QAA at XZ"
@@ -16,4 +10,4 @@ RWBuffer<float> 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 178332d03e6404..34019e5b186931 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<float> Buf;
// CHECK: define linkonce_odr noundef ptr @"??0?$StructuredBuffer at M@hlsl@@QAA at XZ"
diff --git a/clang/test/CodeGenHLSL/builtins/create_handle.hlsl b/clang/test/CodeGenHLSL/builtins/create_handle.hlsl
new file mode 100644
index 00000000000000..61226c2b54e726
--- /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 e7d19c3da7216d..b73c325fae0715 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]+}} <<invalid sloc>> <invalid sloc> class RWBuffer definition implicit_instantiation
// CHECK: -TemplateArgument type 'float'
// CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float'
-// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> implicit h '__hlsl_resource_t
+// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> 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<float> Buffer1;
// CHECK: -TemplateArgument type 'vector<float, 4>'
// CHECK: `-ExtVectorType 0x{{[0-9a-f]+}} 'vector<float, 4>' 4
// CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float'
-// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> implicit h '__hlsl_resource_t
+// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> implicit referenced h 'vector<float *, 4>
// CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]
// CHECK-SAME{LITERAL}: [[hlsl::is_rov]]
// CHECK-SAME{LITERAL}: [[hlsl::contained_type(vector<float, 4>)]]
diff --git a/llvm/include/llvm/IR/IntrinsicsDirectX.td b/llvm/include/llvm/IR/IntrinsicsDirectX.td
index 555877e7aaf0e5..3ce7b8b987ef86 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 a9dbddb13adaa7..7ac479f31386f9 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 0c4af28a2ab57b..a92ffe3cdeb7ec 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 <hekotas at microsoft.com>
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 a9b2042fcff4e3..e4e04bff8b5120 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 311cd58bbcac2c..e088254c566d3e 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 1d8ccdda45573f..ebe76185cbb2d5 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 a7beb9d222c3b5..950bd6db0359d1 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 b2d492d95945c1..1c37d72de86148 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 at +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 at +2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}}
-// expected-error at +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 3b2c12e7a96c5c..15685bd1a3baa7 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 at +1{{attribute 'is_rov' is already applied}}
__hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::is_rov]] [[hlsl::is_rov]] res4;
-
-// expected-error at +2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}}
-// expected-error at +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 77530cbf9e4d92..83273426017ed0 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 at +1{{attribute 'raw_buffer' is already applied}}
__hlsl_resource_t [[hlsl::resource_class(UAV)]] [[hlsl::raw_buffer]] [[hlsl::raw_buffer]] res4;
-
-// expected-error at +2{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}}
-// expected-error at +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 63e39daff949b4..01ff1c007e2b57 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 at +1{{'resource_class' attribute takes one argument}}
__hlsl_resource_t [[hlsl::resource_class(SRV, "aa")]] e5;
-
-// expected-error at +1{{attribute 'resource_class' can be used only on HLSL intangible type '__hlsl_resource_t'}}
-float [[hlsl::resource_class(UAV)]] e6;
More information about the cfe-commits
mailing list