[SYCL] Optimize sycl::detail::memcpy (#13751)

Fixes #7901.

Author: aelovikov-intel

sycl/include/sycl/detail/memcpy.hpp

@@ -8,17 +8,17 @@
 
 #pragma once
 
-#include <cstddef>
+#include <cstring>
 
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
 inline void memcpy(void *Dst, const void *Src, size_t Size) {
-  char *Destination = reinterpret_cast<char *>(Dst);
-  const char *Source = reinterpret_cast<const char *>(Src);
-  for (size_t I = 0; I < Size; ++I) {
-    Destination[I] = Source[I];
-  }
+#ifdef __SYCL_DEVICE_ONLY__
+  __builtin_memcpy(Dst, Src, Size);
+#else
+  std::memcpy(Dst, Src, Size);
+#endif
 }
 } // namespace detail
 } // namespace _V1

sycl/test/.clang-format

@@ -1,2 +1,2 @@
 BasedOnStyle: LLVM
-CommentPragmas: "RUN|FAIL|REQUIRES|UNSUPPORTED|CHECK|expected-"
+CommentPragmas: "RUN|FAIL|REQUIRES|UNSUPPORTED|CHECK|expected-|update_cc_test_checks.py"

sycl/test/check_device_code/vector/vector_as.cpp

@@ -19,36 +19,12 @@ template SYCL_EXTERNAL sycl::vec<int, 4> sycl::vec<float, 4>::as<sycl::vec<int,
 // CHECK-LABEL: define weak_odr dso_local spir_func void @_ZNK4sycl3_V13vecIfLi4EE2asINS1_IiLi4EEEEET_v(
 // CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable sret(%"class.sycl::_V1::vec") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef align 16 dereferenceable_or_null(16) [[THIS:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] comdat align 2 !srcloc [[META5:![0-9]+]] !sycl_fixed_targets [[META6:![0-9]+]] {
 // CHECK-NEXT:  entry:
-// CHECK-NEXT:    br label [[FOR_COND_I:%.*]]
-// CHECK:       for.cond.i:
-// CHECK-NEXT:    [[I_0_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[INC_I:%.*]], [[FOR_BODY_I:%.*]] ]
-// CHECK-NEXT:    [[CMP_I:%.*]] = icmp ult i64 [[I_0_I]], 16
-// CHECK-NEXT:    br i1 [[CMP_I]], label [[FOR_BODY_I]], label [[_ZN4SYCL3_V16DETAIL6MEMCPYEPVPKVM_EXIT:%.*]]
-// CHECK:       for.body.i:
-// CHECK-NEXT:    [[ARRAYIDX_I:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[THIS]], i64 [[I_0_I]]
-// CHECK-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[ARRAYIDX_I]], align 1, !tbaa [[TBAA7:![0-9]+]]
-// CHECK-NEXT:    [[ARRAYIDX1_I:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[AGG_RESULT]], i64 [[I_0_I]]
-// CHECK-NEXT:    store i8 [[TMP0]], ptr addrspace(4) [[ARRAYIDX1_I]], align 1, !tbaa [[TBAA7]]
-// CHECK-NEXT:    [[INC_I]] = add nuw nsw i64 [[I_0_I]], 1
-// CHECK-NEXT:    br label [[FOR_COND_I]], !llvm.loop [[LOOP10:![0-9]+]]
-// CHECK:       _ZN4sycl3_V16detail6memcpyEPvPKvm.exit:
+// CHECK-NEXT:    tail call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) noundef align 16 dereferenceable(16) [[AGG_RESULT]], ptr addrspace(4) noundef align 16 dereferenceable(16) [[THIS]], i64 16, i1 false)
 // CHECK-NEXT:    ret void
 //
 //
 // CHECK-PREVIEW-LABEL: define weak_odr dso_local spir_func void @_ZNK4sycl3_V13vecIfLi4EE2asINS1_IiLi4EEEEET_v(
 // CHECK-PREVIEW-SAME: ptr addrspace(4) dead_on_unwind noalias writable sret(%"class.sycl::_V1::vec") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef align 16 dereferenceable_or_null(16) [[THIS:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] comdat align 2 !srcloc [[META5:![0-9]+]] !sycl_fixed_targets [[META6:![0-9]+]] {
 // CHECK-PREVIEW-NEXT:  entry:
-// CHECK-PREVIEW-NEXT:    br label [[FOR_COND_I:%.*]]
-// CHECK-PREVIEW:       for.cond.i:
-// CHECK-PREVIEW-NEXT:    [[I_0_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[INC_I:%.*]], [[FOR_BODY_I:%.*]] ]
-// CHECK-PREVIEW-NEXT:    [[CMP_I:%.*]] = icmp ult i64 [[I_0_I]], 16
-// CHECK-PREVIEW-NEXT:    br i1 [[CMP_I]], label [[FOR_BODY_I]], label [[_ZN4SYCL3_V16DETAIL6MEMCPYEPVPKVM_EXIT:%.*]]
-// CHECK-PREVIEW:       for.body.i:
-// CHECK-PREVIEW-NEXT:    [[ARRAYIDX_I:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[THIS]], i64 [[I_0_I]]
-// CHECK-PREVIEW-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[ARRAYIDX_I]], align 1, !tbaa [[TBAA7:![0-9]+]]
-// CHECK-PREVIEW-NEXT:    [[ARRAYIDX1_I:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[AGG_RESULT]], i64 [[I_0_I]]
-// CHECK-PREVIEW-NEXT:    store i8 [[TMP0]], ptr addrspace(4) [[ARRAYIDX1_I]], align 1, !tbaa [[TBAA7]]
-// CHECK-PREVIEW-NEXT:    [[INC_I]] = add nuw nsw i64 [[I_0_I]], 1
-// CHECK-PREVIEW-NEXT:    br label [[FOR_COND_I]], !llvm.loop [[LOOP10:![0-9]+]]
-// CHECK-PREVIEW:       _ZN4sycl3_V16detail6memcpyEPvPKvm.exit:
+// CHECK-PREVIEW-NEXT:    tail call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) noundef align 16 dereferenceable(16) [[AGG_RESULT]], ptr addrspace(4) noundef align 16 dereferenceable(16) [[THIS]], i64 16, i1 false)
 // CHECK-PREVIEW-NEXT:    ret void