Skip to content

fix(rocSOLVER) Fix race condition in SYEVJ GPU kernel on gfx11 (LLVM Issue #207883) #12266

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
EdDAzevedo merged 12 commits into from
Oct 15, 2026
Merged

fix(rocSOLVER) Fix race condition in SYEVJ GPU kernel on gfx11 (LLVM Issue #207883) #12266

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
eefec0e
remove assert on info
EdDAzevedo Sep 17, 2026
06991a6
add calls to syncthread to fix potential race condition
EdDAzevedo Sep 17, 2026
a754c81
update copyright year
EdDAzevedo Sep 17, 2026
520be73
update issue addressed
EdDAzevedo Sep 17, 2026
67e796a
add extra tests for small kernel path in syevj
EdDAzevedo Sep 17, 2026
63682af
Merge branch 'develop' into users/EdDAzevedo/rocm_31219_syevj
EdDAzevedo Sep 30, 2026
52c8e48
Merge branch 'develop' into users/EdDAzevedo/rocm_31219_syevj
EdDAzevedo Oct 13, 2026
ef6f21f
Merge branch 'develop' into users/EdDAzevedo/rocm_31219_syevj
EdDAzevedo Oct 14, 2026
a53374c
restore lapack_host_reference
EdDAzevedo Oct 14, 2026
033e2cc
minor update to comments
EdDAzevedo Oct 14, 2026
d837081
Merge branch 'develop' into users/EdDAzevedo/rocm_31219_syevj
mgates3 Oct 15, 2026
a434b52
Merge branch 'develop' into users/EdDAzevedo/rocm_31219_syevj
mgates3 Oct 15, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions projects/rocsolver/CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,9 @@ Full documentation for rocSOLVER is available at the [rocSOLVER documentation](h
* Improved performance of expert eigensolvers SYEVDX/HEEVDX, SYGVDX/HEGVDX

### Resolved issues

* Fixed a race condition in SYEVJ GPU kernel.

### Known issues
### Upcoming changes

Expand Down
16 changes: 0 additions & 16 deletions projects/rocsolver/clients/common/misc/lapack_host_reference.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6264,7 +6264,6 @@ void cpu_sytrs<float>(rocblas_fill uplo,
rocblas_int info = 0;
char uploC = rocblas2char_fill(uplo);
ssytrs_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &info);
Comment thread
EdDAzevedo marked this conversation as resolved.
Resolved
assert(info == 0);
}

template <>
Expand All @@ -6280,7 +6279,6 @@ void cpu_sytrs<double>(rocblas_fill uplo,
rocblas_int info = 0;
char uploC = rocblas2char_fill(uplo);
dsytrs_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &info);
assert(info == 0);
}

template <>
Expand All @@ -6296,7 +6294,6 @@ void cpu_sytrs<rocblas_float_complex>(rocblas_fill uplo,
rocblas_int info = 0;
char uploC = rocblas2char_fill(uplo);
csytrs_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &info);
assert(info == 0);
}

template <>
Expand All @@ -6312,7 +6309,6 @@ void cpu_sytrs<rocblas_double_complex>(rocblas_fill uplo,
rocblas_int info = 0;
char uploC = rocblas2char_fill(uplo);
zsytrs_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &info);
assert(info == 0);
}

// sytrs2
Expand All @@ -6333,7 +6329,6 @@ void cpu_sytrs2<float>(rocblas_fill uplo,

char uploC = rocblas2char_fill(uplo);
ssytrs2_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &(work[0]), &info);
assert(info == 0);
}

template <>
Expand All @@ -6353,7 +6348,6 @@ void cpu_sytrs2<double>(rocblas_fill uplo,

char uploC = rocblas2char_fill(uplo);
dsytrs2_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &(work[0]), &info);
assert(info == 0);
}

template <>
Expand All @@ -6373,7 +6367,6 @@ void cpu_sytrs2<rocblas_float_complex>(rocblas_fill uplo,

char uploC = rocblas2char_fill(uplo);
csytrs2_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &(work[0]), &info);
assert(info == 0);
}

template <>
Expand All @@ -6393,7 +6386,6 @@ void cpu_sytrs2<rocblas_double_complex>(rocblas_fill uplo,

char uploC = rocblas2char_fill(uplo);
zsytrs2_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &(work[0]), &info);
assert(info == 0);
}
// -----------------------------------------------------------------------
// getrs_npvt
Expand Down Expand Up @@ -8657,7 +8649,6 @@ void cpu_sytf2<float>(rocblas_fill uplo,
{
char uploC = rocblas2char_fill(uplo);
ssytf2_(&uploC, &n, A, &lda, ipiv, info);
assert(*info == 0);
}

template <>
Expand All @@ -8670,7 +8661,6 @@ void cpu_sytf2<double>(rocblas_fill uplo,
{
char uploC = rocblas2char_fill(uplo);
dsytf2_(&uploC, &n, A, &lda, ipiv, info);
assert(*info == 0);
}

template <>
Expand All @@ -8683,7 +8673,6 @@ void cpu_sytf2<rocblas_float_complex>(rocblas_fill uplo,
{
char uploC = rocblas2char_fill(uplo);
csytf2_(&uploC, &n, A, &lda, ipiv, info);
assert(*info == 0);
}

template <>
Expand All @@ -8696,7 +8685,6 @@ void cpu_sytf2<rocblas_double_complex>(rocblas_fill uplo,
{
char uploC = rocblas2char_fill(uplo);
zsytf2_(&uploC, &n, A, &lda, ipiv, info);
assert(*info == 0);
}

// sytrf
Expand All @@ -8712,7 +8700,6 @@ void cpu_sytrf<float>(rocblas_fill uplo,
{
char uploC = rocblas2char_fill(uplo);
ssytrf_(&uploC, &n, A, &lda, ipiv, work, &lwork, info);
assert(*info == 0);
}

template <>
Expand All @@ -8727,7 +8714,6 @@ void cpu_sytrf<double>(rocblas_fill uplo,
{
char uploC = rocblas2char_fill(uplo);
dsytrf_(&uploC, &n, A, &lda, ipiv, work, &lwork, info);
assert(*info == 0);
}

template <>
Expand All @@ -8742,7 +8728,6 @@ void cpu_sytrf<rocblas_float_complex>(rocblas_fill uplo,
{
char uploC = rocblas2char_fill(uplo);
csytrf_(&uploC, &n, A, &lda, ipiv, work, &lwork, info);
assert(*info == 0);
}

template <>
Expand All @@ -8757,7 +8742,6 @@ void cpu_sytrf<rocblas_double_complex>(rocblas_fill uplo,
{
char uploC = rocblas2char_fill(uplo);
zsytrf_(&uploC, &n, A, &lda, ipiv, work, &lwork, info);
assert(*info == 0);
}

// bdsvdx
Expand Down
6 changes: 6 additions & 0 deletions projects/rocsolver/clients/gtest/lapack/syevj_heevj_gtest.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,6 +56,12 @@ const vector<vector<int>> size_range = {
{12, 12},
{23, 30},
{40, 45},
// small-kernel path (n <= SYEVJ_BLOCKED_SWITCH) with a 2D thread array exceeding
// ~512 threads: previously untested window that exposed an eigenvector/eigenvalue
// corruption on gfx11
{48, 50},
{56, 56},
{58, 60},
Comment thread
EdDAzevedo marked this conversation as resolved.
Resolved
{67, 70},
{70, 70},
};
Expand Down
15 changes: 14 additions & 1 deletion projects/rocsolver/library/src/lapack/roclapack_syevj_heevj.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@
* and
* Hari & Kovac (2019). On the Convergence of Complex Jacobi Methods.
* Linear and Multilinear Algebra 69(3), p. 489-514.
* Copyright (C) 2021-2025 Advanced Micro Devices, Inc.
* Copyright (C) 2021-2026 Advanced Micro Devices, Inc.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions
Expand Down Expand Up @@ -169,6 +169,14 @@ __device__ void run_syevj(const rocblas_int dimx,
local_res += cosines_res[i];
local_diag += std::real(sines_diag[i]);
}

// -----------------------------------------------------------------------------
// this read of cosines_res[]/sines_diag[] must complete before the first
// sweep overwrites them with the rotation (c,s1). Without it tolerance is not
Comment thread
EdDAzevedo marked this conversation as resolved.
Resolved
// workgroup-uniform, so the loop may run a different number of barriers per wave
// -----------------------------------------------------------------------------
__syncthreads();

S tolerance = (local_res + local_diag) * abstol * abstol;
S small_num = get_safemin<S>() / eps;

Expand Down Expand Up @@ -320,6 +328,11 @@ __device__ void run_syevj(const rocblas_int dimx,
for(i = 0; i < dimx; i++)
local_res += cosines_res[i];

// ---------------------------------------------------------------
// make sure it is safe for next sweep to over-write cosines_res[]
// ---------------------------------------------------------------
__syncthreads();

sweeps++;
}

Expand Down
Loading