Why memcpy kernels are submitted to command queue one by one?

[GHGmc2]

With ze_tracer tool, I noticed that memcpy kernels in the command list are submited to command queue one by one(I set SYCL_PI_LEVEL_ZERO_BATCH_SIZE=32 and did not call wait() after the memcpy), and there are bubbles between the pipeline.

Why can't we submit them in mini-batch, and is there any change to minimize the bubble between them?

• code of host2device memcpy, the same to d2d and d2h memcpy.

#include "CL/sycl.hpp"

constexpr int ALIGN = 64;
constexpr int SIZE = 256;

int main() {
  sycl::queue q(sycl::gpu_selector{});

  for (int i = 0; i < 20; ++i) {
    void *src = sycl::aligned_alloc_host(ALIGN, SIZE, q);
    void *dst = sycl::aligned_alloc_device(ALIGN, SIZE, q);
    q.memcpy(dst, src, SIZE);
    sycl::free(dst, q);
    sycl::free(src, q);
  }

  return 0;
}

• the timeline the ze_tracer generates(w/ flag --chrome-device-stages)
  

[AlexeySachkov]

Tagging @smaslov-intel and @kbsmith-intel here to help answering this question

[smaslov-intel]

That's right, we are not currently batching memory transfers, and this is WIP to start doing so.
This would save on the submissions time, since there would be a single submission then, and is likely to get rid of bubbles too.

[GHGmc2]

Thanks, will memcpy be submitted with other normal kernels in one batch, or only batch of memcpy can be submitted together?