[CUDA] NVPTX expands an aligned 16- or 32-byte memcpy to ld/st.global.b64

[zjin-lcf]

Describe the bug

On the CUDA target, an aligned sycl::double4 or sycl::uint4 assignment, and an assignment of an alignas(32) struct of four doubles, becomes llvm.memcpy of 16 or 32 bytes. For sm_80, that memcpy is expanded to interleaved ld.global.b64 / st.global.b64.

A Clang <4 x i32> load and store in the same translation unit is lowered to ld.global.v4.b32 / st.global.v4.b32. sycl::vec hits the memcpy path because its storage is a plain array (__SYCL_USE_LIBSYCL8_VEC_IMPL is 0).

The copy is correct. A 32 MB A100 peer put reaches 20.3 GB/s with the 8-byte stores and 86.5 GB/s when both 16-byte vector loads are issued before the two vector stores.

To reproduce

#include <sycl/sycl.hpp>

using u4native = unsigned int __attribute__((ext_vector_type(4)));

int main() {
  sycl::queue q(sycl::gpu_selector_v);
  constexpr size_t n = 1024;
  auto *d4_dst = sycl::malloc_device<sycl::double4>(n, q);
  auto *d4_src = sycl::malloc_device<sycl::double4>(n, q);
  auto *nv_dst = static_cast<u4native *>(sycl::malloc_device(n * sizeof(u4native), q));
  auto *nv_src = static_cast<u4native *>(sycl::malloc_device(n * sizeof(u4native), q));

  q.parallel_for(sycl::nd_range<1>(256, 256), [=](sycl::nd_item<1> it) {
    const size_t i = it.get_local_id(0);
    d4_dst[i] = d4_src[i];
  }).wait();

  q.parallel_for(sycl::nd_range<1>(256, 256), [=](sycl::nd_item<1> it) {
    const size_t i = it.get_local_id(0);
    const u4native a = nv_src[i * 2];
    const u4native b = nv_src[i * 2 + 1];
    nv_dst[i * 2] = a;
    nv_dst[i * 2 + 1] = b;
  }).wait();
  return 0;
}

clang++ -std=c++17 -O3 -fsycl -fsycl-targets=nvptx64-nvidia-cuda   -Xsycl-target-backend --cuda-gpu-arch=sm_80 -save-temps=obj -c repro.cpp
llvm-dis *-sycl-nvptx64-nvidia-cuda-sm_80.bc -o repro.ll
llc -march=nvptx64 -mcpu=sm_80 repro.ll -o repro.ptx

The sycl::double4 kernel is a 32-byte llvm.memcpy and this PTX:

ld.global.b64 [%rd+24], ...
st.global.b64 [%rd+24], ...
ld.global.b64 [%rd+16], ...
st.global.b64 [%rd+16], ...
ld.global.b64 [%rd+8],  ...
st.global.b64 [%rd+8],  ...
ld.global.b64 [%rd],    ...
st.global.b64 [%rd],    ...

The native <4 x i32> kernel is:

ld.global.v4.b32 {%r3, %r4, %r5, %r6}, [%rd9+-16];
ld.global.v4.b32 {%r7, %r8, %r9, %r10}, [%rd9];
st.global.v4.b32 [%rd8+-16], {%r3, %r4, %r5, %r6};
st.global.v4.b32 [%rd8], {%r7, %r8, %r9, %r10};

Expected: the aligned 32-byte copy uses 16-byte vector loads and stores, with both loads before the two stores. sycl::uint4 (16-byte memcpy) is expanded to two b64 pairs in the same way.

Environment

• OS: Linux 4.18.0-553.150.1.el8_10.x86_64
• Target device and vendor: NVIDIA A100-SXM4-40GB
• DPC++ version: 7.1.0 pre-release, clang 23.0.0git (https://github.com/intel/llvm.git fbf7d1fd2cbf40b346e8109a81b3a13cf87d7ebd)
• Dependencies version: driver 610.43.02, CUDA 13.3 (sycl-ls: NVIDIA CUDA BACKEND, NVIDIA A100-SXM4-40GB 8.0 [CUDA 13.3])

Additional context

CUDA 13.3 nvcc lowers the same 32-byte struct to two ld.global.v4.u32 followed by two st.global.v4.u32. This is separate from #6583.