Сокращение количества регистров, используемых в ядре CUDA

Question

У меня есть ядро, которое использует 17 регистров, сокращение его до 16 принесло бы мне 100% загрузки. У меня вопрос: есть ли методы, которые можно использовать для уменьшения числа или используемых регистров, исключая полное переписывание моих алгоритмов другим способом. Я всегда предполагал, что компилятор намного умнее меня, поэтому, например, для ясности я часто использую дополнительные переменные. Я ошибаюсь в этом мышлении?

Обратите внимание: я знаю о флаге --max_registers (или каков бы ни был синтаксис), но использование локальной памяти было бы более вредным, чем 25% -ная занятость (я должен проверить это)

Answer 1

Занятость может немного вводить в заблуждение, и 100% -ая занятость не должна быть вашей основной целью.Если вы можете получить полностью объединенный доступ к глобальной памяти, то на высокопроизводительном графическом процессоре 50% занятости будет достаточно, чтобы скрыть задержку для глобальной памяти (для чисел с плавающей запятой, даже ниже для двойных).Посмотрите презентацию Advanced CUDA C от GTC в прошлом году для получения дополнительной информации по этой теме.

В вашем случае вы должны измерять производительность как с maxrregcount, так и без него, установленным на 16. Задержка длялокальная память должна быть скрыта из-за наличия достаточного количества потоков, при условии, что вы не используете произвольный доступ к локальным массивам (что может привести к не слитному доступу).

Чтобы ответить на конкретный вопрос о сокращении регистров, опубликуйте код для более подробных ответов!Понимание того, как работают компиляторы в целом, может помочь, но помните, что nvcc - это оптимизирующий компилятор с большим пространством параметров, поэтому минимизация количества регистров должна быть сбалансирована с общей производительностью.

Answer 2

Трудно сказать, компилятор nvcc, на мой взгляд, не очень умен.
Вы можете попробовать очевидные вещи, например, использовать short вместо int, передавать и использовать переменные по ссылке (например, & variable), развертывать циклы, использовать шаблоны (как в C ++). Если у вас есть деления, трансцендентные функции, применяемые последовательно, попробуйте сделать их в виде цикла. Попробуйте избавиться от условных выражений, возможно, заменив их избыточными вычислениями.

Если вы отправите код, возможно, вы получите конкретные ответы.

Answer 3

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

Считайте, что ядро ​​вычисляет некоторые значения, и эти вычисленные значения используются всеми потоками,

__global__ void kernel(...) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int id0 = blockDim.x * blockIdx.x;

    int reg = id0 * ...;
    int reg0 = reg * a / x + y;


    ...

    int val =  reg + reg0 + 2 * idx;

    output[idx] = val > 10;
}

Таким образом, вместо сохранения reg и reg0 в качестве регистров и возможности их вытекания в локальную память (глобальную память), мы можем использовать разделяемую память.

__global__ void kernel(...) {
    __shared__ int cache[10];

    int idx = threadIdx.x + blockDim.x * blockIdx.x;

    if (threadIdx.x == 0) {
      int id0 = blockDim.x * blockIdx.x;

      cache[0] = id0 * ...;
      cache[1] = cache[0] * a / x + y;
    }
    __syncthreads();


    ...

    int val =  cache[0] + cache[1] + 2 * idx;

    output[idx] = val > 10;
}

Посмотрите на эту бумагу для получения дополнительной информации ..

Answer 4

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

Как это работает, когда уменьшение регистров вызвало замедление скорости

Скорее всего, компилятору пришлось вылить недостаточно регистровых данных в «локальную» память, которая по сути такая же, как глобальная память, и, следовательно, очень медленная

В целях оптимизации я бы рекомендовал использовать ключевые слова, такие как const, volatile и т. Д., Где это необходимо, чтобы помочь компилятору на этапе оптимизации.

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

Answer 5

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