Поэлементные операции в OpenCL (Cuda)

Question

Я строю ядро ​​для поэлементного умножения двух матриц, но по крайней мере с моими конфигурациями мое ядро ​​OpenCL работает быстрее только тогда, когда каждая матрица больше 2 ГБ. Поэтому мне было интересно, если это из-за моего наивного ядра (см. Ниже) или из-за природы поэлементных операций, то есть, что элементарные операции не выигрывают от использования графических процессоров.

Спасибо за ваш вклад!

Ядро:

KERNEL_CODE = """
// elementwise multiplication: C = A .* B.
__kernel void matrixMul(
        __global float* C,
        __global float* A,
        __global float* B,
        int width, int height)
{
    // ID
    int x = get_global_id(0);
    int y = get_global_id(1);

    // Multiplying
    C[y * height + x ] = A[y * height + x] * B[y * height + x];
}
"""

p.s. Я читал, что некоторые эксперты считают, что CUDA слишком отличается от OpenCL, чтобы отвечать на оба вопроса в одном и том же вопросе, но не смог удалить его из заголовка и тегов.

Answer 1

Операции такого типа имеют N FLOP, но 3N транзакций памяти, поэтому она будет полностью ограничена пропускной способностью памяти.Нет возможности для повторного использования данных, поэтому верхняя граница ускорения по сравнению с эталонной версией ЦП является отношением GPU к пропускной способности ЦП.Это число редко превышает 10 раз и может довольно быстро исчезнуть из-за стоимости перемещения данных в память графического процессора и обратно.Вообще говоря, этот тип операции лучше всего «сливать» с другими операциями O (N) для повышения производительности.Обычно вы никогда не будете просто вычислять продукт Hadamard в одном ядре, вы будете делать это как часть серии операций O (N) в одном ядре.Так что нет, это не лучший кандидат на ускорение, даже если ядро ​​было оптимальным.

А ваше ядро ​​определенно не так.Вы делаете 3 IOP для каждого FLOP, что является огромным штрафом.Вы можете определенно сделать что-то, чтобы улучшить это, но что будет зависеть полностью от того, на каком оборудовании это будет работать.

Answer 2

Говоря об элементарных операциях: это зависит от устройства. Например, в графических процессорах NVidia используются скалярные процессоры (со скалярными инструкциями), векторизация не требуется. Напротив, ATI имеет 5d (или 4d) VLIW процессоры, и для них крайне важна векторизация. Однако иногда он может выполняться компилятором, а не использовать векторные типы данных непосредственно в коде, но это первое, что нужно сделать при оптимизации для графических процессоров ATI.

Тем не менее, как указали talonmies, приведенный выше алгоритм вряд ли ограничен пропускной способностью памяти, и вы не можете ожидать большого ускорения, используя только GPU для него.

Answer 3

Ядро, которое вы разместили, должно быть по крайней мере так же быстро, как и процессорное.Но вы вообще не используете объединенный доступ к памяти!

Это убивает вашу производительность.

Однако, как заявляет @talonmiesЭто не хороший случай для графического процессора.Вы теряете все свое время в памяти копии.