[ESIMD] Make sure metadata references new function in ESIMDLowerVecArg… (#8496)

In ESIMDLowerVecArg, we change the type of a function and fix uses to
use the new one. However, we forgot to handle metadata, so any metadata
referencing the old function would end up as null after this pass, as we
also delete the function. Make sure to fixup metadata uses too.

Note this pass only runs with opaque pointers off, so we may not need it
much longer.

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDVecArg.cpp

@@ -213,6 +213,13 @@ Function *ESIMDLowerVecArgPass::rewriteFunc(Function &F) {
     ReplaceInstWithInst(OldInst, NewInst);
   }
 
+  // Make sure to update any metadata as well
+  if(F.isUsedByMetadata()) {
+    // The old function is about to be destroyed, so
+    // just change its type so all replacement works.
+    F.mutateType(NF->getType());
+    ValueAsMetadata::handleRAUW(&F, NF);
+  }
   F.eraseFromParent();
 
   return NF;

llvm/test/SYCLLowerIR/ESIMD/lower_vec_arg_fp_metadata.ll

@@ -0,0 +1,20 @@
+; RUN: opt < %s -passes=ESIMDLowerVecArg -S | FileCheck %s
+
+; Check that we correctly update metadata to reference the new function
+
+%"class.sycl::_V1::vec" = type { <2 x double> }
+
+$foo = comdat any
+
+define weak_odr dso_local spir_kernel void @foo(%"class.sycl::_V1::vec" addrspace(1)* noundef align 16 %_arg_out) local_unnamed_addr comdat {
+entry:
+  ret void
+}
+
+;CHECK: !genx.kernels = !{![[GenXMD:[0-9]+]]}
+!genx.kernels = !{!0}
+
+;CHECK: ![[GenXMD]] = !{void (<2 x double> addrspace(1)*)* @foo, {{.*}}}
+!0 = !{void (%"class.sycl::_V1::vec" addrspace(1)*)* @foo, !"foo", !1, i32 0, i32 0, !1, !2, i32 0, i32 0}
+!1 = !{i32 0}
+!2 = !{!"svmptr_t"}