Оптимизирует ли nvcc использование регистров?

У меня есть следующее ядро:

    void version1(float *X, float *Y, int N) {
        int n;
        float x,y;

        n = blockIdx.x * blockDim.x + threadIdx.x;
        if (n >= N) return;

        x=X[n];
       x=x+1;
       X[n]=x;

       y=Y[n];
       y=y+1;
       Y[n]=y;
    }

и вторая версия

    void version2(float *X, float *Y, int N) {
        int n;
        float Xb47w;

        n = blockIdx.x * blockDim.x + threadIdx.x;
        if(n >= N) return;

        Xb47w=X[n];
        Xb47w=Xb47w+1;
        X[n]=Xb47w;

        Xb47w=Y[n];
        Xb47w=Xb47w+1;
        Y[n]=Xb47w;
    }

Они дают тот же результат. Однако версию 1 проще читать, а версию 2 сложнее, потому что Xb47w используется как для X, так и для Y. Поэтому я бы предпочел версию 1, но в ней два регистра x y вместо 1 Xb47w для версии 2. У меня есть много ядер, где я сохраняю регистры таким образом, но есть более сложные для чтения и обслуживания.

x больше не используется после X[n]=x, поэтому мне интересно, понимает ли это компилятор CUDA и делает ли версию 1 почти идентичной версии 2, тем самым сохраняя один регистр?


person YLS    schedule 31.07.2020    source источник


Ответы (2)


Оптимизирует ли nvcc использование регистров?

Да, nvcc пытается скомпилировать ваш код, чтобы использовать меньше регистров (хотя минимальное использование регистров само по себе не является целью).

Интересно, понимает ли это компилятор CUDA и делает ли версию 1 почти идентичной версии 2, тем самым экономя один регистр?

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

Таким образом, обе версии вашей функции компилируются в один и тот же код PTX (GodBolt.org)

person einpoklum    schedule 31.07.2020

Внутри nvcc для оптимизации кода используется компилятор C++ (я слишком упрощаю ссылка) Таким образом, возникает вопрос, сохранит ли компилятор C++ регистр?

А ответ - используйте godbolt и сравните сборку ваших двух программ!

Изменить: это не вся история, вы увидите PTX-представление вашей программы (которое вы также можете получить с помощью nvcc). Следующим шагом будет рассмотрение самой сборки графического процессора, называемой SASS (которая зависит от карты).

person Thomas Caissard    schedule 31.07.2020
comment
С GodBolt можно сравнить скомпилированный PTX-код ядер. Это не совсем язык ассемблера; это промежуточное представление, близкое к языку ассемблера графических процессоров NVIDIA и общее для всех них, но само по себе не являющееся сборкой ни одного из них. - person einpoklum; 31.07.2020
comment
Ну да, таким образом, это слишком просто. Но это дало бы вам приблизительное представление о том, что он делает. - person Thomas Caissard; 31.07.2020
comment
Я так понимаю, что не могу быть уверен, что nvcc компилирует версию1 как версию. Это означает, что я должен придерживаться этого управления регистрами, которое может сбивать с толку. - person YLS; 31.07.2020
comment
@YLS: Нет, ты неправильно понял. - person einpoklum; 31.07.2020