Каков наилучший способ доступа к памяти в этой проблеме с N-корпусом, решенной на AMD Radeon RX580?

Question

Я вычисляю траектории N частиц, которые движутся в своем поле силы тяжести. Я написал следующее ядро ​​OpenCL:

#define G 100.0f
#define EPS 1.0f

float2 f (float2 r_me, __constant float *m, __global float2 *r, size_t s, size_t n)
{
    size_t i;
    float2 res = (0.0f, 0.0f);

    for (i=1; i<n; i++) {
        size_t idx = i;
//        size_t idx = (i + s) % n;
        float2 dir = r[idx] - r_me;
        float dist = length (dir);
        res += G*m[idx]/pown(dist + EPS, 3) * dir;
    }

    return res;
}

__kernel void take_step_rk2 (__constant float *m,
                             __global float2 *r,
                             __global float2 *v,
                             float delta)
{
    size_t n = get_global_size(0);
    size_t s = get_global_id(0);


    float2 mv = f(r[s], m, r, s, n);
    float2 mr = v[s];

    float2 vpred1 = v[s] + mv * delta;
    float2 rpred1 = r[s] + mr * delta;

    float2 nv = f(rpred1, m, r, s, n);
    float2 nr = vpred1;

    barrier (CLK_GLOBAL_MEM_FENCE);

    r[s] += (mr + nr) * delta / 2;
    v[s] += (mv + nv) * delta / 2;
}

Затем я запускаю это ядро ​​несколько раз как одномерную задачу с глобальным рабочим размером = [количество тел]:

void take_step (struct cl_state *state)
{
    size_t n = state->nbodies;
    clEnqueueNDRangeKernel (state->queue, state->step, 1, NULL, &n, NULL, 0, NULL, NULL);
    clFinish (state->queue);
}

Это цитата из руководства по оптимизации AMD OpenCL (2015 год):

При определенных условиях одним из неожиданных случаев конфликта каналов является то, что чтение с одного и того же адреса является конфликтом, даже в FastPath. Этого не происходит с памятью только для чтения, такой как буферы констант, текстуры или представление ресурсов шейдера (SRV); но это возможно для чтения / записи памяти UAV или глобальной памяти OpenCL.

Все рабочие элементы в моей очереди пытаются получить доступ к одной и той же памяти в этом цикле, поэтому должен быть конфликт каналов:

for (i=1; i<n; i++) {
        size_t idx = i;
//        size_t idx = (i + s) % n;
        float2 dir = r[idx] - r_me;
        float dist = length (dir);
        res += G*m[idx]/pown(dist + EPS, 3) * dir;
    }

Я заменил

        size_t idx = i;
//        size_t idx = (i + s) % n;

с участием

//        size_t idx = i;
        size_t idx = (i + s) % n;

поэтому первый рабочий элемент (с глобальным идентификатором 0) сначала получить доступ к первому элементу в массиве r, второй рабочий элемент обращается ко второму элементу и так далее.

Я ожидал, что это изменение должно привести к повышению производительности, но, напротив, оно привело к значительному снижению производительности (примерно в 2 раза). Что мне не хватает? Почему в этой ситуации лучше обращаться к одной и той же памяти?

Если у вас есть другие советы по повышению производительности, поделитесь со мной. Руководство по оптимизации OpenCL очень сбивает с толку.

Answer 1

Цикл функции f не имеет препятствий для повторной конвергенции для объединенного доступа. Как только некоторые элементы получают свои r-данные, они начинают вычисления, но те, кто не может ждать своих данных, теряют целостность слияния. Чтобы перегруппировать их, добавьте 1 барьер как минимум на 10 или 2 итерации, а может быть, даже на каждую итерацию. Но доступ к global имеет большую задержку. Барьер + задержка плохо сказываются на производительности. Здесь вам нужна локальная память, поскольку она имеет низкую задержку и возможность широковещательной передачи, что позволяет ей терять объединенность только на гранулах, превышающих номер локального потока (64?), Что также неплохо для доступа к глобальной памяти (вам нужно заполнять локальную память из глобальной в каждом K-я итерация, где N разделено на группы размером K).

Источник от 2013 года (http://developer.amd.com/wordpress/media/2013/07/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide-rev-2.7.pdf):

Таким образом, ключом к эффективному использованию LDS является управление шаблоном доступа, чтобы доступы, созданные в одном цикле, отображались в разные банки в LDS. Одним примечательным исключением является то, что доступ к одному и тому же адресу (даже если у них одинаковые биты 6:2) может транслироваться всем запрашивающим и не генерировать конфликт банков.

Использование LDS(__local) для этого даст хорошие характеристики. Поскольку LDS небольшой, вы должны делать это небольшими участками, например, по 256 частиц за раз.

Кроме того, использование i в качестве idx очень удобно для кеширования, но модульная версия - очень неприятный кеш. Если данные могут существовать в кеше, не имеет значения, выполнено ли N запросов. Сейчас они из кеша. Но с модулем вы уничтожаете ингредиенты кеша до того, как они будут повторно использованы, в зависимости от N. Для малых N это должно быть быстрее, как вы и предполагали. Для больших N и с маленьким кешем GPU было бы намного хуже. Как только 1 глобальный запрос за цикл по сравнению с глобальными запросами N-cache_size за цикл.

Я предполагаю, что с таким мощным графическим процессором у вас было высокое значение N, такое как 64k тел, для которых требовалось 2 переменных на тело и 4 байта на переменную на общую сумму 512 КБ, что не может соответствовать L1. Может быть, только L2, который медленнее, чем idx=i через L1.

Ответ:

• все в тот же L1 cache adr быстрее, чем all to global и L2 cache adr

• использовать локальную память в алгоритме "блокировки / исправления" для достижения высокой скорости