Есть ли способ избежать CUDA atomicAdd в моей ситуации?

Question

Я делаю операцию, как показано на рисунке ниже.

Вот мое ядро.

Как показано на рисунке, я делаю небольшую матрицу, используя около миллиона векторов, и накапливаю ее в большой подготовленной матрице.

Мне нужна идея, которая может улучшить производительность, не превышая 8 ГБ графического процессораглобальная память.

Как мне избежать атомарных операций?Я использую GTX1080.Существующие ядра занимают около 250 мс.

__global__ void buildMatrixKernel(const CostJacobianCT *src, const int num, const int st, const int mw_width, double *A, double *b)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < num)
    {
        if (src[idx].mask == 1)
        {
            // matrix width
            int cols = 6 * (mw_width + 1);

            // calc position for insert
            int idx0 = (src[idx].fid0 - st);
            if (idx0 == mw_width - 2)
            {
                idx0 = idx0 - 1;
            }
            else if (idx0 == mw_width - 1)
            {
                idx0 = idx0 - 2;
            }

            int idx1 = (src[idx].fid1 - st);
            if (idx1 == mw_width - 2)
            {
                idx1 = idx1 - 1;
            }
            else if (idx1 == mw_width - 1)
            {
                idx1 = idx1 - 2;
            }
            int pos0 = idx0 * 6;
            int pos1 = idx1 * 6;

            // set tempolar matrix
            double _A00[24 * 24];
            double _A11[24 * 24];
            double _A01[24 * 24];
            double _b0[24];
            double _b1[24];
            for (int y = 0; y < 24; y++)
            {
                for (int x = 0; x < 24; x++)
                {
                    _A00[y * 24 + x] = src[idx].w * src[idx].J0[y] * src[idx].J0[x];
                    _A11[y * 24 + x] = src[idx].w * src[idx].J1[y] * src[idx].J1[x];
                    _A01[y * 24 + x] = src[idx].w * src[idx].J0[y] * src[idx].J1[x];
                }
                _b0[y] = src[idx].w * src[idx].c * src[idx].J0[y];
                _b1[y] = src[idx].w * src[idx].c * src[idx].J1[y];
            }

            // set final matrix
            for (int i = 0; i < 24; i++)
            {
                for (int j = 0; j < 24; j++)
                {
                    atomicAdd(&A[(i + pos0) * cols + (j + pos0)], _A00[i * 24 + j]); // 00
                    atomicAdd(&A[(i + pos1) * cols + (j + pos1)], _A11[i * 24 + j]); // 11
                    atomicAdd(&A[(i + pos0) * cols + (j + pos1)], _A01[i * 24 + j]); // 01
                    atomicAdd(&A[(i + pos1) * cols + (j + pos0)], _A01[j * 24 + i]); // 10
                }
                atomicAdd(&b[i + pos0], _b0[i]); // 0
                atomicAdd(&b[i + pos1], _b1[i]); // 1
            }
        }
    }
}

---

2019.3.6.Я изменил код ниже, чтобы увидеть некоторые улучшения производительности.250 мс -> 95 мс

__global__ void buildMatrixKernel(const CostJacobianCT *src, const int num, const int st, const int mw_width, double *A, double *b)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < num)
    {
        int src_idx = idx / 576;
        if (src[src_idx].mask == 1)
        {
            int cols = 6 * (mw_width + 1);
            int pos0 = src[src_idx].pos0;
            int pos1 = src[src_idx].pos1;
            double w = src[src_idx].w;
            double c = src[src_idx].c;

            int sub_idx = idx % 576;
            int i = sub_idx / 24;
            int j = sub_idx % 24;

            double J0_i = src[src_idx].J0[i];
            double J0_j = src[src_idx].J0[j];
            double J1_i = src[src_idx].J1[i];
            double J1_j = src[src_idx].J1[j];

            atomicAdd(&A[(i + pos0) * cols + (j + pos0)], w * J0_i * J0_j); // 00
            atomicAdd(&A[(i + pos1) * cols + (j + pos1)], w * J1_i * J1_j); // 11
            atomicAdd(&A[(i + pos0) * cols + (j + pos1)], w * J0_i * J1_j); // 01
            atomicAdd(&A[(i + pos1) * cols + (j + pos0)], w * J1_i * J0_j); // 10

            if (j == 0)
            {
                atomicAdd(&b[i + pos0], w * c * J0_i); // 0
                atomicAdd(&b[i + pos1], w * c * J1_i); // 1
            }
        }
    }
}