Merge remote-tracking branch 'upstream/master' (including PR #809 and previous ones) into patchmad_nofileeeeeeeeeeeeeeeeeeeee

Author: valassi

MG5aMC/mg5amcnlo

@@ -1,5 +1,5 @@
 diff --git b/epochX/cudacpp/gg_tt.mad/SubProcesses/P1_gg_ttx/auto_dsig1.f a/epochX/cudacpp/gg_tt.mad/SubProcesses/P1_gg_ttx/auto_dsig1.f
-index 880769442..5a3da931f 100644
+index 0a3dfa449..8f4aaecd6 100644
 --- b/epochX/cudacpp/gg_tt.mad/SubProcesses/P1_gg_ttx/auto_dsig1.f
 +++ a/epochX/cudacpp/gg_tt.mad/SubProcesses/P1_gg_ttx/auto_dsig1.f
 @@ -484,23 +484,140 @@ C
@@ -77,9 +77,9 @@ index 880769442..5a3da931f 100644
 +          STOP
 +        ENDIF
 +        IF ( FIRST ) THEN ! exclude first pass (helicity filtering) from timers (#461)
-+          CALL FBRIDGESEQUENCE(FBRIDGE_PBRIDGE, P_MULTI, ALL_G,
-+     &      HEL_RAND, COL_RAND, 0, OUT2,
-+     &      SELECTED_HEL2, SELECTED_COL2 ) ! 0: multi channel disabled for helicity filtering
++          CALL FBRIDGESEQUENCE_NOMULTICHANNEL( FBRIDGE_PBRIDGE, ! multi channel disabled for helicity filtering
++     &      P_MULTI, ALL_G, HEL_RAND, COL_RAND, OUT2,
++     &      SELECTED_HEL2, SELECTED_COL2 )
 +          FIRST = .FALSE.
 +c         ! This is a workaround for https://github.com/oliviermattelaer/mg5amc_test/issues/22 (see PR #486)
 +          IF( FBRIDGE_MODE .EQ. 1 ) THEN ! (CppOnly=1 : SMATRIX1 is not called at all)
@@ -96,9 +96,9 @@ index 880769442..5a3da931f 100644
 +        ENDIF
 +        call counters_smatrix1multi_start( 0, VECSIZE_USED ) ! cudacpp=0
 +        IF ( .NOT. MULTI_CHANNEL ) THEN
-+          CALL FBRIDGESEQUENCE(FBRIDGE_PBRIDGE, P_MULTI, ALL_G,
-+     &      HEL_RAND, COL_RAND, 0, OUT2,
-+     &      SELECTED_HEL2, SELECTED_COL2 ) ! 0: multi channel disabled
++          CALL FBRIDGESEQUENCE_NOMULTICHANNEL( FBRIDGE_PBRIDGE, ! multi channel disabled
++     &      P_MULTI, ALL_G, HEL_RAND, COL_RAND, OUT2,
++     &      SELECTED_HEL2, SELECTED_COL2 )
 +        ELSE
 +          IF( SDE_STRAT.NE.1 ) THEN
 +            WRITE(6,*) 'ERROR! The cudacpp bridge requires SDE=1' ! multi channel single-diagram enhancement strategy
@@ -284,7 +284,7 @@ index 71fbf2b25..0f1d199fc 100644
        open(unit=lun,file=tempname,status='old',ERR=20)
        fopened=.true.
 diff --git b/epochX/cudacpp/gg_tt.mad/SubProcesses/P1_gg_ttx/matrix1.f a/epochX/cudacpp/gg_tt.mad/SubProcesses/P1_gg_ttx/matrix1.f
-index 3ac962688..daea73a6d 100644
+index 817af778b..0c2ce6ec4 100644
 --- b/epochX/cudacpp/gg_tt.mad/SubProcesses/P1_gg_ttx/matrix1.f
 +++ a/epochX/cudacpp/gg_tt.mad/SubProcesses/P1_gg_ttx/matrix1.f
 @@ -72,7 +72,10 @@ C

epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/MG5aMC_patches/PROD/patch.P1

@@ -38,7 +38,7 @@ index 617f10b93..00c73099a 100644
 +cleanall: cleanSource # THIS IS THE ONE
 +	for i in `ls -d ../SubProcesses/P*`; do cd $$i; make cleanavxs; cd -; done;
 diff --git b/epochX/cudacpp/gg_tt.mad/SubProcesses/makefile a/epochX/cudacpp/gg_tt.mad/SubProcesses/makefile
-index 348c283be..65369d610 100644
+index 348c283be..d572486c2 100644
 --- b/epochX/cudacpp/gg_tt.mad/SubProcesses/makefile
 +++ a/epochX/cudacpp/gg_tt.mad/SubProcesses/makefile
 @@ -1,6 +1,28 @@
@@ -297,10 +297,10 @@ index 348c283be..65369d610 100644
 +distclean: cleanall # Clean all fortran and cudacpp builds as well as the googletest installation
 +	$(MAKE) -f $(CUDACPP_MAKEFILE) distclean
 diff --git b/epochX/cudacpp/gg_tt.mad/bin/internal/gen_ximprove.py a/epochX/cudacpp/gg_tt.mad/bin/internal/gen_ximprove.py
-index ebbc1ac1d..a88d60b28 100755
+index fb7efa87c..5fd170d18 100755
 --- b/epochX/cudacpp/gg_tt.mad/bin/internal/gen_ximprove.py
 +++ a/epochX/cudacpp/gg_tt.mad/bin/internal/gen_ximprove.py
-@@ -385,8 +385,20 @@ class gensym(object):
+@@ -391,8 +391,20 @@ class gensym(object):
                          done = True
                  if not done:
                      raise Exception('Parsing error in gensym: %s' % stdout)
@@ -324,7 +324,7 @@ index ebbc1ac1d..a88d60b28 100755
                  self.submit_to_cluster(job_list)
                  job_list = {}
 diff --git b/epochX/cudacpp/gg_tt.mad/bin/internal/madevent_interface.py a/epochX/cudacpp/gg_tt.mad/bin/internal/madevent_interface.py
-index 389b93ab8..d72270289 100755
+index 8c509e83f..cb6bf4ca5 100755
 --- b/epochX/cudacpp/gg_tt.mad/bin/internal/madevent_interface.py
 +++ a/epochX/cudacpp/gg_tt.mad/bin/internal/madevent_interface.py
 @@ -3614,8 +3614,20 @@ Beware that this can be dangerous for local multicore runs.""")

epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/MG5aMC_patches/PROD/patch.common

@@ -10,6 +10,7 @@
 #include <cassert>
 
 #ifndef MGONGPU_HAS_NO_CURAND /* clang-format off */
+// NB This must come AFTER mgOnGpuConfig.h which contains our definition of __global__ when MGONGPUCPP_GPUIMPL is not defined
 #include "curand.h"
 #define checkCurand( code ){ assertCurand( code, __FILE__, __LINE__ ); }
 inline void assertCurand( curandStatus_t code, const char *file, int line, bool abort = true )

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

@@ -0,0 +1,145 @@
+// Copyright (C) 2020-2024 CERN and UCLouvain.
+// Licensed under the GNU Lesser General Public License (version 3 or later).
+// Created by: A. Valassi (Jan 2024) for the MG5aMC CUDACPP plugin.
+// Further modified by: A. Valassi (2024) for the MG5aMC CUDACPP plugin.
+
+#include "mgOnGpuConfig.h"
+
+#include "GpuRuntime.h"
+#include "MemoryBuffers.h"
+#include "RandomNumberKernels.h"
+
+#include <cassert>
+
+#ifndef MGONGPU_HAS_NO_HIPRAND /* clang-format off */
+#ifndef __HIP_PLATFORM_AMD__
+#define __HIP_PLATFORM_AMD__ 1 // enable hiprand for AMD (rocrand)
+#endif
+#include <hiprand/hiprand.h>
+#define checkHiprand( code ){ assertHiprand( code, __FILE__, __LINE__ ); }
+inline void assertHiprand( hiprandStatus_t code, const char *file, int line, bool abort = true )
+{
+  if ( code != HIPRAND_STATUS_SUCCESS )
+  {
+    printf( "HiprandAssert: %s:%d code=%d\n", file, line, code );
+    if ( abort ) assert( code == HIPRAND_STATUS_SUCCESS );
+  }
+}
+#endif /* clang-format on */
+
+#ifdef MGONGPUCPP_GPUIMPL
+namespace mg5amcGpu
+#else
+namespace mg5amcCpu
+#endif
+{
+  //--------------------------------------------------------------------------
+#ifndef MGONGPU_HAS_NO_HIPRAND
+  HiprandRandomNumberKernel::HiprandRandomNumberKernel( BufferRndNumMomenta& rnarray, const bool onDevice )
+    : RandomNumberKernelBase( rnarray )
+    , m_isOnDevice( onDevice )
+  {
+    if( m_isOnDevice )
+    {
+#ifdef MGONGPUCPP_GPUIMPL
+      if( !m_rnarray.isOnDevice() )
+        throw std::runtime_error( "HiprandRandomNumberKernel on device with a host random number array" );
+#else
+      throw std::runtime_error( "HiprandRandomNumberKernel does not support HiprandDevice on CPU host" );
+#endif
+    }
+    else
+    {
+      if( m_rnarray.isOnDevice() )
+        throw std::runtime_error( "HiprandRandomNumberKernel on host with a device random number array" );
+    }
+    createGenerator();
+  }
+
+  //--------------------------------------------------------------------------
+
+  HiprandRandomNumberKernel::~HiprandRandomNumberKernel()
+  {
+    destroyGenerator();
+  }
+
+  //--------------------------------------------------------------------------
+
+  void HiprandRandomNumberKernel::seedGenerator( const unsigned int seed )
+  {
+    if( m_isOnDevice )
+    {
+      destroyGenerator(); // workaround for #429
+      createGenerator();  // workaround for #429
+    }
+    //printf( "seedGenerator: seed %d\n", seed );
+    checkHiprand( hiprandSetPseudoRandomGeneratorSeed( m_rnGen, seed ) );
+  }
+
+  //--------------------------------------------------------------------------
+
+  void HiprandRandomNumberKernel::createGenerator()
+  {
+    //const hiprandRngType_t type = HIPRAND_RNG_PSEUDO_DEFAULT;
+    //const hiprandRngType_t type = HIPRAND_RNG_PSEUDO_XORWOW;
+    //const hiprandRngType_t type = HIPRAND_RNG_PSEUDO_MRG32K3A;
+    const hiprandRngType_t type = HIPRAND_RNG_PSEUDO_MTGP32; // same as curand; not implemented yet (code=1000) in host code
+    //const hiprandRngType_t type = HIPRAND_RNG_PSEUDO_MT19937;
+    //const hiprandRngType_t type = HIPRAND_RNG_PSEUDO_PHILOX4_32_10;
+    if( m_isOnDevice )
+    {
+      checkHiprand( hiprandCreateGenerator( &m_rnGen, type ) );
+    }
+    else
+    {
+      // See https://github.com/ROCm/hipRAND/issues/76
+      throw std::runtime_error( "HiprandRandomNumberKernel on host is not supported yet (hiprandCreateGeneratorHost is not implemented yet)" );
+      //checkHiprand( hiprandCreateGeneratorHost( &m_rnGen, type ) ); // ALWAYS FAILS WITH CODE=1000
+    }
+    // FIXME: hiprand ordering is not implemented yet
+    // See https://github.com/ROCm/hipRAND/issues/75
+    /*
+    //checkHiprand( hiprandSetGeneratorOrdering( *&m_rnGen, HIPRAND_ORDERING_PSEUDO_LEGACY ) );
+    checkHiprand( hiprandSetGeneratorOrdering( *&m_rnGen, HIPRAND_ORDERING_PSEUDO_BEST ) );
+    //checkHiprand( hiprandSetGeneratorOrdering( *&m_rnGen, HIPRAND_ORDERING_PSEUDO_DYNAMIC ) );
+    //checkHiprand( hiprandSetGeneratorOrdering( *&m_rnGen, HIPRAND_ORDERING_PSEUDO_SEEDED ) );
+    */
+  }
+
+  //--------------------------------------------------------------------------
+
+  void HiprandRandomNumberKernel::destroyGenerator()
+  {
+    checkHiprand( hiprandDestroyGenerator( m_rnGen ) );
+  }
+
+  //--------------------------------------------------------------------------
+
+  void HiprandRandomNumberKernel::generateRnarray()
+  {
+#if defined MGONGPU_FPTYPE_DOUBLE
+    checkHiprand( hiprandGenerateUniformDouble( m_rnGen, m_rnarray.data(), m_rnarray.size() ) );
+#elif defined MGONGPU_FPTYPE_FLOAT
+    checkHiprand( hiprandGenerateUniform( m_rnGen, m_rnarray.data(), m_rnarray.size() ) );
+#endif
+    /*
+    printf( "\nHiprandRandomNumberKernel::generateRnarray size = %d\n", (int)m_rnarray.size() );
+    fptype* data = m_rnarray.data();
+#ifdef MGONGPUCPP_GPUIMPL
+    if( m_rnarray.isOnDevice() )
+    {
+      data = new fptype[m_rnarray.size()]();
+      checkCuda( cudaMemcpy( data, m_rnarray.data(), m_rnarray.bytes(), cudaMemcpyDeviceToHost ) );
+    }
+#endif
+    for( int i = 0; i < ( (int)m_rnarray.size() / 4 ); i++ )
+      printf( "[%4d] %f %f %f %f\n", i * 4, data[i * 4], data[i * 4 + 2], data[i * 4 + 2], data[i * 4 + 3] );
+#ifdef MGONGPUCPP_GPUIMPL
+    if( m_rnarray.isOnDevice() ) delete[] data;
+#endif
+    */
+  }
+
+  //--------------------------------------------------------------------------
+#endif
+}

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

@@ -1,21 +1,22 @@
-// Copyright (C) 2020-2023 CERN and UCLouvain.
+// Copyright (C) 2020-2024 CERN and UCLouvain.
 // Licensed under the GNU Lesser General Public License (version 3 or later).
 // Created by: A. Valassi (Dec 2021) for the MG5aMC CUDACPP plugin.
-// Further modified by: J. Teig, A. Valassi (2021-2023) for the MG5aMC CUDACPP plugin.
+// Further modified by: J. Teig, A. Valassi (2021-2024) for the MG5aMC CUDACPP plugin.
 
 #ifndef RANDOMNUMBERKERNELS_H
 #define RANDOMNUMBERKERNELS_H 1
 
 #include "mgOnGpuConfig.h"
 
-// NB This must come AFTER mgOnGpuConfig.h which contains our definition of __global__ when MGONGPUCPP_GPUIMPL is not defined
-#ifndef MGONGPU_HAS_NO_CURAND
-//#include "curand.h"
-struct curandGenerator_st; // forward definition from curand.h
-#endif
-
 #include "MemoryBuffers.h"
 
+// Forward definition from curand.h (the full header is only needed in CurandRandomKernel.cc)
+struct curandGenerator_st;
+
+// Forward definition from hiprand.h (the full header is only needed in HiprandRandomKernel.cc)
+struct rocrand_generator_base_type;
+typedef rocrand_generator_base_type hiprandGenerator_st;
+
 #ifdef MGONGPUCPP_GPUIMPL
 namespace mg5amcGpu
 #else
@@ -107,7 +108,6 @@ namespace mg5amcCpu
 
   //--------------------------------------------------------------------------
 
-#ifndef MGONGPU_HAS_NO_CURAND
   // A class encapsulating CURAND random number generation on a CPU host or on a GPU device
   class CurandRandomNumberKernel final : public RandomNumberKernelBase
   {
@@ -142,11 +142,49 @@ namespace mg5amcCpu
     const bool m_isOnDevice;
 
     // The curand generator
-    // (NB: curand.h defines typedef generator_t as a pointer to forward-defined 'struct curandGenerator_st')
+    // (NB: curand.h defines typedef curandGenerator_t as a pointer to forward-defined 'struct curandGenerator_st')
     curandGenerator_st* m_rnGen;
   };
 
-#endif
+  //--------------------------------------------------------------------------
+
+  // A class encapsulating HIPRAND random number generation on a CPU host or on a GPU device
+  class HiprandRandomNumberKernel final : public RandomNumberKernelBase
+  {
+  public:
+
+    // Constructor from an existing output buffer
+    HiprandRandomNumberKernel( BufferRndNumMomenta& rnarray, const bool onDevice );
+
+    // Destructor
+    ~HiprandRandomNumberKernel();
+
+    // Seed the random number generator
+    void seedGenerator( const unsigned int seed ) override final;
+
+    // Generate the random number array
+    void generateRnarray() override final;
+
+    // Is this a host or device kernel?
+    bool isOnDevice() const override final { return m_isOnDevice; }
+
+  private:
+
+    // Create the generator (workaround for #429: do this in every seedGenerator call rather than only in the ctor)
+    void createGenerator();
+
+    // Destroy the generator (workaround for #429: do this in every seedGenerator call rather than only in the ctor)
+    void destroyGenerator();
+
+  private:
+
+    // Is this a host or device kernel?
+    const bool m_isOnDevice;
+
+    // The hiprand generator
+    // (NB: hiprand.h defines typedef hiprandGenerator_t as a pointer to forward-defined 'struct hiprandGenerator_st')
+    hiprandGenerator_st* m_rnGen;
+  };
 
   //--------------------------------------------------------------------------
 }