[ROCm] HIP port for AMD GPUs - #186
Conversation
Ports PopSift (GPU SIFT) to ROCm/HIP. A single compatibility header src/popsift/cuda_to_hip.h aliases the CUDA spellings to HIP under USE_HIP and is a no-op include of the CUDA headers otherwise; the .cu sources build as HIP, and the CUDA build path is byte-for-byte unchanged (every change is wrapped in USE_HIP / if(NOT USE_HIP)). README.md documents the USE_HIP build option, the build recipe, and the supported AMD targets. Getting correct, deterministic SIFT output required addressing five distinct HIP/ROCm fault classes, each behind a USE_HIP guard. 1. Hardware linear-filter textures. HIP rejects a cudaFilterModeLinear + element-read float texture at creation, so the linear fetches do manual bilinear interpolation in software (point-sample the 4 neighbors and lerp, CUDA's -0.5 texel-center convention). 2. Layered cudaArray dimension collapse. A multi-layer cudaArrayLayered array written one layer per kernel launch collapses to a single, last-written layer on any later-launch read on this ROCm (tex2DLayered, surf2DLayeredread, and even host hipMemcpy3D all return the last layer for every index; see ROCm/clr#275, which PopSift motivated, and the partial fix ROCm/rocm-systems#6683 covering only surf2DLayered). The Gaussian pyramid and DoG arrays are converted to non-layered 3D arrays (surf3Dwrite/surf3Dread/tex3D with the level as a real z coordinate, done once in the compat header's write wrapper plus the texture-fetch helper), which is fully coherent across launches. 3. Wavefront-width-generic warp collectives. The bitonic sort, exclusive prefix sums, descriptor reductions and orientation ballots operate per warpSize-lane group rather than assuming a 32-lane warp. The extrema counter in s_extrema.cu is the critical case: the in-wavefront lane is the block's linear thread index modulo the actual warpSize, the ballot/popc/shuffle widths follow warpSize, and the leader is lane 0 of the real wavefront. On wave64 this reduces exactly to a 64-bit ballot with an exclusive-prefix slot assignment; on wave32 it is correct where a fixed 64-lane geometry would have lost atomicAdds and read phantom lanes. 4. RootSift and L2 NaN. The default RootSift normalization sqrt(bin/L1) and the L2 path are guarded against sqrt of a slightly-negative bin and the 0*inf of an all-zero descriptor. 5. Example linkage. The PopSift library links roc::rocthrust PRIVATE and exposes hip::host PUBLIC, so the plain-C++ examples (main.cpp/match.cpp/pgmread.cpp, which have no device code) compile as CXX and just link the HIP runtime, rather than being forced through the HIP compiler by a leaked hip::device interface. The compat header's forward-compat <hip/hip_bf16.h> include is gated on __clang__ so a gcc host consumer never parses that clang-only header; it is pulled in before the __shfl_*_sync compat macros so newer ROCm's real __shfl_*_sync<...> functions are defined before the macros (guarded with __has_include so older ROCm, where the header has no such functions, is a harmless no-op). Tested on three AMD targets. PopSift ships no CPU reference, so the gates are run-to-run determinism, value sanity (finite, correctly normalized descriptors; in-bounds keypoints), and cross-architecture agreement between gfx90a and gfx1100 on the same images. - gfx90a (AMD Instinct MI250X, Linux, ROCm 7.2.1): popsift-demo gives 895 features / 1494 descriptors, identical across 5/5 runs; 0 NaN/Inf in all 191232 descriptor values; per-descriptor L2 in [0.9993, 1.0009], mean 1.0000; a rotated image gives 896/1484, also deterministic. popsift-match produces real finite distances and a sane ratio test. - gfx1100 (AMD Radeon Pro W7800, Linux, RDNA3 wave32, ROCm 7.2.1): the Oxford feature-detection images reproduce the gfx90a reference feature counts across all six images, deterministic across runs, 0 NaN/Inf across all descriptor modes. - gfx1151 (AMD Radeon 8060S APU, Windows, TheRock ROCm): SIFT on a real image gives non-zero, deterministic features (187/221 across runs), 0 NaN/Inf, RootSift L2 ~1.0, keypoints in bounds. Authored with the assistance of Claude (Anthropic).
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Code Review
This pull request adds support for AMD GPUs via ROCm/HIP to the PopSift library, enabling compilation and execution on AMD hardware while preserving the existing CUDA path. It introduces a CUDA-to-HIP compatibility shim, addresses warp-size differences (32-lane vs. 64-lane) in reduction and shuffle operations, and works around ROCm-specific limitations regarding layered image coherency and hardware bilinear filtering. The review feedback highlights two key areas for improvement: addressing a potential NaN corruption in s_desc_norm_rs.h when sum is extremely small, and initializing write_index in s_extrema.cu to prevent compiler or static analysis warnings regarding uninitialized variables.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
Addresses automated review on the AMD support PR. s_desc_norm_rs.h: gate the RootSift divisor at a small threshold so a degenerate near-zero descriptor sum is treated as all-zero. The 0*inf=NaN case was already mapped to 0 by the existing fmaxf(.,0), but a tiny subnormal sum makes 1/sum overflow to +inf and a positive bin would then normalize to +inf; the threshold removes that. No effect on any descriptor with real gradient content. s_extrema.cu: initialize write_index. __shfl reads only lane 0, so the prior uninitialized value on other lanes was never used, but the initializer silences -Wmaybe-uninitialized. Both changes are HIP-guarded; the CUDA path is unchanged. Authored with the assistance of Claude (Anthropic).
|
On On |
Authored with the assistance of Claude (Anthropic).
The documented Linux HIP configure command set -DCMAKE_HIP_COMPILER by absolute path but omitted any ROCm prefix hint. On a clean ROCm container where /opt/rocm is not on PATH, CMake then fails to locate the hip and rocThrust packages (hip_DIR / rocThrust reported NOTFOUND), so the documented command does not configure out of the box. Passing -DCMAKE_HIP_COMPILER by absolute path and setting the ROCM_PATH environment variable do not resolve this; only adding the ROCm install prefix to CMAKE_PREFIX_PATH (or putting /opt/rocm/bin on PATH) does. Add -DCMAKE_PREFIX_PATH=/opt/rocm to the Linux example and a short note on pointing it at the ROCm install. Documentation only; no build or source changes. The Windows block already supplies its own prefix and is untouched. Authored with the assistance of an AI coding agent.
The block that defaulted CMAKE_HIP_ARCHITECTURES to gfx90a in src/CMakeLists.txt was unreachable dead code: it ran from add_subdirectory(src), which the top-level CMakeLists reaches only after project(... LANGUAGES CXX HIP). By that point CMake has already resolved the HIP architecture -- auto-detecting the host GPU when no arch is given, honoring an explicit -DCMAKE_HIP_ARCHITECTURES, and erroring on a host with no GPU -- so the guard was always false and the gfx90a default never took effect. The set_target_properties() call that consumes CMAKE_HIP_ARCHITECTURES is kept; it now references the value CMake resolved at project(). Removing the dead block drops the misleading appearance of a gfx90a default. No functional change: the removed code never executed. Authored with assistance from Claude (Anthropic).
|
Hi Jeff, I'm really thankful that you have made the HIP port! I'm working abroad this year and I have only a Mac with me. I'll try to get my fingers of an NVidia and an AMD machine from my hosts, but I don't know if I can get it. The develop branch doesn't have a CPU-only version as you write, but there is the "sycl" branch, which can support CPU only including Mac via SYCL. However, that version is stripped down to one or at most two algorithms in every step of the pipeline. But in any case, our reference is always vlfeat. For a very long time, I've been annoyed that the vlfeat command line has stopped worked. But a few months ago, I finally found the bug in vlfeat. The README for fixing the bug and the python/opencv scripts for running the tests are 3 commits in the sycl branch. Another note on SYCL: Intel's SYCL compiler doesn't do very well with the PopSift port, while U.Stuttgart's compiler doesn't support hardware interpolation via sycl::image. CUDA and HIP still win, although the SYCL port shows that PopSift's DoG and normalize code could be better. |
|
Thanks Carsten, and thanks for the vlfeat pointers. I ran your vlfeat comparison on an AMD GPU (gfx90a) using your own setup: vlfeat built from source with the vl/pgm.c fix from your README, popsift-demo with On the AMD side this is already running on gfx90a, gfx1100 (RDNA3), and gfx1201 (RDNA4) through the same SIFT pipeline. The AMD support is entirely behind USE_HIP and the CUDA path is unchanged; I also confirmed the NVIDIA build still compiles clean under CUDA 12.6 (the full library, sm_86), so existing CUDA users are unaffected. You don't need to source an AMD machine to take this. If it's useful I can also send the side-by-side keypoint image that openCVMatches.py produces. One unrelated note while I was in there: develop does not build under CUDA 13, because the newer Thrust no longer transitively includes <thrust/tuple.h> and s_filtergrid.cu then fails on thrust::get/make_tuple. It is pre-existing (the unmodified file fails the same way) and trivial to fix with an explicit include. |
griwodz
left a comment
griwodz
left a comment
There was a problem hiding this comment.
I still don't have access to an AMD machine, but at least I can accress my NVidia machine again.
These are my current comments.
|
|
||
| add_executable(popsift-demo main.cpp pgmread.cpp pgmread.h) | ||
|
|
||
| set_property(TARGET popsift-demo PROPERTY CXX_STANDARD 11) |
There was a problem hiding this comment.
I've proposed in PR #190 to remove support for CUDA 10 and earlier. C++17 can be the default for CUDA as well.
| * textures are created with point filtering and we reproduce CUDA's bilinear | ||
| * filter in software here. This is an empirical limitation on this device/ROCm; | ||
| * it may not hold on other arches (re-verify on RDNA). | ||
| * CUDA's unnormalized linear filter on tex2DLayered(c) samples at index c-0.5, |
There was a problem hiding this comment.
I've always found that the CUDA docs are missing a good explanation for this +0.5f offset. I believe that the weird CUDA addressing exists because normalized addressing requires that the left-/topmost corner or the left-/topmost pixel have index 0 (instead of the pixel center). If index 0 were the center of the pixel, normalized addressing would have to work on negative coordinates.
| */ | ||
| const float fx = floorf( x ); | ||
| const float fy = floorf( y ); | ||
| const float ax = x - fx; |
There was a problem hiding this comment.
Since x is always positive, x - floorf(x) is always the same as x-static_cast(x)
For some reason, CUDA doesn't have a floor intrinsic and the cast is faster. Maybe HIP is better?
| const float t10 = tex2DLayered<float>( tex, fx + 1.5f, fy + 0.5f, z ); | ||
| const float t01 = tex2DLayered<float>( tex, fx + 0.5f, fy + 1.5f, z ); | ||
| const float t11 = tex2DLayered<float>( tex, fx + 1.5f, fy + 1.5f, z ); | ||
| const float top = t00 + ax * ( t10 - t00 ); |
There was a problem hiding this comment.
Have you considered using fmaf() ?
It was a real time-saver in the s_desc* functions.
Or perhaps the equivalent of cuda::std::lerp() exists for HIP.
| * blurred value a producer wrote at (x,y,level) is exactly what the consumer reads. | ||
| * tex is carried for parity with the CUDA signature but is unused on HIP. | ||
| * On CUDA this struct is unused (callers pass the texture object straight to the | ||
| * tex2DLayered readTex above) and the read path is byte-for-byte unchanged. |
There was a problem hiding this comment.
That is very unfortunate. The main reason for going through the hassle with textures and surfaces was to benefit from the hardware interpolation. If you have to do that manually, wouldn't it be more efficient to drop textures altogether and use "normal" cudaMalloc?
Also getting rid of the half-pixel offset.
| _data_ext.height = _h; | ||
| _data_ext.depth = _levels; | ||
|
|
||
| // Observed on gfx90a (ROCm 7.2.1): layered images are incoherent across kernel |
There was a problem hiding this comment.
I'm not clear how this happens.
The CUDA code uses one layer of one texture backed by an array for reading, and it uses a different layer of the same surface backed by the same array for writing.
Layers are never touched by one another. Since you don't have hardware interpolation, you could use a simple 3D point texture.
Does the problem appear because the reading from a texture and writing to a surface backed by the same Array is undefined in ROCm even when the cache lines never touch? That is inconvenient.
But I was wondering: I want a cudaArray for CUDA because I can only interpolate in X/Y while remaining discrete in Z when I used a Layered2D instead of 3D. And I can only define Layered2D over a cudaArray. If linear interpolation is not working anyway, you could use plain 3D memory, and the caching would happen in L1. Although that would cost read performance unless you transpose the Intermediate layers.
There was a problem hiding this comment.
PS: It could be a really good idea to create an alternative Octave class that works on flat memory (3D or Pitched). That alternative would have transposed Intermediate planes and need no texture at all. It might make the HIP code easier to read, and it would overcome the CUDA limit for large images, which are today constrained by the maximum Texture dimensions.
| tex_desc.addressMode[2] = cudaAddressModeClamp; | ||
| tex_desc.readMode = cudaReadModeElementType; // read as float | ||
| tex_desc.filterMode = cudaFilterModeLinear; // no interpolation | ||
| tex_desc.filterMode = cudaFilterModeLinear; // hardware bilinear (CUDA) |
| cudaSurfaceObject_t tex; | ||
| // Holds a texture object (assigned from cudaCreateTextureObject and consumed | ||
| // by tex2DLayered). On CUDA texture and surface handles are both unsigned | ||
| // long long so the original cudaSurfaceObject_t typing compiled; HIP uses |
There was a problem hiding this comment.
My original code was buggy. The final commit shouldn't keep the comment because it would confuse future readers.
| target_include_directories(popsift PRIVATE | ||
| ${CMAKE_CURRENT_SOURCE_DIR}/popsift/hip_compat) | ||
| else() | ||
| set_target_properties(popsift PROPERTIES CUDA_SEPARABLE_COMPILATION ON) |
There was a problem hiding this comment.
The need for rdc / CUDA_SEPARABLE_COMPILATION is an artifact of a mistake I made a long time ago: the host code that calls a global function should always be in the same compile unit but it isn't.
There is a branch that removes a multitude of alternatives that don't work well, where this mistake is also fixed.
I'm afraid it has to stay in the CUDA case for now.
| # The CC list for Tegras and Jetson will require manual updates | ||
| set(CMAKE_CUDA_ARCHITECTURES "53;62;72;87" | ||
| CACHE | ||
| STRING "Which CUDA CCs to support: native, all, all-major or an explicit list delimited by semicolons") |
There was a problem hiding this comment.
My comment string for the Tegra case is wrong. CMake is unable to choose the right CCs for Tegra. If we don't force it, it will configure for discrete-GPU CCs also on ARM hosts. But Tegra-family machines cannot host discrete NVidia GPUs.
Every warp shuffle in PopSift operates on a group of 32 threads that comes from the launch configuration (a descriptor tile, an image row, one row of the normalization block, the bitonic network), not from the hardware warp size. Passing that width explicitly states the intent, is correct when a wavefront is wider than the group, and removes the per-platform branch that used to spell the same reduction twice. The lane zero broadcast that picks the strongest orientation is one of those groups: ori_par launches a (32,1) block, so the width names the block row there as well. The extremum counter is the exception, because it counts over a whole wavefront, and its shuffles keep the hardware width. On NVIDIA the width equals the warp size, so the generated code and the results are unchanged. The compatibility header no longer redefines the __shfl_*_sync builtins. HIP provides the mask-free __shfl/__ballot/__any/__all spelling with a width parameter, which is the same spelling assist.h selects for CUDA before 9.0, so the AMD build now simply takes that branch through PopSift_HAVE_SHFL_DOWN_SYNC. That also drops the <hip/hip_bf16.h> include ordering workaround, which only existed because those macros used to rewrite the header's own declarations. This work was done with the assistance of an AI coding agent. Test Plan: ``` cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release \ -DPopSift_BUILD_EXAMPLES=ON -DBUILD_SHARED_LIBS=ON -DPopSift_USE_TEST_CMD=ON \ -DPopSift_TESTFILE_PATH=<path to the Oxford datasets> cmake --build build-hip -j cmake --build build-hip --target run-test-boat ``` On a Radeon Pro W7800 (gfx1100) the six Oxford boat images give 8351/9874, 7946/9452, 6158/7280, 4802/5799, 4618/5476 and 3855/4618 features/descriptors, identical to the counts before this change, and img1 features.txt is md5-identical.
A descriptor whose norm is zero cannot come from a real extremum, so both normalizers now reject that case in one place instead of scaling by an infinite factor and repairing the result afterwards. L2 keeps the inverse norm at zero when the sum of squares is zero, and RootSift writes an all-zero descriptor when the bin sum is not positive. The per-bin test in RootSift replaces the previous clamp. It also covers a bin that came out marginally negative, which can happen where the weight accumulation has no round-toward-positive-infinity intrinsic and falls back to round-to-nearest, and keeps such a bin out of the square root. The shuffle width in both files is now passed explicitly on every platform, for the same reason as the surrounding kernels: the reduction and the lane-zero broadcast span one row of the (32,32) normalization block, which is one descriptor, and not the hardware warp. This work was done with the assistance of an AI coding agent. Test Plan: ``` cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release \ -DPopSift_BUILD_EXAMPLES=ON -DBUILD_SHARED_LIBS=ON -DPopSift_USE_TEST_CMD=ON \ -DPopSift_TESTFILE_PATH=<path to the Oxford datasets> cmake --build build-hip -j cmake --build build-hip --target run-test-boat ``` The run-test-boat target uses --root-sift, so it exercises the RootSift path. On a Radeon Pro W7800 (gfx1100) the six Oxford boat images give the same counts as before this change and img1 features.txt is md5-identical, with no NaN or infinity in any descriptor.
The read handle a kernel takes for a pyramid array is now declared in src/popsift/sift_textures.h, together with LinearTexture, rather than inside common/assist.h, because the type appears in kernel signatures all over the library and is not an implementation detail of the assist helpers. The header is the place for the texture and surface handle types. Octave hands out the read handle itself. It knows the array, the texture, the surface and the level dimensions, so a call site no longer assembles a handle out of four accessors: it asks for the handle it wants, point filtered or linear filtered, over the blur data, over the intermediate data or over the DoG. The accessors that only returned the raw texture object are gone, since a caller that wants to read has a better one to use. The helper that fetches a single clamped texel is now called texFetchClamped: it emulates a point filtered texture fetch, including the clamping and the byte offset of the surface x coordinate, so its name should say texture rather than surface. It clamps with std::clamp. The manual bilinear interpolation uses fmaf, as it does elsewhere in the descriptor code, and the comments about hardware linear filtering now record that gfx1100 accepts what gfx90a rejects, which is why the software path stays in the shared build. This work was done with the assistance of an AI coding agent. Test Plan: ``` cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release \ -DPopSift_BUILD_EXAMPLES=ON -DBUILD_SHARED_LIBS=ON -DPopSift_USE_TEST_CMD=ON \ -DPopSift_TESTFILE_PATH=<path to the Oxford datasets> cmake --build build-hip -j cmake --build build-hip --target run-test-boat ``` On a Radeon Pro W7800 (gfx1100) the six Oxford boat images give the same counts as before this change and img1 features.txt is md5-identical.
The compatibility header described ROCm/rocm-systems#6683 as a partial fix covering only the layered surface read. A controlled experiment on a standalone reproducer showed the opposite: the collapse comes from the write, which passed the layer index in the mipmap level slot, so every layer landed in the same slot, and once the write is corrected all three read paths (layered surface read, layered texture fetch and host copy) return the right per-layer data. That result was posted on ROCm/clr#275. Comment only. The pyramid arrays stay non-layered, because the fix is not in ROCm 7.2.x. This work was done with the assistance of an AI coding agent. Test Plan: ``` cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release \ -DPopSift_BUILD_EXAMPLES=ON -DBUILD_SHARED_LIBS=ON -DPopSift_USE_TEST_CMD=ON \ -DPopSift_TESTFILE_PATH=<path to the Oxford datasets> cmake --build build-hip -j cmake --build build-hip --target run-test-boat ``` On a Radeon Pro W7800 (gfx1100) the six Oxford boat images give the same counts as before this change.
Grid filtering is the only Thrust user today, but the setup it needs is not specific to it: the stream-bound execution policy lives in a different namespace on each platform. Putting it in common/thrust_setup.h lets a second user pick it up without repeating the distinction. This work was done with the assistance of an AI coding agent. Test Plan: ``` cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release \ -DPopSift_BUILD_EXAMPLES=ON -DBUILD_SHARED_LIBS=ON -DPopSift_USE_TEST_CMD=ON \ -DPopSift_TESTFILE_PATH=<path to the Oxford datasets> cmake --build build-hip -j cmake --build build-hip --target run-test-boat ``` Grid filtering is compiled in by default, so the Thrust translation unit is built. On a Radeon Pro W7800 (gfx1100) the six Oxford boat images give the same counts as before this change.
Both creations of a linear filtered texture reported a point texture in their error message, copied from the point filtered creation above them. The comment next to the point filtering fallback asked for the behaviour to be re-verified on RDNA. It has been: gfx1100 creates the hardware linear filtered texture that gfx90a rejects, so the fallback is per device and the software interpolation in readTex is what keeps one build correct on either of them. This work was done with the assistance of an AI coding agent. Test Plan: ``` cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release \ -DPopSift_BUILD_EXAMPLES=ON -DBUILD_SHARED_LIBS=ON -DPopSift_USE_TEST_CMD=ON \ -DPopSift_TESTFILE_PATH=<path to the Oxford datasets> cmake --build build-hip -j cmake --build build-hip --target run-test-boat ``` Message and comment text only. On a Radeon Pro W7800 (gfx1100) the six Oxford boat images give the same counts as before this change.
The commit that simplified the normalization guards reduced the RootSift test to a comparison of the bin sum against zero. That reopens an overflow this branch had already closed: a subnormal sum makes the reciprocal overflow to infinity, so a positive bin normalizes to infinity instead of to zero. Measured on an Instinct MI300X, with the test against zero a bin of 1e-41 with a sum of 1e-40 gives infinity, and with the threshold it gives zero. A sum that small is a degenerate window rather than a real descriptor, so treating it like an all-zero one is what the surrounding code already does for a sum of exactly zero. The threshold applies on both platforms now; before the simplification it was on the AMD path only, while the NVIDIA path divided by the sum per bin. This work was done with the assistance of an AI coding agent. Test Plan: ``` cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_HIP_COMPILER=<rocm>/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release \ -DPopSift_BUILD_EXAMPLES=ON -DBUILD_SHARED_LIBS=ON -DPopSift_USE_TEST_CMD=ON \ -DPopSift_TESTFILE_PATH=<path to the Oxford datasets> cmake --build build-hip -j cmake --build build-hip --target run-test-boat ``` The run-test-boat target passes --root-sift, so it exercises this path. The threshold only changes the outcome for a bin sum below 1e-20, which no real image produces, so the Oxford boat feature and descriptor counts are unchanged. The NVIDIA path was compile-checked with nvcc 12.8 for sm_86.
The comment above the rocThrust lookup said that the Thrust code in the grid filter compiles unchanged on AMD. That was never accurate: the stream-bound execution policy is in a different namespace there, which is the reason common/thrust_setup.h exists. The comment now points at that header instead of claiming there is nothing to do. Comment only, no change to what is built. This work was done with the assistance of an AI coding agent. Test Plan: ``` cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_HIP_COMPILER=<rocm>/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release \ -DPopSift_BUILD_EXAMPLES=ON -DBUILD_SHARED_LIBS=ON cmake --build build-hip -j ``` Grid filtering is on by default, so the Thrust translation unit is compiled.
|
Thanks for the careful review. All comments are addressed on the branch; here is a summary. Every shuffle now passes its width explicitly, on both platforms, so the The shuffle compatibility now goes through The texture and surface types now live in one header, Octave now returns filled read handles ( The RootSift normalization has the structure you wrote: the early The L2 normalization guard is one line on both paths: Also done: the copy-pasted "point texture" error string on the linear-texture creation, the stale comments (including recording that gfx1100 accepts hardware linear filtering, so the software bilinear is a per-device workaround kept so one binary serves both), a shared Thrust setup header ( One suggestion I did not take: On the other points: agreed that rdc stays; the CUDA 10 / C++17 CMake modernization makes sense in #190 rather than here; and replacing the texture path with flat 3D memory and a leaner Octave class reads like the right long-term direction, which this change deliberately does not start. Since your review the branch was validated again end to end on an MI250X (gfx90a) and a Radeon 8060S (gfx1151, Windows), with earlier passes of the same changes on a Radeon Pro W7800 (gfx1100) and two more Windows machines (gfx1201 and gfx1101); the CUDA arm compile-checks clean with nvcc 12.8. On gfx90a one of the six benchmark images yields 9451 descriptors where the other architectures give 9452, stable across runs -- a single boundary keypoint on the 64-wide wavefront, noted here for completeness. |
3 participants
Ports PopSift (GPU SIFT) to ROCm/HIP. A single compatibility header src/popsift/cuda_to_hip.h aliases the CUDA spellings to HIP under USE_HIP and is a no-op include of the CUDA headers otherwise; the .cu sources build as HIP, and the CUDA build path is byte-for-byte unchanged (every change is wrapped in USE_HIP / if(NOT USE_HIP)). README.md documents the USE_HIP build option, the build recipe, and the supported AMD targets.
Getting correct, deterministic SIFT output required addressing five distinct HIP/ROCm fault classes, each behind a USE_HIP guard.
Hardware linear-filter textures. HIP rejects a cudaFilterModeLinear + element-read float texture at creation, so the linear fetches do manual bilinear interpolation in software (point-sample the 4 neighbors and lerp, CUDA's -0.5 texel-center convention).
Layered cudaArray dimension collapse. A multi-layer cudaArrayLayered array written one layer per kernel launch collapses to a single, last-written layer on any later-launch read on this ROCm (tex2DLayered, surf2DLayeredread, and even host hipMemcpy3D all return the last layer for every index; see Layered cudaArray collapses to one layer: surf2DLayeredread/tex2DLayered return the last-written layer for every layer index (gfx90a) ROCm/clr#275, which PopSift motivated, and the partial fix [HIP] Fix incorrect ockl binding for layered surface functions JIRA ID : SWDEV-597321 ROCm/rocm-systems#6683 covering only surf2DLayered). The Gaussian pyramid and DoG arrays are converted to non-layered 3D arrays (surf3Dwrite/surf3Dread/tex3D with the level as a real z coordinate, done once in the compat header's write wrapper plus the texture-fetch helper), which is fully coherent across launches.
Wavefront-width-generic warp collectives. The bitonic sort, exclusive prefix sums, descriptor reductions and orientation ballots operate per warpSize-lane group rather than assuming a 32-lane warp. The extrema counter in s_extrema.cu is the critical case: the in-wavefront lane is the block's linear thread index modulo the actual warpSize, the ballot/popc/shuffle widths follow warpSize, and the leader is lane 0 of the real wavefront. On wave64 this reduces exactly to a 64-bit ballot with an exclusive-prefix slot assignment; on wave32 it is correct where a fixed 64-lane geometry would have lost atomicAdds and read phantom lanes.
RootSift and L2 NaN. The default RootSift normalization sqrt(bin/L1) and the L2 path are guarded against sqrt of a slightly-negative bin and the 0*inf of an all-zero descriptor.
Example linkage. The PopSift library links roc::rocthrust PRIVATE and exposes hip::host PUBLIC, so the plain-C++ examples (main.cpp/match.cpp/pgmread.cpp, which have no device code) compile as CXX and just link the HIP runtime, rather than being forced through the HIP compiler by a leaked hip::device interface. The compat header's forward-compat <hip/hip_bf16.h> include is gated on clang so a gcc host consumer never parses that clang-only header; it is pulled in before the _shfl_sync compat macros so newer ROCm's real _shfl_sync<...> functions are defined before the macros (guarded with __has_include so older ROCm, where the header has no such functions, is a harmless no-op).
Tested on three AMD targets. PopSift ships no CPU reference, so the gates are run-to-run determinism, value sanity (finite, correctly normalized descriptors; in-bounds keypoints), and cross-architecture agreement between gfx90a and gfx1100 on the same images.
Authored with the assistance of Claude (Anthropic).