OpenCL AES распараллеливание

Question

Я пытаюсь написать код, который выполняет расшифровку AES для сервера SSL. Чтобы ускорить его, я пытаюсь объединить несколько пакетов для одновременной расшифровки на графическом процессоре.

Если я просто перебираю каждый пакет и отправляю каждое ядро ​​в gpu, а затем выполняю чтение, которое использует событие kernel для своего ожидания. Затем я собираю события для всех операций чтения и жду их одновременно, но кажется, что запускается только один блок за раз, а затем выполняется следующий блок. Это не то, что я ожидал. Я ожидаю, что если я поставлю все ядра в очередь, то надеюсь, что драйверы попытаются выполнить как можно больше работы параллельно.

Я что-то упустил? Нужно ли указывать глобальный рабочий размер, равный размеру всех блоков пакета, и указывать локальный размер ядра, равный размеру блоков каждого пакета?

Это мой код для моего ядра OpenCL.

__kernel void decryptCBC( __global const uchar *rkey, const uint rounds, 
    __global const uchar* prev, __global const uchar *data, 
    __global uchar *result, const uint blocks ) {

    const size_t id = get_global_id( 0 );
    if( id > blocks ) return;

    const size_t startPos = BlockSize * id;

    // Create Block
    uchar block[BlockSize];
    for( uint i = 0; i < BlockSize; i++) block[i] = data[startPos+i];

    // Calculate Result
    AddRoundKey( rkey, block, rounds );

    for( uint j = 1; j < rounds; ++j ){
        const uint round = rounds - j;
        InverseShiftRows( block );
        InverseSubBytes( block );
        AddRoundKey( rkey, block, round );
        InverseMixColumns( block );
    }

    InverseSubBytes( block );
    InverseShiftRows( block );
    AddRoundKey( rkey, block, 0 );

    // Store Result
    for( uint i = 0; i < BlockSize; i++ ) {
        result[startPos+i] = block[i] ^ prev[startPos+i];
    }
}

С этим ядром я могу превзойти 8-ядерный процессор с 125 блоками данных в одном пакете. Чтобы ускорить несколько пакетов, я попытался объединить все элементы данных. Это включало объединение входных данных в единый вектор, а затем возникли сложности из-за необходимости, чтобы каждое ядро ​​знало, где получить доступ к ключу, что привело к двум дополнительным массивам, содержащим число раундов и смещение раундов. Это оказалось даже медленнее, чем отдельное выполнение ядра для каждого пакета.

Answer 1

Рассматривайте ваше ядро ​​как функцию, выполняющую работу CBC.Как вы выяснили, его цепная природа означает, что сама задача CBC принципиально сериализована.Кроме того, графический процессор предпочитает запускать 16 потоков с одинаковыми рабочими нагрузками.По сути, это размер одной задачи в многопроцессорном ядре, которых, как правило, десятки;но система управления может выполнять только некоторые из этих задач в целом, и система памяти редко справляется с ними.Кроме того, циклы являются одним из худших вариантов использования ядра, поскольку графические процессоры не предназначены для выполнения большого потока управления.

Итак, глядя на AES, он работает с 16-байтовыми блоками, но только в байтовых операциях.Это будет ваше первое измерение - каждый блок должен быть обработан 16 потоками (вероятно, локальный размер работы на языке opencl).Обязательно перенесите блок в локальную память, где все потоки могут работать в режиме блокировки, делая произвольный доступ с очень низкой задержкой.Разверните все в операции блока AES, используя get_local_id (0), чтобы узнать, с каким байтом работает каждый поток.Синхронизировать с барьером (CLK_LOCAL_MEM_FENCE) в случае, если рабочая группа работает на процессоре, который может выйти из режима блокировки.Ключ, вероятно, может быть сохранен в постоянной памяти, поскольку он может быть кэширован.Цепочка блоков может быть подходящим уровнем для цикла, хотя бы для того, чтобы избежать перезагрузки предыдущего зашифрованного текста блока из глобальной памяти.Также может помочь асинхронное сохранение завершенного зашифрованного текста с использованием async_work_group_copy ().Возможно, вы можете заставить поток выполнять больше работы, используя векторы, но это, вероятно, не поможет из-за таких шагов, как shiftRows.

В принципе, если какой-либо поток в группе из 16 потоков (может отличаться в зависимости от архитектуры) получает какой-либо другой поток управления, ваш графический процессор останавливается.И если таких групп недостаточно для заполнения конвейеров и мультипроцессоров, ваш графический процессор бездействует.Пока вы не очень тщательно оптимизируете доступ к памяти, он не приблизится к скорости процессора, и даже после этого вам нужно будет обрабатывать десятки пакетов одновременно, чтобы избежать создания слишком маленьких рабочих групп для GPU.Проблема заключается в том, что, хотя графический процессор может выполнять тысячи потоков, его структура управления в любое время обрабатывает только несколько рабочих групп.

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

Answer 2

Из вашего описания не совсем понятно, но я думаю, что есть некоторая концептуальная путаница.

Не перебирайте каждый пакет и не запускайте новое ядро.Вам не нужно указывать OpenCL запустить несколько ядер.Вместо этого загрузите как можно больше пакетов в графический процессор, а затем запустите ядро ​​только один раз.Когда вы указываете размер рабочей группы, это количество ядер, которые GPU пытается запустить одновременно.

Вам нужно будет запрограммировать ваши ядра для каждого просмотра в разных местах данных, которые вы загрузили, чтобы найти их пакет.Например, если вы собираетесь добавить два массива в третий массив, ваше ядро ​​будет выглядеть так:

__kernel void vectorAdd(__global const int* a,
                        __global const int* b,
                        __global int* c) {
  int idx = get_global_id(0);
  c[idx] = a[idx] + b[idx];
}

Важной частью является то, что каждое ядро ​​знает индекс в массиве, используя свой глобальный идентификатор.Вы хотите сделать что-то подобное.