Я пытаюсь предварительно извлечь некоторые данные.Обычно для этого я полагаюсь на компилятор, так как над ним работают многие, многие тысячи людей, и я всего лишь одно скромное существо.
Как говорится, иногда, если я не делаю что-то наНа алгоритмическом уровне у компилятора нет шансов.В этом случае мне нужно, чтобы графический процессор запускал предварительную выборку значений из основной памяти графического процессора за один или два выполнения цикла до того, как это значение действительно понадобится.Вы можете увидеть мою реализацию для этого в приведенном ниже коде.
Около 10% примеров указателей инструкций приходится на строку 382, поэтому оптимизация этого внесла бы существенную разницу в производительность (я в длинном хвосте)).Я удалил много отвлекающих слов (и заменил их пробелами), чтобы код было легче читать.
I 'На самом деле эта техника работала в более сложном коде (который я здесь не включаю).Это была похожая ситуация, когда происходило так много всего, что компилятору не было никакого разумного способа выполнить предварительную выборку - поэтому мне пришлось вмешаться на алгоритмическом уровне.Однако в этой ситуации это не помогает.
Мое основное подозрение в том, что компилятор выяснил, что ему на самом деле не нужны curr_operand_value_p1, curr_operand_value_p2 или curr_operand_value_p3, и что он оптимизировал ихдалеко.Действительно, я не смог найти код ptx, соответствующий этим строкам (я не эксперт по чтению кода ptx, но переход к строкам 370 и 371 в режиме Nsight Eclipse Profiling не показывает никакого соответствия с ptxмонтаж).Из-за моего небольшого размера пакета (512) и большого количества регистров графического процессора я не ограничен регистром (см. Примечание в конце).Поэтому я бы предпочел, чтобы компилятор был менее умным.
Независимо от этого, как я могу заставить компилятор выдавать эти загрузки раньше времени? В качестве альтернативы, есть что-тоболее высокий уровень, который я должен рассмотреть?
Я также попытался написать несколько ptx для предварительной выборки значения, но это не повлияло на производительность.В частности, я написал:
__device__ void prefetch_l1(unsigned long long addr) {
asm(" prefetch.global.L1 [ %1 ];": "=l"(addr) : "l"(addr));
}
Затем я вызвал эту функцию с адресом текущего значения операнда (который инкапсулирован в определении get_operand_value), но существенной разницы в производительности не было.Я также назвал его адресом следующих трех операндов, но опять же, без существенной разницы в производительности.Доступы для get_operand_value сильно слиты (тип данных bool operands[a large number][512]), поэтому я не думаю, что я глубоко перебиваю свой L1, выпуская эти предварительные выборки.
Некоторые связанные ссылки:
Примечание: я мог бы работать с большим размером пакета, и у меня было бы лучшее покрытие задержки - за исключением того, что я в настоящее время ограничен памятью.Кроме того, наличие меньшего размера пакета помогает удовлетворить наши требования к задержке - я знаю, что графические процессоры не известны тем, что они уменьшают задержку, но если я смогу заставить его работать эффективно на графическом процессоре и снизить затраты, то я мог бы также использоватьGPU вместо FPGA или CPU.