From 419bc4ba6d1b6ff01c8f2f8ac2306307d9022cc9 Mon Sep 17 00:00:00 2001 From: David Neto Date: Thu, 17 Oct 2024 03:29:12 -0400 Subject: [PATCH] Update LLVM (#1406) - Rename Intrinsic::getDeclaration to Intrinsic::getOrInsertDeclaration - Update one test to avoid disabling int8 Bug: crbug.com/373696543 --- deps.json | 2 +- lib/LongVectorLoweringPass.cpp | 54 +++++++-------- lib/ReplaceLLVMIntrinsicsPass.cpp | 7 +- lib/ReplaceOpenCLBuiltinPass.cpp | 68 ++++++++++--------- .../in_storage_buffer_descriptor_map.cl | 2 +- test/UBO/transform_global.cl | 7 +- 6 files changed, 71 insertions(+), 69 deletions(-) diff --git a/deps.json b/deps.json index 2b886a387..00e6e907c 100644 --- a/deps.json +++ b/deps.json @@ -6,7 +6,7 @@ "subrepo" : "llvm/llvm-project", "branch" : "main", "subdir" : "third_party/llvm", - "commit" : "125aa10b3d645bd26523a1bc321bb2e6b1cf04e1" + "commit" : "30deb76d46053c243561c6fa072c5a30407241cb" }, { "name" : "SPIRV-Headers", diff --git a/lib/LongVectorLoweringPass.cpp b/lib/LongVectorLoweringPass.cpp index 0ce00fd1b..51c30cf45 100644 --- a/lib/LongVectorLoweringPass.cpp +++ b/lib/LongVectorLoweringPass.cpp @@ -31,9 +31,9 @@ #include "Constants.h" #include "clspv/Passes.h" +#include "BitcastUtils.h" #include "Builtins.h" #include "LongVectorLoweringPass.h" -#include "BitcastUtils.h" #include #include @@ -120,7 +120,8 @@ Function *getIntrinsicScalarVersion(Function &Intrinsic) { Param = Param->getScalarType(); } - return Intrinsic::getDeclaration(Intrinsic.getParent(), id, ParamTys); + return Intrinsic::getOrInsertDeclaration(Intrinsic.getParent(), id, + ParamTys); break; } } @@ -154,30 +155,29 @@ std::string getSpirvCompliantName(const clspv::Builtins::FunctionInfo &IInfo) { Type *getScalarPointerType(Function &Builtin) { const auto &Info = clspv::Builtins::Lookup(&Builtin); switch (Info.getType()) { - case clspv::Builtins::kSincos: - case clspv::Builtins::kModf: - case clspv::Builtins::kFract: - return Builtin.getReturnType()->getScalarType(); - case clspv::Builtins::kFrexp: - case clspv::Builtins::kRemquo: - case clspv::Builtins::kLgammaR: - return Type::getInt32Ty(Builtin.getParent()->getContext()); - case clspv::Builtins::kVloadHalf: - case clspv::Builtins::kVloadaHalf: - case clspv::Builtins::kVstoreHalf: - case clspv::Builtins::kVstoreaHalf: - return Type::getHalfTy(Builtin.getParent()->getContext()); - case clspv::Builtins::kVstore: - return Builtin.getArg(0)->getType()->getScalarType(); - case clspv::Builtins::kVload: - return Builtin.getReturnType()->getScalarType(); - default: - // What about llvm intrinsics (e.g. memcpy) or other OpenCL builtins? - return nullptr; + case clspv::Builtins::kSincos: + case clspv::Builtins::kModf: + case clspv::Builtins::kFract: + return Builtin.getReturnType()->getScalarType(); + case clspv::Builtins::kFrexp: + case clspv::Builtins::kRemquo: + case clspv::Builtins::kLgammaR: + return Type::getInt32Ty(Builtin.getParent()->getContext()); + case clspv::Builtins::kVloadHalf: + case clspv::Builtins::kVloadaHalf: + case clspv::Builtins::kVstoreHalf: + case clspv::Builtins::kVstoreaHalf: + return Type::getHalfTy(Builtin.getParent()->getContext()); + case clspv::Builtins::kVstore: + return Builtin.getArg(0)->getType()->getScalarType(); + case clspv::Builtins::kVload: + return Builtin.getReturnType()->getScalarType(); + default: + // What about llvm intrinsics (e.g. memcpy) or other OpenCL builtins? + return nullptr; } } - /// Get the scalar overload for the given OpenCL builtin function @p Builtin. Function *getBIFScalarVersion(Function &Builtin) { assert(!Builtin.isIntrinsic()); @@ -520,7 +520,8 @@ Function *createFunctionWithMappedTypes(Function &F, B.CreateRetVoid(); } else { auto *EquivalentReturnTy = EquivalentFunctionTy->getReturnType(); - Value *ReturnValue = convertEquivalentValue(B, Call, EquivalentReturnTy, DL); + Value *ReturnValue = + convertEquivalentValue(B, Call, EquivalentReturnTy, DL); B.CreateRet(ReturnValue); } @@ -916,8 +917,7 @@ Value *clspv::LongVectorLoweringPass::visitCastInst(CastInst &I) { V = B.CreatePtrToInt(EquivalentValue, EquivalentDestTy, I.getName()); break; } - case Instruction::IntToPtr: - { + case Instruction::IntToPtr: { IRBuilder<> B(&I); V = B.CreateIntToPtr(EquivalentValue, EquivalentDestTy, I.getName()); break; @@ -1049,7 +1049,7 @@ void clspv::LongVectorLoweringPass::reworkIndices( for (unsigned i = 1; i < Idxs.size(); i++) { Indices.push_back(Idxs[i]); // Get original indices up to ith element for below: - auto CumulativeOldIdxs = ArrayRef(Idxs.begin(), Idxs.begin() + i); + auto CumulativeOldIdxs = ArrayRef(Idxs.begin(), Idxs.begin() + i); auto IndexedTy = GetElementPtrInst::getIndexedType(Ty, CumulativeOldIdxs); if (getEquivalentType(IndexedTy)) { auto Idx = Indices.pop_back_val(); diff --git a/lib/ReplaceLLVMIntrinsicsPass.cpp b/lib/ReplaceLLVMIntrinsicsPass.cpp index 58ea74206..6e4da9eb9 100644 --- a/lib/ReplaceLLVMIntrinsicsPass.cpp +++ b/lib/ReplaceLLVMIntrinsicsPass.cpp @@ -95,8 +95,9 @@ bool clspv::ReplaceLLVMIntrinsicsPass::runOnFunction(Function &F) { return replaceAddSubSat(F, true, true); case Intrinsic::is_fpclass: return replaceIsFpClass(F); - // SPIR-V OpAssumeTrueKHR requires ExpectAssumeKHR capability in SPV_KHR_expect_assume extension. - // Vulkan doesn't support that, so remove assume declaration. + // SPIR-V OpAssumeTrueKHR requires ExpectAssumeKHR capability in + // SPV_KHR_expect_assume extension. Vulkan doesn't support that, so remove + // assume declaration. case Intrinsic::assume: // SPIR-V OpLifetimeStart and OpLifetimeEnd require Kernel capability. // Vulkan doesn't support that, so remove all lifteime bounds declarations. @@ -465,7 +466,7 @@ bool clspv::ReplaceLLVMIntrinsicsPass::replaceCountZeroes(Function &F, IRBuilder<> builder(Call); auto ty = Call->getType()->getWithNewBitWidth(32); auto c32 = ConstantInt::get(ty, 32); - auto func_32bit = Intrinsic::getDeclaration( + auto func_32bit = Intrinsic::getOrInsertDeclaration( F.getParent(), leading ? Intrinsic::ctlz : Intrinsic::cttz, ty); if (bitwidth < 32) { // Extend the input to 32-bits and perform a clz/ctz. diff --git a/lib/ReplaceOpenCLBuiltinPass.cpp b/lib/ReplaceOpenCLBuiltinPass.cpp index accb17d20..fb52370f2 100644 --- a/lib/ReplaceOpenCLBuiltinPass.cpp +++ b/lib/ReplaceOpenCLBuiltinPass.cpp @@ -878,7 +878,8 @@ Value *ReplaceOpenCLBuiltinPass::InsertOpMulExtended(Instruction *InsertPoint, res_neg = Builder.CreateICmpSLT(a_b_xor, ConstantInt::get(Ty, 0, true)); auto F = InsertPoint->getFunction(); - auto abs = Intrinsic::getDeclaration(F->getParent(), Intrinsic::abs, Ty); + auto abs = + Intrinsic::getOrInsertDeclaration(F->getParent(), Intrinsic::abs, Ty); a = Builder.CreateCall(abs, {a, Builder.getInt1(false)}); b = Builder.CreateCall(abs, {b, Builder.getInt1(false)}); } @@ -1190,7 +1191,7 @@ bool ReplaceOpenCLBuiltinPass::replaceCopysign(Function &F) { return replaceCallsWithValue(F, [&F](CallInst *Call) { const auto x = Call->getArgOperand(0); const auto y = Call->getArgOperand(1); - auto intrinsic = Intrinsic::getDeclaration( + auto intrinsic = Intrinsic::getOrInsertDeclaration( F.getParent(), Intrinsic::copysign, Call->getType()); return CallInst::Create(intrinsic->getFunctionType(), intrinsic, {x, y}, "", Call); @@ -1308,8 +1309,8 @@ bool ReplaceOpenCLBuiltinPass::replaceLog1p(Function &F) { auto ArgP1 = BinaryOperator::Create( Instruction::FAdd, ConstantFP::get(Arg->getType(), 1.0), Arg, "", CI); - auto log = - Intrinsic::getDeclaration(F.getParent(), Intrinsic::log, CI->getType()); + auto log = Intrinsic::getOrInsertDeclaration(F.getParent(), Intrinsic::log, + CI->getType()); return CallInst::Create(log, ArgP1, "", CI); }); } @@ -1728,8 +1729,8 @@ bool ReplaceOpenCLBuiltinPass::replaceRotate(Function &F) { // Replace with LLVM's funnel shift left intrinsic because it is more // generic than rotate. - Function *intrinsic = - Intrinsic::getDeclaration(F.getParent(), Intrinsic::fshl, SrcType); + Function *intrinsic = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::fshl, SrcType); return CallInst::Create(intrinsic->getFunctionType(), intrinsic, {SrcValue, SrcValue, RotAmount}, "", CI); }); @@ -2883,7 +2884,8 @@ llvm::Value *ReplaceOpenCLBuiltinPass::createVstoreHalf(llvm::Module &M, CallInst::Create(NewF, Params, "store_halfword_xor_trick", CI); // Return a Nop so the old Call is removed - Function *donothing = Intrinsic::getDeclaration(&M, Intrinsic::donothing); + Function *donothing = + Intrinsic::getOrInsertDeclaration(&M, Intrinsic::donothing); V = CallInst::Create(donothing, {}, "", CI); } @@ -3633,8 +3635,8 @@ bool ReplaceOpenCLBuiltinPass::replaceAddSubSat(Function &F, bool is_signed, : (is_add ? Intrinsic::uadd_sat : Intrinsic::usub_sat); auto a = Call->getArgOperand(0); auto b = Call->getArgOperand(1); - auto intrinsic = Intrinsic::getDeclaration(F.getParent(), intrinsic_type, - Call->getType()); + auto intrinsic = Intrinsic::getOrInsertDeclaration( + F.getParent(), intrinsic_type, Call->getType()); return CallInst::Create(intrinsic->getFunctionType(), intrinsic, {a, b}, "", Call); }); @@ -3821,7 +3823,7 @@ bool ReplaceOpenCLBuiltinPass::replaceCountZeroes(Function &F, bool leading) { return false; return replaceCallsWithValue(F, [&F, leading](CallInst *Call) { - Function *intrinsic = Intrinsic::getDeclaration( + Function *intrinsic = Intrinsic::getOrInsertDeclaration( F.getParent(), leading ? Intrinsic::ctlz : Intrinsic::cttz, Call->getType()); const auto c_false = ConstantInt::getFalse(Call->getContext()); @@ -4104,13 +4106,13 @@ bool ReplaceOpenCLBuiltinPass::replaceRound(Function &F) { clspv_fract_fn->setCallingConv(CallingConv::SPIR_FUNC); } - auto ceil = Intrinsic::getDeclaration(F.getParent(), Intrinsic::ceil, - Call->getType()); - auto floor = Intrinsic::getDeclaration(F.getParent(), Intrinsic::floor, - Call->getType()); - auto fabs = Intrinsic::getDeclaration(F.getParent(), Intrinsic::fabs, - Call->getType()); - auto copysign = Intrinsic::getDeclaration( + auto ceil = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::ceil, Call->getType()); + auto floor = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::floor, Call->getType()); + auto fabs = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::fabs, Call->getType()); + auto copysign = Intrinsic::getOrInsertDeclaration( F.getParent(), Intrinsic::copysign, {Call->getType(), Call->getType()}); IRBuilder<> builder(Call); @@ -4138,21 +4140,21 @@ bool ReplaceOpenCLBuiltinPass::replaceTrigPi(Function &F, auto mul = builder.CreateFMul(x, pi); switch (type) { case Builtins::kSinpi: { - auto func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin, - x->getType()); + auto func = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::sin, x->getType()); return builder.CreateCall(func->getFunctionType(), func, {mul}); } case Builtins::kCospi: { - auto func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos, - x->getType()); + auto func = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::cos, x->getType()); return builder.CreateCall(func->getFunctionType(), func, {mul}); } case Builtins::kTanpi: { - auto sin = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin, - x->getType()); + auto sin = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::sin, x->getType()); auto sin_call = builder.CreateCall(sin->getFunctionType(), sin, {mul}); - auto cos = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos, - x->getType()); + auto cos = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::cos, x->getType()); auto cos_call = builder.CreateCall(cos->getFunctionType(), cos, {mul}); return builder.CreateFDiv(sin_call, cos_call); } @@ -4166,10 +4168,10 @@ bool ReplaceOpenCLBuiltinPass::replaceTrigPi(Function &F, bool ReplaceOpenCLBuiltinPass::replaceSincos(Function &F) { return replaceCallsWithValue(F, [&F](CallInst *Call) { - auto sin_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin, - Call->getType()); - auto cos_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos, - Call->getType()); + auto sin_func = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::sin, Call->getType()); + auto cos_func = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::cos, Call->getType()); IRBuilder<> builder(Call); auto sin = builder.CreateCall(sin_func->getFunctionType(), sin_func, @@ -4183,8 +4185,8 @@ bool ReplaceOpenCLBuiltinPass::replaceSincos(Function &F) { bool ReplaceOpenCLBuiltinPass::replaceExpm1(Function &F) { return replaceCallsWithValue(F, [&F](CallInst *Call) { - auto exp_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::exp, - Call->getType()); + auto exp_func = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::exp, Call->getType()); IRBuilder<> builder(Call); auto exp = builder.CreateCall(exp_func->getFunctionType(), exp_func, @@ -4195,8 +4197,8 @@ bool ReplaceOpenCLBuiltinPass::replaceExpm1(Function &F) { bool ReplaceOpenCLBuiltinPass::replacePown(Function &F) { return replaceCallsWithValue(F, [&F](CallInst *Call) { - auto pow_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::pow, - Call->getType()); + auto pow_func = Intrinsic::getOrInsertDeclaration( + F.getParent(), Intrinsic::pow, Call->getType()); IRBuilder<> builder(Call); auto conv = builder.CreateSIToFP(Call->getArgOperand(1), Call->getType()); diff --git a/test/ProgramScopeConstants/in_storage_buffer_descriptor_map.cl b/test/ProgramScopeConstants/in_storage_buffer_descriptor_map.cl index 2a12d00cb..3568bbf43 100644 --- a/test/ProgramScopeConstants/in_storage_buffer_descriptor_map.cl +++ b/test/ProgramScopeConstants/in_storage_buffer_descriptor_map.cl @@ -1,4 +1,4 @@ -// RUN: clspv %target %s -o %t.spv -module-constants-in-storage-buffer -int8=0 +// RUN: clspv %target %s -o %t.spv -module-constants-in-storage-buffer // RUN: clspv-reflection %t.spv -o %t.map // RUN: FileCheck -check-prefix=MAP %s < %t.map // RUN: spirv-val --target-env vulkan1.0 %t.spv diff --git a/test/UBO/transform_global.cl b/test/UBO/transform_global.cl index 3851c91df..b96870374 100644 --- a/test/UBO/transform_global.cl +++ b/test/UBO/transform_global.cl @@ -54,11 +54,10 @@ __kernel void foo(__global data_type *data, __constant data_type *c_arg, // CHECK-DAG: [[c_var_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Private [[array]] // CHECK-DAG: [[c_var_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Private [[int]] // CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK-DAG: [[undef:%[0-9a-zA-Z_]+]] = OpUndef [[int]] -// CHECK-DAG: [[zero_undef:%[0-9a-zA-Z_]+]] = OpConstantComposite [[data_type]] [[zero]] [[undef]] +// CHECK-DAG: [[zero_zero:%[0-9a-zA-Z_]+]] = OpConstantNull [[data_type]] // CHECK-DAG: [[one:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1 -// CHECK-DAG: [[one_undef:%[0-9a-zA-Z_]+]] = OpConstantComposite [[data_type]] [[one]] [[undef]] -// CHECK-DAG: [[array_const:%[0-9a-zA-Z_]+]] = OpConstantComposite [[array]] [[zero_undef]] [[one_undef]] +// CHECK-DAG: [[one_zero:%[0-9a-zA-Z_]+]] = OpConstantComposite [[data_type]] [[one]] [[zero]] +// CHECK-DAG: [[array_const:%[0-9a-zA-Z_]+]] = OpConstantComposite [[array]] [[zero_zero]] [[one_zero]] // CHECK-DAG: [[c_var:%[0-9a-zA-Z_]+]] = OpVariable [[c_var_ptr]] Private [[array_const]] // CHECK-DAG: [[data]] = OpVariable [[data_ptr]] StorageBuffer // CHECK: [[c_arg]] = OpVariable [[c_arg_ptr]] Uniform