Оптимизация загрузки глобальной памяти в CUDA - PullRequest
Новая
       26

Оптимизация загрузки глобальной памяти в CUDA

2 голосов
/ 25 февраля 2020

Моя задача: у меня есть две матрицы: A - (18 x 4194304); Б - (18 х 1024).

Я должен взять каждый вектор длины 18 из A и вычислить расстояние с каждым вектором длины 18 из B и найти минимальное расстояние и индекс.

Мой код: 

__device__
void GetMin(float &dist, int &idx)
{
    float dist2;
    int idx2;
    dist2 = __shfl_down_sync(0xFFFFFFFF, dist, 16, 32);
    idx2 = __shfl_down_sync(0xFFFFFFFF, idx, 16);
    if (dist > dist2)
    {
        dist = dist2;
        idx = idx2;
    }

    dist2 = __shfl_down_sync(0xFFFFFFFF, dist, 8, 32);
    idx2 = __shfl_down_sync(0xFFFFFFFF, idx, 8);
    if (dist > dist2)
    {
        dist = dist2;
        idx = idx2;
    }

    dist2 = __shfl_down_sync(0xFFFFFFFF, dist, 4, 32);
    idx2 = __shfl_down_sync(0xFFFFFFFF, idx, 4);
    if (dist > dist2)
    {
        dist = dist2;
        idx = idx2;
    }

    dist2 = __shfl_down_sync(0xFFFFFFFF, dist, 2, 32);
    idx2 = __shfl_down_sync(0xFFFFFFFF, idx, 2);
    if (dist > dist2)
    {
        dist = dist2;
        idx = idx2;
    }

    dist2 = __shfl_down_sync(0xFFFFFFFF, dist, 1, 32);
    idx2 = __shfl_down_sync(0xFFFFFFFF, idx, 1);
    if (dist > dist2)
    {
        dist = dist2;
        idx = idx2;
    }
}

__global__
void CalcMinDist_kernel(const float *A, const float *B, float *output, const int nNumPixels, int nNumImages)
{
    int tx = threadIdx.x + blockIdx.x * blockDim.x;
    int ty = threadIdx.y;

    int lane_id = tx % 32;

    float dist = 0;
    int idx = 0;

    float fMin = 99999999;
    int nMinIdx = -1;

    for(int i = lane_id; i < 1024; i += 32)
    {
        dist = 0;
        for(int  j = 0; j < nNumImages; ++j)
        {
            int img_idx = blockIdx.x * ty + j * nNumPixels;
            dist += (A[img_idx] - B[i * nNumImages + j]) * 
                    (A[img_idx] - B[i * nNumImages + j]);
        }
        idx = i;
        GetMin(dist, idx);

        if(threadIdx.x == 0)
        {
            if(fMin > dist)
            {
                fMin = dist;
                nMinIdx = idx;
            }
        }
    }

    if(threadIdx.x == 0)
    {
        output[blockIdx.x * ty] = nMinIdx;
    }
}

Глядя на профилировщик, я ограничен в памяти и занимаю ~ 90%. Есть ли способ ускорить эту операцию?

Дайте мне знать, если мне нужно предоставить какую-либо другую информацию.

1 Ответ

1 голос
/ 25 февраля 2020

На самом деле, я бы сначала посмотрел на алгоритм . Это проблема геометрии c - рассматривайте ее так.

Вы должны представлять данные B, используя другую структуру данных, например, путем кластеризации или построения структуры разделов (например, kd tree ). Это позволит вам избежать фактического вычисления расстояния от большинства элементов B. (Вы могли бы также рассмотреть проект с меньшим количеством измерений, но выгода от этого может быть более неуловимой.)

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

(я также думаю, что GetMin() мог бы избежать некоторых свопов с индексами, но это несущественно, поскольку вы выполняете только очень немногие из них.)

...