Add LLVM memmove intrinsic support

[EwanC]

The llvm.memmove intrinsic can exist in input LLVM IR generated from the following OpenCL-C code, but is not currently handled by clspv.

$ cat /tmp/memmove.cl
__kernel void foo(__global int* a, int size) {
  for (int i=0; i < size-1; i++) {
    a[i] = a[i+1];
  }
}

$ clang-19 -S -emit-llvm /tmp/memmove.cl -o /tmp/memmove.ll && cat /tmp/memmove.ll
; ModuleID = '/tmp/memmove.cl'
source_filename = "/tmp/memmove.cl"
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
target triple = "x86_64-pc-linux-gnu"

; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite) uwtable
define dso_local spir_kernel void @foo(ptr nocapture noundef align 4 %0, i32 noundef %1) local_unnamed_addr #0 !kernel_arg_addr_space !7 !kernel_arg_access_qual !8 !kernel_arg_type !9 !kernel_arg_base_type !9 !kernel_arg_type_qual !10 {
  %3 = icmp sgt i32 %1, 1
  br i1 %3, label %4, label %9

4:                                                ; preds = %2
  %5 = add nsw i32 %1, -1
  %6 = getelementptr i8, ptr %0, i64 4
  %7 = zext nneg i32 %5 to i64
  %8 = shl nuw nsw i64 %7, 2
  tail call void @llvm.memmove.p0.p0.i64(ptr align 4 %0, ptr align 4 %6, i64 %8, i1 false), !tbaa !11
  br label %9

9:                                                ; preds = %4, %2
  ret void
}

; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)
declare void @llvm.memmove.p0.p0.i64(ptr nocapture writeonly, ptr nocapture readonly, i64, i1 immarg) #1

This patch expands the memmove intrinsic in ReplaceLLVMIntrinsicPass by copying to an intermediate alloca byte array when the number of bytes is known at compile time. When the number of bytes to copy isn't known at compile time, then we fallback to LLVMs utility for creating loops to do the copy. The same strategy as SPIRVLowerMemmove.cpp in the SPIRV-LLVM-Translator.

[rjodinchr]

That alloca could be huge, are we sure we want it whatever the size?

[EwanC]

That's a great point, I've added a command-line option that controls the limit and if that's exceeded then we fallback to the dynamic implemented. I've defaulted that to 128 but that was an arbitrary choice and happy to change that (or this solution if you don't think a command-line option is appropriate)

[rjodinchr]

Looking good, thank you.
For the default value, I guess whatever works for you right now would be the best start.

[EwanC]

Thanks for rerunning the CI @rjodinchr I can see the clang-format fail log but not the amber-ubuntu-gcc build log, could paste the relevant compiler error so I can address it?

[rjodinchr]

It seems to be an infra error, but somehow I'm not able to re-trigger it at the moment.
We will see when the CI will run with the fixed formatting.