Нет выгоды от использования локальной памяти на Intel UHD 620

Question

Я новичок в этих вещах GPGPU и теперь экспериментирую с OpenCL для кластеризации k-средних. В образовательных целях я пытаюсь добиться максимальной производительности для расчета таблицы расстояний в квадрате с N точками и только одним центроидом.

Я использую для этой цели свой интегрированный графический процессор Intel UHD 620 и сделал синтетическую настройку для упрощения тестирования (N — общее количество входных точек, D — размерность каждой точки):

N: 131072D: 256WorkItemDimensions: 1LocalWorkSize: 256

Моя наивная реализация использует объединенное чтение входных точек (матрица D на N) и дает мне 20,44 MFLOPS (что выглядит очень плохо, если предположить, что производительность просмотра для моего графического процессора составляет ~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% пиковой производительности для этой задачи.

Заранее спасибо!