Skip to content

[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

Merged
merged 1 commit into from
Sep 26, 2024

Conversation

hekota
Copy link
Member

@hekota hekota commented Sep 25, 2024

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.

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.
@hekota hekota closed this Sep 25, 2024
@hekota hekota reopened this Sep 26, 2024
@hekota hekota marked this pull request as ready for review September 26, 2024 01:48
@llvmbot llvmbot added clang Clang issues not falling into any other category clang:frontend Language frontend issues, e.g. anything involving "Sema" backend:DirectX HLSL HLSL Language Support llvm:ir labels Sep 26, 2024
@llvmbot
Copy link
Member

llvmbot commented Sep 26, 2024

@llvm/pr-subscribers-llvm-ir
@llvm/pr-subscribers-backend-directx
@llvm/pr-subscribers-hlsl

@llvm/pr-subscribers-clang

Author: Helena Kotas (hekota)

Changes

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.


Full diff: https://github.com/llvm/llvm-project/pull/109910.diff

11 Files Affected:

  • (modified) clang/include/clang/Basic/Builtins.td (-6)
  • (modified) clang/lib/Sema/HLSLExternalSemaSource.cpp (+2-30)
  • (modified) clang/test/AST/HLSL/RWBuffer-AST.hlsl (+1-1)
  • (modified) clang/test/AST/HLSL/StructuredBuffer-AST.hlsl (+1-1)
  • (modified) clang/test/CodeGenHLSL/builtins/RWBuffer-constructor.hlsl (+7-1)
  • (modified) clang/test/CodeGenHLSL/builtins/StructuredBuffer-constructor.hlsl (+7)
  • (removed) clang/test/CodeGenHLSL/builtins/create_handle.hlsl (-7)
  • (modified) clang/test/ParserHLSL/hlsl_resource_handle_attrs.hlsl (+2-2)
  • (modified) llvm/include/llvm/IR/IntrinsicsDirectX.td (-3)
  • (modified) llvm/include/llvm/IR/IntrinsicsSPIRV.td (-2)
  • (modified) llvm/unittests/IR/IntrinsicsTest.cpp (-1)
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},

Copy link
Contributor

@damyanp damyanp left a 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.

@hekota
Copy link
Member Author

hekota commented Sep 26, 2024

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.

@hekota hekota merged commit 90b7fe4 into llvm:main Sep 26, 2024
21 checks passed
Sterling-Augustine pushed a commit to Sterling-Augustine/llvm-project that referenced this pull request Sep 27, 2024
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.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
backend:DirectX clang:frontend Language frontend issues, e.g. anything involving "Sema" clang Clang issues not falling into any other category HLSL HLSL Language Support llvm:ir
Projects
Archived in project
Development

Successfully merging this pull request may close these issues.

4 participants