llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-hlsl @llvm/pr-subscribers-clang Author: Helena Kotas (hekota) <details> <summary>Changes</summary> 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 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}, `````````` </details> https://github.com/llvm/llvm-project/pull/109910 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
