Experimental/hip - #3721

The clang-resource-version scan in the HIP backend enumerated
rocmRoot/lib/llvm/lib/clang using POSIX dirent.h (opendir/readdir/
closedir) to discover the versioned clang include directory. Those
headers do not exist on Windows, so the HIP backend failed to compile
there.

Replace the directory scan with std::filesystem::directory_iterator
under #if defined(_WIN32). The original POSIX path is preserved
byte-identical under #else, so Linux builds are unaffected. The scan
semantics are unchanged: pick the first non-dot entry as the version
directory and fall back to the bare include path when none is found.

This closes a gap where the Windows fix had been validated locally but
was never committed to the branch, so the build could only succeed with
an uncommitted working-tree edit.

Authored with assistance from Claude.

Test Plan:
  cmake --build build-gfx1201 -j24
  HIP_VISIBLE_DEVICES=0 ctest --test-dir build-gfx1201 -R "cuda" -j1

The previous default pinned CMAKE_HIP_ARCHITECTURES to gfx90a before
enable_language(HIP), which preempted CMake's own host-GPU detection. A
user on a non-gfx90a AMD GPU who omitted -DCMAKE_HIP_ARCHITECTURES would
silently build gfx90a code objects that then fail to load on their card
at runtime ("no kernel image is available for execution on the device").

Remove the pin and rely on enable_language(HIP): it honors an explicit
-DCMAKE_HIP_ARCHITECTURES when given, otherwise auto-detects the host
GPU(s), and errors out when no GPU is found rather than guessing. Builds
that pass an explicit architecture are unchanged.

Authored with assistance from Claude.

Two Windows-only build fixes for the HIP backend, reported by @villekf
while building on Windows with a clang toolchain.

1. Force shared-library link mode on afcuda on Windows. The clang driver
   does not always infer the DLL subsystem for this backend (seen with the
   Visual Studio generator selecting clang), which surfaces as a
   missing-subsystem link error. Adding -shared to the link makes the DLL
   build regardless of how the generator selects the driver.

2. Define BOOST_USE_WINDOWS_H on Windows so Boost.Stacktrace's windbg
   backend includes <windows.h> directly instead of self-declaring the
   Win32 API through Boost.WinAPI. Without it the default
   AF_STACKTRACE_TYPE=Windbg fails to compile with conflicting-types errors
   for CreateFileA, ReadFile, WriteFile and other Win32 entry points (the
   self-declarations clash with the real Windows SDK headers that other
   translation units pull in), forcing users to fall back to
   AF_STACKTRACE_TYPE=None. With the define the default stacktrace builds.

Both changes are guarded by if(WIN32); Linux and macOS builds are byte
unchanged.

Test Plan: built the HIP backend for gfx1201 (Radeon RX 9070 XT) with
TheRock ROCm 7.14 clang, Ninja generator.

```
# Fix 1: afcuda relinks cleanly with -shared (default None stacktrace)
cmake --build build-gfx1201 --target afcuda    # links bin/afcuda.dll, exit 0

# Fix 2: prove the default Windbg stacktrace now compiles
cmake -S . -B build-gfx1201 -DAF_STACKTRACE_TYPE=Windbg
cmake --build build-gfx1201 --target afcuda    # 0 errors, links afcuda.dll
# (without BOOST_USE_WINDOWS_H this fails in cufft.cu.obj with the
#  conflicting-types errors described above)
```

Authored with the assistance of an AI coding assistant.

The HIP backend's hipBLAS GemmEx call uses hipblasGemmEx with hipDataType
and hipblasComputeType_t (blas.cu, cudaDataType.hpp). That signature was
introduced in hipBLAS 3, which ships with ROCm 7.0. ROCm 6.x ships
hipBLAS 2, whose hipblasGemmEx takes hipblasDatatype_t and the older
compute-type enum, so the backend does not compile there (reported on
ROCm 6.4.4; ROCm 7.0.3 builds). State the minimum next to the
AF_BUILD_HIP wiring where the backend is documented.

Documentation only; no source or device code changes.

Authored with the assistance of Claude (Anthropic).

Test Plan:
Documentation-only change (a CMake comment). No build or GPU run needed;
verified inert with the MOAT regression classifier (comment/doc-only,
carries prior validation forward on every platform).

Several comments in the HIP backend explained a hazard by naming other
porting work and an internal document instead of the problem itself.
Those references mean nothing to a reader of this repository, so each
comment now states the hazard directly.

The substance is unchanged. ROCm typedefs several logically distinct
library handles to the same void*, so a wrapper keyed on the handle type
cannot separate them, which is why hip_unique_handle.hpp keys its RAII on
a tag type instead. HIP device compilation defines __HIP_DEVICE_COMPILE__
rather than __CUDA_ARCH__. clang enforces matching __host__ __device__
attributes between a specialization and its primary template where nvcc
does not. AMD wavefronts are 64 lanes on CDNA and 32 on RDNA, so
warp-staged kernels must use kWarpSize rather than a literal 32.

Comment text only: no declaration, definition or preprocessor condition
changes.

Authored with the assistance of Claude, an AI assistant made by Anthropic.

Test Plan
=========

Rebuilt the HIP backend on an MI250X (gfx90a) with ROCm 7.2.1 to confirm
the reworded comment blocks still compile:

```
cmake -S . -B build-hip -GNinja -DCMAKE_BUILD_TYPE=Release \
  -DAF_BUILD_HIP=ON -DAF_BUILD_CUDA=OFF -DAF_BUILD_CPU=OFF \
  -DAF_BUILD_OPENCL=OFF -DAF_BUILD_ONEAPI=OFF -DAF_BUILD_UNIFIED=OFF \
  -DAF_BUILD_EXAMPLES=OFF -DAF_BUILD_FORGE=OFF -DAF_WITH_CUDNN=OFF \
  -DAF_WITH_IMAGEIO=OFF -DAF_BUILD_DOCS=OFF \
  -DCMAKE_HIP_ARCHITECTURES=gfx90a \
  -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++
cmake --build build-hip -j 32 --target afcuda
```

Built clean, 425/425 steps, no errors.

Confirmed the change reaches no code, only comments:

```
git diff -U0 | grep -E '^[+-]' | grep -vE '^(\+\+\+|---)' | grep -vE '^[+-]\s*//'
```

Printed nothing.

The HIP backend's FAST, ORB and connected-components kernels read their lookup tables through a texture object, which does not build on CDNA3 and newer devices. ROCm's HIP headers define __HIP_NO_IMAGE_SUPPORT whenever __gfx94plus_clr__ is set, and that marks every tex1Dfetch overload in texture_indirect_functions.h unavailable, so the call is a compile error rather than a runtime fallback: 'tex1Dfetch<unsigned char, nullptr>' is unavailable: The image/texture API not supported on the device. The same sources compile on gfx90a and on gfx1100, so this appears only when the offload target is gfx942 or gfx950.

All three sites are plain point lookups. fast.cu and orb.cu bind a fixed 16 to 64 entry corner-test table through LookupTable1D, and regions.cu binds the equivalency map it is about to relabel; none of them use filtering, normalized coordinates or address clamping, and the regions kernel already read the map directly for the double instantiation. So the texture object buys nothing on any device and is replaced everywhere by an indexed read through a const pointer: LookupTable1D keeps its shape and its Array copy for the reference count but now hands out the device pointer, the two lookup helpers and the regions fetch helper index that pointer, and regions no longer creates or destroys a texture object at all. The now-unused texture entries in hip_compat.h go with it. This is unconditional rather than guarded on the architecture, because the fix is correct everywhere and an architecture-specific variant would leave the CDNA3 path untested on the machines people build on.

The CUDA backend is untouched, so builds on NVIDIA GPUs keep the texture path exactly as before.

Authored with the assistance of Claude (Anthropic).

Test Plan:
- Build for gfx942 (MI300X, ROCm 7.14), which previously failed to compile:

```
cmake -S . -B build-hip -GNinja -DCMAKE_BUILD_TYPE=Release \
  -DAF_BUILD_HIP=ON -DAF_BUILD_CUDA=OFF -DAF_BUILD_CPU=ON \
  -DAF_BUILD_OPENCL=OFF -DAF_BUILD_ONEAPI=OFF -DAF_BUILD_UNIFIED=ON \
  -DAF_BUILD_EXAMPLES=OFF -DAF_BUILD_FORGE=OFF -DAF_WITH_CUDNN=OFF \
  -DAF_WITH_IMAGEIO=OFF -DAF_BUILD_DOCS=OFF -DAF_BUILD_TESTS=ON \
  -DCMAKE_HIP_ARCHITECTURES=gfx942
cmake --build build-hip -j 32
```

  The three previously failing translation units (fast.cu, orb.cu, regions.cu) now compile, and the full library and all test binaries build.

- Run the affected feature-detection and labelling tests on gfx942:

```
HIP_VISIBLE_DEVICES=0 ctest --test-dir build-hip --output-on-failure \
  -R 'test_(fast|orb|regions|harris|sift|homography)_cuda'
```

  100% tests passed, 0 tests failed out of 6.

The comment added by the previous commit, and that commit's own message, describe the two lookup tables wrongly. The comment calls them "a few dozen entries" and the message calls them a "fixed 16 to 64 entry corner-test table". Neither is right: FAST_LUT in kernel/fast_lut.hpp is 65536 unsigned char entries (64 KiB), as it must be, because kernel/fast.hpp builds the bright and dark responses as 16 bit masks and indexes the table with them directly, and d_ref_pat in kernel/orb_patch.hpp is REF_PAT_SAMPLES * REF_PAT_COORDS = 1024 int entries (4 KiB).

Table size was never what made reading them through a plain pointer sound, so the comment now states the property that does: they are only ever point sampled, with no filtering, no normalized coordinates and no address modes, so an indexed read returns exactly what the texture fetch returned.

The previous message also claimed the texture object "buys nothing on any device". That was an assertion, not a measurement, and it is withdrawn here. What is established is that the results are unchanged and that the texture fetch builtins are unavailable, so the code does not compile at all, when the offload target is gfx942 or gfx950. That is the reason the tables are read directly on every target rather than under an architecture guard.

Authored with the assistance of Claude (Anthropic).

Test Plan:
- This changes a comment only, so the behaviour of the previous commit is unaffected and its gfx942 feature-detection test run still applies. Rebuild the two translation units that include the header, on gfx942 (MI300X, ROCm 7.14), to confirm they still compile:

```
cmake -S . -B build-hip -GNinja -DCMAKE_BUILD_TYPE=Release \
  -DAF_BUILD_HIP=ON -DAF_BUILD_CUDA=OFF -DAF_BUILD_CPU=ON \
  -DAF_BUILD_OPENCL=OFF -DAF_BUILD_ONEAPI=OFF -DAF_BUILD_UNIFIED=ON \
  -DAF_BUILD_EXAMPLES=OFF -DAF_BUILD_FORGE=OFF -DAF_WITH_CUDNN=OFF \
  -DAF_WITH_IMAGEIO=OFF -DAF_BUILD_DOCS=OFF -DAF_BUILD_TESTS=ON \
  -DCMAKE_HIP_ARCHITECTURES=gfx942
ninja -C build-hip src/backend/hip/CMakeFiles/afcuda.dir/fast.cu.o \
  src/backend/hip/CMakeFiles/afcuda.dir/orb.cu.o
```

  Both objects build.