0

我是这些 GPGPU 的新手，现在正在尝试使用 OpenCL 进行 k-means 聚类。出于教育目的，我正在尝试使用N个点和一个质心来实现平方距离表计算的最大性能。

为此，我使用集成的 Intel UHD 620 GPU 并进行了合成设置以简化测试（N 是输入点的总数，D 是每个点的维度）：

N：131072 D：256 WorkItemDimensions：1 LocalWorkSize：256

我的幼稚实现使用输入点的合并读取（D x N 矩阵）并给我20.44 MFLOPS（假设我的 GPU 的 peek 性能为~430 GFLOPS看起来真的很糟糕）。这是我天真的内核：

__kernel void naive(__global const float* points, __global const float* centroid, __global float* distances)
{
    const int gid = get_global_id(0);

    float distance = 0.0f;

    #pragma unroll
    for (int i = 0; i < D; i++) {
        float diff = points[i * N + gid] - centroid[i];
        distance += diff * diff;
    }

    distances[gid] = distance;
}

据我了解，对于此类问题，主要瓶颈是内存带宽，因此我决定尝试通过将其放入本地内存来减少对全局质心缓冲区的访问。我的工作组大小是 256，因此我希望将单个质心向量复制到本地内存会提高性能，但没有运气 - 下面的代码给了我19.78 MFLOPS。

__kernel void local_memory(__global const float* points, __global const float* centroid, __global float* distances)
{
    const int gid = get_global_id(0);
    const int lid = get_local_id(0);

    // Work group size equals to a centroid vector dimensionality
    __local float l_centroid[D];
    l_centroid[lid] = centroid[lid];
    barrier(CLK_LOCAL_MEM_FENCE);

    float distance = 0.0f;

    #pragma unroll
    for (int i = 0; i < D; i++) {
        float diff = points[i * N + gid] - l_centroid[i];
        distance += diff * diff;
    }

    distances[gid] = distance;
}

现在我真的很想弄清楚这里发生了什么，以及如何为这项任务获得至少 50% 的峰值性能。

提前致谢！

4

0 回答 0