[SYCL][CUDA][HIP] warp misaligned address on CUDA and results mismatch on HIP

[zjin-lcf]

Running the example https://github.com/zjin-lcf/HeCBench/blob/master/aop-sycl/main.cpp built with the CUDA support on a P100 GPU
shows warp misaligned address may be caused by the shared local memory "double4 lsums" in the kernel prepare_svd_kernel<256, PayoffPut>. The SYCL program runs successfully on an Intel GPU.

Did you encounter warp misaligned address when porting a CUDA program ?

Running the example built with the HIP support shows the result does not match the HIP/CUDA version:
To reproduce

make HIP=yes
./main

==============
Num Timesteps         : 100
Num Paths             : 32K
Num Runs              : 1
T                     : 1.000000
S0                    : 3.600000
K                     : 4.000000
r                     : 0.060000
sigma                 : 0.200000
Option Type           : American Put
==============
GPU Longstaff-Schwartz: 0.39776070   (the expected is 0.44783124)

[npmiller]

Hello @zjin-lcf,

I've looked at this on CUDA and there is a bug with the way local kernel arguments are laid out in memory which causes the double4 argument address to be misaligned.

I'm investigating this for a proper fix however as a workaround you can simply re-order the kernel arguments, for example as follows:

diff --git a/aop-sycl/main.cpp b/aop-sycl/main.cpp
index 6a9c9edc..2b8830f0 100644
--- a/aop-sycl/main.cpp
+++ b/aop-sycl/main.cpp
@@ -1046,7 +1046,7 @@ void do_run(queue &q,
     accessor<double4, 1, sycl_read_write, access::target::local> lsums (1, cgh);
     accessor<int, 1, sycl_read_write, access::target::local> lsum (1, cgh);
     accessor<double, 1, sycl_read_write, access::target::local> smem_svds (R_W_MATRICES_SMEM_SLOTS, cgh);
-    cgh.parallel_for<class prepare_svd<Payoff>>(nd_range<1>(gws_prepare_svd, lws_prepare_svd), [=] (nd_item<1> item) {
+    cgh.parallel_for<class prepare_svd<Payoff>>(nd_range<1>(gws_prepare_svd, lws_prepare_svd), [num_paths, payoff, paths, all_out_of_the_money, svds, lsums, smem_svds, scan_input, scan_output, lsum] (nd_item<1> item) {
       prepare_svd_kernel<NUM_THREADS_PER_BLOCK1>(
           item,
           num_paths,

The local memory arguments seem to be placed in a shared memory buffer one after the other, so for example with an int and a double 4 you would have [int, double4], with no padding in between this breaks the alignment, so this patch just swaps them around [double4, int] so double4 gets the correct alignment and it still works for int because it requires a smaller alignment.

I'll update this ticket when I have a proper fix for this.

[zjin-lcf]

The function is called in the following way for the three pointers,

      lsums.get_pointer(),
      lsum.get_pointer(),
      smem_svds.get_pointer(),

Does the compiler have its own way to determine the order ?
Does [...] override the order ? Is it documented ?

Thanks

[npmiller]

The actual "kernel" is the lambda of the parallel for, here there's no issues with the function.

So here the kernel arguments are whatever is captured by the [=] lambda, and the underlying order of the kernel arguments is whatever C++ usually does when capturing variables. Using [a, b] specifies manually what is being captured and in which order, so it allows you to re-order the underlying kernel arguments. This is not really SYCL specific though, just C++ lambdas.

I'm not entirely sure how [=] works but it's possible re-ordering the accessor declaration might also change the order and "fix" it.

In most cases the kernel argument order shouldn't matter and we definitely need to fix this so users don't have to worry about kernel argument order.

[zjin-lcf]

Thank you for explaining the kernel argument order.

[AerialMantis]

We believe the issue for CUDA to be address by #5113, there is a further pull request open for genializing this across both the NVPTX and AMDGCN LLVM backends, resolving this for the HIP as well - #5149.

Edit: I originally closed this issue and then re-opened it as it's not yet addressed for the HIP backend.