diff --git a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas.h b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas.h index fcfc4b3153..2e87db1a2a 100644 --- a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas.h +++ b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas.h @@ -9,148 +9,158 @@ !========================================================================== //-------------------------------------------------------------------------- -#ifdef MGONGPU_INLINE_HELAMPS -#define INLINE inline -#define ALWAYS_INLINE __attribute__( ( always_inline ) ) +#if defined(MGONGPU_INLINE_HELAMPS) && defined(MGONGPU_NOINLINE_HELAMPS) +#error "Conflicting HELAS inlining options" +#elif defined(MGONGPU_INLINE_HELAMPS) +#if defined(MGONGPUCPP_GPUIMPL) +#define MGONGPU_HELAS_INLINE __forceinline__ #else -#define INLINE -#define ALWAYS_INLINE +#define MGONGPU_HELAS_INLINE inline __attribute__( ( always_inline ) ) +#endif +#elif defined(MGONGPU_NOINLINE_HELAMPS) +#if defined(MGONGPUCPP_GPUIMPL) +#define MGONGPU_HELAS_INLINE __noinline__ +#else +#define MGONGPU_HELAS_INLINE __attribute__( ( noinline ) ) +#endif +#else +#define MGONGPU_HELAS_INLINE #endif //-------------------------------------------------------------------------- // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void ixxxxx( const fptype momenta[], // input: momenta const fptype fmass, // input: fermion mass const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PX == PY == 0 and E == +PZ > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void ipzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PX == PY == 0 and E == -PZ > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void imzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PT > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void ixzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction vc[6] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void vxxxxx( const fptype momenta[], // input: momenta const fptype vmass, // input: vector boson mass const int nhel, // input: -1, 0 (only if vmass!=0) or +1 (helicity of vector boson) const int nsv, // input: +1 (final) or -1 (initial) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction sc[3] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void sxxxxx( const fptype momenta[], // input: momenta //const fptype, // WARNING: input "smass" unused (missing in Fortran) - scalar boson mass //const int, // WARNING: input "nhel" unused (missing in Fortran) - scalar has no helicity! const int nss, // input: +1 (final) or -1 (initial) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fo[6] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void oxxxxx( const fptype momenta[], // input: momenta const fptype fmass, // input: fermion mass const int nhel, // input: -1, 0 (only if vmass!=0) or +1 (helicity of vector boson) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fo[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PX == PY == 0 and E == +PZ > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void opzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fo[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PX == PY == 0 and E == -PZ > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void omzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fo[6] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void oxzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //========================================================================== diff --git a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas_fd.h b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas_fd.h index 205708a5ba..9d7cbdb243 100644 --- a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas_fd.h +++ b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/aloha/template_files/gpu/helas_fd.h @@ -9,176 +9,186 @@ !========================================================================== //-------------------------------------------------------------------------- -#ifdef MGONGPU_INLINE_HELAMPS -#define INLINE inline -#define ALWAYS_INLINE __attribute__( ( always_inline ) ) +#if defined(MGONGPU_INLINE_HELAMPS) && defined(MGONGPU_NOINLINE_HELAMPS) +#error "Conflicting HELAS inlining options" +#elif defined(MGONGPU_INLINE_HELAMPS) +#if defined(MGONGPUCPP_GPUIMPL) +#define MGONGPU_HELAS_INLINE __forceinline__ #else -#define INLINE -#define ALWAYS_INLINE +#define MGONGPU_HELAS_INLINE inline __attribute__( ( always_inline ) ) +#endif +#elif defined(MGONGPU_NOINLINE_HELAMPS) +#if defined(MGONGPUCPP_GPUIMPL) +#define MGONGPU_HELAS_INLINE __noinline__ +#else +#define MGONGPU_HELAS_INLINE __attribute__( ( noinline ) ) +#endif +#else +#define MGONGPU_HELAS_INLINE #endif //-------------------------------------------------------------------------- // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void ixxxxx( const fptype momenta[], // input: momenta const fptype fmass, // input: fermion mass const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PX == PY == 0 and E == +PZ > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void ipzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PX == PY == 0 and E == -PZ > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void imzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PT > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void ixzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction vc[6] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void vxxxxx( const fptype momenta[], // input: momenta const fptype vmass, // input: vector boson mass const int nhel, // input: -1, 0 (only if vmass!=0) or +1 (helicity of vector boson) const int nsv, // input: +1 (final) or -1 (initial) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction sc[3] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void sxxxxx( const fptype momenta[], // input: momenta //const fptype, // WARNING: input "smass" unused (missing in Fortran) - scalar boson mass //const int, // WARNING: input "nhel" unused (missing in Fortran) - scalar has no helicity! const int nss, // input: +1 (final) or -1 (initial) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fo[6] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void oxxxxx( const fptype momenta[], // input: momenta const fptype fmass, // input: fermion mass const int nhel, // input: -1, 0 (only if vmass!=0) or +1 (helicity of vector boson) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fo[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PX == PY == 0 and E == +PZ > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void opzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fo[6] from the input momenta[npar*4*nevt] // ASSUMPTIONS: (FMASS == 0) and (PX == PY == 0 and E == -PZ > 0) template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void omzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the output wavefunction fo[6] from the input momenta[npar*4*nevt] template - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void oxzxxx( const fptype momenta[], // input: momenta //const fptype fmass, // [skip: ASSUME fermion mass==0] const int nhel, // input: -1 or +1 (helicity of fermion) const int nsf, // input: +1 (particle) or -1 (antiparticle) fptype wavefunctions[], // output: wavefunctions const int ipar // input: particle# out of npar - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute the direction n[5] of the gauge q[5] - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void define_gauge_dir( const fptype q[], // input: gauge fptype n[] // output: direction - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // Compute a propagator factor d out of gauge q[5] and a mass - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void calculate_propagator_factor( const fptype_sv q[5], // input: gauge const fptype_sv mass, // input: mass fptype_sv *d // output: propagator factor - ) ALWAYS_INLINE; + ); //-------------------------------------------------------------------------- // multiply by propagation factor from m and wawefunctionsin[] and output them // as wavefunctionout[] template< class W_ACCESS> - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void multiply_propagator_factor( const fptype wavefunctionsin[], // input: wavefunctions const fptype m, // input: mass fptype wavefunctionsout[] // output: wavefunctions - ) ALWAYS_INLINE; + ); //========================================================================== // Compute the output wavefunction fi[6] from the input momenta[npar*4*nevt] @@ -947,7 +957,7 @@ //-------------------------------------------------------------------------- // Compute the direction n[5] of the gauge q[5] - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void define_gauge_dir( const cxtype_sv q[5], // input: gauge fptype_sv n[5] ) // output: direction { @@ -992,7 +1002,7 @@ //-------------------------------------------------------------------------- // Compute propagator factor n[5] of the gauge q[5] - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void calculate_propagator_factor( const cxtype_sv q[5], // input: gauge const fptype mass, // input: mass fptype_sv *d ) // output: propagator factor @@ -1005,7 +1015,7 @@ // multiply by propagation factor from m and wawefunctionsin[] and output them // as wavefunctionout[] template< class W_ACCESS> - __host__ __device__ INLINE void + __host__ __device__ MGONGPU_HELAS_INLINE void multiply_propagator_factor( const fptype wavefunctionsin[], // input: wavefunctions const fptype m, // input: mass fptype wavefunctionsout[] ) // output: wavefunctions diff --git a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/cudacpp.mk b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/cudacpp.mk index 78188d3f83..d250f01e95 100644 --- a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/cudacpp.mk +++ b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/cudacpp.mk @@ -624,13 +624,19 @@ else $(error Unknown FPTYPE='$(FPTYPE)': only 'd', 'f' and 'm' are supported) endif -# Set the build flags appropriate to each HELINL choice (example: "make HELINL=1") +# Set the build flags appropriate to each HELINL choice: +# HELINL=1 forces inlining +# HELINL=0 leaves the decision to the compiler +# HELINL=-1 prevents inlining $(info HELINL='$(HELINL)') ifeq ($(HELINL),1) CXXFLAGS += -DMGONGPU_INLINE_HELAMPS GPUFLAGS += -DMGONGPU_INLINE_HELAMPS +else ifeq ($(HELINL),-1) + CXXFLAGS += -DMGONGPU_NOINLINE_HELAMPS + GPUFLAGS += -DMGONGPU_NOINLINE_HELAMPS else ifneq ($(HELINL),0) - $(error Unknown HELINL='$(HELINL)': only '0' and '1' are supported) + $(error Unknown HELINL='$(HELINL)': only '-1', '0' and '1' are supported) endif # Set the build flags appropriate to each HRDCOD choice (example: "make HRDCOD=1") diff --git a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/cudacpp_config.mk b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/cudacpp_config.mk index b57e56d182..6ce993ebe1 100644 --- a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/cudacpp_config.mk +++ b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/cudacpp_config.mk @@ -22,7 +22,7 @@ ifeq ($(FPTYPE),) override FPTYPE = m endif -# Set the default HELINL (inline helicities?) choice +# Set the default HELINL choice (1: inline, 0: compiler default, -1: noinline) ifeq ($(HELINL),) override HELINL = 0 endif @@ -44,7 +44,7 @@ ifneq ($(words $(filter $(FPTYPE), $(SUPPORTED_FPTYPES))),1) $(error Invalid fptype FPTYPE='$(FPTYPE)': supported fptypes are $(foreach fptype,$(SUPPORTED_FPTYPES),'$(fptype)')) endif -override SUPPORTED_HELINLS = 0 1 +override SUPPORTED_HELINLS = -1 0 1 ifneq ($(words $(filter $(HELINL), $(SUPPORTED_HELINLS))),1) $(error Invalid helinl HELINL='$(HELINL)': supported helinls are $(foreach helinl,$(SUPPORTED_HELINLS),'$(helinl)')) endif diff --git a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/process_function_definitions.inc b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/process_function_definitions.inc index cd73d52ed3..a5413a9c41 100644 --- a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/process_function_definitions.inc +++ b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/madgraph/iolibs/template_files/gpu/process_function_definitions.inc @@ -120,7 +120,7 @@ namespace mg5amcCpu //-------------------------------------------------------------------------- #ifdef MGONGPUCPP_GPUIMPL - __device__ INLINE unsigned int + __device__ MGONGPU_HELAS_INLINE unsigned int gpu_channelId( const unsigned int* allChannelIds ) { unsigned int channelId = 0; // disable multichannel single-diagram enhancement unless allChannelIds != nullptr diff --git a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/model_handling.py b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/model_handling.py index b72c6c209a..4a93a71641 100644 --- a/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/model_handling.py +++ b/epochX/cudacpp/CODEGEN/PLUGIN/CUDACPP_SA_OUTPUT/model_handling.py @@ -240,8 +240,8 @@ def get_header_txt(self, name=None, couplings=None,mode='', combined=False): out.write(' %(comment)s\n %(template)s\n %(prefix)s void\n %(name)s( const %(args)s,\n%(indent)s%(output)s )%(suffix)s' % {'comment': comment, # AV - add comment 'template': template, # AV - add template - 'prefix': self.prefix + ( ' INLINE' if 'is_h' in mode else '' ), # AV - add INLINE - 'suffix': ( ' ALWAYS_INLINE' if 'is_h' in mode else '' ), # AV - add ALWAYS_INLINE + 'prefix': self.prefix + ( ' MGONGPU_HELAS_INLINE' if 'is_h' in mode else '' ), # AV - add INLINE + 'suffix': "", 'indent':indent, 'output':output, 'name': name, 'args': (',\n' + indent + 'const ').join(args)}) # AV - add const, add indent if 'is_h' in mode: @@ -1500,9 +1500,12 @@ def get_all_sigmaKin_lines(self, color_amplitudes, class_name): const int ievt00 // input: first event number in current C++ event page (for CUDA, ievt depends on threadid) #endif ) - //ALWAYS_INLINE // attributes are not permitted in a function definition + // attributes are not permitted in a function definition { #ifdef MGONGPUCPP_GPUIMPL +#if defined(__CUDACC__) && (CUDART_VERSION >= 13000) + asm volatile (".pragma \\"enable_smem_spilling\\";"); +#endif using namespace mg5amcGpu; using M_ACCESS = DeviceAccessMomenta; // non-trivial access: buffer includes all events using W_ACCESS = DeviceAccessWavefunctions; // TRIVIAL ACCESS (no kernel splitting yet): buffer for one event