Skip to content

Commit fb4aac1

Browse files
valassivalassi
committed
[hack_ihel] in gg_tt.sa (MemoryAccessMatrixElements.h and various .cc files), fix "CUDA_HOME=none make" builds
This completes the C++ build when moving ME buffers from 1 per event to ncomb+1 per event However runTest.exe for C++ is now failing (output ME is 0)
1 parent c65f089 commit fb4aac1

File tree

5 files changed

+33
-27
lines changed

5 files changed

+33
-27
lines changed

‎epochX/cudacpp/gg_tt.sa/SubProcesses/CrossSectionKernels.cc‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
// Copyright (C) 2020-2023 CERN and UCLouvain.
1+
// Copyright (C) 2020-2024 CERN and UCLouvain.
22
// Licensed under the GNU Lesser General Public License (version 3 or later).
33
// Created by: A. Valassi (Jan 2022) for the MG5aMC CUDACPP plugin.
4-
// Further modified by: J. Teig, A. Valassi (2022-2023) for the MG5aMC CUDACPP plugin.
4+
// Further modified by: J. Teig, A. Valassi (2022-2024) for the MG5aMC CUDACPP plugin.
55

66
#include "CrossSectionKernels.h"
77

@@ -121,7 +121,7 @@ namespace mg5amcCpu
121121
// FIRST PASS: COUNT ALL/ABN/ZERO EVENTS, COMPUTE MIN/MAX, COMPUTE REFS AS MEANS OF SIMPLE SUMS
122122
for( size_t ievt = 0; ievt < nevt(); ++ievt ) // Loop over all events in this iteration
123123
{
124-
const fptype& me = MemoryAccessMatrixElements::ieventAccessConst( m_matrixElements.data(), ievt );
124+
const fptype& me = MemoryAccessMatrixElements::ieventAccessIhelConst( m_matrixElements.data(), ievt, MemoryBuffers::ncomb );
125125
const fptype& wg = MemoryAccessWeights::ieventAccessConst( m_samplingWeights.data(), ievt );
126126
const size_t ievtALL = m_iter * nevt() + ievt;
127127
// The following events are abnormal in a run with "-p 2048 256 12 -d"
@@ -156,7 +156,7 @@ namespace mg5amcCpu
156156
// SECOND PASS: IMPROVE MEANS FROM SUMS OF DIFFS TO PREVIOUS REF, UPDATE REF
157157
for( size_t ievt = 0; ievt < nevt(); ++ievt ) // Loop over all events in this iteration
158158
{
159-
const fptype& me = MemoryAccessMatrixElements::ieventAccessConst( m_matrixElements.data(), ievt );
159+
const fptype& me = MemoryAccessMatrixElements::ieventAccessIhelConst( m_matrixElements.data(), ievt, MemoryBuffers::ncomb );
160160
const fptype& wg = MemoryAccessWeights::ieventAccessConst( m_samplingWeights.data(), ievt );
161161
if( fp_is_abnormal( me ) ) continue;
162162
stats.sumMEdiff += ( me - stats.refME );
@@ -169,7 +169,7 @@ namespace mg5amcCpu
169169
// THIRD PASS: COMPUTE STDDEV FROM SQUARED SUMS OF DIFFS TO REF
170170
for( size_t ievt = 0; ievt < nevt(); ++ievt ) // Loop over all events in this iteration
171171
{
172-
const fptype& me = MemoryAccessMatrixElements::ieventAccessConst( m_matrixElements.data(), ievt );
172+
const fptype& me = MemoryAccessMatrixElements::ieventAccessIhelConst( m_matrixElements.data(), ievt, MemoryBuffers::ncomb );
173173
const fptype& wg = MemoryAccessWeights::ieventAccessConst( m_samplingWeights.data(), ievt );
174174
if( fp_is_abnormal( me ) ) continue;
175175
stats.sqsMEdiff += std::pow( me - stats.refME, 2 );

‎epochX/cudacpp/gg_tt.sa/SubProcesses/MemoryAccessMatrixElements.h‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
// Copyright (C) 2020-2023 CERN and UCLouvain.
1+
// Copyright (C) 2020-2024 CERN and UCLouvain.
22
// Licensed under the GNU Lesser General Public License (version 3 or later).
33
// Created by: A. Valassi (Jan 2022) for the MG5aMC CUDACPP plugin.
4-
// Further modified by: A. Valassi (2022-2023) for the MG5aMC CUDACPP plugin.
4+
// Further modified by: A. Valassi (2022-2024) for the MG5aMC CUDACPP plugin.
55

66
#ifndef MemoryAccessMatrixElements_H
77
#define MemoryAccessMatrixElements_H 1
@@ -107,12 +107,12 @@ namespace mg5amcCpu
107107
MemoryAccessHelper<MemoryAccessMatrixElementsBase>::template decodeRecordConst<int>;
108108

109109
// Locate a field (output) in a memory buffer (input) from the given event number (input) and the given field indexes (input)
110-
// [Signature (non-const) ===> fptype& ieventAccess( fptype* buffer, const ievt, const ihel ) <===]
110+
// [Signature (non-const) ===> fptype& ieventAccessIhel( fptype* buffer, const ievt, const ihel ) <===]
111111
static constexpr auto ieventAccessIhel =
112112
MemoryAccessHelper<MemoryAccessMatrixElementsBase>::template ieventAccessField<int>;
113113

114114
// Locate a field (output) in a memory buffer (input) from the given event number (input) and the given field indexes (input)
115-
// [Signature (const) ===> const fptype& ieventAccessConst( const fptype* buffer, const ievt, const ihel ) <===]
115+
// [Signature (const) ===> const fptype& ieventAccessIhelConst( const fptype* buffer, const ievt, const ihel ) <===]
116116
static constexpr auto ieventAccessIhelConst =
117117
MemoryAccessHelper<MemoryAccessMatrixElementsBase>::template ieventAccessFieldConst<int>;
118118
};
@@ -130,16 +130,16 @@ namespace mg5amcCpu
130130
static constexpr auto ieventAccessRecord = MemoryAccessMatrixElements::ieventAccessRecord;
131131

132132
// Locate a field (output) in a memory buffer (input) from a kernel event-indexing mechanism (internal) and the given field indexes (input)
133-
// [Signature (non-const, SCALAR) ===> fptype& kernelAccess_s( fptype* buffer, const int ihel ) <===]
133+
// [Signature (non-const, SCALAR) ===> fptype& kernelAccessIhel_s( fptype* buffer, const int ihel ) <===]
134134
static constexpr auto kernelAccessIhel_s =
135135
KernelAccessHelper<MemoryAccessMatrixElementsBase, onDevice>::template kernelAccessField<int>; // requires cuda 11.4
136136

137137
// Locate a field (output) in a memory buffer (input) from a kernel event-indexing mechanism (internal)
138-
// [Signature (non const, SCALAR OR VECTOR) ===> fptype_sv& kernelAccess( const fptype* buffer, const int ihel ) <===]
138+
// [Signature (non const, SCALAR OR VECTOR) ===> fptype_sv& kernelAccessIhel( const fptype* buffer, const int ihel ) <===]
139139
static __host__ __device__ inline fptype_sv&
140-
kernelAccessIhel( fptype* buffer )
140+
kernelAccessIhel( fptype* buffer, const int ihel )
141141
{
142-
fptype& out = kernelAccessIhel_s( buffer );
142+
fptype& out = kernelAccessIhel_s( buffer, ihel );
143143
#ifndef MGONGPU_CPPSIMD
144144
return out;
145145
#else
@@ -151,7 +151,7 @@ namespace mg5amcCpu
151151
}
152152

153153
// Locate a field (output) in a memory buffer (input) from a kernel event-indexing mechanism (internal) and the given field indexes (input)
154-
// [Signature (const) ===> const fptype& kernelAccessConst( const fptype* buffer, const int ihel ) <===]
154+
// [Signature (const) ===> const fptype& kernelAccessIhelConst( const fptype* buffer, const int ihel ) <===]
155155
static constexpr auto kernelAccessIhelConst =
156156
KernelAccessHelper<MemoryAccessMatrixElementsBase, onDevice>::template kernelAccessFieldConst<int>; // requires cuda 11.4
157157
};

‎epochX/cudacpp/gg_tt.sa/SubProcesses/P1_Sigma_sm_gg_ttx/CPPProcess.cc‎

Lines changed: 14 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -407,10 +407,10 @@ namespace mg5amcCpu
407407
// *** STORE THE RESULTS ***
408408

409409
// NB: calculate_wavefunctions ADDS |M|^2 for a given ihel to the running sum of |M|^2 over helicities for the given event(s)
410-
fptype_sv& MEs_sv = E_ACCESS::kernelAccess( MEs );
410+
fptype_sv& MEs_sv = E_ACCESS::kernelAccessIhel( MEs, ihel );
411411
MEs_sv += deltaMEs; // fix #435
412412
#if defined MGONGPU_CPPSIMD and defined MGONGPU_FPTYPE_DOUBLE and defined MGONGPU_FPTYPE2_FLOAT
413-
fptype_sv& MEs_sv_previous = E_ACCESS::kernelAccess( MEs_previous );
413+
fptype_sv& MEs_sv_previous = E_ACCESS::kernelAccessIhel( MEs_previous, ihel );
414414
MEs_sv_previous += deltaMEs_previous;
415415
#endif
416416
/*
@@ -849,8 +849,11 @@ namespace mg5amcCpu
849849
{
850850
const int ievt0 = ipagV * neppV;
851851
fptype* MEs = E_ACCESS::ieventAccessRecord( allMEs, ievt0 );
852-
fptype_sv& MEs_sv = E_ACCESS::kernelAccess( MEs );
853-
MEs_sv = fptype_sv{ 0 };
852+
for( int ihel = 0; ihel < ncomb; ihel++ )
853+
{
854+
fptype_sv& MEs_sv = E_ACCESS::kernelAccessIhel( MEs, ihel );
855+
MEs_sv = fptype_sv{ 0 };
856+
}
854857
#ifdef MGONGPU_SUPPORTS_MULTICHANNEL
855858
fptype* numerators = NUM_ACCESS::ieventAccessRecord( allNumerators, ievt0 );
856859
fptype* denominators = DEN_ACCESS::ieventAccessRecord( allDenominators, ievt0 );
@@ -969,9 +972,9 @@ namespace mg5amcCpu
969972
#else
970973
calculate_wavefunctions( ihel, allmomenta, allcouplings, allMEs, jamp2_sv, ievt00 );
971974
#endif
972-
MEs_ighel[ighel] = E_ACCESS::kernelAccess( E_ACCESS::ieventAccessRecord( allMEs, ievt00 ) );
975+
MEs_ighel[ighel] = E_ACCESS::kernelAccessIhel( E_ACCESS::ieventAccessRecord( allMEs, ievt00 ), ihel );
973976
#if defined MGONGPU_CPPSIMD and defined MGONGPU_FPTYPE_DOUBLE and defined MGONGPU_FPTYPE2_FLOAT
974-
MEs_ighel2[ighel] = E_ACCESS::kernelAccess( E_ACCESS::ieventAccessRecord( allMEs, ievt00 + neppV ) );
977+
MEs_ighel2[ighel] = E_ACCESS::kernelAccessIhel( E_ACCESS::ieventAccessRecord( allMEs, ievt00 + neppV ), ihel );
975978
#endif
976979
}
977980
// Event-by-event random choice of helicity #403
@@ -1087,8 +1090,11 @@ namespace mg5amcCpu
10871090
{
10881091
const int ievt0 = ipagV * neppV;
10891092
fptype* MEs = E_ACCESS::ieventAccessRecord( allMEs, ievt0 );
1090-
fptype_sv& MEs_sv = E_ACCESS::kernelAccess( MEs );
1091-
MEs_sv /= helcolDenominators[0];
1093+
for( int ihel = 0; ihel < ncomb; ihel++ )
1094+
{
1095+
fptype_sv& MEs_sv = E_ACCESS::kernelAccessIhel( MEs, ihel );
1096+
MEs_sv /= helcolDenominators[0];
1097+
}
10921098
#ifdef MGONGPU_SUPPORTS_MULTICHANNEL
10931099
if( channelId > 0 )
10941100
{

‎epochX/cudacpp/gg_tt.sa/SubProcesses/P1_Sigma_sm_gg_ttx/check_sa.cc‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -746,7 +746,7 @@ main( int argc, char** argv )
746746
std::cout << std::string( SEP79, '-' ) << std::endl;
747747
// Display matrix elements
748748
std::cout << " Matrix element = "
749-
<< MemoryAccessMatrixElements::ieventAccessIcombConst( hstMatrixElements.data(), ievt, CPPProcess::ncomb )
749+
<< MemoryAccessMatrixElements::ieventAccessIhelConst( hstMatrixElements.data(), ievt, CPPProcess::ncomb )
750750
<< " GeV^" << meGeVexponent << std::endl;
751751
std::cout << std::string( SEP79, '-' ) << std::endl;
752752
}

‎epochX/cudacpp/gg_tt.sa/SubProcesses/runTest.cc‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
// Copyright (C) 2020-2023 CERN and UCLouvain.
1+
// Copyright (C) 2020-2024 CERN and UCLouvain.
22
// Licensed under the GNU Lesser General Public License (version 3 or later).
33
// Created by: S. Hageboeck (Nov 2020) for the MG5aMC CUDACPP plugin.
4-
// Further modified by: S. Hageboeck, O. Mattelaer, S. Roiser, J. Teig, A. Valassi (2020-2023) for the MG5aMC CUDACPP plugin.
4+
// Further modified by: S. Hageboeck, O. Mattelaer, S. Roiser, J. Teig, A. Valassi (2020-2024) for the MG5aMC CUDACPP plugin.
55
//----------------------------------------------------------------------------
66
// Use ./runTest.exe --gtest_filter=*xxx to run only testxxx.cc tests
77
//----------------------------------------------------------------------------
@@ -114,7 +114,7 @@ struct CPUTest : public CUDA_CPU_TestBase
114114

115115
fptype getMatrixElement( std::size_t ievt ) const override
116116
{
117-
return MemoryAccessMatrixElements::ieventAccessConst( hstMatrixElements.data(), ievt );
117+
return MemoryAccessMatrixElements::ieventAccessIhelConst( hstMatrixElements.data(), ievt, CPPProcess::ncomb );
118118
}
119119
};
120120
#endif
@@ -235,7 +235,7 @@ struct CUDATest : public CUDA_CPU_TestBase
235235

236236
fptype getMatrixElement( std::size_t ievt ) const override
237237
{
238-
return MemoryAccessMatrixElements::ieventAccessConst( hstMatrixElements.data(), ievt );
238+
return MemoryAccessMatrixElements::ieventAccessIhelConst( hstMatrixElements.data(), ievt, CPPProcess::ncomb );
239239
}
240240
};
241241
#endif /* clang-format off */

0 commit comments

Comments
 (0)