Crossing: integrate the helicity PRs (#278 #279 #263 #255 #261 #217), MadSpin with crossing

.github/gpu_ci/README.md

@@ -24,18 +24,28 @@ allocation per workflow run, and every CI job of that run executes in it:
 | `pp_jj` | cluster | the same with `generate p p > j j` (`PROCESS`/`TAG` of `pp_ttx_mg7.sh`): its jet cuts act on several objects, which `p p > t t~` has none of |
 | `crossing_folding` | cluster | `p p > w+ j`, `p p > j j` and `p p > w+ j j` generated `--use_crossing=True`, each written folded and `--use_crossing=False` expanded, each run on the GPU and on cpu with one seed, then the cpu-made `p p > w+ j` gridpacks on the GPU (`crossing_folding_mg7.sh`): every run must agree with its counterpart (folded = expanded, GPU = cpu, gridpack = its run) |
 | `madspace_tests` | cluster | the madspace unit tests, `madspace/tests`, with `MADSPACE_TEST_DEVICE=cuda` or `hip` (`madspace_tests_gpu.sh`): the tests put their torch tensors on the GPU, so madspace runs them on its GPU runtime. Host-side tests (marker `cpu_only`) are skipped. pytest and scipy are installed in the cache when the modules lack them, and so is a CUDA torch wheel when the PyTorch module is a CPU build (manneback). The log and `junit.xml` are uploaded as an artifact |
+| `acceptance` | cluster | the acceptance tests that need a GPU, through `tests/test_manager.py` (`acceptance_gpu.sh`): they skip without `nvcc`/`hipcc` or a device, so a skip fails the job. Now `test_standalone_umami_first_call_gpu` (the good helicities chosen at `umami_initialize`, with device buffers). To add a test, append it to `DEFAULT_TESTS` in `acceptance_gpu.sh` |
 | `interference` | cluster | squared split orders on the GPU backend of madmatrix (`interference_gpu.sh`, `interference_checks.py`): standalone `u u~ > u u~` constraints, the helicity choice of an interference matrix element through umami, with `check_sa.exe` (exact and summed), and `p p > u u~ QCD^2==2` with mg7 against its CPU reference. Input `interference`, on by default and on push |
-| `madmatrix` | cluster | the GPU backend of madmatrix against its scalar CPU backend (`madmatrix_gpu.sh`, `madmatrix_checks.py`), with the `check_sa.exe` of each backend fed the same events from files: MSSM `p p > go go` event by event (running flavour couplings), a subprocess with no good helicity (`b b~ > ta+ ta- / z a`, ymtau = 0), and a non-blocking stream handed to umami (`g g > t t~ g`, `--nonblocking-stream`) |
+| `madmatrix` | cluster | the GPU backend of madmatrix against its scalar CPU backend (`madmatrix_gpu.sh`, `madmatrix_checks.py`), with the `check_sa.exe` of each backend fed the same events from files: MSSM `p p > go go` event by event (running flavour couplings), MSSM `p p > go go j` over every crossing of the subprocess that folds them (extended flavour ids, helicities and colours included), the same for `u g > u z` with `u u~ > z g` (a massive leg crossed into a massless slot) for `p p > z j --axial_gauge=True` (a crossed axial-gauge gluon) and for `p p > j j` and `p p > t t~ j` at `QED^2==2` (a crossed interference, whose helicities are signed: helicity sum and signed draws), a subprocess with no good helicity (`b b~ > ta+ ta- / z a`, ymtau = 0) and |M|^2 = 0 with good helicities (`p p > w+ j` at alpha_s = 0, after a call at 0.118): no helicity or colour, and a non-blocking stream handed to umami (`g g > t t~ g`, `--nonblocking-stream`) |
 | `stop_runner` | cluster | clean up, then create the stop file: `runner_batch.sh` stops the runner and the allocation ends |
 | `watchdog` | vmadgraph | from the start of the runner to the end of the allocation: if the allocation ends before `stop_runner` completed, cancel the run (`ceci-gate watch-runner`, `remote_watch_runner.sh`) |
 
+Without a GPU, `host_hip/` runs the madmatrix GPU code on the host: with `host_hip/bin` first on
+the `PATH`, `make BACKEND=hip` builds the GPU templates as host C++ (its `hipcc` rewrites the kernel
+launches, `include/hip/hip_runtime.h` runs every kernel serially with host memory as device
+memory), so `madmatrix_checks.py --backend hip` compares the logic of the GPU code with the CPU
+backend. The acceptance test `test_gpu_backend_host_emulation` does that for the crossing checks.
+It is no substitute for an nvcc/hipcc build on a device.
+
 To add a CI job, give it `needs: [start_runner, ...]`,
 `if: needs.start_runner.outputs.attempt == github.run_attempt` (see *Re-runs* below) and
 `runs-on: [self-hosted, "${{ inputs.runner_label }}"]`, and add it to the `needs` of `stop_runner`. It runs in the same allocation, after the others, so
 `time_limit` must cover all the jobs together. The first madspace build (OpenBLAS + GPU code)
 is the slow part. The default is kept short so that Slurm can backfill the allocation:
-20 minutes on lemaitre4 (successful runs take 1-12 minutes) and 30 minutes on manneback
-(usually 2-17 minutes, but the runner set-up and the checkout are sometimes slow there).
+30 minutes on both clusters. On lemaitre4 a run takes about 19 minutes with the madspace
+build cached (the madmatrix and acceptance jobs brought it there from 15), and a madspace
+rebuild adds about 3. On manneback runs usually take 2-17 minutes, but the runner set-up and
+the checkout are sometimes slow there.
 Raise it from the Actions tab for a run that needs more, or in the caller workflow when a
 new job makes every run longer.
 

.github/gpu_ci/acceptance_gpu.sh

@@ -0,0 +1,74 @@
+#!/bin/bash -l
+# Runs on the self-hosted runner (GPU node): job acceptance of gpu_runner_ci.yml.
+# The acceptance tests that need a GPU, through tests/test_manager.py. Such a test skips
+# without nvcc/hipcc or without a device, which on a GPU node is a failure (test_manager
+# exits 1 on a skip, and summary.txt says which).
+# Environment: as pp_ttx_mg7.sh (BACKEND, MODULES, GPU_ARCH, VENV, WORKDIR); no madspace
+# and no PDF set needed (standalone outputs only). GPU_TESTS: the tests to run, space
+# separated (test names, or test_manager expressions), by default the list below.
+set -eo pipefail
+HERE=$(cd "$(dirname "$0")" && pwd)
+REPO=$(cd "$HERE/../.." && pwd)
+section() { echo; echo "=================== $* ($(date +%T))"; }
+
+# the good helicities chosen at umami_initialize, with device buffers
+DEFAULT_TESTS="test_standalone_umami_first_call_gpu"
+
+section "Environment"
+if [ -n "$MODULES" ]; then source "$HERE/load_modules.sh"; fi
+source "$VENV/bin/activate"
+python3 --version
+case $BACKEND in
+    cuda) nvidia-smi
+          GPU_NAME=$(nvidia-smi --query-gpu=name --format=csv,noheader | sed -n 1p) ;;
+    hip)  rocm-smi || true
+          GPU_NAME=$(rocminfo 2> /dev/null | awk -F': *' '/Marketing Name/ {name=$2} /Device Type: *GPU/ {print name; exit}' || true) ;;
+    *)    echo "::error::the GPU acceptance tests need a GPU backend, not '$BACKEND'"; exit 2 ;;
+esac
+export GPU_NAME
+if [ -n "$GPU_ARCH" ]; then
+    export MADGRAPH_CUDA_ARCHITECTURE=$GPU_ARCH MADGRAPH_HIP_ARCHITECTURE=$GPU_ARCH
+fi
+# the tests call make without -j
+export MAKEFLAGS="-j${SLURM_CPUS_PER_TASK:-4}"
+
+# default configuration, as .github/actions/checkout_mg5 does for the other CI jobs
+cp "$REPO/input/.mg7_configuration_default.txt" "$REPO/input/mg7_configuration.txt"
+cp "$REPO/Template/LO/Source/.make_opts" "$REPO/Template/LO/Source/make_opts"
+cat >> "$REPO/input/mg7_configuration.txt" << EOF
+auto_update = 0
+automatic_html_opening = False
+notification_center = False
+nb_core = ${SLURM_CPUS_PER_TASK:-4}
+EOF
+
+rm -rf "$WORKDIR"
+mkdir -p "$WORKDIR"
+cd "$WORKDIR"
+STATUS=0
+: > summary.txt
+for test in ${GPU_TESTS:-$DEFAULT_TESTS}; do
+    section "$test"
+    log=$WORKDIR/${test//[^A-Za-z0-9_.-]/_}.log
+    if (cd "$REPO" && python3 ./tests/test_manager.py -p A -t0 -v 2 "$test") 2>&1 | tee "$log"; then
+        result=ok
+    else
+        result=FAILED
+        STATUS=1
+    fi
+    # unittest's own verdict ("Ran 1 test in 95.2s", "OK"), and why a test skipped: the
+    # verdict of test_manager's runner says OK for a skip
+    ran=$(grep -E '^Ran [0-9]+ tests? in ' "$log" | tail -1 || true)
+    verdict=$(grep -E '^(OK|FAILED)( \(.*\))?$' "$log" | tail -1 || true)
+    skipped=$(grep -oE '\.\.\. skipped .*' "$log" | sed 's/^\.\.\. //' | tail -1 || true)
+    echo "$test=$result (${ran:-no test ran}, ${verdict:-no verdict}${skipped:+, $skipped})" >> summary.txt
+done
+{
+    echo "node=$(hostname)"
+    echo "gpu=${GPU_NAME:-none}"
+    echo "backend=$BACKEND"
+} >> summary.txt
+
+section "Summary"
+cat summary.txt
+exit $STATUS

.github/gpu_ci/host_hip/bin/hipcc

@@ -0,0 +1,69 @@
+#!/usr/bin/env python3
+"""hipcc of the host emulation of HIP (see ../include/hip/hip_runtime.h).
+
+With this directory first on the PATH, 'make BACKEND=hip' in a madmatrix output
+builds the GPU templates (backend/gpu) as host C++: a source is preprocessed,
+its kernel launches kernel<<<...>>>( args ) are rewritten into calls of
+host_hip::Launch, and the result is compiled by the host compiler ($CXX, else
+c++). The GPU-only options of madmatrix.mk (--offload-arch, -target, -x hip) are
+dropped. Only for checking the GPU code without a GPU toolchain.
+"""
+
+import os
+import re
+import subprocess
+import sys
+import tempfile
+
+ROOT = os.path.dirname(os.path.dirname(os.path.abspath(__file__)))
+LAUNCH = re.compile(r'([A-Za-z_][\w:]*(?:<[^<>;]*>)?)\s*<<<(.*?)>>>', re.S)
+
+
+def main(argv):
+    args, sources, out, compile_only = [], [], None, False
+    i = 0
+    while i < len(argv):
+        arg = argv[i]
+        if arg.startswith('--offload-arch') or arg.startswith('-fgpu-rdc'):
+            pass
+        elif arg in ('-target', '-x'):
+            i += 1  # the language is C++ once the launches are rewritten
+        elif arg == '--shared':
+            args.append('-shared')
+        elif (sys.platform == 'darwin' and arg == '-Xlinker' and i + 3 < len(argv)
+              and argv[i + 1] == '-soname'):
+            args.append('-Wl,-install_name,@rpath/' + argv[i + 3])
+            i += 3
+        elif arg == '-o':
+            out = argv[i + 1]
+            i += 1
+        elif arg == '-c':
+            compile_only = True
+        elif not arg.startswith('-') and re.search(r'\.(cc|cpp|cu)$', arg):
+            sources.append(arg)
+        else:
+            args.append(arg)
+        i += 1
+    cxx = os.environ.get('CXX') or 'c++'
+    base = ['-D__HIPCC__', '-D__HIP_PLATFORM_AMD__', '-DMGONGPU_HAS_NO_HIPRAND',
+            '-I' + os.path.join(ROOT, 'include')]
+    if not compile_only:
+        return subprocess.call([cxx] + base + args + sources + (['-o', out] if out else []))
+    (source,) = sources
+    pre = subprocess.run([cxx, '-E', '-x', 'c++'] + base + args + [source],
+                         stdout=subprocess.PIPE, universal_newlines=True)
+    if pre.returncode:
+        return pre.returncode
+    text = LAUNCH.sub(r'host_hip::Launch(\2)(\1)', pre.stdout)
+    fd, expanded = tempfile.mkstemp(suffix='.ii')
+    try:
+        with os.fdopen(fd, 'w') as fsock:
+            fsock.write(text)
+        return subprocess.call([cxx, '-c', '-x', 'c++-cpp-output'] + base + args +
+                               [expanded, '-o', out])
+    finally:
+        os.unlink(expanded)
+
+
+if __name__ == '__main__':
+    sys.exit(main(sys.argv[1:]))

.github/gpu_ci/host_hip/bin/hipconfig

@@ -0,0 +1,6 @@
+#!/bin/sh
+# hipconfig of the host emulation of HIP: madmatrix.mk takes HIP_HOME from it
+case "$1" in
+    --rocmpath|--path) cd "$(dirname "$0")/.." && pwd ;;
+    *) echo "host emulation of HIP (.github/gpu_ci/host_hip)" ;;
+esac

.github/gpu_ci/host_hip/include/hip/hip_runtime.h

@@ -0,0 +1,172 @@
+// Serial host emulation of the HIP runtime, for the madmatrix GPU templates
+// (backend/gpu) on a machine without a GPU toolchain: bin/hipcc compiles them as
+// host C++ with this header, every kernel launch runs its grid x block threads one
+// after the other on the host, and device memory is host memory. It checks the
+// logic of the GPU code (indexing, crossing tables, good helicities, selection)
+// against the CPU backends; it cannot replace a real nvcc/hipcc build and run.
+// The GPU code has no __syncthreads and no shared memory, so a serial run of the
+// threads is a valid schedule; atomicAdd is then a plain addition.
+// Compiled with clang -x cuda (__CUDA__), the same header only declares the
+// device attributes: a device-side semantic check of the sources.
+#ifndef HOST_HIP_RUNTIME_H
+#define HOST_HIP_RUNTIME_H 1
+
+#include <cstddef>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+#ifdef __CUDA__
+#define __device__ __attribute__( ( device ) )
+#define __global__ __attribute__( ( global ) )
+#define __host__ __attribute__( ( host ) )
+#define __constant__ __attribute__( ( constant ) )
+#define __shared__ __attribute__( ( shared ) )
+#define __forceinline__ __attribute__( ( always_inline ) ) inline
+#include <__clang_cuda_builtin_vars.h>
+#define HOST_HIP_HD __attribute__( ( host ) ) __attribute__( ( device ) )
+#else
+#define __device__
+#define __global__
+#define __host__
+#define __constant__
+#define __shared__
+#define __forceinline__ inline
+#define HOST_HIP_HD
+#endif
+
+struct dim3
+{
+  unsigned int x, y, z;
+  HOST_HIP_HD constexpr dim3( unsigned int x_ = 1, unsigned int y_ = 1, unsigned int z_ = 1 )
+    : x( x_ ), y( y_ ), z( z_ ) {}
+};
+
+#ifdef __CUDA__
+extern "C" __host__ int cudaConfigureCall( dim3 grid, dim3 block, size_t shared = 0, void* stream = 0 );
+#else
+inline dim3 threadIdx( 0, 0, 0 ), blockIdx( 0, 0, 0 ), blockDim( 1, 1, 1 ), gridDim( 1, 1, 1 );
+#endif
+
+typedef enum
+{
+  hipSuccess = 0,
+  hipErrorInvalidValue = 1
+} hipError_t;
+typedef enum
+{
+  hipMemcpyHostToHost = 0,
+  hipMemcpyHostToDevice = 1,
+  hipMemcpyDeviceToHost = 2,
+  hipMemcpyDeviceToDevice = 3,
+  hipMemcpyDefault = 4
+} hipMemcpyKind;
+typedef void* hipStream_t;
+#define hipStreamNonBlocking 1
+
+inline const char* hipGetErrorString( hipError_t e ) { return e == hipSuccess ? "hipSuccess" : "hipError"; }
+inline hipError_t hipPeekAtLastError() { return hipSuccess; }
+inline hipError_t hipGetLastError() { return hipSuccess; }
+inline hipError_t hipSetDevice( int ) { return hipSuccess; }
+inline hipError_t hipDeviceSynchronize() { return hipSuccess; }
+inline hipError_t hipDeviceReset() { return hipSuccess; }
+
+// A fresh allocation is filled with 0xff bytes (NaN doubles, -1 ints), so that a
+// read of memory the GPU code never wrote shows up instead of reading zeros
+inline void* host_hip_alloc( size_t n )
+{
+  void* p = nullptr;
+  if( posix_memalign( &p, 256, n ? n : 1 ) != 0 ) return nullptr;
+  memset( p, 0xff, n );
+  return p;
+}
+template<typename T>
+inline hipError_t hipMalloc( T** p, size_t n ) { *p = (T*)host_hip_alloc( n ); return *p ? hipSuccess : hipErrorInvalidValue; }
+template<typename T>
+inline hipError_t hipHostMalloc( T** p, size_t n, unsigned int = 0 ) { return hipMalloc( p, n ); }
+template<typename T>
+inline hipError_t hipMallocAsync( T** p, size_t n, hipStream_t ) { return hipMalloc( p, n ); }
+inline hipError_t hipFree( void* p ) { free( p ); return hipSuccess; }
+inline hipError_t hipHostFree( void* p ) { free( p ); return hipSuccess; }
+inline hipError_t hipFreeAsync( void* p, hipStream_t ) { free( p ); return hipSuccess; }
+inline hipError_t hipMemcpy( void* d, const void* s, size_t n, hipMemcpyKind ) { memmove( d, s, n ); return hipSuccess; }
+inline hipError_t hipMemcpyAsync( void* d, const void* s, size_t n, hipMemcpyKind, hipStream_t = nullptr ) { memmove( d, s, n ); return hipSuccess; }
+inline hipError_t hipMemset( void* d, int v, size_t n ) { memset( d, v, n ); return hipSuccess; }
+inline hipError_t hipMemsetAsync( void* d, int v, size_t n, hipStream_t = nullptr ) { memset( d, v, n ); return hipSuccess; }
+template<typename T>
+inline hipError_t hipMemcpyToSymbol( T& sym, const void* s, size_t n, size_t off = 0, hipMemcpyKind = hipMemcpyHostToDevice )
+{
+  if( off + n > sizeof( T ) )
+  {
+    fprintf( stderr, "host hip: hipMemcpyToSymbol of %zu bytes at %zu into a %zu-byte symbol\n", n, off, sizeof( T ) );
+    abort();
+  }
+  memcpy( (char*)&sym + off, s, n );
+  return hipSuccess;
+}
+template<typename T>
+inline hipError_t hipMemcpyFromSymbol( void* d, const T& sym, size_t n, size_t off = 0, hipMemcpyKind = hipMemcpyDeviceToHost )
+{
+  if( off + n > sizeof( T ) )
+  {
+    fprintf( stderr, "host hip: hipMemcpyFromSymbol of %zu bytes at %zu from a %zu-byte symbol\n", n, off, sizeof( T ) );
+    abort();
+  }
+  memcpy( d, (const char*)&sym + off, n );
+  return hipSuccess;
+}
+template<typename P, typename T>
+inline hipError_t hipGetSymbolAddress( P** p, T& sym ) { *p = (P*)&sym; return hipSuccess; }
+inline hipError_t hipStreamCreate( hipStream_t* s ) { *s = (hipStream_t)0x1; return hipSuccess; }
+inline hipError_t hipStreamCreateWithFlags( hipStream_t* s, unsigned int ) { *s = (hipStream_t)0x2; return hipSuccess; }
+inline hipError_t hipStreamDestroy( hipStream_t ) { return hipSuccess; }
+inline hipError_t hipStreamSynchronize( hipStream_t ) { return hipSuccess; }
+
+template<typename T>
+HOST_HIP_HD inline T atomicAdd( T* a, T v )
+{
+  T old = *a;
+  *a += v;
+  return old;
+}
+
+#ifndef __CUDA__
+namespace host_hip
+{
+  // bin/hipcc rewrites kernel<<<grid, block, shared, stream>>>( args ) into
+  // host_hip::Launch( grid, block, shared, stream )( kernel )( args )
+  struct Launch
+  {
+    dim3 grid, block;
+    Launch( dim3 g, dim3 b, size_t = 0, hipStream_t = nullptr ) : grid( g ), block( b ) {}
+    template<typename K>
+    auto operator()( K kernel ) const
+    {
+      const dim3 g = grid, b = block;
+      return [=]( auto&&... args ) {
+        if( g.x * g.y * g.z == 0 || b.x * b.y * b.z == 0 )
+        {
+          // a real GPU fails the launch (invalid configuration)
+          fprintf( stderr, "host hip: invalid launch configuration\n" );
+          abort();
+        }
+        gridDim = g;
+        blockDim = b;
+        for( unsigned bz = 0; bz < g.z; bz++ )
+          for( unsigned by = 0; by < g.y; by++ )
+            for( unsigned bx = 0; bx < g.x; bx++ )
+              for( unsigned tz = 0; tz < b.z; tz++ )
+                for( unsigned ty = 0; ty < b.y; ty++ )
+                  for( unsigned tx = 0; tx < b.x; tx++ )
+                  {
+                    blockIdx = dim3( bx, by, bz );
+                    threadIdx = dim3( tx, ty, tz );
+                    kernel( args... );
+                  }
+      };
+    }
+  };
+}
+#endif
+
+#endif // HOST_HIP_RUNTIME_H

.github/gpu_ci/interference_checks.py

@@ -7,7 +7,7 @@
 1. standalone ``u u~ > u u~`` with QED^2==0, ==2, ==4 and <=4, FPTYPE=d, at the classic
    RAMBO point (1000 GeV) of check_sa.exe:
    - QED^2==2, the QCD x EW interference alone, is the Fortran split-order value
-     -5.5828746494657265e-02 (an unmasked colour sum would give the total, +2.776);
+     -5.5828747824035893e-02 (an unmasked colour sum would give the total, +2.776);
    - the three components add up to the QED^2<=4 total of the same build: this pins the
      pair loop of the colour sum (all ORDERED pairs of amplitude orders, not a triangle);
    and QED^2==2 again with FPTYPE=m (colour algebra in single precision).
@@ -44,8 +44,9 @@
 GPU_BACKENDS = ('cuda', 'hip')
 
 # u u~ > u u~ QED^2==2 at the check_sa RAMBO point, from the Fortran split-order
-# standalone (madmatrix CPU: -5.5828746494657258e-02)
-UUX_INTERFERENCE = -5.5828746494657265e-02
+# standalone (madmatrix CPU simd_128: -5.5828747824035928e-02); the t-channel Z has
+# no width (zerowidth_tchannel, the default: -5.5828746494657265e-02 with it)
+UUX_INTERFERENCE = -5.5828747824035893e-02
 # mg7 p p > u u~ QCD^2==2, fixed scale 91.188 GeV, NNPDF23_lo_as_0130_qed, default cuts:
 # three CPU runs of 200k events (tests/acceptance_tests/test_mg7_interference.py)
 UUX_MG7_CROSS, UUX_MG7_ERROR = -12253., 10.