WIP (DOCUMENTATION ONLY): gg to ttgggg (2->6 process) with diagram splitting

epochX/cudacpp/.gitignore

@@ -6,3 +6,5 @@ run_[0-9]*
 events.lhe*
 
 py3_model.pkl
+
+perf.data*

epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas.cu

@@ -2,10 +2,10 @@
 ! Copyright (C) 2010 The MadGraph5_aMC@NLO development team and contributors.
 ! Created by: J. Alwall (Sep 2010) for the MG5aMC CPP backend.
 !==========================================================================
-! Copyright (C) 2020-2024 CERN and UCLouvain.
+! Copyright (C) 2020-2025 CERN and UCLouvain.
 ! Licensed under the GNU Lesser General Public License (version 3 or later).
 ! Modified by: O. Mattelaer (Mar 2020) for the MG5aMC CUDACPP plugin.
-! Further modified by: O. Mattelaer, A. Valassi (2020-2024) for the MG5aMC CUDACPP plugin.
+! Further modified by: O. Mattelaer, A. Valassi (2020-2025) for the MG5aMC CUDACPP plugin.
 !==========================================================================
 
   //==========================================================================

epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/cpp_model_parameters_cc.inc

@@ -1,10 +1,10 @@
 // Copyright (C) 2010 The MadGraph5_aMC@NLO development team and contributors.
 // Created by: J. Alwall (Oct 2010) for the MG5aMC CPP backend.
 //==========================================================================
-// Copyright (C) 2020-2024 CERN and UCLouvain.
+// Copyright (C) 2020-2025 CERN and UCLouvain.
 // Licensed under the GNU Lesser General Public License (version 3 or later).
 // Modified by: A. Valassi (Sep 2021) for the MG5aMC CUDACPP plugin.
-// Further modified by: J. Teig, A. Valassi (2021-2024) for the MG5aMC CUDACPP plugin.
+// Further modified by: J. Teig, A. Valassi (2021-2025) for the MG5aMC CUDACPP plugin.
 //==========================================================================
 // This file has been automatically generated for CUDA/C++ standalone by
 %(info_lines)s

epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/cpp_model_parameters_h.inc

@@ -1,10 +1,10 @@
 // Copyright (C) 2010 The MadGraph5_aMC@NLO development team and contributors.
 // Created by: J. Alwall (Oct 2010) for the MG5aMC CPP backend.
 //==========================================================================
-// Copyright (C) 2020-2024 CERN and UCLouvain.
+// Copyright (C) 2020-2025 CERN and UCLouvain.
 // Licensed under the GNU Lesser General Public License (version 3 or later).
 // Modified by: A. Valassi (Sep 2021) for the MG5aMC CUDACPP plugin.
-// Further modified by: A. Valassi (2021-2024) for the MG5aMC CUDACPP plugin.
+// Further modified by: A. Valassi (2021-2025) for the MG5aMC CUDACPP plugin.
 //==========================================================================
 // This file has been automatically generated for CUDA/C++ standalone by
 %(info_lines)s
@@ -217,7 +217,7 @@ namespace mg5amcCpu
 #pragma GCC diagnostic ignored "-Wunused-but-set-variable" // e.g. <<warning: variable ‘couplings_sv’ set but not used [-Wunused-but-set-variable]>>
 #endif
   // Compute the output couplings (e.g. gc10 and gc11) from the input gs
-  template<class G_ACCESS, class C_ACCESS>
+  template<class G_ACCESS, class CD_ACCESS>
   __device__ inline void
   G2COUP( const fptype gs[],
           fptype couplings[],

epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/GpuAbstraction.h

@@ -74,6 +74,18 @@
 #define GPUBLAS_OP_N CUBLAS_OP_N
 #define GPUBLAS_OP_T CUBLAS_OP_T
 
+#define gpuGraph_t cudaGraph_t
+#define gpuGraphExec_t cudaGraphExec_t
+#define gpuGraphNode_t cudaGraphNode_t
+#define gpuKernelNodeParams cudaKernelNodeParams
+
+//#define gpuGraphAddDependencies cudaGraphAddDependencies
+#define gpuGraphAddKernelNode cudaGraphAddKernelNode
+#define gpuGraphCreate cudaGraphCreate
+#define gpuGraphExecKernelNodeSetParams cudaGraphExecKernelNodeSetParams
+#define gpuGraphInstantiate cudaGraphInstantiate
+#define gpuGraphLaunch cudaGraphLaunch
+
 //--------------------------------------------------------------------------
 
 #elif defined __HIPCC__
@@ -138,6 +150,18 @@
 #define GPUBLAS_OP_N HIPBLAS_OP_N
 #define GPUBLAS_OP_T HIPBLAS_OP_T
 
+#define gpuGraph_t hipGraph_t
+#define gpuGraphExec_t hipGraphExec_t
+#define gpuGraphNode_t hipGraphNode_t
+#define gpuKernelNodeParams hipKernelNodeParams
+
+//#define gpuGraphAddDependencies hipGraphAddDependencies
+#define gpuGraphAddKernelNode hipGraphAddKernelNode
+#define gpuGraphExecKernelNodeSetParams hipGraphExecKernelNodeSetParams
+#define gpuGraphCreate hipGraphCreate
+#define gpuGraphInstantiate hipGraphInstantiate
+#define gpuGraphLaunch hipGraphLaunch
+
 #endif
 
 //--------------------------------------------------------------------------

epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/MatrixElementKernels.cc

@@ -139,6 +139,44 @@ namespace mg5amcCpu
   }
 
   //--------------------------------------------------------------------------
+
+  bool precomputedGoodHelicities( bool* isGoodHel ) // output: isGoodHel[ncomb] - host array
+  {
+    static bool first = true;
+    static bool allGoodHel = false;
+    // Analyse environment variable CUDACPP_RUNTIME_GOODHELICITIES
+    if( first )
+    {
+      first = false;
+      const char* ghelEnv = getenv( "CUDACPP_RUNTIME_GOODHELICITIES" );
+      if( ghelEnv && std::string( ghelEnv ) != "" )
+      {
+        std::string ghelStr = std::string( ghelEnv );
+        // Case "ALL": treat all ncomb helicities as good helicities
+        if( ghelStr == "ALL" )
+        {
+          std::cout << "INFO: Env variable CUDACPP_RUNTIME_GOODHELICITIES equals \"ALL\": keep all helicities" << std::endl;
+          allGoodHel = true;
+        }
+        // TODO Case "DUMP": encode and dump the input mask isGoodHel as a string "xxxx" (e.g. in binary or hex format)
+        // TODO Case "xxxx": decode string "xxxx" into the output mask isGoodHel of good helicities
+        else
+        {
+          std::cout << "WARNING: Env variable CUDACPP_RUNTIME_GOODHELICITIES is set to an unknown value \"" << ghelStr << "\" and will be ignored" << std::endl;
+        }
+      }
+    }
+    // Case "ALL": treat all ncomb helicities as good helicities
+    if( allGoodHel )
+    {
+      for( int ihel = 0; ihel < CPPProcess::ncomb; ihel++ ) isGoodHel[ihel] = true;
+      return true;
+    }
+    else
+      return false;
+  }
+
+  //--------------------------------------------------------------------------
 }
 
 //============================================================================
@@ -201,13 +239,16 @@ namespace mg5amcCpu
   int MatrixElementKernelHost::computeGoodHelicities()
   {
     HostBufferHelicityMask hstIsGoodHel( CPPProcess::ncomb );
-    // ... 0d1. Compute good helicity mask on the host
-    computeDependentCouplings( m_gs.data(), m_couplings.data(), m_gs.size() );
+    if( !precomputedGoodHelicities( hstIsGoodHel.data() ) )
+    {
+      // ... 0d1. Compute good helicity mask on the host
+      computeDependentCouplings( m_gs.data(), m_couplings.data(), m_gs.size() );
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
-    sigmaKin_getGoodHel( m_momenta.data(), m_couplings.data(), m_matrixElements.data(), m_numerators.data(), m_denominators.data(), hstIsGoodHel.data(), nevt() );
+      sigmaKin_getGoodHel( m_momenta.data(), m_couplings.data(), m_matrixElements.data(), m_numerators.data(), m_denominators.data(), hstIsGoodHel.data(), nevt() );
 #else
-    sigmaKin_getGoodHel( m_momenta.data(), m_couplings.data(), m_matrixElements.data(), hstIsGoodHel.data(), nevt() );
+      sigmaKin_getGoodHel( m_momenta.data(), m_couplings.data(), m_matrixElements.data(), hstIsGoodHel.data(), nevt() );
 #endif
+    }
     // ... 0d2. Copy good helicity list to static memory on the host
     // [FIXME! REMOVE THIS STATIC THAT BREAKS MULTITHREADING?]
     return sigmaKin_setGoodHel( hstIsGoodHel.data() );
@@ -316,6 +357,7 @@ namespace mg5amcGpu
     , m_couplings( this->nevt() )
     , m_pHelMEs()
     , m_pHelJamps()
+    , m_pHelWfs()
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
     , m_pHelNumerators()
     , m_pHelDenominators()
@@ -354,6 +396,11 @@ namespace mg5amcGpu
     }
     // Create the "one-helicity" jamp buffer that will be used for helicity filtering
     m_pHelJamps.reset( new DeviceBufferSimple( CPPProcess::ncolor * mgOnGpu::nx2 * this->nevt() ) );
+#ifndef MGONGPU_RDC_DIAGRAMS
+    // Create the "one-helicity" wavefunction buffer that will be used for helicity filtering
+    if constexpr( CPPProcess::ndiagramgroups > 1 )
+      m_pHelWfs.reset( new DeviceBufferSimple( CPPProcess::nwf * CPPProcess::nw6 * mgOnGpu::nx2 * this->nevt() ) );
+#endif
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
     // Create the "one-helicity" numerator and denominator buffers that will be used for helicity filtering
     m_pHelNumerators.reset( new DeviceBufferSimple( this->nevt() ) );
@@ -442,15 +489,19 @@ namespace mg5amcGpu
 
   int MatrixElementKernelDevice::computeGoodHelicities()
   {
-    PinnedHostBufferHelicityMask hstIsGoodHel( CPPProcess::ncomb );
-    // ... 0d1. Compute good helicity mask (a host variable) on the device
-    gpuLaunchKernel( computeDependentCouplings, m_gpublocks, m_gputhreads, m_gs.data(), m_couplings.data() );
     const int nevt = m_gpublocks * m_gputhreads;
+    PinnedHostBufferHelicityMask hstIsGoodHel( CPPProcess::ncomb );
+    if( !precomputedGoodHelicities( hstIsGoodHel.data() ) )
+    {
+      // ... 0d1. Compute good helicity mask (a host variable) on the device
+      gpuLaunchKernel( computeDependentCouplings, m_gpublocks, m_gputhreads, m_gs.data(), m_couplings.data() );
+      fptype* helWfsData = ( m_pHelWfs ? m_pHelWfs->data() : nullptr );
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
-    sigmaKin_getGoodHel( m_momenta.data(), m_couplings.data(), m_matrixElements.data(), m_pHelJamps->data(), m_pHelNumerators->data(), m_pHelDenominators->data(), hstIsGoodHel.data(), nevt );
+      sigmaKin_getGoodHel( m_momenta.data(), m_couplings.data(), m_matrixElements.data(), m_pHelNumerators->data(), m_pHelDenominators->data(), m_pHelJamps->data(), helWfsData, hstIsGoodHel.data(), nevt );
 #else
-    sigmaKin_getGoodHel( m_momenta.data(), m_couplings.data(), m_matrixElements.data(), m_pHelJamps->data(), hstIsGoodHel.data(), nevt );
+      sigmaKin_getGoodHel( m_momenta.data(), m_couplings.data(), m_matrixElements.data(), m_pHelJamps->data(), helWfsData, hstIsGoodHel.data(), nevt );
 #endif
+    }
     // ... 0d3. Set good helicity list in host static memory
     int nGoodHel = sigmaKin_setGoodHel( hstIsGoodHel.data() );
     assert( nGoodHel > 0 ); // SANITY CHECK: there should be at least one good helicity
@@ -470,9 +521,15 @@ namespace mg5amcGpu
 #endif
     // ... Create the "many-helicity" super-buffer of nGoodHel ME buffers (dynamically allocated because nGoodHel is determined at runtime)
     m_pHelMEs.reset( new DeviceBufferSimple( nGoodHel * nevt ) );
-    // ... Create the "many-helicity" super-buffer of nGoodHel ME buffers (dynamically allocated because nGoodHel is determined at runtime)
+    // ... Create the "many-helicity" super-buffer of nGoodHel jamp buffers (dynamically allocated because nGoodHel is determined at runtime)
     // ... (calling reset here deletes the previously created "one-helicity" buffers used for helicity filtering)
     m_pHelJamps.reset( new DeviceBufferSimple( nGoodHel * CPPProcess::ncolor * mgOnGpu::nx2 * nevt ) );
+#ifndef MGONGPU_RDC_DIAGRAMS
+    // ... Create the "many-helicity" super-buffer of nGoodHel wavefunction buffers (dynamically allocated because nGoodHel is determined at runtime)
+    // ... (calling reset here deletes the previously created "one-helicity" buffers used for helicity filtering)
+    if constexpr( CPPProcess::ndiagramgroups > 1 )
+      m_pHelWfs.reset( new DeviceBufferSimple( nGoodHel * CPPProcess::nwf * CPPProcess::nw6 * mgOnGpu::nx2 * nevt ) );
+#endif
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
     // ... Create the "many-helicity" super-buffers of nGoodHel numerator and denominator buffers (dynamically allocated)
     // ... (calling reset here deletes the previously created "one-helicity" buffers used for helicity filtering)
@@ -505,20 +562,21 @@ namespace mg5amcGpu
     fptype2* ghelAllBlasTmp = nullptr;
     gpuBlasHandle_t* pBlasHandle = nullptr;
 #endif
+    fptype* helWfsData = ( m_pHelWfs ? m_pHelWfs->data() : nullptr );
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
     const unsigned int* pChannelIds = ( useChannelIds ? m_channelIds.data() : nullptr );
-    sigmaKin( m_momenta.data(), m_couplings.data(), m_rndhel.data(), m_rndcol.data(), pChannelIds, m_matrixElements.data(), m_selhel.data(), m_selcol.data(), m_colJamp2s.data(), m_pHelNumerators->data(), m_pHelDenominators->data(), m_pHelMEs->data(), m_pHelJamps->data(), ghelAllBlasTmp, pBlasHandle, m_helStreams, m_gpublocks, m_gputhreads );
+    sigmaKin( m_momenta.data(), m_couplings.data(), m_rndhel.data(), m_rndcol.data(), pChannelIds, m_matrixElements.data(), m_selhel.data(), m_selcol.data(), m_colJamp2s.data(), m_pHelNumerators->data(), m_pHelDenominators->data(), m_pHelMEs->data(), m_pHelJamps->data(), helWfsData, ghelAllBlasTmp, pBlasHandle, m_helStreams, m_gpublocks, m_gputhreads );
 #else
     assert( useChannelIds == false );
-    sigmaKin( m_momenta.data(), m_couplings.data(), m_rndhel.data(), m_matrixElements.data(), m_selhel.data(), m_pHelMEs->data(), m_pHelJamps->data(), ghelAllBlasTmp, pBlasHandle, m_helStreams, m_gpublocks, m_gputhreads );
+    sigmaKin( m_momenta.data(), m_couplings.data(), m_rndhel.data(), m_matrixElements.data(), m_selhel.data(), m_pHelMEs->data(), m_pHelJamps->data(), helWfsData, ghelAllBlasTmp, pBlasHandle, m_helStreams, m_gpublocks, m_gputhreads );
 #endif
 #ifdef MGONGPU_CHANNELID_DEBUG
     //std::cout << "DEBUG: MatrixElementKernelDevice::computeMatrixElements " << this << " " << ( useChannelIds ? "T" : "F" ) << " " << nevt() << std::endl;
     copyHostFromDevice( m_hstChannelIds, m_channelIds ); // FIXME?!
     const unsigned int* pHstChannelIds = ( useChannelIds ? m_hstChannelIds.data() : nullptr );
     MatrixElementKernelBase::updateNevtProcessedByChannel( pHstChannelIds, nevt() );
 #endif
-    checkGpu( gpuPeekAtLastError() );   // is this needed?
+    checkGpu( gpuPeekAtLastError() );   // not strictly needed, but useful if previous calls were not wrapped in checkGpu
     checkGpu( gpuDeviceSynchronize() ); // probably not needed? but it avoids errors in sigmaKin above from appearing later on in random places...
   }
 

epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/MatrixElementKernels.h

@@ -200,6 +200,9 @@ namespace mg5amcCpu
     // The super-buffer of nGoodHel jamp buffers (dynamically allocated because nGoodHel is determined at runtime)
     std::unique_ptr<DeviceBufferSimple> m_pHelJamps;
 
+    // The super-buffer of nGoodHel wavefunction buffers (dynamically allocated because nGoodHel is determined at runtime)
+    std::unique_ptr<DeviceBufferSimple> m_pHelWfs;
+
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
     // The super-buffer of nGoodHel numerator buffers (dynamically allocated because nGoodHel is determined at runtime)
     std::unique_ptr<DeviceBufferSimple> m_pHelNumerators;