[schur-complement] validate the arguments and report allocation and launch failures - #316
Open
zjin-lcf wants to merge 5 commits into
Open
[schur-complement] validate the arguments and report allocation and launch failures#316zjin-lcf wants to merge 5 commits into
zjin-lcf wants to merge 5 commits into
Conversation
…SYCL and OpenMP kernel launches The CSR row pointers and column indices are 32-bit, so reject non-positive sizes and sizes whose nonzero count would overflow an int instead of silently generating a corrupt matrix. Also share the block size macros in the SYCL version and set the team size explicitly in the OpenMP version. Co-authored-by: Cursor <cursoragent@cursor.com>
The examples themselves are already on master, so keep only the argument validation and the kernel-launch alignment here, and take the updated Makefiles (LAUNCHER, ARCH and SM overridable from the environment) from master. Co-authored-by: Cursor <cursoragent@cursor.com>
…n and launch failures - Move the command-line checks into reference.h so the four versions cannot drift apart, and bound the number of rows as well: m + 1 row pointers and rounding m up to the block size have to stay within an int. - Report a failed host allocation instead of terminating on bad_alloc, and a failed device allocation in the SYCL version instead of dereferencing null in the kernel. - Check the CUDA and HIP kernel launches, so an invalid launch configuration is reported rather than showing up as a verification failure. - Fix the two SYCL kernel comments that still described the original HiOp one-work-item-per-row mapping. Co-authored-by: Cursor <cursoragent@cursor.com>
zjin-lcf
changed the title
There was a problem hiding this comment.
Warning
Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.
Pull request overview
Hardens the schur-complement benchmark variants by validating input sizes, improving allocation/launch failure reporting, and aligning parallel execution parameters across backends.
Changes:
- Added centralized argument validation (
valid_problem_size) and applied it in SYCL/OMP/CUDA/HIP mains. - Added host allocation failure handling (catching
std::bad_alloc) and device allocation failure reporting (SYCL). - Added CUDA/HIP kernel launch error checks and updated SYCL/OpenMP configuration/comments for the 2D (i,j) mapping.
Reviewed changes
Copilot reviewed 5 out of 5 changed files in this pull request and generated 4 comments.
Show a summary per file
| File | Description |
|---|---|
| src/schur-complement-sycl/main.cpp | Adds size validation, host/device allocation failure reporting, and hoists SYCL block geometry constants. |
| src/schur-complement-omp/main.cpp | Adds size validation, host allocation failure reporting, and adjusts OpenMP target threading parameters. |
| src/schur-complement-hip/main.cu | Adds size validation, host allocation failure reporting, and checks HIP kernel launches. |
| src/schur-complement-cuda/reference.h | Introduces valid_problem_size to prevent CSR-related 32-bit overflows. |
| src/schur-complement-cuda/main.cu | Adds size validation, host allocation failure reporting, and checks CUDA kernel launches. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+22
to
+28
| // row pointers and rounding m up to a multiple of the block size. | ||
| static bool valid_problem_size(int m, int nnz_row, int repeat) | ||
| { | ||
| return m > 0 && nnz_row > 0 && repeat > 0 && | ||
| m <= INT_MAX - 1024 && // m + 1, and the launch range | ||
| nnz_row <= (INT_MAX - 1024) / 8 && // nx = 8 * nnz_row + 1024 | ||
| (long long)m * nnz_row <= INT_MAX; // number of nonzeros |
|
|
||
| if (!valid_problem_size(m, nnz_row, repeat)) { | ||
| printf("Invalid arguments: <rows>, <nnz per row> and <repeat> must be " | ||
| "positive, and the number of nonzeros must fit in a 32-bit int\n"); |
Comment on lines
69
to
+78
| double *d_D = sycl::malloc_device<double>(nx, q); | ||
| double *d_W = sycl::malloc_device<double>(w_elems, q); | ||
|
|
||
| if (d_rs1 == nullptr || d_jc1 == nullptr || d_v1 == nullptr || | ||
| d_rs2 == nullptr || d_jc2 == nullptr || d_v2 == nullptr || | ||
| d_D == nullptr || d_W == nullptr) { | ||
| printf("Failed to allocate the device buffers: the dense block alone needs " | ||
| "%zu bytes\n", w_bytes); | ||
| return 1; | ||
| } |
|
|
||
| // host/device correctness check (run once, verify against reference) before timing | ||
| #pragma omp target teams distribute parallel for thread_limit(128) | ||
| #pragma omp target teams distribute parallel for |
Collaborator
Author
There was a problem hiding this comment.
num_threads(256) is optional here
Co-authored-by: Cursor <cursoragent@cursor.com>
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.
The examples themselves are now on master, so what is left here is the hardening of the four
schur-complement-*versions. Merged master to resolve the conflicts, keeping master's Makefiles (LAUNCHER,ARCH,SMandCUDA_ARCH/HIP_ARCHoverridable from the environment).Summary of the changes:
<rows>,<nnz per row>and<repeat>. The CSR row pointers and column indices are 32-bit, as in HiOp, so reject non-positive values, a nonzero count that would not fit an int, and row counts for whichm + 1or roundingmup to the block size would overflow. The check lives inreference.h(valid_problem_size) so the four versions cannot drift apart.bad_alloc, and a failed device allocation in the SYCL version instead of dereferencing a null pointer inside the kernel.BLOCK_X/BLOCK_Yas in the other versions.Verified with
./main 1024 32 20(all PASS): CUDA on a Tesla M40, HIP and OpenMP (AOMP) on an MI100, and SYCL on the M40 through the CUDA backend; the SYCL version also compiles withicpx. Invalid arguments (zero, negative, and out-of-range sizes) are rejected with exit status 1.