search

icl-utk-edu / heffte

BSD 3-Clause "New" or "Revised" License
20 stars 15 forks source link

Fix Stock FFT for CUDA 12.2 on Perlmutter #54

Closed ax3l closed 3 weeks ago

ax3l commented 3 weeks ago

Relying on the compiler intrinsic macro __AVX__ leads to compile errors when including the public heFFTe header in CUDA code of WarpX of the form:

/usr/lib64/gcc/x86_64-suse-linux/12/include/avx512fp16intrin.h(38): error: vector_size attribute requires an arithmetic or enum type
  typedef __half __v8hf __attribute__ ((__vector_size__ (16)));
                                        ^
...
/usr/lib64/gcc/x86_64-suse-linux/12/include/avx512fp16intrin.h(62): error: more than one conversion function from "__half" to "<error-type>" applies:
            function "__half::operator __half_raw() const" (declared at line 309 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator float() const" (declared at line 337 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator signed char() const" (declared at line 436 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator unsigned char() const" (declared at line 443 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator char() const" (declared at line 451 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator short() const" (declared at line 476 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator unsigned short() const" (declared at line 483 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator int() const" (declared at line 490 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator unsigned int() const" (declared at line 497 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator long() const" (declared at line 504 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator unsigned long() const" (declared at line 528 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator long long() const" (declared at line 554 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator unsigned long long() const" (declared at line 561 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
            function "__half::operator bool() const" (declared at line 593 of /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/include/cuda_fp16.hpp)
    return __extension__ (__m128h)(__v8hf){ __A0, __A1, __A2, __A3,
                                            ^
...

Since heFFTe has explicit control for AVX via macros, we can rely on those. Most GPU users will disable AVX support anyway when building for GPU.

System details

Perlmutter (NERSC) HPE machine with:

Currently Loaded Modules:
  1) libfabric/1.15.2.0                      4) gcc-native/12.3          7) gpu/1.0                  10) cray-libsci/23.12.5 (math)  13) craype-x86-milan      (cpe)  16) cmake/3.24.3                (buildtools)
  2) craype-network-ofi                      5) perftools-base/23.12.0   8) cray-dsmml/0.2.2         11) PrgEnv-gnu/8.5.0    (cpe)   14) craype-accel-nvidia80        17) cray-hdf5-parallel/1.12.2.9 (io)
  3) xpmem/2.6.2-2.5_2.38__gd067c3f.shasta   6) cpe/23.12                9) cray-mpich/8.1.28 (mpi)  12) craype/2.7.30       (c)     15) cudatoolkit/12.2      (g)    18) cray-python/3.11.5          (dev)

Kudos

Thanks to @aeriforme and @Haavaan for reporting this. First seen in https://github.com/ECP-WarpX/WarpX/pull/4937

ax3l commented 3 weeks ago

attn @mkstoyanov

mkstoyanov commented 3 weeks ago

I'll test and merge some time tomorrow. Thanks for fixing this.

ax3l commented 3 weeks ago

Fully understand, thanks a lot 😊 have a nice evening!

ax3l commented 3 weeks ago

Thank you for merging! :pray: