[GCC14] Avoid std::clamp in device code #47266

iarspider · 2025-02-05T10:07:13Z

PR description:

When testing CMSSW with gcc14, the following error is emitted:

/cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/bits/stl_algo.h:3626:7: error: reference to __host__ function '__glibcxx_assert_fail' in __host__ __device__ function
 3626 |       __glibcxx_assert(!(__hi < __lo));
      |       ^
/cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/x86_64-redhat-linux-gnu/bits/c++config.h:614:12: note: expanded from macro '__glibcxx_assert'
  614 |       std::__glibcxx_assert_fail();                                     \
      |            ^
src/RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksByDensity.h:79:17: note: called by 'clusterTracksByDensity'
   79 |       iz = std::clamp(iz, INT8_MIN, INT8_MAX);
      |                 ^

This issue was reported to LLVM and GCC. Fix inspired by pytorch's fix of the same problem.

PR validation:

Bot tests

iarspider · 2025-02-05T10:07:34Z

please test

cmsbuild · 2025-02-05T10:07:34Z

cms-bot internal usage

cmsbuild · 2025-02-05T10:10:28Z

+code-checks

Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-47266/43568

cmsbuild · 2025-02-05T10:10:53Z

A new Pull Request was created by @iarspider for master.

It involves the following packages:

RecoVertex/PixelVertexFinding (reconstruction)

@jfernan2, @mandrenguyen can you please review it and eventually sign? Thanks.
@GiacomoSguazzoni, @VinInn, @VourMa, @dgulhan, @fabiocos, @martinamalberti, @missirol, @mmusich, @mtosi, @rovere this is something you requested to watch as well.
@antoniovilela, @mandrenguyen, @rappoccio, @sextonkennedy you are the release manager for this.

cms-bot commands are listed here

iarspider · 2025-02-05T10:17:49Z

please test for el8_amd64_gcc14

VinInn · 2025-02-05T10:32:32Z

RecoVertex/PixelVertexFinding/plugins/alpaka/clusterTracksByDensity.h

+      // which doesn't compile with gcc14 due to reference to __glibcxx_assert
+      // See https://github.com/llvm/llvm-project/issues/95183
+      int tmp_max = std::max<int>(iz, INT8_MIN);
+      iz = std::min<int>(iz, INT8_MAX);


tmp_max is unused....

maybe we should use NVIDIA implementation for cuda backends
https://nvidia.github.io/cccl/libcudacxx/

https://github.com/NVIDIA/cccl/blob/3786a088f986179b30bff006bd0d78853e5013dc/libcudacxx/include/cuda/std/__algorithm/clamp.h

True, but we will also need to implement workaround for ROCm (I think it has the same issue)

did you check with hipcc? My understanding was that llvm took care of those

did you check with hipcc? My understanding was that llvm took care of those

>> Compiling edm plugin src/RecoLocalTracker/SubCollectionProducers/src/SeedClusterRemover.cc In file included from <built-in>:1: In file included from /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/rocm/6.2.4-4e7b4b5933de5a6448978654268b9b7f/lib/llvm/lib/clang/18/include/__clang_hip_runtime_wrapper.h:145: In file included from /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/rocm/6.2.4-4e7b4b5933de5a6448978654268b9b7f/lib/llvm/lib/clang/18/include/cuda_wrappers/algorithm:55: In file included from /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/algorithm:61: /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/bits/stl_algo.h:3626:7: error: reference to __host__ function '__glibcxx_assert_fail' in __host__ __device__ function 3626 | __glibcxx_assert(!(__hi < __lo)); | ^ /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/x86_64-redhat-linux-gnu/bits/c++config.h:614:12: note: expanded from macro '__glibcxx_assert' 614 | std::__glibcxx_assert_fail(); \ | ^ src/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforDevice.h:382:19: note: called by 'errorFromDB<pixelTopology::Phase1>' 382 | int jx = std::clamp(bin_value, low_value, high_value); | ^ src/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHits.h:183:34: note: called by 'operator()' 183 | pixelCPEforDevice::errorFromDB<TrackerTraits>( | ^ /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/alpaka/1.2.0-bfbd7b77995d3af9fba1eba156781229/include/alpaka/kernel/TaskKernelGpuUniformCudaHipRt.hpp:75:13: note: called by 'gpuKernel<alpaka_rocm_async::pixelRecHits::GetHits<pixelTopology::Phase1>, alpaka::ApiHipRt, alpaka::AccGpuUniformCudaHipRt<alpaka::ApiHipRt, std::integral_constant<unsigned long, 1>, unsigned int>, std::integral_constant<unsigned long, 1>, unsigned int, const pixelCPEforDevice::ParamsOnDeviceT<pixelTopology::Phase1> *, const BeamSpotPOD *, SiPixelDigisLayout<>::ConstViewTemplateFreeParams<128, false, true, true>, unsigned int, unsigned int, SiPixelClustersLayout<>::ConstViewTemplateFreeParams<128, false, true, true>, TrackingRecHitSoA<pixelTopology::Phase1>::Layout<>::ViewTemplateFreeParams<128, false, true, true>>' 75 | kernelFnObj(const_cast<TAcc const&>(acc), args...); | ^ /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/x86_64-redhat-linux-gnu/bits/c++config.h:608:3: note: '__glibcxx_assert_fail' declared here 608 | __glibcxx_assert_fail() | ^ In file included from <built-in>:1: In file included from /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/rocm/6.2.4-4e7b4b5933de5a6448978654268b9b7f/lib/llvm/lib/clang/18/include/__clang_hip_runtime_wrapper.h:145: In file included from /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/rocm/6.2.4-4e7b4b5933de5a6448978654268b9b7f/lib/llvm/lib/clang/18/include/cuda_wrappers/algorithm:55: In file included from /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/algorithm:61: /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/bits/stl_algo.h:3626:7: error: reference to __host__ function '__glibcxx_assert_fail' in __host__ __device__ function 3626 | __glibcxx_assert(!(__hi < __lo)); | ^ /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/x86_64-redhat-linux-gnu/bits/c++config.h:614:12: note: expanded from macro '__glibcxx_assert' 614 | std::__glibcxx_assert_fail(); \ | ^ src/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforDevice.h:382:19: note: called by 'errorFromDB<pixelTopology::HIonPhase1>' 382 | int jx = std::clamp(bin_value, low_value, high_value); | ^ src/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHits.h:183:34: note: called by 'operator()' 183 | pixelCPEforDevice::errorFromDB<TrackerTraits>( | ^ /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/alpaka/1.2.0-bfbd7b77995d3af9fba1eba156781229/include/alpaka/kernel/TaskKernelGpuUniformCudaHipRt.hpp:75:13: note: called by 'gpuKernel<alpaka_rocm_async::pixelRecHits::GetHits<pixelTopology::HIonPhase1>, alpaka::ApiHipRt, alpaka::AccGpuUniformCudaHipRt<alpaka::ApiHipRt, std::integral_constant<unsigned long, 1>, unsigned int>, std::integral_constant<unsigned long, 1>, unsigned int, const pixelCPEforDevice::ParamsOnDeviceT<pixelTopology::HIonPhase1> *, const BeamSpotPOD *, SiPixelDigisLayout<>::ConstViewTemplateFreeParams<128, false, true, true>, unsigned int, unsigned int, SiPixelClustersLayout<>::ConstViewTemplateFreeParams<128, false, true, true>, TrackingRecHitSoA<pixelTopology::HIonPhase1>::Layout<>::ViewTemplateFreeParams<128, false, true, true>>' 75 | kernelFnObj(const_cast<TAcc const&>(acc), args...); | ^ /cvmfs/cms-ib.cern.ch/sw/x86_64/nweek-02874/el8_amd64_gcc14/external/gcc/14.2.1-60e8d7f953e2e96ab45420ba74bbf3ae/lib/gcc/x86_64-redhat-linux-gnu/14.2.1/../../../../include/c++/14.2.1/x86_64-redhat-linux-gnu/bits/c++config.h:608:3: note: '__glibcxx_assert_fail' declared here 608 | __glibcxx_assert_fail() | ^ 2 errors generated when compiling for gfx1030.

So it seems to fail for ROCm, but not for CUDA (at least I couldn't find any failures for CUDA in the output)

iarspider · 2025-02-05T10:35:26Z

please abort

iarspider · 2025-02-05T10:36:37Z

please test

cmsbuild · 2025-02-05T10:38:25Z

+code-checks

Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-47266/43570

iarspider · 2025-02-05T10:40:27Z

please test

cmsbuild · 2025-02-05T10:41:25Z

+code-checks

Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-47266/43572

cmsbuild · 2025-02-05T10:41:50Z

Pull request #47266 was updated. @jfernan2, @mandrenguyen can you please check and sign again.

VinInn · 2025-02-05T11:30:50Z

is cuda that is messing up with the gcc config...
to trigger the assert one needs either _GLIBCXX_ASSERTIONS or _GLIBCXX_DEBUG to be set
in gcc "x86: this is NOT default
see https://godbolt.org/z/d8jM1fo9n
std::clamp do even vectorize at -O2!

iarspider · 2025-02-05T12:31:17Z

std::clamp do even vectorize at -O2!

... as does the manual implementation: https://godbolt.org/z/j1dGET7Pc

cmsbuild · 2025-02-05T13:19:09Z

+1

Size: This PR adds an extra 28KB to repository
Summary: https://cmssdt.cern.ch/SDT/jenkins-artifacts/pull-request-integration/PR-06436c/44214/summary.html
COMMIT: 9291eee
CMSSW: CMSSW_15_0_X_2025-02-04-2300/el8_amd64_gcc12
User test area: For local testing, you can use /cvmfs/cms-ci.cern.ch/week1/cms-sw/cmssw/47266/44214/install.sh to create a dev area with all the needed externals and cmssw changes.

The following merge commits were also included on top of IB + this PR after doing git cms-merge-topic:

@mmusich add fillDescriptions to several plugins used at HLT (5/N) #47224

You can see more details here:
https://cmssdt.cern.ch/SDT/jenkins-artifacts/pull-request-integration/PR-06436c/44214/git-recent-commits.json
https://cmssdt.cern.ch/SDT/jenkins-artifacts/pull-request-integration/PR-06436c/44214/git-merge-result

Comparison Summary

Summary:

You potentially added 2 lines to the logs
Reco comparison results: 36 differences found in the comparisons
DQMHistoTests: Total files compared: 51
DQMHistoTests: Total histograms compared: 4130882
DQMHistoTests: Total failures: 1283
DQMHistoTests: Total nulls: 0
DQMHistoTests: Total successes: 4129579
DQMHistoTests: Total skipped: 20
DQMHistoTests: Total Missing objects: 0
DQMHistoSizes: Histogram memory added: 0.0 KiB( 50 files compared)
Checked 222 log files, 194 edm output root files, 51 DQM output files
TriggerResults: no differences found

jfernan2 · 2025-02-05T17:42:12Z

+1

cmsbuild · 2025-02-05T17:42:30Z

This pull request is fully signed and it will be integrated in one of the next master IBs (tests are also fine). This pull request will now be reviewed by the release team before it's merged. @sextonkennedy, @antoniovilela, @mandrenguyen, @rappoccio (and backports should be raised in the release meeting by the corresponding L2)

mandrenguyen · 2025-02-05T18:47:18Z

+1

cmsbuild added this to the CMSSW_15_0_X milestone Feb 5, 2025

cmsbuild added reconstruction-pending pending-signatures tests-pending orp-pending code-checks-pending tracking labels Feb 5, 2025

cmsbuild added tests-started and removed tests-pending labels Feb 5, 2025

cmsbuild added code-checks-approved and removed code-checks-pending labels Feb 5, 2025

VinInn reviewed Feb 5, 2025

View reviewed changes

cmsbuild added tests-pending and removed tests-started labels Feb 5, 2025

[GCC14] Avoid std::clamp in device code

04aa605

iarspider force-pushed the no-std-clamp branch from 4d90f6a to 04aa605 Compare February 5, 2025 10:36

cmsbuild added code-checks-pending and removed code-checks-approved labels Feb 5, 2025

cmsbuild added tests-started and removed tests-pending labels Feb 5, 2025

cmsbuild added code-checks-approved and removed code-checks-pending labels Feb 5, 2025

Replace std::clamp in pixelCPEforDevice.h

9291eee

cmsbuild added tests-pending code-checks-pending trk and removed tests-started code-checks-approved labels Feb 5, 2025

cmsbuild added tests-started and removed tests-pending labels Feb 5, 2025

cmsbuild added code-checks-approved and removed code-checks-pending labels Feb 5, 2025

iarspider mentioned this pull request Feb 5, 2025

[GCC14] template-id not allowed for constructor in C++20 #47265

Merged

cmsbuild added tests-approved and removed tests-started labels Feb 5, 2025

cmsbuild added reconstruction-approved fully-signed and removed reconstruction-pending pending-signatures labels Feb 5, 2025

cmsbuild added orp-approved and removed orp-pending labels Feb 5, 2025

cmsbuild merged commit 1cd573e into cms-sw:master Feb 5, 2025
11 checks passed

iarspider deleted the no-std-clamp branch February 7, 2025 10:17

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[GCC14] Avoid std::clamp in device code #47266

[GCC14] Avoid std::clamp in device code #47266

iarspider commented Feb 5, 2025

iarspider commented Feb 5, 2025

cmsbuild commented Feb 5, 2025 •

edited

Loading

cmsbuild commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

iarspider commented Feb 5, 2025

VinInn Feb 5, 2025 •

edited

Loading

iarspider Feb 5, 2025

VinInn Feb 5, 2025

VinInn Feb 5, 2025

iarspider Feb 5, 2025

VinInn Feb 5, 2025

iarspider Feb 5, 2025 •

edited

Loading

iarspider commented Feb 5, 2025

iarspider commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

iarspider commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

VinInn commented Feb 5, 2025

iarspider commented Feb 5, 2025 •

edited

Loading

cmsbuild commented Feb 5, 2025

jfernan2 commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

mandrenguyen commented Feb 5, 2025

[GCC14] Avoid std::clamp in device code #47266

[GCC14] Avoid std::clamp in device code #47266

Conversation

iarspider commented Feb 5, 2025

PR description:

PR validation:

iarspider commented Feb 5, 2025

cmsbuild commented Feb 5, 2025 • edited Loading

cmsbuild commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

iarspider commented Feb 5, 2025

VinInn Feb 5, 2025 • edited Loading

Choose a reason for hiding this comment

iarspider Feb 5, 2025

Choose a reason for hiding this comment

VinInn Feb 5, 2025

Choose a reason for hiding this comment

VinInn Feb 5, 2025

Choose a reason for hiding this comment

iarspider Feb 5, 2025

Choose a reason for hiding this comment

VinInn Feb 5, 2025

Choose a reason for hiding this comment

iarspider Feb 5, 2025 • edited Loading

Choose a reason for hiding this comment

iarspider commented Feb 5, 2025

iarspider commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

iarspider commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

VinInn commented Feb 5, 2025

iarspider commented Feb 5, 2025 • edited Loading

cmsbuild commented Feb 5, 2025

Comparison Summary

jfernan2 commented Feb 5, 2025

cmsbuild commented Feb 5, 2025

mandrenguyen commented Feb 5, 2025

cmsbuild commented Feb 5, 2025 •

edited

Loading

VinInn Feb 5, 2025 •

edited

Loading

iarspider Feb 5, 2025 •

edited

Loading

iarspider commented Feb 5, 2025 •

edited

Loading