[llvm] 90b7fe4 - [HLSL] Remove `__builtin_hlsl_create_handle` (#109910)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Sep 26 13:56:52 PDT 2024
Author: Helena Kotas
Date: 2024-09-26T13:56:49-07:00
New Revision: 90b7fe42d8e6f8647ce9279d6d026c36ccfcbb8f
URL: https://github.com/llvm/llvm-project/commit/90b7fe42d8e6f8647ce9279d6d026c36ccfcbb8f
DIFF: https://github.com/llvm/llvm-project/commit/90b7fe42d8e6f8647ce9279d6d026c36ccfcbb8f.diff
LOG: [HLSL] Remove `__builtin_hlsl_create_handle` (#109910)
The `__builtin_hlsl_create_handle` called from the constructor of
resource buffer class was supposed to initialize the resource handle
based on resource type and registry binding information. It is not
possible to do though that because the registry binding information is
not accessible from the constructor during codegen.
Instead, the handle should be initialized to an empty or null handle
with something like `__builtin_hlsl_create_null_handle`. This PR is
removing `__builtin_hlsl_create_handle` first and the
`__builtin_hlsl_create_null_handle` will be added to the constructor
once the handle type changes to `__hlsl_resource_t` and
HLSLAttributeResourceType is updated to be a canonical type, which will
allow the initialization assignment.
The actual handle initialization based on the registry binding will be
implemented part 2/2 of llvm/llvm-project#105076 once the dependent
tasks are completed.
Part 1/2 of llvm/llvm-project#105076.
Added:
Modified:
clang/include/clang/Basic/Builtins.td
clang/lib/Sema/HLSLExternalSemaSource.cpp
clang/test/AST/HLSL/RWBuffer-AST.hlsl
clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl
clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl
clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl
llvm/include/llvm/IR/IntrinsicsDirectX.td
llvm/include/llvm/IR/IntrinsicsSPIRV.td
llvm/unittests/IR/IntrinsicsTest.cpp
Removed:
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 d19f79b6ddefcd..ca521dc0bcd26b 100644
--- a/clang/lib/Sema/HLSLExternalSemaSource.cpp
+++ b/clang/lib/Sema/HLSLExternalSemaSource.cpp
@@ -193,36 +193,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/RWBuffer-AST.hlsl b/clang/test/AST/HLSL/RWBuffer-AST.hlsl
index c3ba520e0f68e4..a95be63da5dc17 100644
--- a/clang/test/AST/HLSL/RWBuffer-AST.hlsl
+++ b/clang/test/AST/HLSL/RWBuffer-AST.hlsl
@@ -66,7 +66,7 @@ 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 'float *
// CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]]
// CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]]
// CHECK-SAME: ':'float *'
diff --git a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
index 1a3deba5830fa7..a186779870c263 100644
--- a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
+++ b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl
@@ -70,7 +70,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 '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 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 301d61c0e906ed..5e4ed96561a300 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 'float *
+// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <<invalid sloc>> <invalid sloc> implicit h 'float *
// CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)]]
// CHECK-SAME{LITERAL}: [[hlsl::contained_type(float)]]
// CHECK-SAME: ':'float *'
@@ -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 '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 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 c5c60963ed6fda..7ff3d58690ba75 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -59,8 +59,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},
More information about the llvm-commits
mailing list