diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td index 8c5d7ad763bf9..33791270800c9 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 d19f79b6ddefc..ca521dc0bcd26 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 c3ba520e0f68e..a95be63da5dc1 100644 --- a/clang/test/AST/HLSL/RWBuffer-AST.hlsl +++ b/clang/test/AST/HLSL/RWBuffer-AST.hlsl @@ -66,7 +66,7 @@ RWBuffer Buffer; // CHECK: TemplateArgument type 'float' // CHECK-NEXT: BuiltinType 0x{{[0-9A-Fa-f]+}} 'float' // CHECK-NEXT: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit referenced h 'float * +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> 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 1a3deba5830fa..a186779870c26 100644 --- a/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl +++ b/clang/test/AST/HLSL/StructuredBuffer-AST.hlsl @@ -70,7 +70,7 @@ StructuredBuffer Buffer; // CHECK: TemplateArgument type 'float' // CHECK-NEXT: BuiltinType 0x{{[0-9A-Fa-f]+}} 'float' // CHECK-NEXT: FinalAttr 0x{{[0-9A-Fa-f]+}} <> Implicit final -// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> implicit referenced h 'float * +// CHECK-NEXT: FieldDecl 0x{{[0-9A-Fa-f]+}} <> 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 174f4c3eaaad2..19699dcf14d9f 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 Buf; // CHECK: define linkonce_odr noundef ptr @"??0?$RWBuffer@M@hlsl@@QAA@XZ" @@ -10,4 +16,4 @@ RWBuffer 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 34019e5b18693..178332d03e640 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 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 61226c2b54e72..0000000000000 --- 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 301d61c0e906e..5e4ed96561a30 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]+}} <> class RWBuffer definition implicit_instantiation // CHECK: -TemplateArgument type 'float' // CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float' -// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit referenced h 'float * +// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> 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 Buffer1; // CHECK: -TemplateArgument type 'vector' // CHECK: `-ExtVectorType 0x{{[0-9a-f]+}} 'vector' 4 // CHECK: `-BuiltinType 0x{{[0-9a-f]+}} 'float' -// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit referenced h 'vector +// CHECK: -FieldDecl 0x{{[0-9a-f]+}} <> implicit h 'vector // CHECK-SAME{LITERAL}: [[hlsl::resource_class(UAV)] // CHECK-SAME{LITERAL}: [[hlsl::is_rov]] // CHECK-SAME{LITERAL}: [[hlsl::contained_type(vector)]] diff --git a/llvm/include/llvm/IR/IntrinsicsDirectX.td b/llvm/include/llvm/IR/IntrinsicsDirectX.td index 3ce7b8b987ef8..555877e7aaf0e 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 7ac479f31386f..a9dbddb13adaa 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 5916a194f76d4..66e1e432ddf64 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},