[tgrogers]

Builds and runs the whole suite on a GB300 (Blackwell, sm_103, aarch64) under
CUDA 13.2. make goes from failing immediately on dev to completing with
472 binaries and zero errors.

Why this is a re-port

This work was originally done on a branch cut from release. That made a direct
merge into dev produce 36 conflicts, so the diff was re-applied onto dev
instead. Everything below is the dev version of that work, plus the extra
issues that only dev's own code exposed.

Architecture selection

setup_environment previously carried a hand-maintained per-CUDA-version arch
table that stopped at compute_100. It now intersects a desired-arch wish-list
with whatever nvcc --list-gpu-arch reports, with a static fallback for
pre-CUDA-11 toolkits. On CUDA 13.2 this resolves to 75/80/86/90/100/103.

Gencode uses the bracket form code=[sm_X,compute_X] rather than
code="sm_X,compute_X": identical SASS and PTX, but it contains no spaces, so
passing it through MAKE_ARGS cannot split into a stray make goal. PTX stays
embedded so GPGPU-Sim can still run these binaries.

Multi-arch is deliberate. The suite is built once and simulated as many
different machines, so each target wants the PTX a real build for that arch
would emit. Narrowing to the local GPU is available via GPUAPPS_ARCHS="103"
to cut build time.

CUDA 13 API removals

• Legacy texture-reference API (texture<>, cudaBindTexture,
  cudaBindTextureToArray, cudaUnbindTexture) is gone. Migrated ~68 files to
  cudaTextureObject_t, declaring the handles __constant__ under their
  original names so every fetch site and dispatch macro is unchanged; the host
  fills them with cudaMemcpyToSymbol. Only the explicit template arguments
  (tex1Dfetch<float>, tex2D<uint4>) differ.
• cudaDeviceProp::clockRate / ::memoryClockRate / ::computeMode are gone;
  replaced with cudaDeviceGetAttribute. Five sites.
• cudaThreadSynchronize() is gone (lonestargpu pta).
• thrust::unary_function is gone, and using namespace thrust now collides
  with ::cuda/::cub in nvcc-generated stubs.

Blackwell / newer-toolkit build breaks

• arch_filter.mk (new) intersects any hardcoded ARCH list with
  nvcc --list-gpu-arch, so the five ubenches pinning sm_101 and friends stop
  being fatal. A MIN_ARCH mechanism in common.mk drops gencodes below a
  benchmark's floor (cp.async needs 80).
• Submodule bumps: cutlass-bench v4.1.0 to v4.8.1 (4.1.0 warns
  unsupported or deprecated compute capabilities 103a), cuda-samples v12.8
  to v13.2 (12.8 pins compute_50, fatal under CUDA 13).
• heterosync hardcoded compute_60; cleared so CUDA_CPPFLAGS decides.
• ispass-2009/WP linked -lGL although no source includes a GL header.
• MapSMtoCores tables stopped at SM 7.0 and SM 3.0; a GB300 fell off the end
  and reported 152x64 cores, or -1. Both now cover SM 3.2 to 12.1.

aarch64

• -m64 is x86-only; made conditional in pathfinder and pathfinder_k1,
  matching the pattern already used in huffman and mummergpu.
• char is unsigned by default, so getopt_long returning -1 compared as 255.
  Fixed in dwt2d along with a missing {0,0,0,0} longopts sentinel.

Correctness bugs found by running, not building

These are pre-existing and would silently corrupt results:

• rodinia-3.1 kmeans used the pre-CUDA-5 cudaMemcpyToSymbol("c_clusters",...)
  string-literal form, which passes the address of the string. Unquoted.
• pathfinder was missing two __syncthreads(), a real data race.
• parboil sgemm declared readColMajorMatrixFile as bool and fell off the
  end, giving SIGTRAP under a modern compiler.
• mummergpu has a if (this) null check inside a member function, which is
  UB that modern GCC deletes; built with -fno-delete-null-pointer-checks.

CUTLASS build

Examples 10 and 11 link cutlass_library, which drags in all 56,823 generated
kernels spanning sm_50 to sm_103; a serial sub-make then spent over six minutes
in a single ptxas on an FP4 kernel for an arch with no FP4 hardware. Filtering
to CUTLASS_LIBRARY_KERNELS=*planar_complex* keeps all 97 examples building
while cutting generated sources to 1,363. The profiler is no longer built, since
the install step only ever copied from build/examples. Arches are a
CUTLASS_ARCHS variable instead of a hardcoded 80;90;90a.

Makefile reconciliation

dev had drifted from the sub-Makefiles: 24 stale mv lines in the
rodinia-3.1 target relocated binaries that now link straight into BINDIR.
Removed.

Re-enabled

dev disabled several apps that release builds. They build and run clean
under CUDA 13, so they are back: ispass-2009 AES and MUM, rodinia-2.0-ft
kmeans, lonestar bfs-wlc and sssp-wlc (which needed uncommenting in two
places -- the VARIANTS list and the install mv), and parboil in all:.

Verification

Full make from clean: exit 0, zero errors, 472 binaries.

All nine suites run under the accel-sim run_hw.py harness at exit 0:
rodinia_2.0-ft, rodinia-3.1, ispass-2009, parboil, polybench, pannotia,
proxy-apps-doe, lonestargpu-2.0, GPU_Microbenchmark -- 94 apps, 130 run
configurations, with no app missing a binary.

Known pre-existing failures, unchanged by this PR: lonestar bh deadlocks and
pta goes out of bounds under Volta+ independent thread scheduling (both are
pre-Volta warp-synchronous code), mnistCUDNN needs APIs removed in cuDNN 9,
and shoc-Stability tries a 256GB allocation.

Hardware counter collection could not be exercised here: this machine has
RmProfilingAdminOnly: 1, so every ncu invocation fails ERR_NVGPUCTRPERM.
The applications themselves still ran to completion under the harness.

Deliberately out of scope

src/cuda/gpuwattch-ubench still calls cudaThreadSynchronize(), which CUDA 13
removed, so it will not compile on a CUDA 13 toolkit. This PR does not fix it.

An earlier revision did -- a mechanical cudaThreadSynchronize() ->
cudaDeviceSynchronize() rename across 181 files -- but no Makefile target
references that directory (the only wattch target points at
accelwattch-ubench), so nothing compiled the result and the change could not
be verified. It was 40% of this PR's file count for zero tested benefit, so it
has been dropped. The rename is almost certainly still correct and belongs in
its own commit, alongside a target that actually builds the suite.

Also reverted to dev's content: src/cuda/shoc-master/build-aux/config.guess
and config.sub. Those were regenerated as a byproduct of running autotools
during a shoc build attempt, not a deliberate change, and shoc is not re-enabled
here. They were ~16% of the line count.

Together these took the diff from 482 files / +15,277 -3,825 to 298 files /
+12,471 -1,457.

Companion changes not in this PR

• Stop missing binaries aborting hw runs; dedupe the runner scripts; aarch64 NVBit accel-sim-framework#575 — harness fixes: a missing binary no longer
  aborts a whole suite, the three runner scripts are deduplicated, NVBit is
  fetched for the host architecture, and the app list is corrected.
• Fix the build on aarch64 and CUDA 13 gpgpu-sim_distribution#150 — the three aarch64/CUDA-13 fixes needed
  to build accel-sim.out here: x86-only -msse2/-mfpmath=sse, an x86-only
  asm("int $03"), and cudaDeviceProp::clockRate.
• Ubuntu 24.04 CUDA 13.2.1 minimal Dockerfile#12 — a Ubuntu-24.04-cuda-13.2-minimal image recipe, so
  the CUDA 13 CI job here can eventually drop its apt step and start from a
  published tag.
• The data tarball behind get_data.sh is missing
  lonestargpu-2.0/lonestar-bfs-wla/data -> ../inputs, so that one app cannot
  find its graph. Fixing it needs access to the published archive.

🤖 Generated with Claude Code