在CUDA中优化全局内存负载

Question

我的任务:我有两个矩阵:A-(18x1024)；B-(18x1024)。

我必须从A中提取每个18长度的向量，并与B中的每个18长度向量计算距离，然后找到最小距离和索引。

我的代码：

__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%的内存。有什么方法可以加速这个操作吗？

如果我需要提供任何其他信息，请告诉我。

Answer 1

实际上，我会先看看算法。这是一个几何问题--就像这样对待它。

您应该使用不同的数据结构来表示B数据，例如，通过集群或构建分区结构(例如k-d tree)。这将使您避免实际计算到大多数B元素的距离。(您也可以考虑使用较少维度的项目，但这样做的好处可能更难以捉摸。)

关于访问模式-让连续的线程处理18个元素长的向量的连续元素，而不是让线程单独处理完整的18个元素长的向量，这可能会使您受益。这将更好地适合内存布局-现在，曲折读取是由许多元素组成的，这些元素彼此之间的距离为18。如果我能正确理解代码的话。

(我还认为GetMin()可以避免一些指数掉期，但这并不重要，因为您只执行了很少的几个。)