Update to cuDSS 0.8 - #12
Merged
Merged
Conversation
📝 WalkthroughWalkthroughSwitches cuDSS from optional local-source builds to downloaded prebuilt archives (CUDSS_VERSION), removes LOCAL_CUDSS_BUILD conditionals, adds CUDSS_NEW_API gating from header version macros, and refactors helper and linear-solver flows to support new and legacy cuDSS APIs. ChangescuDSS Prebuilt Archive Migration and API Versioning
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~25 minutes Poem
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches📝 Generate docstrings
🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 2
🧹 Nitpick comments (1)
cmake/AddCUDSS.cmake (1)
59-65: 💤 Low valueConsider pinning a
URL_HASHfor the fetched archive.
FetchContent_Declarefor a remote binary without an integrity hash leaves the build vulnerable to corrupted or tampered downloads and is non-reproducible. Since the version is fixed, aURL_HASH SHA256=...per archive would harden the supply chain.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cmake/AddCUDSS.cmake` around lines 59 - 65, Add an integrity hash to the FetchContent declaration to pin the downloaded archive: compute the SHA256 of the file at CUDSS_URL (or obtain the vendor-provided checksum) and add a URL_HASH SHA256=<checksum> entry to the FetchContent_Declare call in AddCUDSS.cmake (the block that declares "cudss" and uses CUDSS_URL and FetchContent_MakeAvailable). This ensures reproducible, tamper-evident fetches; update the stored checksum when intentionally changing versions.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cmake/AddCUDSS.cmake`:
- Around line 51-65: The FetchContent_Declare block for the cudss dependency
currently uses URL ${CUDSS_URL} without integrity verification; compute the
SHA256 of the expected .tar.xz archive for each supported release and add an
URL_HASH entry (e.g., URL_HASH SHA256=<computed-hash>) to the
FetchContent_Declare(cudss ...) call so FetchContent will verify the downloaded
artifact; update the FetchContent_Declare invocation in AddCUDSS.cmake alongside
the existing URL and keep the CUDSS_URL logic unchanged.
In `@cunls/common/cudss_helper.cpp`:
- Around line 245-254: The CUDSS_NEW_API branch of cudssMatrixCreateCsr is
calling the wrong signature and using non-public CUDSS_R_* enums; update the
CUDSS_NEW_API call in cudss_helper.cpp to match the public cuDSS CSR signature
used elsewhere by passing the same parameter order and types as the non-NEW
branch (use a single index type and a single value type), replacing
CUDSS_R_32I/CUDSS_R_32F with the public CUDA data-type enums (CUDA_R_32I,
CUDA_R_32F) and ensure the NULL/cols/values parameters match the non-NEW_API
invocation so the call mirrors the working cudssMatrixCreateCsr(&mat,
matrix_size, matrix_size, num_nonzeros, rows_ptr, NULL, cols_ptr, values_ptr,
CUDA_R_32I, CUDA_R_32F, CUDSS_MTYPE_SYMMETRIC, CUDSS_MVIEW_FULL,
CUDSS_BASE_ZERO).
---
Nitpick comments:
In `@cmake/AddCUDSS.cmake`:
- Around line 59-65: Add an integrity hash to the FetchContent declaration to
pin the downloaded archive: compute the SHA256 of the file at CUDSS_URL (or
obtain the vendor-provided checksum) and add a URL_HASH SHA256=<checksum> entry
to the FetchContent_Declare call in AddCUDSS.cmake (the block that declares
"cudss" and uses CUDSS_URL and FetchContent_MakeAvailable). This ensures
reproducible, tamper-evident fetches; update the stored checksum when
intentionally changing versions.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Enterprise
Run ID: 5e7fe0a8-0e83-4c84-a43e-264ef769a7f5
📒 Files selected for processing (6)
CMakeLists.txtcmake/AddCUDSS.cmakecunls/common/CMakeLists.txtcunls/common/cudss_helper.cppcunls/linear_solver/CMakeLists.txtcunls/linear_solver/cudss_sparse_linear_solver.cpp
💤 Files with no reviewable changes (2)
- cunls/common/CMakeLists.txt
- cunls/linear_solver/CMakeLists.txt
ZeroScalarKernel and Zero2ScalarKernel were never launched (the solver zeros scalar slots with cudaMemsetAsync), producing nvcc #177-D "declared but never referenced" warnings. Remove them and update the two comments that referenced them. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cunls/linear_solver/block_sparse_pcg_solver.cu (1)
650-664:⚠️ Potential issue | 🔴 Critical | ⚡ Quick winFix grid-wide race in
PcgDirectionKernelupdatingrz_old_ptr[0]
PcgDirectionKernelhas CTAs readingrz_old_ptr[0](fordenom) whileblockIdx.x==0writesrz_old_ptr[0]=numwithin the same kernel launch;__syncthreads()doesn’t order across CTAs, sobetacan vary nondeterministically.🐛 Proposed fix
__global__ void PcgDirectionKernel(const float *__restrict__ rz_new_ptr, - float *__restrict__ rz_old_ptr, const float *__restrict__ z, + const float *__restrict__ rz_old_ptr, + const float *__restrict__ z, float *__restrict__ p, int n) { __shared__ float b; if (threadIdx.x == 0) { float num = rz_new_ptr[0]; float denom = rz_old_ptr[0]; b = (fabsf(denom) > 0.f) ? num / denom : 0.f; - // Only the first CTA's thread 0 writes rz_old; all other CTAs - // skip the write (the rz_old slot is read by the *next* - // iteration's PcgUpdate, which happens after this kernel - // completes, so any one writer is enough). - if (blockIdx.x == 0) { - rz_old_ptr[0] = num; - } } __syncthreads(); int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= n) { return; } p[i] = z[i] + b * p[i]; }PcgDirectionKernel<<<blocks, threads, 0, stream>>>( d_scratch_.data() + kRzNew, d_scratch_.data() + kRzOld, z_.data(), p_.data(), n); + THROW_ON_CUDA_ERROR(cudaMemcpyAsync(d_scratch_.data() + kRzOld, + d_scratch_.data() + kRzNew, + sizeof(float), + cudaMemcpyDeviceToDevice, + stream));🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/linear_solver/block_sparse_pcg_solver.cu` around lines 650 - 664, PcgDirectionKernel races because multiple CTAs read rz_old_ptr[0] (denom) while blockIdx.x==0 writes rz_old_ptr[0]=num in the same kernel; remove the in-kernel write to rz_old_ptr[0] inside PcgDirectionKernel to eliminate the cross-CTA race and instead perform the update to rz_old_ptr[0] from the host (or in a separate single-block kernel launched after PcgDirectionKernel completes) so denom is read-only inside the kernel; keep the rest of the logic (reading rz_new_ptr[0] into num, computing b, and updating p/z) unchanged and reference PcgDirectionKernel, rz_new_ptr, rz_old_ptr, num, denom, and blockIdx.x when applying the change.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Outside diff comments:
In `@cunls/linear_solver/block_sparse_pcg_solver.cu`:
- Around line 650-664: PcgDirectionKernel races because multiple CTAs read
rz_old_ptr[0] (denom) while blockIdx.x==0 writes rz_old_ptr[0]=num in the same
kernel; remove the in-kernel write to rz_old_ptr[0] inside PcgDirectionKernel to
eliminate the cross-CTA race and instead perform the update to rz_old_ptr[0]
from the host (or in a separate single-block kernel launched after
PcgDirectionKernel completes) so denom is read-only inside the kernel; keep the
rest of the logic (reading rz_new_ptr[0] into num, computing b, and updating
p/z) unchanged and reference PcgDirectionKernel, rz_new_ptr, rz_old_ptr, num,
denom, and blockIdx.x when applying the change.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Enterprise
Run ID: 6c632087-a817-4ab0-b637-9833f915a17a
📒 Files selected for processing (1)
cunls/linear_solver/block_sparse_pcg_solver.cu
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
1 participant
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.
Summary by CodeRabbit
Chores
Refactor