CUDA with C++20, main branch (2024.03.21.) (!70017) · Merge requests · atlas / athena

Attila Krasznahorkay requested to merge akraszna/athena:CUDACxx20-main-20240320 into main Mar 21, 2024

As I mentioned in some meetings, our current nightlies still build all CUDA source files with the C++14 standard. 😱 This MR is meant to fix that.

Currently we only build CUDA code as part of Athena, so I only added a CMAKE_CUDA_STANDARD setting for that project.

To demonstrate the advantages that this update brings, I decided to simplify one of the example CUDA algorithms. Since now that the host and CUDA code are both being built with the C++20 standard, we don't need to strictly separate "Athena" and "CUDA" code.

However to do this, I had to do 2 things:

I added the EIGEN_NO_CUDA definition to the build of AthExCUDA. This flag makes it impossible to use Eigen structures in CUDA device code. (Which the example doesn't currently do anyway.) But it also silences a ton of warnings coming from CUDA, because of the mistakes in the Eigen version that we use. Warnings like:

...
[ 91%] Building CUDA object Control/AthenaExamples/AthExCUDA/CMakeFiles/AthExCUDA.dir/src/TrackParticleCalibratorExampleAlg.cu.o
/cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/eigen/3.4.0-6ce89/x86_64-el9-gcc13-opt/include/eigen3/Eigen/src/Core/DenseStorage.h(223): warning #20012-D: __host__ annotation is ignored on a function("DenseStorage") that is explicitly defaulted on its first declaration
     __attribute__((host)) __attribute__((device)) 
                    ^

Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

/cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/eigen/3.4.0-6ce89/x86_64-el9-gcc13-opt/include/eigen3/Eigen/src/Core/DenseStorage.h(223): warning #20012-D: __device__ annotation is ignored on a function("DenseStorage") that is explicitly defaulted on its first declaration
     __attribute__((host)) __attribute__((device)) 
                                          ^
...

CUDA really didn't like the shenanigans going on in Tracking/TrkEvent/TrkParametersBase. Clearly it is trying to instantiate templates a little differently than GCC and Clang do. Leading to:

...
[17/50] Building CUDA object Control/AthenaExamples/AthExCUDA/CMakeFiles/AthExCUDA.dir/src/TrackParticleCalibratorExampleAlg.cu.o
FAILED: Control/AthenaExamples/AthExCUDA/CMakeFiles/AthExCUDA.dir/src/TrackParticleCalibratorExampleAlg.cu.o 
/cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/cuda/12.4/x86_64-el9-gcc13-opt/bin/nvcc -forward-unknown-to-host-compiler -ccbin=/cvmfs/sft.cern.ch/lcg/releases/gcc/13.1.0-b3d18/x86_64-el9/bin/g++ -DATLAS_PACKAGE_NAME=\"AthExCUDA\" -DAthExCUDA_EXPORTS -DBOOST_FILESYSTEM_DYN_LINK -DBOOST_FILESYSTEM_NO_LIB -DBOOST_REGEX_DYN_LINK -DBOOST_REGEX_NO_LIB -DBOOST_SYSTEM_DYN_LINK -DBOOST_SYSTEM_NO_LIB -DBOOST_THREAD_DYN_LINK -DBOOST_THREAD_NO_LIB -DCLHEP_ABS_DEFINED -DCLHEP_MAX_MIN_DEFINED -DCLHEP_SQR_DEFINED -DEIGEN_NO_CUDA -DFMT_SHARED -DGAUDI_V20_COMPAT -DJSON_DIAGNOSTICS=0 -DJSON_USE_IMPLICIT_CONVERSIONS=1 -DVECMEM_DEBUG_MSG_LVL=0 -DVECMEM_HAVE_PMR_MEMORY_RESOURCE -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthenaExamples/AthExCUDA -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthenaBaseComps -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthenaKernel -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/CxxUtils -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/DataModelRoot -I/cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/Boost/1.82.0-fbfc9/x86_64-el9-gcc13-opt/include -I/cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/fmt/10.0.0/x86_64-el9-gcc13-opt/include -I/cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/jsonmcpp/3.10.5/x86_64-el9-gcc13-opt/include -I/home/krasznaa/ATLAS/projects/cuda/athena/Event/xAOD/xAODEventInfo -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthContainers -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthContainersInterfaces -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthLinks -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/SGTools -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthAllocators -I/home/krasznaa/ATLAS/projects/cuda/athena/Event/xAOD/xAODCore -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/RootUtils -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/StoreGate -I/home/krasznaa/ATLAS/projects/cuda/athena/Database/PersistentDataModel -I/home/krasznaa/ATLAS/projects/cuda/athena/Event/xAOD/xAODTracking -I/home/krasznaa/ATLAS/projects/cuda/athena/DetectorDescription/GeoPrimitives -I/home/krasznaa/ATLAS/projects/cuda/athena/Event/EventPrimitives -I/home/krasznaa/ATLAS/projects/cuda/athena/Event/xAOD/xAODBase -I/home/krasznaa/ATLAS/projects/cuda/athena/Event/xAOD/xAODMeasurementBase -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkNeutralParameters -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkDetDescr/TrkSurfaces -I/home/krasznaa/ATLAS/projects/cuda/athena/DetectorDescription/Identifier -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkParametersBase -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkEventPrimitives -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkDetDescr/TrkDetDescrUtils -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkDetDescr/TrkDetElementBase -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkParameters -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkTrack -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkMaterialOnTrack -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkMeasurementBase -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkTrackLink -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkTrackSummary -I/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/VxVertex -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthCUDA/AthCUDAInterfaces -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthCUDA/AthCUDAKernel -I/home/krasznaa/ATLAS/projects/cuda/athena/Control/AthCUDA/AthCUDACore -isystem /cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/2024-03-20T2101/AthenaExternals/25.0.2/InstallArea/x86_64-el9-gcc13-opt/include -isystem /cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/Boost/1.82.0/x86_64-el9-gcc13-opt/include -isystem /cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/ROOT/6.28.12/x86_64-el9-gcc13-opt/include -isystem /cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/tbb/2020_U2/x86_64-el9-gcc13-opt/include -isystem /cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/cppgsl/3.1.0/x86_64-el9-gcc13-opt/include -isystem /cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/rangev3/0.11.0/x86_64-el9-gcc13-opt/include -isystem /cvmfs/atlas.cern.ch/repo/sw/tdaq/tdaq-common/tdaq-common-11-02-01/installed/external/x86_64-el9-gcc13-opt/include -isystem /cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/2024-03-20T2101/AthenaExternals/25.0.2/InstallArea/x86_64-el9-gcc13-opt/lib/CLHEP-2.4.1.3/../../include -isystem /cvmfs/atlas-nightlies.cern.ch/repo/sw/main_Athena_x86_64-el9-gcc13-opt/sw/lcg/releases/LCG_104d_ATLAS_9/eigen/3.4.0/x86_64-el9-gcc13-opt/include/eigen3 -O3 -DNDEBUG -std=c++20 "--generate-code=arch=compute_52,code=[compute_52,sm_52]" -Xcompiler=-fPIC -Xcudafe --diag_suppress=177 -MD -MT Control/AthenaExamples/AthExCUDA/CMakeFiles/AthExCUDA.dir/src/TrackParticleCalibratorExampleAlg.cu.o -MF Control/AthenaExamples/AthExCUDA/CMakeFiles/AthExCUDA.dir/src/TrackParticleCalibratorExampleAlg.cu.o.d -x cu -c /home/krasznaa/ATLAS/projects/cuda/athena/Control/AthenaExamples/AthExCUDA/src/TrackParticleCalibratorExampleAlg.cu -o Control/AthenaExamples/AthExCUDA/CMakeFiles/AthExCUDA.dir/src/TrackParticleCalibratorExampleAlg.cu.o
/home/krasznaa/ATLAS/projects/cuda/athena/Tracking/TrkEvent/TrkParametersBase/TrkParametersBase/ParametersBase.icc(152): error: no operator "<<" matches these operands
            operand types are: std::ostream << const Trk::Surface
      sl << this->associatedSurface() << std::endl;
         ^
/cvmfs/sft.cern.ch/lcg/releases/gcc/13.1.0-b3d18/x86_64-el9/include/c++/13.1.0/ostream(330): note #3326-D: function "std::basic_ostream<_CharT, _Traits>::operator<<(std::basic_ostream<_CharT, _Traits>::__streambuf_type *) [with _CharT=char, _Traits=std::char_traits<char>]" does not match because argument #1 does not match parameter
        operator<<(__streambuf_type* __sb);
        ^
/cvmfs/sft.cern.ch/lcg/releases/gcc/13.1.0-b3d18/x86_64-el9/include/c++/13.1.0/ostream(297): note #3326-D: function "std::basic_ostream<_CharT, _Traits>::operator<<(std::nullptr_t) [with _CharT=char, _Traits=std::char_traits<char>]" does not match because argument #1 does not match parameter
        operator<<(nullptr_t)
        ^
...

After some amount of experimentation, I came up with the included update. This is a formalism that nvcc accepts, and doesn't require too much of a change in our code. 🤔 Though it's worth noting that just removing "this->", and making that line just:

sl << associatedSurface() << std::endl;

, is also acceptable to nvcc. 😕 Still, I thought the formalism I added now, could be the easiest for compilers to understand. 🤔

CUDA with C++20, main branch (2024.03.21.)

Merge request reports