[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 https://github.com/intel/llvm/pull/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 - https://github.com/intel/llvm/pull/5149.

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

[zjin-lcf]

Running the sycl program on an nvidia gpu produces the right result after "ulimit -s unlimited'. However, running the program on an amd gpu (gfx908) causes segfault. Could you reproduce that ? Thanks.

[npmiller]

I'm not having any issues with aop-sycl on gfx908 with the latest:

• intel/llvm: 028055893dc16605d966592f97843173bc2221d7
• HeCBench: 13c7d1d55bcad1575d584b6f1ed4c2a663677b5b

Build command:

make GPU=yes HIP=yes HIP_ARCH=gfx908

Run command:

SYCL_DEVICE_FILTER=hip ./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.44783124
Binonmial             : 0.44880498
European Price        : 0.38443078
==============
elapsed time for each run         : 22.000ms
==============

[zjin-lcf]

Could you run: ./main -runs 100 ?

Thanks

[npmiller]

Oh, that's interesting, I am getting a segfault with that:

SYCL_DEVICE_FILTER=hip ./main -runs 100                                                       
==============                                                                                                                                                               
Num Timesteps         : 100                                                                                                                                                  
Num Paths             : 32K                                                                                                                                                  
Num Runs              : 100               
T                     : 1.000000
S0                    : 3.600000                                                                                                                                             
K                     : 4.000000                                                                                                                                             
r                     : 0.060000                                                                                                                                             
sigma                 : 0.200000                                                                                                                                             
Option Type           : American Put                                                                                                                                         
[1]    40238 segmentation fault  SYCL_DEVICE_FILTER=hip ./main -runs 100
SYCL_DEVICE_FILTER=hip ./main -runs 100  19.99s user 0.21s system 103% cpu 19.450 total                                                                                      

[zjin-lcf]

I remember segfault occurs on an nvidia p100 gpu without "ulimit -s unlimited".

[npmiller]

So looking a bit closer at this I think there might be a bug and/or race condition in the SYCL runtime with regards to command deletion which triggers the segfault. I'll have to dig into it further but it seems like the environment variable SYCL_DISABLE_POST_ENQUEUE_CLEANUP works around the problem:

SYCL_DEVICE_FILTER=hip SYCL_DISABLE_POST_ENQUEUE_CLEANUP=1 ./main -runs 100
==============
Num Timesteps         : 100
Num Paths             : 32K
Num Runs              : 100
T                     : 1.000000
S0                    : 3.600000
K                     : 4.000000
r                     : 0.060000
sigma                 : 0.200000
Option Type           : American Put
==============
GPU Longstaff-Schwartz: 0.44747600
Binonmial             : 0.44880498
European Price        : 0.38443078
==============
elapsed time for each run         : 20.040ms
==============
SYCL_DEVICE_FILTER=hip SYCL_DISABLE_POST_ENQUEUE_CLEANUP=1 ./main -runs 100  37.71s user 0.27s system 104% cpu 36.477 total

With this it will likely use more memory but at least it shouldn't segfault while trying to access a deleted command.

[bader]

@sergey-semenov, FYI.

[zjin-lcf]

Thank you for pointing out the issue.

[pvchupin]

@sergey-semenov, can you confirm please that this is general runtime issue rather than CUDA specific?

[sergey-semenov]

It certainly appears so. There's one known segfault problem related to post-enqueue cleanup and it has to do with execution graph leaf handling. @npmiller Could you please check if this fix (https://github.com/intel/llvm/pull/5417) takes care of this segfault too? @KseniyaTikhomirova FYI

[npmiller]

@sergey-semenov I've just tried it and as far as I can tell it does seem that #5417 fixes the segfault in this test. I've ran the sample successfully several times with 100 iterations and a couple times with 500 iterations with no issues.

@zjin-lcf I suspect this may also fix it for CUDA even without ulimit -s unlimited

[zjin-lcf]

I would like to run the program after #5417 is merged. I hope that is fine. Thank you for your updates and tests.

[npmiller]

@zjin-lcf #5417 has been merged, have you had the time to test this again? Can we close this ticket?

[zjin-lcf]

Yes. I ran the example. Thanks.