Update LLVM (#1406)

- Rename Intrinsic::getDeclaration to Intrinsic::getOrInsertDeclaration - Update one test to avoid disabling int8 Bug: crbug.com/373696543
google · Oct 17, 2024 · 419bc4b · 419bc4b
1 parent 40c590a
commit 419bc4b
Show file tree

Hide file tree

Showing 6 changed files with 71 additions and 69 deletions.
diff --git 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
@@ -31,9 +31,9 @@
 #include "Constants.h"
 #include "clspv/Passes.h"
 
+#include "BitcastUtils.h"
 #include "Builtins.h"
 #include "LongVectorLoweringPass.h"
-#include "BitcastUtils.h"
 
 #include <array>
 #include <functional>
@@ -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<Value*>(Idxs.begin(), Idxs.begin() + i);
+    auto CumulativeOldIdxs = ArrayRef<Value *>(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
@@ -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
@@ -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
@@ -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
@@ -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