-
Notifications
You must be signed in to change notification settings - Fork 14.3k
[HLSL] Remove __builtin_hlsl_create_handle
#109910
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Conversation
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#105076 once the dependent tasks are completed. Part 1/2 of llvm#105076.
@llvm/pr-subscribers-llvm-ir @llvm/pr-subscribers-clang Author: Helena Kotas (hekota) ChangesThe Instead, the handle should be initialized to an empty or null handle with something like 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. Full diff: https://github.com/llvm/llvm-project/pull/109910.diff 11 Files Affected:
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@M@hlsl@@QAA@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@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 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 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 5916a194f76d48..66e1e432ddf641 100644
--- a/llvm/unittests/IR/IntrinsicsTest.cpp
+++ b/llvm/unittests/IR/IntrinsicsTest.cpp
@@ -92,7 +92,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},
|
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
It's a shame we don't have a better way to review changes like this, so I'm a little uncomfortable about a commit leaving things in a transitional state, but I think my appreciation of being able to review this separately from the next step outweighs that.
Yeah, I had all of the resource codegen changes in one place, but it was huge. So I am splitting it to about 5 separate PRs for easier reviews. |
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#105076 once the dependent tasks are completed. Part 1/2 of llvm#105076.
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 #105076 once the dependent tasks are completed.
Part 1/2 of #105076.