[SYCL][NFC] Modernize use of StringRef::startswith (#12406)

The method was deprecated in favor of `StringRef::starts_with` in
llvm/llvm-project#75491

Author: AlexeySachkov

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -309,7 +309,7 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
     return std::nullopt;
   StringRef AttrKindStr = Attr.getKindAsString();
   // Early exit if it is not a sycl-* attribute.
-  if (!AttrKindStr.startswith("sycl-"))
+  if (!AttrKindStr.starts_with("sycl-"))
     return std::nullopt;
 
   auto AddFPControlMetadataForWidth = [&](int32_t SPIRVFPControl,

llvm/lib/SYCLLowerIR/ESIMD/ESIMDRemoveHostCode.cpp

@@ -50,8 +50,8 @@ PreservedAnalyses ESIMDRemoveHostCodePass::run(Module &M,
     id::OutputBuffer NameBuf;
     NameNode->print(NameBuf);
     StringRef Name(NameBuf.getBuffer(), NameBuf.getCurrentPosition());
-    if (!Name.startswith("sycl::_V1::ext::intel::esimd::") &&
-        !Name.startswith("sycl::_V1::ext::intel::experimental::esimd::"))
+    if (!Name.starts_with("sycl::_V1::ext::intel::esimd::") &&
+        !Name.starts_with("sycl::_V1::ext::intel::experimental::esimd::"))
       continue;
     SmallVector<BasicBlock *> BBV;
     for (BasicBlock &BB : F) {

llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp

@@ -33,31 +33,31 @@ constexpr char SLM_ALLOCATOR_DTOR_SUFFIX[] = "EED2Ev";
 
 bool isSlmAllocatorConstructor(const Function &F) {
   auto Name = F.getName();
-  return Name.startswith(SLM_ALLOCATOR_CTOR_DTOR_PREFIX) &&
+  return Name.starts_with(SLM_ALLOCATOR_CTOR_DTOR_PREFIX) &&
          Name.endswith(SLM_ALLOCATOR_CTOR_SUFFIX);
 }
 
 bool isSlmAllocatorDestructor(const Function &F) {
   auto Name = F.getName();
-  return Name.startswith(SLM_ALLOCATOR_CTOR_DTOR_PREFIX) &&
+  return Name.starts_with(SLM_ALLOCATOR_CTOR_DTOR_PREFIX) &&
          Name.endswith(SLM_ALLOCATOR_DTOR_SUFFIX);
 }
 
 bool isSlmInit(const Function &F) {
-  return F.getName().startswith(SLM_INIT_PREFIX);
+  return F.getName().starts_with(SLM_INIT_PREFIX);
 }
 
 bool isSlmAlloc(const Function &F) {
-  return F.getName().startswith(SLM_ALLOC_PREFIX);
+  return F.getName().starts_with(SLM_ALLOC_PREFIX);
 }
 
 bool isSlmFree(const Function &F) {
-  return F.getName().startswith(SLM_FREE_PREFIX);
+  return F.getName().starts_with(SLM_FREE_PREFIX);
 }
 
 bool isAssertFail(const Function &F) {
-  return F.getName().startswith("__assert_fail") ||
-         F.getName().startswith("__devicelib_assert_fail");
+  return F.getName().starts_with("__assert_fail") ||
+         F.getName().starts_with("__devicelib_assert_fail");
 }
 
 bool isESIMD(const Function &F) {

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -150,10 +150,10 @@ class ESIMDVerifierImpl {
 
           // We are interested in functions defined in SYCL namespace, but
           // outside of ESIMD namespaces.
-          if (!Name.startswith("sycl::_V1::") ||
-              Name.startswith("sycl::_V1::detail::") ||
-              Name.startswith("sycl::_V1::ext::intel::esimd::") ||
-              Name.startswith("sycl::_V1::ext::intel::experimental::esimd::"))
+          if (!Name.starts_with("sycl::_V1::") ||
+              Name.starts_with("sycl::_V1::detail::") ||
+              Name.starts_with("sycl::_V1::ext::intel::esimd::") ||
+              Name.starts_with("sycl::_V1::ext::intel::experimental::esimd::"))
             continue;
 
           // Check if function name matches any allowed SYCL function name.

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

@@ -704,17 +704,17 @@ static std::string mangleFunction(StringRef FunctionName) {
   // These functions are defined as extern "C" which Demangler that is used
   // fails to handle properly.
   if (isDevicelibFunction(FunctionName)) {
-    if (FunctionName.startswith("__devicelib_ConvertFToBF16INTEL")) {
+    if (FunctionName.starts_with("__devicelib_ConvertFToBF16INTEL")) {
       return (Twine("_Z31") + FunctionName + "RKf").str();
     }
-    if (FunctionName.startswith("__devicelib_ConvertBF16ToFINTEL")) {
+    if (FunctionName.starts_with("__devicelib_ConvertBF16ToFINTEL")) {
       return (Twine("_Z31") + FunctionName + "RKt").str();
     }
   }
   // Every inserted vstore gets its own function with the same name,
   // so they are mangled with ".[0-9]+". Just use the
   // raw name to pass through the demangler.
-  if (FunctionName.startswith(ESIMD_INSERTED_VSTORE_FUNC_NAME))
+  if (FunctionName.starts_with(ESIMD_INSERTED_VSTORE_FUNC_NAME))
     return ESIMD_INSERTED_VSTORE_FUNC_NAME;
   return FunctionName.str();
 }
@@ -1968,36 +1968,36 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F,
       // process ESIMD builtins that go through special handling instead of
       // the translation procedure
 
-      if (Name.startswith("__esimd_svm_block_ld") ||
-          Name.startswith("__esimd_slm_block_ld")) {
-        translateBlockLoad(*CI, Name.startswith("__esimd_slm_block_ld"));
+      if (Name.starts_with("__esimd_svm_block_ld") ||
+          Name.starts_with("__esimd_slm_block_ld")) {
+        translateBlockLoad(*CI, Name.starts_with("__esimd_slm_block_ld"));
         ToErase.push_back(CI);
         continue;
       }
-      if (Name.startswith("__esimd_svm_block_st") ||
-          Name.startswith("__esimd_slm_block_st")) {
-        translateBlockStore(*CI, Name.startswith("__esimd_slm_block_st"));
+      if (Name.starts_with("__esimd_svm_block_st") ||
+          Name.starts_with("__esimd_slm_block_st")) {
+        translateBlockStore(*CI, Name.starts_with("__esimd_slm_block_st"));
         ToErase.push_back(CI);
         continue;
       }
-      if (Name.startswith("__esimd_gather_ld") ||
-          Name.startswith("__esimd_slm_gather_ld")) {
-        translateGatherLoad(*CI, Name.startswith("__esimd_slm_gather_ld"));
+      if (Name.starts_with("__esimd_gather_ld") ||
+          Name.starts_with("__esimd_slm_gather_ld")) {
+        translateGatherLoad(*CI, Name.starts_with("__esimd_slm_gather_ld"));
         ToErase.push_back(CI);
         continue;
       }
 
-      if (Name.startswith("__esimd_nbarrier_init")) {
+      if (Name.starts_with("__esimd_nbarrier_init")) {
         translateNbarrierInit(*CI);
         ToErase.push_back(CI);
         continue;
       }
-      if (Name.startswith("__esimd_pack_mask")) {
+      if (Name.starts_with("__esimd_pack_mask")) {
         translatePackMask(*CI);
         ToErase.push_back(CI);
         continue;
       }
-      if (Name.startswith("__esimd_unpack_mask")) {
+      if (Name.starts_with("__esimd_unpack_mask")) {
         translateUnPackMask(*CI);
         ToErase.push_back(CI);
         continue;
@@ -2006,21 +2006,21 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F,
       // those globals marked as genx_volatile, We can translate
       // them directly into generic load/store inst. In this way
       // those insts can be optimized by llvm ASAP.
-      if (Name.startswith("__esimd_vload")) {
+      if (Name.starts_with("__esimd_vload")) {
         if (translateVLoad(*CI, GVTS)) {
           ToErase.push_back(CI);
           continue;
         }
       }
-      if (Name.startswith("__esimd_vstore")) {
+      if (Name.starts_with("__esimd_vstore")) {
         if (translateVStore(*CI, GVTS)) {
           ToErase.push_back(CI);
           continue;
         }
       }
 
       if (Name.empty() ||
-          (!Name.startswith(ESIMD_INTRIN_PREF1) && !isDevicelibFunction(Name)))
+          (!Name.starts_with(ESIMD_INTRIN_PREF1) && !isDevicelibFunction(Name)))
         continue;
       // this is ESIMD intrinsic - record for later translation
       ESIMDIntrCalls.push_back(CI);
@@ -2045,7 +2045,7 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F,
       }
 
       if (!isa<GlobalVariable>(SpirvGlobal) ||
-          !SpirvGlobal->getName().startswith(SPIRV_INTRIN_PREF))
+          !SpirvGlobal->getName().starts_with(SPIRV_INTRIN_PREF))
         continue;
 
       auto PrefLen = StringRef(SPIRV_INTRIN_PREF).size();

llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp

@@ -88,7 +88,7 @@ std::pair<Value *, Value *>
 getHelperAndInvokeeIfInvokeSimdCall(const CallInst *CI) {
   Function *F = CI->getCalledFunction();
 
-  if (F && F->getName().startswith(esimd::INVOKE_SIMD_PREF)) {
+  if (F && F->getName().starts_with(esimd::INVOKE_SIMD_PREF)) {
     return {CI->getArgOperand(0), CI->getArgOperand(1)};
   }
   return {nullptr, nullptr};
@@ -426,7 +426,7 @@ PreservedAnalyses SYCLLowerInvokeSimdPass::run(Module &M,
 
   for (Function &F : M) {
     if (!F.isDeclaration() ||
-        !F.getName().startswith(esimd::INVOKE_SIMD_PREF)) {
+        !F.getName().starts_with(esimd::INVOKE_SIMD_PREF)) {
       continue;
     }
     SmallVector<User *, 4> Users(F.users());

llvm/lib/SYCLLowerIR/MutatePrintfAddrspace.cpp

@@ -78,7 +78,7 @@ SYCLMutatePrintfAddrspacePass::run(Module &M, ModuleAnalysisManager &MAM) {
   for (Function &F : M) {
     if (!F.isDeclaration())
       continue;
-    if (!F.getName().startswith("_Z18__spirv_ocl_printf"))
+    if (!F.getName().starts_with("_Z18__spirv_ocl_printf"))
       continue;
     if (F.getArg(0)->getType() == CASLiteralType)
       // No need to replace the literal type and its printf users

llvm/lib/SYCLLowerIR/RenameKernelSYCLNativeCPU.cpp

@@ -23,7 +23,7 @@ static bool isSpirvSyclBuiltin(StringRef FName) {
   // now skip the digits
   FName = FName.drop_while([](char C) { return std::isdigit(C); });
 
-  return FName.startswith("__spirv_") || FName.startswith("__sycl_");
+  return FName.starts_with("__spirv_") || FName.starts_with("__sycl_");
 }
 
 PreservedAnalyses

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -508,7 +508,7 @@ bool isSpirvSyclBuiltin(StringRef FName) {
   // now skip the digits
   FName = FName.drop_while([](char C) { return std::isdigit(C); });
 
-  return FName.startswith("__spirv_") || FName.startswith("__sycl_");
+  return FName.starts_with("__spirv_") || FName.starts_with("__sycl_");
 }
 
 bool isEntryPoint(const Function &F) {

llvm/lib/SYCLLowerIR/SYCLPropagateJointMatrixUsage.cpp

@@ -41,7 +41,7 @@ bool isSpirvSyclBuiltin(StringRef FName) {
   // now skip the digits
   FName = FName.drop_while([](char C) { return std::isdigit(C); });
 
-  return FName.startswith("__spirv_") || FName.startswith("__sycl_");
+  return FName.starts_with("__spirv_") || FName.starts_with("__sycl_");
 }
 
 bool isEntryPoint(const Function &F) {