SYCL Poor FFT Performance Compared to CUDA in Double Precision

yencal · 6 July 2024 02:01

Describe the bug

The double precision performance of fft-sycl lags behind fft-cuda (HeCBench), achieving only 70% of CUDA’s performance. Profiling with nsys (nsys nvprof --print-gpu-trace) reveals that the SYCL code utilizes more registers (255 registers/thread) than its CUDA counterpart (72 registers/thread), leading to register spills as observed with -Xcuda-ptxas --verbose.

Interestingly, the single precision version of the code shows comparable performance between SYCL and CUDA. Although SYCL uses (83 registers/thread) and cuda uses (48 registers/thread), there is no register spill for both.

Seeking recommendations to enhance the performance of the double precision SYCL implementation would be greatly appreciated.

To reproduce

Clone GitHub - zjin-lcf/HeCBench

For sycl (double precission):

cd HeCBench/src/fft-sycl
make CUDA=yes CUDA_ARCH=sm_80 GCC_TOOLCHAIN=""
Run code: ./main 3 1000

For cuda (double precission):

cd HeCBench/src/fft-cuda
make ARCH=sm_80
Run code: ./main 3 1000

To build the single precision versions, ensure the SINGLE_PRECISION preprocessor directive is defined by editing main.cpp or Makefile.

Here is the result of my run on NVIDIA A100:

For Double Precision

SYCL:

~/HeCBench/src/fft-sycl$ ./main 3 1000
used_bytes=268435456, n_cmplx=16777216
FFT PASS
iFFT PASS
Average kernel execution time 0.00117844 (s)

CUDA:

~/HeCBench/src/fft-cuda$ ./main 3 1000
used_bytes=268435456, n_cmplx=1.67772e+07
FFT PASS
iFFT PASS
Average kernel execution time 0.000832436 (s)

For Single Precision

SYCL:

~/HeCBench/src/fft-sycl$ ./main 3 1000
used_bytes=268435456, n_cmplx=33554432
FFT PASS
iFFT PASS
Average kernel execution time 0.000819632 (s)

CUDA:

~/HeCBench/src/fft-cuda$ ./main 3 1000
used_bytes=268435456, n_cmplx=3.35544e+07
FFT PASS
iFFT PASS
Average kernel execution time 0.000798072 (s)

Environment

OS: Ubuntu 23.10
NVIDIA A100
clang++ --version:

$ clang++ --version
clang version 19.0.0git (https://github.com/intel/llvm 666cf66258363ba1c416d054cab38c85c04fe389)
Target: x86_64-unknown-linux-gnu
Thread model: posix
Build config: +assertions

sycl-ls --verbose:

$ sycl-ls --verbose
[opencl:fpga][opencl:0] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2  [2024.17.5.0.08_160000.xmain-hotfix]
[opencl:cpu][opencl:1] Intel(R) OpenCL,            Intel(R) Xeon(R) CPU @ 2.20GHz OpenCL 3.0 (Build 0) [2024.17.5.0.08_160000.xmain-hotfix]
[cuda:gpu][cuda:0] NVIDIA CUDA BACKEND, NVIDIA A100-SXM4-40GB 8.0 [CUDA 12.4]

Platforms: 3
Platform [#1]:
    Version  : OpenCL 1.2 Intel(R) FPGA SDK for OpenCL(TM), Version 20.3
    Name     : Intel(R) FPGA Emulation Platform for OpenCL(TM)
    Vendor   : Intel(R) Corporation
    Devices  : 1
        Device [#0]:
        Type              : fpga
        Version           : OpenCL 1.2 
        Name              : Intel(R) FPGA Emulation Device
        Vendor            : Intel(R) Corporation
        Driver            : 2024.17.5.0.08_160000.xmain-hotfix
        Num SubDevices    : 0
        Num SubSubDevices : 0
        Aspects           : accelerator fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations usm_atomic_host_allocations usm_atomic_shared_allocations ext_oneapi_srgb ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistic_group ext_oneapi_tangle_group ext_intel_fpga_task_sequence ext_oneapi_private_alloca
        info::device::sub_group_sizes: 4 8 16 32 64
        Architecture: unknown
Platform [#2]:
    Version  : OpenCL 3.0 LINUX
    Name     : Intel(R) OpenCL
    Vendor   : Intel(R) Corporation
    Devices  : 1
        Device [#1]:
        Type              : cpu
        Version           : OpenCL 3.0 (Build 0)
        Name              :            Intel(R) Xeon(R) CPU @ 2.20GHz
        Vendor            : Intel(R) Corporation
        Driver            : 2024.17.5.0.08_160000.xmain-hotfix
        Num SubDevices    : 0
        Num SubSubDevices : 0
        Aspects           : cpu fp16 fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations usm_system_allocations usm_atomic_host_allocations usm_atomic_shared_allocations atomic64 ext_oneapi_srgb ext_oneapi_native_assert ext_intel_legacy_image ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistic_group ext_oneapi_tangle_group ext_oneapi_private_alloca
        info::device::sub_group_sizes: 4 8 16 32 64
        Architecture: x86_64
Platform [#3]:
    Version  : CUDA 12.4
    Name     : NVIDIA CUDA BACKEND
    Vendor   : NVIDIA Corporation
    Devices  : 1
        Device [#0]:
        Type              : gpu
        Version           : 8.0
        Name              : NVIDIA A100-SXM4-40GB
        Vendor            : NVIDIA Corporation
        Driver            : CUDA 12.4
        UUID              : 9031191727913913712018114216122172201180135
        Num SubDevices    : 0
        Num SubSubDevices : 0
        Aspects           : gpu fp16 fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations ext_intel_pci_address usm_atomic_shared_allocations atomic64 ext_intel_device_info_uuid ext_oneapi_bfloat16_math_functions ext_intel_free_memory ext_intel_device_id ext_intel_memory_clock_rate ext_intel_memory_bus_widthImages are not fully supported by the CUDA BE, their support is disabled by default. Their partial support can be activated by setting SYCL_PI_CUDA_ENABLE_IMAGE_SUPPORT environment variable at runtime.
 ext_oneapi_bindless_images ext_oneapi_bindless_images_shared_usm ext_oneapi_bindless_images_2d_usm ext_oneapi_interop_memory_import ext_oneapi_interop_semaphore_import ext_oneapi_mipmap ext_oneapi_mipmap_anisotropy ext_oneapi_mipmap_level_reference ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistic_group ext_oneapi_graph ext_oneapi_limited_graph ext_oneapi_cubemap ext_oneapi_cubemap_seamless_filtering ext_oneapi_bindless_sampled_image_fetch_1d_usm ext_oneapi_bindless_sampled_image_fetch_2d_usm ext_oneapi_bindless_sampled_image_fetch_2d ext_oneapi_bindless_sampled_image_fetch_3d ext_oneapi_queue_profiling_tag
        info::device::sub_group_sizes: 32
        Architecture: nvidia_gpu_sm_80
default_selector()      : gpu, NVIDIA CUDA BACKEND, NVIDIA A100-SXM4-40GB 8.0 [CUDA 12.4]
accelerator_selector()  : fpga, Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2  [2024.17.5.0.08_160000.xmain-hotfix]
cpu_selector()          : cpu, Intel(R) OpenCL,            Intel(R) Xeon(R) CPU @ 2.20GHz OpenCL 3.0 (Build 0) [2024.17.5.0.08_160000.xmain-hotfix]
gpu_selector()          : gpu, NVIDIA CUDA BACKEND, NVIDIA A100-SXM4-40GB 8.0 [CUDA 12.4]
custom_selector(gpu)    : gpu, NVIDIA CUDA BACKEND, NVIDIA A100-SXM4-40GB 8.0 [CUDA 12.4]
custom_selector(cpu)    : cpu, Intel(R) OpenCL,            Intel(R) Xeon(R) CPU @ 2.20GHz OpenCL 3.0 (Build 0) [2024.17.5.0.08_160000.xmain-hotfix]
custom_selector(acc)    : fpga, Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2  [2024.17.5.0.08_160000.xmain-hotfix]

rbielski · 8 July 2024 12:13

Hi @yencal,
I could reproduce the issue and we will be looking into why it’s happening. Right now, I can offer a workaround forcing the max number of registers in the PTX->assembly compilation. Try adding:

-Xcuda-ptxas --maxrregcount=64

in the compiler invocation. This recovered the performance for me on the A100 GPU. Please let us know if you see the same.

Thanks,
Rafal

rbielski · 8 July 2024 15:38

Hi @yencal,
we found the reason for the difference and it’s because the code is different. These functions:

github.com

zjin-lcf/HeCBench/blob/4c0ab5cb9716b29ac3b9f79d880adae0864f093c/src/fft-sycl/main.cpp#L29-L36


      
          inline T2 exp_i( T phi ) {
            return (T2){ sycl::cos(phi), sycl::sin(phi) };
          }
          
          inline T2 cmplx_mul( T2 a, T2 b ) { return (T2){ a.x()*b.x()-a.y()*b.y(), a.x()*b.y()+a.y()*b.x() }; }
          inline T2 cm_fl_mul( T2 a, T  b ) { return (T2){ b*a.x(), b*a.y() }; }
          inline T2 cmplx_add( T2 a, T2 b ) { return (T2){ a.x() + b.x(), a.y() + b.y() }; }
          inline T2 cmplx_sub( T2 a, T2 b ) { return (T2){ a.x() - b.x(), a.y() - b.y() }; }

are declared inline in the SYCL version but are not inlined in the CUDA version:

github.com

zjin-lcf/HeCBench/blob/4c0ab5cb9716b29ac3b9f79d880adae0864f093c/src/fft-cuda/main.cu#L29-L41


      
          __host__ __device__
          T2 exp_i( T phi ) {
            return (T2){ cos(phi), sin(phi) };
          }
          
          __host__ __device__
          T2 cmplx_mul( T2 a, T2 b ) { return (T2){ a.x*b.x-a.y*b.y, a.x*b.y+a.y*b.x }; }
          __host__ __device__
          T2 cm_fl_mul( T2 a, T  b ) { return (T2){ b*a.x, b*a.y }; }
          __host__ __device__
          T2 cmplx_add( T2 a, T2 b ) { return (T2){ a.x + b.x, a.y + b.y }; }
          __host__ __device__
          T2 cmplx_sub( T2 a, T2 b ) { return (T2){ a.x - b.x, a.y - b.y }; }

When I changed the inline to [[clang::noinline]], I see the SYCL version matching the CUDA performance.

Hope this helps!
Thanks,
Rafal

yencal · 9 July 2024 16:31

Hi @rbielski,
Thanks for the insight. I did notice the inline functions in the sycl code. I removed them but that did not work. But using [[clang::noinline]] to force the compiler not to inline does the trick (90% cuda performance even with sycl::ext::oneapi::property::queue::discard_events{}). It would be nice if the sycl compiler could make better decisions about when to inline or not inline. Instead of having to manually force [[clang::noinline]].

Thanks for your help.

Caleb Yenusah

yencal · 9 July 2024 20:55

Hi @rbielski
I actually noticed that only the exp_i function should not be inlined ([[clang::noinline]]) to match the cuda performance. If [[clang::noinline]] is applied to all the other functions, the performance drops substantially (~1% cuda performance).

Here is what gave the best performance:

[[clang::noinline]] T2 exp_i( T phi ) { 
  return (T2){ sycl::cos(phi), sycl::sin(phi) };
}

inline T2 cmplx_mul( T2 a, T2 b ) { return (T2){ a.x()*b.x()-a.y()*b.y(), a.x()*b.y()+a.y()*b.x() }; }

inline T2 cm_fl_mul( T2 a, T  b ) { return (T2){ b*a.x(), b*a.y() }; }

inline T2 cmplx_add( T2 a, T2 b ) { return (T2){ a.x() + b.x(), a.y() + b.y() }; }

inline T2 cmplx_sub( T2 a, T2 b ) { return (T2){ a.x() - b.x(), a.y() - b.y() }; }

yencal · 10 July 2024 19:38

Hi @rbielski,

Upon further analysis, although the CUDA versions of the functions did not explicitly specify inline, these functions were indeed inlined, as evident from the PTX. Similarly, in the SYCL versions, the functions are also inlined regardless of whether they are explicitly marked as inline. Therefore, the difference in performance between the two versions is not due to the inline keyword.

The issue seems to lie with the sycl::cos and sycl::sin functions. When these functions are inlined, they significantly increase register usage, which adversely affects performance. That is why when [[clang::noinline]] is applied to the exp_i function, register usage is reduced, improving performance in the SYCL version.

Interestingly, when cos and sin are inlined in the CUDA version, the register usage does not increase dramatically as it does in the SYCL version.

Thank you,
Caleb

zjin-lcf · 15 October 2024 21:11

Thank you for your analysis of the benchmark.

zjin-lcf · 15 October 2024 21:42

@rbielski
Any comments on the impacts of the sycl::sin()/sycl::cos() functions ?

Topic		Replies	Views
Poor performance on matrix multiplication oneAPI for NVIDIA GPUs	7	954	8 October 2024
Working example on Windows 10 + Nvidia SYCL development	11	2018	19 April 2019
Poor memory performance? SYCL development	4	1295	22 May 2019
Performance FAQ and best practices SYCL performance	3	36	15 July 2024
Utilizing Preallocated Device Memory oneAPI for NVIDIA GPUs	6	89	11 July 2024