GV

GeezerVison

The World Wide Web, restored to 800 pixels of glory.

Home|About|Gaming cabinet|Archives|How it works|Starter sites

On this site

  • Sign in
  • ActionsCI/CD and automation
  • CodespacesCloud development environments
  • All Features
  • Documentation
  • Training
  • Certifications
  • Build with Copilot
  • Customize Copilot
  • Integrations
  • Trending repos
  • GitHub Sponsors
  • Developer articles
  • Enterprises

This view

  • Original page
  • New address

GeezerVison is showing github.com — [ROCm] Merge upstream master (3.10.1), fix HIP reduce-by-key keys by jeffdaily · Pull Request #3741 · arrayfire/arrayfire

Conversation

@jeffdaily

@jeffdaily jeffdaily commented Nov 5, 2026

Copy link
Copy Markdown

My agent came up with this analysis, but I've read and approved it.

This brings experimental/hip up to date with master at 64c88d9 (the 23 commits after v3.10.0, including the 3.10.1 release prep and the clang-format 15 pass) and adds one fix to the HIP reduce-by-key kernels. It is a merge, not a rebase, so the commits already on experimental/hip stay as they are, and it fast-forwards the branch.

The merge has no textual conflicts. Three files changed on both sides:

  • CMakeLists.txt: upstream adds the CUDA compiler/toolkit check, the af_fmt variable (fmt::fmt-header-only when that target exists, otherwise fmt::fmt), the Windows standalone runtime collection and the 3.10.1 version bump. The AF_BUILD_HIP option, its block and conditional_directory(AF_BUILD_HIP ...) are untouched. The HIP backend builds as target afcuda, so it is covered by built_backends and the new standalone install step, and it takes fmt through afcommon_interface, which now links ${af_fmt}.
  • README.md: upstream's new "Experimental backends" table (pointing at experimental/hip) and the port's feature-list line both stay.
  • src/backend/common/half.hpp: upstream's only change is clang-format collapsing the __half2uint_rn else branch in half2int. It lands inside the port's __DH__ version of that function; the hipRTC traits block and the device conversion paths are unchanged.

The other upstream changes to the CUDA backend (NVRTC arch query, SASS instead of PTX, cuDNN 9, CUDA 13 thrust and cuFFT guards, compute-capability tables) are NVIDIA-specific and have no HIP counterpart: the HIP backend compiles kernels with hipRTC, and the ROCm Thrust here still provides thrust::unary_function and thrust/system/hip/detail/par.h. The merge itself changes nothing under src/backend/hip.

The commit after the merge, 83be197, changes only src/backend/hip/kernel/reduce_by_key.hpp:

  • Initialize the keys in reduce_blocks_by_key, reduce_blocks_dim_by_key and test_needs_reduction. Threads past the end of the input left k uninitialized and then passed it to the shuffle wrappers, which is undefined behaviour. The ROCm 7.16 compiler acts on it: SimplifyCFG turns the tidx < n branch around the key load into an assumption, so out-of-range threads load past the end of the input and count as unique keys, and every input whose last block is partial gets wrong reduce-by-key sizes (479 failures in reduce_cuda on gfx1100). ROCm 7.14 kept the branch. With the keys value-initialized the branch is kept and out-of-range threads keep the reduction's identity value.
  • Use the wavefront size, not 32, for the first index of the next warp in the warp-boundary check. On 64-wide wavefronts the old index let a block whose valid elements end in its first wavefront compare a key against shared memory that block never wrote. 32-wide wavefronts compile to the same code as before.

The CUDA backend's copy of these kernels is left as upstream has it.

Test data: sparse.tamu.edu now refuses downloads (HTTP 403). The build used upstream's own replacement, .github/scripts/check-matrix-data.py plus cmake -C .github/scripts/matrixmarket.cmake, which restores the same nine Matrix Market files from the pinned arrayfire-data commit.

Test plan

gfx1100 (Radeon Pro W7800), ROCm 7.16:

python3 -I .github/scripts/check-matrix-data.py . build-hip-gfx1100
AF_CI_MTX_FILES=ON cmake -C .github/scripts/matrixmarket.cmake -S . -B build-hip-gfx1100 -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=gfx1100
cmake --build build-hip-gfx1100
build-hip-gfx1100/test/reduce_cuda
ctest --test-dir build-hip-gfx1100 -R '_cuda$' -j1 --output-on-failure

Result at 83be197: reduce_cuda 1062/1062 in two runs (583/1062 at the merge alone), and 131/132 test binaries pass. The one failure is confidence_connected, because this build has no Image IO (AF_WITH_IMAGEIO=OFF): its 21 TEMP_FORMAT cases call af::loadImage and fail with "ArrayFire compiled without Image IO", 12 cases skip and 3 pass, the same as every build with that setting. The HIP runtime compiler needs ROCM_PATH set when ROCm is not installed in /opt/rocm.

Tested at 83be197 on:

  • MI250X (gfx90a, wave64), Linux: 131/132, reduce_cuda 1062/1062 (583/1062 at the merge alone). Same single confidence_connected failure.
  • MI300X (gfx942, wave64), Linux: 131/132, reduce_cuda 1062/1062. Same single failure.
  • Radeon Pro W7800 (gfx1100, wave32), Linux: 131/132 as above.
  • Radeon 8060S (gfx1151), Windows: 128 of 130 run pass, reduce_cuda 1062/1062. sparse_arith was not run (the known Windows rocsparse csrgeam2 crash, as at 950dcdd). Besides confidence_connected, test_threading_cuda hung twice in Threading.FFT_ALL after FFT_R2C and FFT_C2C passed; it passed in 14 s at 950dcdd on this machine, FFT_ALL, BLAS and Sparse pass when run alone, and we have seen the same timeout on two other Windows GPUs before this change. It looks like an intermittent multi-threaded hipFFT hang that this change does not touch; it is not root-caused.
  • CUDA: nvcc 12.8 builds the CUDA backend from this branch and from upstream 64c88d9 with identical results: 707/707 steps, the same 438 upstream warnings, and identical kernel code (no NVIDIA GPU, so compile only).

Sorry, something went wrong.

Uh oh!

There was an error while loading. Please reload this page.

All reactions
melonakos and others added 25 commits September 10, 2026 14:11
@melonakos
README: point to the experimental Metal and HIP backend branches
9ff5a08
@melonakos
Merge pull request arrayfire#3716 from arrayfire/docs/experimental-ba…
ab4a506
…ckends

README: link the experimental Metal and HIP backend branches
@roaffix @syurkevi
Replace the old CI workflows with ci.sh and its scripts
2d799c6
The Unix and Windows workflows downloaded their own CMake, Ninja and
Doxygen, ran partly on a retired runner image (Ubuntu 20.04), repeated
their build steps in YAML and reported to a CDash dashboard instead of
the pull request. They are removed here; their
jobs come back in the following commits as calls to one script. Doing
both in one commit keeps every workflow in the tree passing actionlint.

.github/scripts/ci.sh runs one CI job: it installs the job's
dependencies, then runs the job's CMake workflow preset or check, in
the job's pinned Ubuntu container (--docker) or on the macOS and
Windows hosts, so CI and a developer run the same command; --list
shows the jobs. With --docker it chooses the build parallelism from the
container's memory, keeps ccache in a volume or in a host directory CI
can cache, and builds either this checkout or another commit's tree
(AF_CI_SOURCE_DIR) with this repository's scripts and settings, also as
a copy without .git. It must run on macOS's bash 3.2.

install-deps-ubuntu.sh, install-deps-macos.sh and
install-deps-windows.ps1 install each job's packages: the compilers and
BLAS, the CUDA, OpenCL and oneAPI toolkits by version (with
nvJitLink's development package, which the installers' library bundling
looks up by name), the runtime libraries alone for GPU test containers,
and vcpkg at a pinned release (2026.10.16), whose MSYS2 helpers still
find their packages. lint.sh checks the formatting of the changed lines
with clang-format 15. checkout-commit.sh fetches one commit
anonymously. count-vcpkg-archives.sh tells whether a run added vcpkg
archives worth saving. .gitattributes keeps the shell scripts' LF
endings, which Git Bash on Windows runners needs.

The release source workflow builds the tarball with
make-source-tarball.sh and uploads it with gh release upload instead of
the archived upload-release-asset action; only the upload step can
write, and a manual run builds the tarball as an artifact without
publishing it. It no longer fetches the SuiteSparse matrices, which the
tarball never included.
@roaffix @syurkevi
Fix the presets' LAPACK value, the clFFT pin, Doxygen and .gitignore
871a3ce
- The non-MKL presets set AF_COMPUTE_LIBRARY to "FFTW/LAPCK/BLAS",
  which matches neither supported value, so configure skipped the
  FFTW, CBLAS and LAPACK checks.
- clFFT followed a moving branch, the only way the OpenCL build could
  change without a commit here. It is pinned to the commit its
  arrayfire-release branch points to, like every other OpenCL
  dependency.
- NUM_PROC_THREADS = 0 runs Doxygen's experimental parallel input
  processing on every core. With Doxygen 1.9.7 it intermittently
  corrupted the heap and aborted: 4 of 10 runs on a 24-core machine.
  On one thread every run passed, in the same time.
- The docs build writes docs/details/examples.dox into the source tree
  on every configure, which left the tree dirty; it is ignored.
@roaffix @syurkevi
Add CMake presets for the CI configurations
7e0dc03
The ci-* configure, build, test and workflow presets hold each CI
job's settings, so ci.sh runs "cmake --workflow --preset ci-<job>" and
a developer can run the same:

- CPU on Linux with GCC and Clang and OpenBLAS, with icx and oneMKL,
  and on the newest Ubuntu; macOS with Apple Clang and Accelerate,
  built with spdlog's bundled fmt (Homebrew's fmt 12 breaks it) and
  without the mean/var cases that FMA contraction breaks there; Windows
  with MSVC 2022 and vcpkg, at /W3 (/Wall flooded the log with
  thousands of system-header warnings) and with MSBuild's multi-tool
  task compiling in parallel;
- compile-only CUDA, OpenCL and oneAPI builds. CUDA builds for compute
  capabilities 7.5, 8.9 and 9.0, plus Blackwell's 10.0 and 12.0 through
  nvcc flags, because CUDA_architecture_build_targets goes through
  CMake's own FindCUDA copy of cuda_select_nvcc_arch_flags on this
  tree, which rejects two-digit majors; and a Windows CUDA build;
- the documentation.

The test presets write JUnit, cap the BLAS threads at four and fail
when no test runs. The Matrix Market tests follow AF_CI_MTX_FILES,
with FetchContent pointed at the data that ci.sh downloads from
arrayfire-data and restores into the build tree (added in a
following commit).

The Visual Studio 2019 presets become 2022 ones, since no GitHub runner
has 2019 any more, and the file moves to format 6 (CMake 3.25) for the
workflow presets.
@roaffix @syurkevi
Add the light checks: lint, docs and workflows
23ffbdb
checks.yml runs on every push to a pull request or master and in the
merge queue. It is cheap, so the CPU builds can run on a schedule of
their own:

- lint: clang-format 15 on the lines changed since the base. The tree
  has unformatted files, and library code isn't reformatted here.
- docs: the Doxygen build.
- workflows: actionlint (which also runs shellcheck on the workflows'
  run: steps), shellcheck on the CI scripts, the CI scripts' Python
  test suites, and check-bash32.sh.

check-bash32.sh runs the scripts that the macOS job runs on the host
under bash 3.2, the macOS runner's bash, in a bash:3.2 image pinned by
digest and without network: their syntax, ci.sh --list, a short SHA
through checkout-commit.sh, and a grep for bash-4-only syntax. Nothing
else caught bash 4 syntax until the macOS job failed.
@roaffix @syurkevi
Take the LAPACK overlay and triplet from the fix/restore-ci branch
4ce7858
vcpkg's current blas and lapack ports collide with the lapack-reference
overlay port: the overlay's lapack.pc, its blas.pc links for the
noblas feature and its wrapper in share/lapack belong to them now, and
the Windows build stopped at the file conflict. ArrayFire's own CI
pull request, the fix/restore-ci branch, fixes the overlay, and its
Windows run passes with it. Its portfile and triplet are taken as they
are, so the two branches don't conflict:

- the overlay installs lapack-reference.pc and its wrapper in
  share/lapack-reference, and no blas.pc links;
- vcpkg_cmake_config_fixup uses PACKAGE_NAME ${PORT};
- the triplet sets VCPKG_PROVIDED_FORTRAN, since vcpkg 2026.07.29 and
  later install MinGW gfortran for the port only when the triplet asks
  for it, as vcpkg's own x64-windows triplet does.
@roaffix @syurkevi
Get the Matrix Market test data from a pinned arrayfire-data commit
c6f333c
The sparse tests' matrices come from the SuiteSparse collection, whose
storage refuses every download ("403 AccessDenied"). ArrayFire's
arrayfire-data repository holds the same nine files, unchanged, under
matrixmarket/, with their CC BY 4.0 attribution. CI downloads them from
arrayfire-data at commit 2723164, which matrixmarket.json pins, so the
Matrix Market tests run again; an empty commit would turn them off.

The manifest of each matrix's size and SHA-256 and the idea of checking
the data before a build come from ArrayFire's fix/restore-ci branch
(arrayfire#3735), whose check_matrix_data.py is adapted as check-matrix-data.py.
Before every build, ci.sh runs it. It downloads each matrix from the
pinned commit on raw.githubusercontent.com, checks it against the
manifest and the tested tree's list of matrices, fails naming every
file that differs or doesn't download, restores the matrices into the
build tree and prints ON; without a pinned commit it prints OFF.

ci.sh passes the answer on as AF_CI_MTX_FILES. The tests download each
matrix through af_dep_check_and_populate, whose FetchContent_Populate
takes FETCHCONTENT_SOURCE_DIR_<name> as the source when it is set; the
ci preset and matrixmarket.cmake (for the GPU builds and the CMake 3.16
build) set those to the restored matrices, and AF_TEST_WITH_MTX_FILES
to AF_CI_MTX_FILES. Nothing is downloaded from SuiteSparse, and nothing
under test/ changes. The GPU build archives carry the matrices to the
test jobs.
@roaffix @syurkevi
Add the CUDA release checks and the Windows CUDA scripts
bb1bed7
check-cuda-release.sh configures the CUDA backend as users and the
installers get it: ArrayFire's default architecture list and
AF_INSTALL_STANDALONE=ON. It prints the architecture flags and
compiles one CUDA object with them, which catches a default list with
an architecture the toolkit dropped (with CUDA 13, this tree's list
still has compute_53). It installs the cuda_dependencies component and
checks that every library the installers bundle on Linux is there,
under the toolkit's soname, the name libafcuda loads it by (with CUDA
13, this tree bundles cuFFT, cuSOLVER and cuSPARSE under the wrong
major). It needs only a configure and one compile.

install-cuda-windows.ps1 installs a CUDA toolkit on Windows with
NVIDIA's network installer, pinned by URL and SHA-256 per version
(12.9, 13.1, 13.4), with only the subpackages the build needs. From
CUDA 13.0 on, nvcc no longer includes crt (crt/host_config.h) and nvvm
(cicc): NVIDIA's redistrib manifests list them apart and the
installers carry crt_<version> and nvvm_<version>, though the
installation guides list them only from 13.2 on. msvc-env-windows.ps1
sets up Visual Studio's x64 developer environment for the Ninja
generator.
@roaffix @syurkevi
Add a GPU test workflow with one approval per request
aef5361
gpu.yml tests one commit on GPU runners. The scripts and settings come
from the ref it runs on, and only the source tree from the commit.
Each leg, a CUDA version from AF_CI_CUDA_VERSIONS or OpenCL, is built
on a free standard runner into an archive of the test programs,
libraries and test data, then tested on a GPU runner (AF_CI_GPU_RUNNER)
in a container with the runtime libraries alone, which installs in
about a third of the time of the toolkit. One approve job holds the
gpu-tests environment, so one reviewer's approval covers every leg, and
it starts with the builds. A failed build fails only its own leg; a
summary job reports every leg's passed and failed tests.

The CUDA test job compares the driver with the toolkit (cuda-compat.sh)
and, when the driver is older, installs NVIDIA's forward-compatibility
package for data-center GPUs such as the T4: ArrayFire's JIT has the
driver load PTX from NVRTC, which minor-version compatibility doesn't
cover. The OpenCL test job registers NVIDIA's ICD when the container
toolkit doesn't, prints the platforms and fails clearly without an
NVIDIA one. linux-cuda-gpu.cmake and linux-opencl-gpu.cmake hold the
builds' trusted settings for any commit.

The workflow is off unless AF_CI_GPU_TESTS is true.
@roaffix @syurkevi
Run the CPU tests in the merge queue, on master and nightly
b6b985f
cpu.yml runs the CPU builds and tests on Linux (GCC and Clang, in
pinned containers), macOS and Windows, in the merge queue and on every
push to master, and for a given commit when started by hand or by
another workflow; such a commit is built on its own with the trusted
scripts and presets. Actions are pinned to commit SHAs. Only master's
own builds save the ccache, the vcpkg cache and the downloaded Matrix
Market data, because a build of another commit could plant entries in
them; the vcpkg cache is saved
whenever a run added archives, even if the build failed later. The
macOS and Windows jobs build outside a container, so a request runs
them only when its caller asks for them.

With AF_CI_ACCELERATOR_BUILDS set to true, pushes and requests also
compile the CUDA (once per AF_CI_CUDA_VERSIONS version, with Blackwell
targets and the release checks), OpenCL and oneAPI backends, and the
CUDA backend on Windows. The matrix job has a fixed name: GitHub never
expands a skipped matrix job's name, which showed up as a check named
"matrix.job".

cpu-placeholders.yml gives a pull request instant checks with the same
names, so that it can join a merge queue that requires them. The
nightly runs the icx and oneMKL job, the job on the minimum CMake
version and the job on the newest Ubuntu, which may fail without
failing the run.

A merge-queue entry or a push to master that only adds or changes
non-code files (code-changes.py: Markdown at the top level, docs/, the
issue and pull-request templates, GitHub's citation files, .gitignore,
.mailmap) builds nothing. The changes job lists the changed files with
the compare API and the read-only token; a merge-queue entry is
compared with the branch it merges into, since the group's parent may
be the entry ahead. None of those files is read by a CPU build, an
installer rule or the release tarball. A removed or renamed file runs
the tests in full, since CPack checks at every Linux configure that
README.md exists; so does any other file and every request. The CPU
jobs still run and skip their steps, so the required checks keep their
names (a skipped matrix job's name isn't expanded) and pass in seconds,
each saying why.
@roaffix @syurkevi
Run CPU and GPU tests on request with /cpu-tests and /gpu-tests
54e96d2
A /cpu-tests or /gpu-tests comment on a pull request asks for the
tests of its current head. test-requests.yml runs from the default
branch, so the pull request can't change it, and nothing from the pull
request is checked out or run in it. test-requests.py decides, pins
the head SHA, and cpu.yml or gpu.yml test exactly that commit; the
result goes to the cpu-tests or gpu-tests commit status and a comment.

- /cpu-tests (all adds the accelerator builds): maintainers with write
  access on any pull request; the pull request's author on their own,
  once the light checks have run on it, on Linux only.
- /gpu-tests [cuda|opencl|all] [cuda=X.Y]: maintainers with write
  access only. Without a backend the changed files choose, and a
  change that can't affect a GPU backend gets a reply and starts
  nothing; CUDA versions come from AF_CI_CUDA_VERSIONS (12.9 when
  unset). Every reply says what runs and why.
- An accepted reply says when the branch is behind its base, since CI
  fixes made there since aren't in the test.

Failed API calls become annotations of the run, readable without
logging in; a failed comment after the status is set is only a
warning. The tests run the decisions on fixtures, the commands with
failing writes, and the CUDA driver check.
@roaffix @syurkevi
Describe CI in CONTRIBUTING.md
4053b1b
The "Running CI" section says what runs when, which checks the light
checks are, how to request the CPU and GPU tests and what they cost,
how to run any CI job locally with ci.sh, and where the Matrix Market
data comes from.
@roaffix @syurkevi
Let Dependabot keep the pinned actions current
19513a5
The workflows pin every action to a commit SHA; Dependabot proposes the
updates weekly.
@syurkevi
allow /cpu-tests and /gpu-tests to specify SHA
1e74154
@syurkevi
build windows CI vcpkg dependencies for release only
9887d73
@syurkevi
keep CMakePresets.json at format 2
33719e5
@syurkevi
build updates for cuda 13
ab7aa82
@syurkevi
fix fp64 aspect errors on devices without double support
e3455bc
@syurkevi
fix doxygen warnings and windows docs build
dc7dac0
@syurkevi
cuda: load JIT kernels on drivers older than the toolkit
6c1939f
@syurkevi
prepare 3.10.1 release with cuda 12.9 and 13.4 support
df6d9f4
@syurkevi
apply clang-format 15 to src, test and examples
64c88d9
@jeffdaily
Merge branch 'master' of https://github.com/arrayfire/arrayfire
5d09b5c
@jeffdaily
[ROCm] Fix undefined key reads in the reduce-by-key kernels
83be197
Threads past the end of the input leave their key uninitialized in reduce_blocks_by_key, reduce_blocks_dim_by_key and test_needs_reduction, then pass it to the shuffle wrappers and compare it. The wrappers take the value as a by-value parameter, which clang marks noundef, so passing an indeterminate value is undefined behaviour and the compiler may assume it never happens.

The HIP compiler in ROCm 7.16 acts on it: after inlining, SimplifyCFG replaces the `tidx < n` branch around the key load with an assumption that it is always taken, so the key and value are loaded unconditionally and the `tidx < n` term disappears from the unique-key flag. Every block with fewer valid elements than threads then counts its out-of-range lanes as keys, and reduce-by-key returns wrong sizes or throws `seq.end < parentDim`. ROCm 7.14 compiles the same source with the branch intact, so the bug did not show before: an older compiler is free not to exploit the undefined behaviour, and it did not.

Value-initialize the keys. Out-of-range lanes then keep the reduction's identity value, and the existing bounds checks keep them out of the unique-key flags and the segmented reduction, so the key value given to them does not change any result.

The same kernels compute the first index of the next warp in the warp-boundary check as `(warpid + 1) * 32`. On a 64-wide wavefront that index lies inside the current warp, so when a block's valid elements end in its first wavefront the bounds check passes and the key is compared with the next warp's left-boundary key in shared memory, which that block never wrote. Use kWarpSize, as the rest of the kernel does; 32-wide wavefronts compile to the same code.

The CUDA backend's copy of these kernels has the same uninitialized keys, but nvcc does not exploit them, and it is left as upstream has it.

Authored with Claude.

Test Plan:

```
cmake --build build-hip-gfx1100
build-hip-gfx1100/test/reduce_cuda
ctest --test-dir build-hip-gfx1100 -R '_cuda$' -j1 --output-on-failure
```

On a Radeon Pro W7800 (gfx1100) with ROCm 7.16, reduce_cuda goes from 583/1062 to 1062/1062, and 131 of 132 test binaries pass. The one failure is confidence_connected, which needs Image IO and this build leaves it out: 21 of its 36 cases fail, 12 skip and 3 pass.
syurkevi
syurkevi approved these changes Nov 11, 2026
View reviewed changes

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@syurkevi syurkevi syurkevi approved these changes

Uh oh!

There was an error while loading. Please reload this page.

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

Uh oh!

There was an error while loading. Please reload this page.

4 participants

@jeffdaily @syurkevi @melonakos @roaffix
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
You are visitor 01402503 — GeezerVison does not store the pages you look up.
Best viewed with Netscape Navigator 4.7 or Internet Explorer 6 at 800×600. No cookies required.