Правильная реализация треугольного решателя для разреженной матрицы с OpenACC

В настоящее время я реализую треугольный решатель для разреженной матрицы, и я пытаюсь ускорить его, используя директивы OpenACC. Учитывая мои матричные факторы LU в разреженном формате CSR, OpenACC удалось правильно решить L-фактор, но U-фактор дает совершенно неверные результаты по сравнению с истинным решением приложения. Вот код ускоренного ядра для задачи обратной подстановки:

#pragma acc kernels deviceptr( ia, ja, factorValsU, y, x )
{
    for ( int idx = size; idx > 0; idx-- )
    {
        double temp = 0.0;
        int rowInit = ia[ idx - 1];
        int rowEnd  = ia[ idx ];
        #pragma acc loop vector reduction( + : temp)
        for ( int k = rowInit + 1; k < rowEnd; k++ )
        {
            temp += factorValsU[ k ] * x[ ja[ k ] ];
        }
        x[ idx ] = (y[ idx ] - temp) / factorValsU[ rowInit ];
    }
 }

Я не понимаю, почему эти ядра дают неправильный результат. Я уже пробовал другую версию ядра, в которой матрица сохраняется в обратном порядке, то есть сверху вниз, что в принципе можно решить с помощью следующего ядра:

#pragma acc kernels deviceptr( ia, ja, factorValsU, y, x )
{
    for ( int idx = 0; idx < size; idx++ )
    {
        double temp = 0.0;
        int rowInit = ia[ idx ];
        int rowEnd  = ia[ idx + 1 ];
        #pragma acc loop vector reduction( + : temp)
        for ( int k = rowInit + 1; k < rowEnd; k++ )
        {
            temp += factorValsU[ k ] * x[ ja[ k ] ];
        }
        x[ size - idx ] = (y[ size - idx ] - temp) / factorValsU[ rowInit ];
    }
 }

Но результат всегда неверный. Пропустил ли я что-то фундаментальное в украшении обычного кода директивами OpenACC для достижения правильных результатов?

Как упоминалось ранее, прямая подстановка фактора L работает правильно, поэтому для полноты картины я публикую здесь код.

#pragma acc kernels deviceptr( ia, ja, factorValsL, y, x )
{
    for ( int idx = 0; idx < size; idx++ )
    {
        double temp = 0.0;
        int rowInit = ia[ idx ];
        int rowEnd  = ia[ idx + 1 ];
        #pragma acc loop vector reduction( + : temp)
        for ( int k = rowInit; k < rowEnd; k++ )
        {
            temp += factorValsL[ k ] * x[ ja[ k ] ];
        }
        x[ idx ] = y[ idx ] - temp;
    }
 }

Обратите внимание на тонкую разницу между ядром для прямой подстановки (работает) и обратной подстановки (оба не работают), это область памяти, в которой сохраняется результат:

x[ idx ] = y[ idx ] - temp     for the L factor 
x[ size - idx ] = (y[ size - idx ] - temp) / factorValsU[ rowInit ] for the U factor;

Есть ли причина, по которой решатель U-фактора вычисляет ошибочные результаты из-за порядка, в котором выполняется присвоение (и лекция) в памяти?

Для полноты информации, предоставленной компилятором pgi18.4 о ядре:

triangularSolverU_acc(int, const int *, const int *, const double *, const double *, double *, bool):
614, Complex loop carried dependence of y->,x->,factorVals-> prevents parallelization
     Loop carried dependence of x-> prevents parallelization
     Loop carried backward dependence of x-> prevents vectorization
     Accelerator kernel generated
     Generating Tesla code
    614, #pragma acc loop seq
    621, #pragma acc loop vector(128) /* threadIdx.x */
         Generating reduction(+:temp)
621, Loop is parallelizable

Это показывает, что внешний цикл был сериализован, а внутренний цикл сокращен.