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

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

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

Ответ 1

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

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

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

Ответ 2

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

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

Ответ 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;
}

Итак, вместо того, чтобы сохранять регистры и 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;
}

Взгляните на для получения дополнительной информации..

Ответ 4

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

Ответ 5

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

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

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

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

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