Repository navigation
Test the device pointer mode - #569
Merged
Merged
Conversation
#564 bound the scalars that rocBLAS/hipBLAS/rocSPARSE/hipSPARSE read per pointer mode as type(c_ptr),value, so they can live in host OR device memory. Nothing tested the device half: all 152 set_pointer_mode calls in the suite selected host mode, so the suite proved the new binding compiled and nothing about whether it works. Add a _devptr variant of every f2003 test that passes a dual-mode scalar: alpha and beta are staged in device memory with c_sizeof(), the handle is switched to the library's device enum, and the call sites take the device pointer. The host declarations and every verification expression are left alone, so a library that read the scalar from host memory produces a wrong answer rather than passing. 166 tests across the four libraries that have a pointer mode. f2003 only. The f2008/f2018 array wrappers take a host scalar and call c_loc() themselves, so a device scalar cannot reach them -- there is no overload mixing a c_ptr scalar with a Fortran array. Passing is not evidence on its own: a test that quietly kept its host scalar would also pass. Two things guard against that. test/check_devptr_tests.sh, registered as the hipfort_devptr_audit test, asserts per file that the device enum is selected, that no host spelling survives, that each d_alpha/d_beta is allocated, filled from its host twin and freed, that no call site still passes c_loc(alpha) outside the staging memcpy, and that the host scalar is still read by the verification. And the negative control: flipping the enum back to host mode makes rocblas/sscal, hipblas/sgemm, rocsparse/rocsparse_sspmv and hipsparse/hipsparse_sspmv all segfault, because the device address is then dereferenced on the host. Not covered, for want of a host test to mirror: the rot/rotg family and rocblas_set_solution_fitness_query. rocSOLVER and hipSOLVER are out of scope -- neither has a pointer-mode setter, and their .ptrmode entries are always-device out-pointers rather than dual-mode scalars. Verified on ROCm 10.2 amdflang and gfortran 13, gfx1101: the suite builds clean under both and all 166 pass.
The eight rocsparse csrgeam/csrgemm tests handed rocsparse_csrgeam_nnz and rocsparse_csrgemm_nnz a DEVICE pointer for nnz_C while never calling set_pointer_mode, so the handle stayed in the default HOST mode. rocSPARSE reads and writes nnzTotalDevHostPtr according to that mode, so it dereferenced a device address on the host. All sixteen (f2003 + f2008) died without output; they were the bulk of the suite's standing failures and had been taken for environment noise. Switching the handle to device mode would not fix it: alpha and beta in these same files are host pointers, so that trades one mismatch for another. Pass nnz_C on the host instead, which is what the hipSPARSE csrgemm tests already do -- and why they do not crash. The device staging around it is dropped, and the f2008 nnz_C gains the TARGET attribute c_loc() requires. This surfaced while writing the device-pointer-mode tests in the previous commit: their _devptr variants put alpha, beta and nnz_C all on the device and pass, which is what isolated the mismatch as the cause. With this the suite goes from 19 failures to 3 on ROCm 10.2 / gfx1101; the remaining three are hip mem_advise and module_kernel, unrelated and unchanged.
This was referenced Oct 1, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
No description provided.