Merge pull request #809 from valassi/rocrand

add support for ROCRAND (via HIPRAND)

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

@@ -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/HiprandRandomNumberKernel.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/RandomNumberKernels.h

@@ -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;
+  };
 
   //--------------------------------------------------------------------------
 }