CUDA Каким образом Kernel Fusion повышает производительность приложений с ограниченным объемом памяти на графическом процессоре?

Question

Я проводил исследования потоковых наборов данных, превышающих объем памяти, доступный на GPU, устройству для базовых вычислений.Одним из основных ограничений является тот факт, что шина PCIe обычно ограничена 8 ГБ / с, а объединение ядра может помочь в повторном использовании данных и в том, что они могут использовать общую память и локальность в графическом процессоре.Большинство исследовательских работ, которые я обнаружил, очень трудно понять, и большинство из них реализуют слияние в сложных приложениях, таких как https://ieeexplore.ieee.org/document/6270615.Я прочитал много статей, и они ВСЕ НЕ МОГУТ ОБЪЯСНИТЬ несколько простых шагов по слиянию двух ядер.

Мой вопрос: как на самом деле работает синтез? .Какие шаги нужно пройти, чтобы изменить нормальное ядро ​​на слитое ядро?Кроме того, необходимо ли иметь более одного ядра, чтобы объединить его, поскольку слияние - это просто причудливый термин для устранения некоторых проблем, связанных с памятью, и использования локальности и общей памяти.

Мне нужно понять, как ядроFusion используется для базовой программы CUDA, например, для умножения матриц или для сложения и вычитания ядер.Действительно простой пример (код не верный, но должен дать представление) как:

int *device_A;
int *device_B;
int *device_C;

cudaMalloc(device_A,sizeof(int)*N);

cudaMemcpyAsync(device_A,host_A, N*sizeof(int),HostToDevice,stream);

KernelAdd<<<block,thread,stream>>>(device_A,device_B); //put result in C
KernelSubtract<<<block,thread,stream>>>(device_C);

cudaMemcpyAsync(host_C,device_C, N*sizeof(int),DeviceToHost,stream); //send final result through the PCIe to the CPU

Answer 1

Основная идея слияния ядер заключается в том, что 2 или более ядер будут преобразованы в 1 ядро. Операции объединены. Изначально может быть не очевидно, в чем выгода. Но это может обеспечить два связанных вида преимуществ:

1. путем повторного использования данных, которые ядро ​​могло заполнить либо в регистрах, либо в разделяемой памяти
2. путем уменьшения (то есть устранения) «избыточных» нагрузок и накоплений

Давайте воспользуемся примером, подобным вашему, где у нас есть ядро ​​Add и ядро ​​умножения, и предположим, что каждое ядро ​​работает над вектором, а каждый поток выполняет следующее:

1. Загрузить мой элемент вектора A из глобальной памяти
2. Добавление константы или умножение на константу моего векторного элемента
3. Сохранить мой элемент обратно в вектор A (в глобальной памяти)

Для этой операции требуется одно чтение на поток и одна запись на поток. Если бы мы выполнили оба из них вплотную, последовательность операций выглядела бы так:

Добавить ядро:

1. Загрузить мой элемент вектора A из глобальной памяти
2. Добавить значение в мой векторный элемент
3. Сохранить мой элемент обратно в вектор A (в глобальной памяти)

Умножить ядро:

1. Загрузить мой элемент вектора A из глобальной памяти
2. Умножить мой векторный элемент на значение
3. Сохранить мой элемент обратно в вектор A (в глобальной памяти)

Мы видим, что шаг 3 в первом ядре и шаг 1 во втором ядре делают вещи, которые на самом деле не нужны для достижения конечного результата, но они необходимы из-за структуры этих (независимых) ядер. Одно ядро ​​не может передавать результаты другому ядру, кроме как через глобальную память.

Но если мы объединим два ядра вместе, мы могли бы написать ядро ​​так:

1. Загрузить мой элемент вектора A из глобальной памяти
2. Добавить значение в мой векторный элемент
3. Умножить мой векторный элемент на значение
4. Сохранить мой элемент обратно в вектор A (в глобальной памяти)

Это объединенное ядро ​​выполняет обе операции, выдает один и тот же результат, но вместо 2 операций глобальной загрузки памяти и 2 операций глобального хранения памяти требуется только 1 каждая.

Эта экономия может быть очень значительной для операций с памятью (например, таких) в графическом процессоре. За счет уменьшения количества требуемых загрузок и хранилищ общая производительность улучшается, обычно пропорционально уменьшению количества операций загрузки / хранения.