[willend]

This PR has two parts:

1. Single_crystal selects the reflection during the search and keeps no list by default (numerical change, below). -DMCSX_REFL_SLIST brings the list back on CPU for SPLIT.
2. A new OpenMP target offload backend, -DOMPTARGET, next to OpenACC. It needs no unified memory, so it suits LUMI, where HSA_XNACK is unavailable. mcstas NMX_simplified.instr now gives a .c that builds for it with no hand edits.

OMPTARGET backend

• Macros: -DOPENACC and -DOMPTARGET are mutually exclusive. MCGPU means "device build", and OPENACC stays for the OpenACC API. The guard rename covers the runtime and the ported components only; other components keep OPENACC, which still works for OpenACC builds.
• Code generator: omp_twin() emits the OpenMP counterpart of each #pragma acc line behind #ifdef OMPTARGET, so components need no pragma edits.
• Transfers: a component uploads its heap data at the end of INITIALIZE and downloads at the start of SAVE, keeping device addresses in *_d members. mcgpu_* helpers in mccode-r.c check every transfer, list it on the master rank, and stop on failure.
• Ported components, those NMX_simplified uses: Arm, DiskChopper, DivLambda_monitor, Divergence_monitor, ESS_butterfly, Elliptic_guide_gravity, Guide_m, Monitor_nD, PSD_monitor, Progress_bar, Single_crystal, Slit, TOFLambda_monitor.
• Also ported, for wider test coverage: Guide_gravity, Guide, Guide_channeled, Source_simple, Source_div, Source_gen, Al_window, Beamstop, Collimator_linear, Monochromator_flat, Incoherent, L_monitor, E_monitor, TOF_monitor, Monitor, PSD_monitor_4PI, Brilliance_monitor. reflect in Guide and Guide_gravity and the table files in Source_gen are refused. E_monitor does not print its <E> line under OMPTARGET. Monitor now keeps its sums in an array. OFF geometry stops the build with #error under OMPTARGET.
• Everything else is refused at compile time: a component without #define MCGPU_PORTED_<name> gives #error under OMPTARGET. FUNNEL is refused too, and so are Monitor_nD auto limits and n1D mode, Guide_m reflectivity tables and Elliptic_guide_gravity segments, at run time.

mcrun --openmp builds it, the OpenMP twin of --openacc; the two are mutually exclusive. New config entries OMP and OMPFLAGS default to cc and -O2 -fopenmp -DOMPTARGET, the LUMI build. The Cray wrapper gets the target from craype-accel-amd-gfx90a; other compilers need their own values in the user config. Older user configs keep working. mctest --openmp and the mcgui run menu follow; mctest reports any instrument with an unported component as a compile failure, which is expected for now.

OMPTARGET testing

All on macOS arm64 with clang 21. Not yet run on a real GPU.

• No change for existing builds: old and new generators on 8 instruments, including Monitor_nD_noacc and FUNNEL. CPU and -DOPENACC preprocessed output is identical apart from the new *_d members, and #pragma acc lines are unchanged for the three instruments compared.
• -DOMPTARGET NMX on a stand-in device: host memory acts as the device, and host arrays are overwritten after upload so that reading a host pointer shows up. At -n 1e8 -s 1234, all 67 events, six monitor totals and three monitor arrays match the reference, with real threads.
• mctest --openmp vs CPU on a stand-in device, 57 instruments, 75 test values: 49 ran and all are identical to the CPU. The other 26 are refusals: Monitor_nD n1D mode (20), unported Monochromator_curved (4), OFF geometry (2).
• gfx90a: device compile succeeds with clang, and the image references no libc, only the OpenMP device runtime and math functions.
• Refusals: OPENACC + OMPTARGET, FUNNEL, and unported components all stop the build as intended.

Still to do: a LUMI run with CCE, including whether CCE defines __AMDGCN__ in its device pass; the matching generator changes in mccode-antlr.

Single_crystal: what changes

McStas Single_crystal now picks the reflection during the search and stores no reflection list, on CPU and GPU alike.

• hkl_search_select replaces hkl_search + hkl_select. A weighted reservoir draw keeps one reflection with P = refl_j / Σrefl.
• The T[MCSX_REFL_SLIST_SIZE] array is gone from each GPU thread's stack, about 172 kB per thread at size 1024.
• Reachable reflections are no longer capped. On NMX with a list of 1024, the cap was exceeded in 40% of searches, losing about 23% of coh_refl in those searches.
• Reflections sort on a 1e-9 grid in tau, then h, k, l. Rounding differences such as FMA used to reorder symmetry equivalents, so CPU and GPU picked different, mirrored reflections from the same random draw.

Removed

• By default, the CPU SPLIT cache and its statistics, including the "use SPLIT" hint at the end of a run. Each SPLIT iteration now searches and draws its own reflection.
• The USE_OPENCL branch. It could not link: hkl_search_opencl is not defined anywhere and nothing initialises oclContext_SX. opencl-lib and hkl_search.cl are left for separate cleanup.

Optional list on CPU: -DMCSX_REFL_SLIST

The list and the SPLIT reuse of the search are still worth having on CPU, so this define brings them back, with the cache, its statistics and warnings. Size comes from MCSX_REFL_SLIST_SIZE (default 128), and deltak stays as a parameter, unused otherwise. Both modes share one search body. GPU builds ignore the define.

The switch is a new define on purpose. MCSX_REFL_SLIST_SIZE alone has no effect, because instruments such as NMX_simplified already set it, and using it as the switch would make their CPU and GPU builds select differently.

Numerical change

Fixed-seed results change for every McStas instrument using Single_crystal, because each reachable reflection now draws a random number and the sort order changed. Statistics match the old version, except that searches which overflowed the old list now keep the reflections it dropped. This needs a line in CHANGES_McStas at release time.

Single_crystal testing

All on macOS arm64 with clang, on the CPU.

• NMX_simplified, -n 1e8 -s 1234, generated from the .instr: all 67 detector events and all monitor totals checked match the reference run row for row. That reference is the hand-edited EPICURE version with the same algorithm, whose CPU and OpenACC runs agreed with each other.
• Old vs new component, 11 parameter sets, 54 monitors: Test_SX (including SPLITS=10, order=1), Unittest_SPLIT_sample (SPLITS=1, 100), Test_Monochromators (Mono=6, PG=1, powder=1), templateNMX, templateLaue (2 files). Largest |z| is 1.30.
• -DMCSX_REFL_SLIST against the old component with only the sort changed: bit-identical output at fixed seed, 15 files across Test_SX (order=1 SPLITS=10, SPLITS=5) and templateNMX.
• Run time: Test_SX with order=1 SPLITS=10 takes 7.5 s by default and 4.1 s with -DMCSX_REFL_SLIST, as with the old component. Test_SX with SPLITS=5 and templateNMX show no difference.

Not tested:

• This exact component on GPU. The OpenACC evidence comes from the hand-edited version, and the code here was tidied slightly since (T is a single struct, hkl_search_select returns void).
• Linux, Windows/MSVC, or gcc.
• The full mctest suite.

Scope

McXtrace is deliberately left out. Its Single_crystal and FluoCrystal keep the list and the old sort, no McXtrace component is changed, and McXtrace was not built or tested, so the shared generator and runtime changes could affect it.

🤖 Generated with Claude Code