[helas] regenerate all processes with support for HELINL=L

epochX/cudacpp/ee_mumu.mad/CODEGEN_mad_ee_mumu_log.txt

@@ -62,7 +62,7 @@ generate e+ e- > mu+ mu-
 No model currently active, so we import the Standard Model
 INFO: load particles 
 INFO: load vertices 
-DEBUG: model prefixing  takes 0.005307912826538086 
+DEBUG: model prefixing  takes 0.0057065486907958984 
 INFO: Restrict model sm with file models/sm/restrict_default.dat . 
 DEBUG: Simplifying conditional expressions 
 DEBUG: remove interactions: u s w+ at order: QED=1 
@@ -154,7 +154,7 @@ INFO: Checking for minimal orders which gives processes.
 INFO: Please specify coupling orders to bypass this step. 
 INFO: Trying process: e+ e- > mu+ mu- WEIGHTED<=4 @1  
 INFO: Process has 2 diagrams 
-1 processes with 2 diagrams generated in 0.004 s
+1 processes with 2 diagrams generated in 0.005 s
 Total: 1 processes with 2 diagrams
 output madevent_simd ../TMPOUT/CODEGEN_mad_ee_mumu --hel_recycling=False --vector_size=32
 Load PLUGIN.CUDACPP_OUTPUT
@@ -176,8 +176,8 @@ INFO: Organizing processes into subprocess groups
 INFO: Generating Helas calls for process: e+ e- > mu+ mu- WEIGHTED<=4 @1 
 INFO: Processing color information for process: e+ e- > mu+ mu- @1 
 INFO: Creating files in directory P1_epem_mupmum 
-DEBUG:  kwargs[prefix] = 0 [model_handling.py at line 1152] 
-DEBUG:  process_exporter_cpp =  <PLUGIN.CUDACPP_OUTPUT.model_handling.PLUGIN_OneProcessExporter object at 0x7f09ed66e490> [export_v4.py at line 6261] 
+DEBUG:  kwargs[prefix] = 0 [model_handling.py at line 1217] 
+DEBUG:  process_exporter_cpp =  <PLUGIN.CUDACPP_OUTPUT.model_handling.PLUGIN_OneProcessExporter object at 0x7fa457e6f550> [export_v4.py at line 6261] 
 INFO: Creating files in directory . 
 FileWriter <class 'PLUGIN.CUDACPP_OUTPUT.model_handling.PLUGIN_CPPWriter'> for ././CPPProcess.h
 FileWriter <class 'PLUGIN.CUDACPP_OUTPUT.model_handling.PLUGIN_CPPWriter'> for ././CPPProcess.cc
@@ -194,22 +194,22 @@ INFO: Created files CPPProcess.h and CPPProcess.cc in directory ./.
 INFO: Generating Feynman diagrams for Process: e+ e- > mu+ mu- WEIGHTED<=4 @1 
 INFO: Finding symmetric diagrams for subprocess group epem_mupmum 
 DEBUG:  os.getcwd() =  /data/avalassi/GPU2023/madgraph4gpuX/MG5aMC/TMPOUT/CODEGEN_mad_ee_mumu/SubProcesses/P1_epem_mupmum [export_v4.py at line 6438] 
-DEBUG:  len(subproc_diagrams_for_config) =  2 [model_handling.py at line 1520] 
-DEBUG:  iconfig_to_diag =  {1: 1, 2: 2} [model_handling.py at line 1544] 
-DEBUG:  diag_to_iconfig =  {1: 1, 2: 2} [model_handling.py at line 1545] 
+DEBUG:  len(subproc_diagrams_for_config) =  2 [model_handling.py at line 1589] 
+DEBUG:  iconfig_to_diag =  {1: 1, 2: 2} [model_handling.py at line 1613] 
+DEBUG:  diag_to_iconfig =  {1: 1, 2: 2} [model_handling.py at line 1614] 
 Generated helas calls for 1 subprocesses (2 diagrams) in 0.004 s
-Wrote files for 8 helas calls in 0.112 s
+Wrote files for 8 helas calls in 0.119 s
 ALOHA: aloha starts to compute helicity amplitudes
 ALOHA: aloha creates FFV1 routines
 ALOHA: aloha creates FFV2 routines
 ALOHA: aloha creates FFV4 routines
-ALOHA: aloha creates 3 routines in  0.198 s
+ALOHA: aloha creates 3 routines in  0.212 s
 ALOHA: aloha starts to compute helicity amplitudes
 ALOHA: aloha creates FFV1 routines
 ALOHA: aloha creates FFV2 routines
 ALOHA: aloha creates FFV4 routines
 ALOHA: aloha creates FFV2_4 routines
-ALOHA: aloha creates 7 routines in  0.253 s
+ALOHA: aloha creates 7 routines in  0.272 s
 <class 'aloha.create_aloha.AbstractRoutine'> FFV1
 <class 'aloha.create_aloha.AbstractRoutine'> FFV1
 <class 'aloha.create_aloha.AbstractRoutine'> FFV2
@@ -220,6 +220,8 @@ ALOHA: aloha creates 7 routines in  0.253 s
 <class 'aloha.create_aloha.AbstractRoutine'> FFV2_4
 FileWriter <class 'PLUGIN.CUDACPP_OUTPUT.model_handling.PLUGIN_CPPWriter'> for /data/avalassi/GPU2023/madgraph4gpuX/MG5aMC/TMPOUT/CODEGEN_mad_ee_mumu/src/./HelAmps_sm.h
 INFO: Created file HelAmps_sm.h in directory /data/avalassi/GPU2023/madgraph4gpuX/MG5aMC/TMPOUT/CODEGEN_mad_ee_mumu/src/. 
+FileWriter <class 'PLUGIN.CUDACPP_OUTPUT.model_handling.PLUGIN_CPPWriter'> for /data/avalassi/GPU2023/madgraph4gpuX/MG5aMC/TMPOUT/CODEGEN_mad_ee_mumu/src/../SubProcesses/./HelAmps.cc
+INFO: Created file HelAmps.cc in directory /data/avalassi/GPU2023/madgraph4gpuX/MG5aMC/TMPOUT/CODEGEN_mad_ee_mumu/src/../SubProcesses/. 
 super_write_set_parameters_onlyfixMajorana (hardcoded=False)
 super_write_set_parameters_onlyfixMajorana (hardcoded=True)
 FileWriter <class 'PLUGIN.CUDACPP_OUTPUT.model_handling.PLUGIN_CPPWriter'> for /data/avalassi/GPU2023/madgraph4gpuX/MG5aMC/TMPOUT/CODEGEN_mad_ee_mumu/src/./Parameters_sm.h
@@ -242,19 +244,17 @@ patching file auto_dsig1.f
 Hunk #1 succeeded at 496 (offset 12 lines).
 patching file driver.f
 patching file matrix1.f
-Hunk #3 succeeded at 230 (offset 9 lines).
-Hunk #4 succeeded at 267 (offset 18 lines).
-Hunk #5 succeeded at 312 (offset 18 lines).
+Hunk #2 succeeded at 229 (offset 9 lines).
 DEBUG:  p.returncode =  0 [output.py at line 242] 
 Output to directory /data/avalassi/GPU2023/madgraph4gpuX/MG5aMC/TMPOUT/CODEGEN_mad_ee_mumu done.
 Type "launch" to generate events from this process, or see
 /data/avalassi/GPU2023/madgraph4gpuX/MG5aMC/TMPOUT/CODEGEN_mad_ee_mumu/README
 Run "open index.html" to see more information about this process.
 quit
 
-real	0m2.067s
-user	0m1.807s
-sys	0m0.251s
+real	0m2.169s
+user	0m1.894s
+sys	0m0.268s
 Code generation completed in 2 seconds
 ************************************************************
 *                                                          *

epochX/cudacpp/ee_mumu.mad/SubProcesses/MemoryAccessGs.h

@@ -10,7 +10,7 @@
 
 #include "MemoryAccessHelpers.h"
 #include "MemoryAccessVectors.h"
-#include "MemoryBuffers.h" // for HostBufferMatrixElements::isaligned
+#include "MemoryBuffers.h" // for HostBufferGs::isaligned
 
 // NB: namespaces mg5amcGpu and mg5amcCpu includes types which are defined in different ways for CPU and GPU builds (see #318 and #725)
 #ifdef MGONGPUCPP_GPUIMPL

epochX/cudacpp/ee_mumu.mad/SubProcesses/MemoryAccessMomenta.h

@@ -8,7 +8,7 @@
 
 #include "mgOnGpuConfig.h"
 
-#include "CPPProcess.h"
+#include "CPPProcess.h" // for CPPProcess::np4 and CPPProcess::npar (NB: npar may differ in different P* subprocess directories!)
 #include "MemoryAccessHelpers.h"
 #include "MemoryAccessVectors.h"
 

epochX/cudacpp/ee_mumu.mad/SubProcesses/P1_epem_mupmum/CPPProcess.cc

@@ -204,7 +204,9 @@ namespace mg5amcCpu
     using M_ACCESS = DeviceAccessMomenta;         // non-trivial access: buffer includes all events
     using E_ACCESS = DeviceAccessMatrixElements;  // non-trivial access: buffer includes all events
     using W_ACCESS = DeviceAccessWavefunctions;   // TRIVIAL ACCESS (no kernel splitting yet): buffer for one event
+#ifndef MGONGPU_LINKER_HELAMPS
     using A_ACCESS = DeviceAccessAmplitudes;      // TRIVIAL ACCESS (no kernel splitting yet): buffer for one event
+#endif
     using CD_ACCESS = DeviceAccessCouplings;      // non-trivial access (dependent couplings): buffer includes all events
     using CI_ACCESS = DeviceAccessCouplingsFixed; // TRIVIAL access (independent couplings): buffer for one event
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
@@ -216,7 +218,9 @@ namespace mg5amcCpu
     using M_ACCESS = HostAccessMomenta;         // non-trivial access: buffer includes all events
     using E_ACCESS = HostAccessMatrixElements;  // non-trivial access: buffer includes all events
     using W_ACCESS = HostAccessWavefunctions;   // TRIVIAL ACCESS (no kernel splitting yet): buffer for one event
+#ifndef MGONGPU_LINKER_HELAMPS
     using A_ACCESS = HostAccessAmplitudes;      // TRIVIAL ACCESS (no kernel splitting yet): buffer for one event
+#endif
     using CD_ACCESS = HostAccessCouplings;      // non-trivial access (dependent couplings): buffer includes all events
     using CI_ACCESS = HostAccessCouplingsFixed; // TRIVIAL access (independent couplings): buffer for one event
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
@@ -328,10 +332,10 @@ namespace mg5amcCpu
 
       oxxxxx<M_ACCESS, W_ACCESS>( momenta, 0., cHel[ihel][3], +1, w_fp[3], 3 );
 
-      FFV1P0_3<W_ACCESS, CI_ACCESS>( w_fp[1], w_fp[0], COUPs[ndcoup + 0], 1.0, 0., 0., w_fp[4] );
+      helas_FFV1P0_3( w_fp[1], w_fp[0], COUPs[ndcoup + 0], 1.0, 0., 0., w_fp[4] );
 
       // Amplitude(s) for diagram number 1
-      FFV1_0<W_ACCESS, A_ACCESS, CI_ACCESS>( w_fp[2], w_fp[3], w_fp[4], COUPs[ndcoup + 0], 1.0, &amp_fp[0] );
+      helas_FFV1_0( w_fp[2], w_fp[3], w_fp[4], COUPs[ndcoup + 0], 1.0, &amp_fp[0] );
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
       if( channelId == 1 ) numerators_sv += cxabs2( amp_sv[0] );
       if( channelId != 0 ) denominators_sv += cxabs2( amp_sv[0] );
@@ -341,10 +345,10 @@ namespace mg5amcCpu
       // *** DIAGRAM 2 OF 2 ***
 
       // Wavefunction(s) for diagram number 2
-      FFV2_4_3<W_ACCESS, CI_ACCESS>( w_fp[1], w_fp[0], COUPs[ndcoup + 1], 1.0, COUPs[ndcoup + 2], 1.0, cIPD[0], cIPD[1], w_fp[4] );
+      helas_FFV2_4_3( w_fp[1], w_fp[0], COUPs[ndcoup + 1], 1.0, COUPs[ndcoup + 2], 1.0, cIPD[0], cIPD[1], w_fp[4] );
 
       // Amplitude(s) for diagram number 2
-      FFV2_4_0<W_ACCESS, A_ACCESS, CI_ACCESS>( w_fp[2], w_fp[3], w_fp[4], COUPs[ndcoup + 1], 1.0, COUPs[ndcoup + 2], 1.0, &amp_fp[0] );
+      helas_FFV2_4_0( w_fp[2], w_fp[3], w_fp[4], COUPs[ndcoup + 1], 1.0, COUPs[ndcoup + 2], 1.0, &amp_fp[0] );
 #ifdef MGONGPU_SUPPORTS_MULTICHANNEL
       if( channelId == 2 ) numerators_sv += cxabs2( amp_sv[0] );
       if( channelId != 0 ) denominators_sv += cxabs2( amp_sv[0] );

epochX/cudacpp/ee_mumu.mad/SubProcesses/P1_epem_mupmum/check_sa.cc

@@ -958,6 +958,8 @@ main( int argc, char** argv )
               << " [" << process.getCompiler() << "]"
 #ifdef MGONGPU_INLINE_HELAMPS
               << " [inlineHel=1]"
+#elif defined MGONGPU_LINKER_HELAMPS
+              << " [inlineHel=L]"
 #else
               << " [inlineHel=0]"
 #endif

epochX/cudacpp/ee_mumu.mad/SubProcesses/cudacpp.mk

@@ -551,8 +551,11 @@ $(info HELINL='$(HELINL)')
 ifeq ($(HELINL),1)
   CXXFLAGS += -DMGONGPU_INLINE_HELAMPS
   GPUFLAGS += -DMGONGPU_INLINE_HELAMPS
+else ifeq ($(HELINL),L)
+  CXXFLAGS += -DMGONGPU_LINKER_HELAMPS
+  GPUFLAGS += -DMGONGPU_LINKER_HELAMPS
 else ifneq ($(HELINL),0)
-  $(error Unknown HELINL='$(HELINL)': only '0' and '1' are supported)
+  $(error Unknown HELINL='$(HELINL)': only 'L,', '0' and '1' are supported)
 endif
 
 # Set the build flags appropriate to each HRDCOD choice (example: "make HRDCOD=1")
@@ -647,7 +650,6 @@ override RUNTIME =
 #=== Makefile TARGETS and build rules below
 #===============================================================================
 
-
 ifeq ($(GPUCC),)
   cxx_checkmain=$(BUILDDIR)/check_cpp.exe
   cxx_fcheckmain=$(BUILDDIR)/fcheck_cpp.exe
@@ -776,6 +778,14 @@ gpu_objects_lib=$(BUILDDIR)/CPPProcess_$(GPUSUFFIX).o $(BUILDDIR)/MatrixElementK
 gpu_objects_exe=$(BUILDDIR)/CommonRandomNumberKernel_$(GPUSUFFIX).o $(BUILDDIR)/RamboSamplingKernels_$(GPUSUFFIX).o
 endif
 
+# Add object files and special build flags only for the HELINL=L mode
+ifeq ($(HELINL),L)
+cxx_objects_lib+=$(BUILDDIR)/HelAmps_cpp.o
+gpu_objects_lib+=$(BUILDDIR)/HelAmps_$(GPUSUFFIX).o
+$(BUILDDIR)/CPPProcess_$(GPUSUFFIX).o: GPUFLAGS += -rdc true # compilation fails if this is not added (ptxas fatal: Unresolved extern function)
+$(BUILDDIR)/HelAmps_$(GPUSUFFIX).o: GPUFLAGS += -rdc true # runtime fails if this is not added ('invalid device symbol' in CPPProcess.cc cHel to tHel copy)
+endif
+
 # Target (and build rules): C++ and CUDA/HIP shared libraries
 $(LIBDIR)/lib$(MG5AMC_CXXLIB).so: $(BUILDDIR)/fbridge_cpp.o
 $(LIBDIR)/lib$(MG5AMC_CXXLIB).so: cxx_objects_lib += $(BUILDDIR)/fbridge_cpp.o
@@ -786,12 +796,12 @@ ifneq ($(GPUCC),)
 $(LIBDIR)/lib$(MG5AMC_GPULIB).so: $(BUILDDIR)/fbridge_$(GPUSUFFIX).o
 $(LIBDIR)/lib$(MG5AMC_GPULIB).so: gpu_objects_lib += $(BUILDDIR)/fbridge_$(GPUSUFFIX).o
 $(LIBDIR)/lib$(MG5AMC_GPULIB).so: $(LIBDIR)/lib$(MG5AMC_COMMONLIB).so $(gpu_objects_lib)
-	$(GPUCC) --shared -o $@ $(gpu_objects_lib) $(GPULIBFLAGSRPATH2) -L$(LIBDIR) -l$(MG5AMC_COMMONLIB)
+	$(GPUCC) --shared -o $@ $(gpu_objects_lib) $(GPUARCHFLAGS) $(GPULIBFLAGSRPATH2) -L$(LIBDIR) -l$(MG5AMC_COMMONLIB)
 # Bypass std::filesystem completely to ease portability on LUMI #803
 #ifneq ($(findstring hipcc,$(GPUCC)),)
-#	$(GPUCC) --shared -o $@ $(gpu_objects_lib) $(GPULIBFLAGSRPATH2) -L$(LIBDIR) -l$(MG5AMC_COMMONLIB) -lstdc++fs
+#	$(GPUCC) --shared -o $@ $(gpu_objects_lib) $(GPUARCHFLAGS) $(GPULIBFLAGSRPATH2) -L$(LIBDIR) -l$(MG5AMC_COMMONLIB) -lstdc++fs
 #else
-#	$(GPUCC) --shared -o $@ $(gpu_objects_lib) $(GPULIBFLAGSRPATH2) -L$(LIBDIR) -l$(MG5AMC_COMMONLIB)
+#	$(GPUCC) --shared -o $@ $(gpu_objects_lib) $(GPUARCHFLAGS) $(GPULIBFLAGSRPATH2) -L$(LIBDIR) -l$(MG5AMC_COMMONLIB)
 #endif
 endif
 
@@ -962,6 +972,7 @@ $(cxx_testmain): LIBFLAGS += $(CXXLIBFLAGSRPATH) # avoid the need for LD_LIBRARY
 $(cxx_testmain): $(LIBDIR)/lib$(MG5AMC_COMMONLIB).so $(cxx_objects_lib) $(cxx_objects_exe) $(GTESTLIBS)
 	$(CXX) -o $@ $(cxx_objects_lib) $(cxx_objects_exe) -ldl -pthread $(LIBFLAGS)
 else # link only runTest_$(GPUSUFFIX).o (new: in the past, this was linking both runTest_cpp.o and runTest_$(GPUSUFFIX).o)
+$(gpu_testmain): LIBFLAGS += $(GPUARCHFLAGS) # avoid "nvlink warning: SM Arch not found" when using rdc
 ###$(gpu_testmain): LIBFLAGS += $(GPULIBFLAGSASAN)
 $(gpu_testmain): LIBFLAGS += $(GPULIBFLAGSRPATH) # avoid the need for LD_LIBRARY_PATH
 $(gpu_testmain): $(LIBDIR)/lib$(MG5AMC_COMMONLIB).so $(gpu_objects_lib) $(gpu_objects_exe) $(GTESTLIBS)

epochX/cudacpp/ee_mumu.mad/src/HelAmps_sm.h

@@ -1207,8 +1207,128 @@ namespace mg5amcCpu
     return;
   }
 
+  //==========================================================================
+
+#ifndef MGONGPU_LINKER_HELAMPS
+
+#define helas_FFV1_0 FFV1_0<W_ACCESS, A_ACCESS, CD_ACCESS>
+#define helas_FFV1P0_3 FFV1P0_3<W_ACCESS, CD_ACCESS>
+#define helas_FFV2_0 FFV2_0<W_ACCESS, A_ACCESS, CD_ACCESS>
+#define helas_FFV2_3 FFV2_3<W_ACCESS, CD_ACCESS>
+#define helas_FFV4_0 FFV4_0<W_ACCESS, A_ACCESS, CD_ACCESS>
+#define helas_FFV4_3 FFV4_3<W_ACCESS, CD_ACCESS>
+#define helas_FFV2_4_0 FFV2_4_0<W_ACCESS, A_ACCESS, CD_ACCESS>
+#define helas_FFV2_4_3 FFV2_4_3<W_ACCESS, CD_ACCESS>
+
+#else
+
+#define helas_FFV1_0 linker_FFV1_0
+#define helas_FFV1P0_3 linker_FFV1P0_3
+#define helas_FFV2_0 linker_FFV2_0
+#define helas_FFV2_3 linker_FFV2_3
+#define helas_FFV4_0 linker_FFV4_0
+#define helas_FFV4_3 linker_FFV4_3
+#define helas_FFV2_4_0 linker_FFV2_4_0
+#define helas_FFV2_4_3 linker_FFV2_4_3
+
+  //--------------------------------------------------------------------------
+
+  // Compute the output amplitude 'vertex' from the input wavefunctions F1[6], F2[6], V3[6]
+  __device__ void
+  linker_FFV1_0( const fptype allF1[],
+                 const fptype allF2[],
+                 const fptype allV3[],
+                 const fptype allCOUP[],
+                 const double Ccoeff,
+                 fptype allvertexes[] );
+
+  //--------------------------------------------------------------------------
+
+  // Compute the output wavefunction 'V3[6]' from the input wavefunctions F1[6], F2[6]
+  __device__ void
+  linker_FFV1P0_3( const fptype allF1[],
+                   const fptype allF2[],
+                   const fptype allCOUP[],
+                   const double Ccoeff,
+                   const fptype M3,
+                   const fptype W3,
+                   fptype allV3[] );
+
   //--------------------------------------------------------------------------
 
+  // Compute the output amplitude 'vertex' from the input wavefunctions F1[6], F2[6], V3[6]
+  __device__ void
+  linker_FFV2_0( const fptype allF1[],
+                 const fptype allF2[],
+                 const fptype allV3[],
+                 const fptype allCOUP[],
+                 const double Ccoeff,
+                 fptype allvertexes[] );
+
+  //--------------------------------------------------------------------------
+
+  // Compute the output wavefunction 'V3[6]' from the input wavefunctions F1[6], F2[6]
+  __device__ void
+  linker_FFV2_3( const fptype allF1[],
+                 const fptype allF2[],
+                 const fptype allCOUP[],
+                 const double Ccoeff,
+                 const fptype M3,
+                 const fptype W3,
+                 fptype allV3[] );
+
+  //--------------------------------------------------------------------------
+
+  // Compute the output amplitude 'vertex' from the input wavefunctions F1[6], F2[6], V3[6]
+  __device__ void
+  linker_FFV4_0( const fptype allF1[],
+                 const fptype allF2[],
+                 const fptype allV3[],
+                 const fptype allCOUP[],
+                 const double Ccoeff,
+                 fptype allvertexes[] );
+
+  //--------------------------------------------------------------------------
+
+  // Compute the output wavefunction 'V3[6]' from the input wavefunctions F1[6], F2[6]
+  __device__ void
+  linker_FFV4_3( const fptype allF1[],
+                 const fptype allF2[],
+                 const fptype allCOUP[],
+                 const double Ccoeff,
+                 const fptype M3,
+                 const fptype W3,
+                 fptype allV3[] );
+
+  //--------------------------------------------------------------------------
+
+  // Compute the output amplitude 'vertex' from the input wavefunctions F1[6], F2[6], V3[6]
+  __device__ void
+  linker_FFV2_4_0( const fptype allF1[],
+                   const fptype allF2[],
+                   const fptype allV3[],
+                   const fptype allCOUP[],
+                   const double Ccoeff,
+                   fptype allvertexes[] );
+
+  //--------------------------------------------------------------------------
+
+  // Compute the output wavefunction 'V3[6]' from the input wavefunctions F1[6], F2[6]
+  __device__ void
+  linker_FFV2_4_3( const fptype allF1[],
+                   const fptype allF2[],
+                   const fptype allCOUP[],
+                   const double Ccoeff,
+                   const fptype M3,
+                   const fptype W3,
+                   fptype allV3[] );
+
+  //--------------------------------------------------------------------------
+
+#endif
+
+  //==========================================================================
+
 } // end namespace
 
 #endif // HelAmps_sm_H

epochX/cudacpp/ee_mumu.mad/src/cudacpp_config.mk

@@ -40,7 +40,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 = L 0 1
 ifneq ($(words $(filter $(HELINL), $(SUPPORTED_HELINLS))),1)
   $(error Invalid helinl HELINL='$(HELINL)': supported helinls are $(foreach helinl,$(SUPPORTED_HELINLS),'$(helinl)'))
 endif