Skip to content

Use CUDA- or HIP-specific memcpy when compiling for GPU

Lucas Meyer Garcia requested to merge memcpy_gpu_lgarcia into master

Compiling Allen with Allen!683 (merged) for HIP results in an error in the LHCb::ODIN constructor related to a std::memcpy call:

[147/276] Building HIPCC object device/selections/CMakeFiles/Selections.dir/lines/photon/src/Selections_generated_SingleCaloCluster.cu.o
FAILED: device/selections/CMakeFiles/Selections.dir/lines/photon/src/Selections_generated_SingleCaloCluster.cu.o 
cd /builds/lhcb/Allen/build_x86_64-centos7-clang12-opt_HIP_RelWithDebInfo_all_/device/selections/CMakeFiles/Selections.dir/lines/photon/src && /cvmfs/sft.cern.ch/lcg/releases/CMake/3.20.0-790a8/x86_64-centos7-clang12-opt/bin/cmake -E make_directory /builds/lhcb/Allen/build_x86_64-centos7-clang12-opt_HIP_RelWithDebInfo_all_/device/selections/CMakeFiles/Selections.dir/lines/photon/src/. && /cvmfs/sft.cern.ch/lcg/releases/CMake/3.20.0-790a8/x86_64-centos7-clang12-opt/bin/cmake -D verbose:BOOL=OFF -D build_configuration:STRING=RELWITHDEBINFO -D generated_file:STRING=/builds/lhcb/Allen/build_x86_64-centos7-clang12-opt_HIP_RelWithDebInfo_all_/device/selections/CMakeFiles/Selections.dir/lines/photon/src/./Selections_generated_SingleCaloCluster.cu.o -P /builds/lhcb/Allen/build_x86_64-centos7-clang12-opt_HIP_RelWithDebInfo_all_/device/selections/CMakeFiles/Selections.dir/lines/photon/src/Selections_generated_SingleCaloCluster.cu.o.cmake
In file included from /builds/lhcb/Allen/device/selections/lines/photon/src/SingleCaloCluster.cu:5:
In file included from /builds/lhcb/Allen/device/selections/lines/photon/include/SingleCaloCluster.cuh:6:
In file included from /builds/lhcb/Allen/device/selections/line_types/include/Line.cuh:9:
/builds/lhcb/Allen/build_x86_64-centos7-clang12-opt_HIP_RelWithDebInfo_all_/imported_include/Event/ODIN.h:187:14: error: reference to __host__ function 'memcpy' in __host__ __device__ function
        std::memcpy( data.data(), buffer.data(), sizeof( data ) );
             ^
/builds/lhcb/Allen/device/selections/line_types/include/Line.cuh:178:20: note: called by 'process_line<single_calo_cluster_line::single_calo_cluster_line_t, single_calo_cluster_line::Parameters>'
  const LHCb::ODIN odin {
                   ^

When compiling for gpu, this constructor is declared as a __host__ __device__ function, so the CUDA- or HIP-specific memcpy should be called instead. This MR changes that.

FYI @rangel @jbaptist

Merge request reports