CUDA with C++20, main branch (2024.03.21.)
As I mentioned in some meetings, our current nightlies still build all CUDA source files with the C++14 standard.
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_CUDAdefinition to the build ofAthExCUDA. 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. this->", and making that line just:
sl << associatedSurface() << std::endl;, is also acceptable to nvcc.