Почему целочисленное деление и модуль не оптимизированы в NVRTC

Question

Я собрал ядро ​​в NVRTC:

__global__ void kernel_A(/* args */) {
    unsigned short idx = threadIdx.x;
    unsigned char warp_id = idx / 32;
    unsigned char lane_id = idx % 32;
    /* ... */
}

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

__global__ void kernel_B(/* args */) {
    unsigned short idx = threadIdx.x;
    unsigned char warp_id = idx >> 5;
    unsigned char lane_id = idx & 31;
    /* ... */
}

похоже на то kernel_B просто бежит быстрее. При пропуске всех других кодов в ядре, начиная с 1024 блоков размером 1024, nvprof шоу kernel_A работает в среднем за 15,2 с, а kernel_B работает 7.4us в среднем. Я предполагаю, что NVRTC не оптимизировал целочисленное деление и по модулю.

Результат получается на GeForce 750 Ti, CUDA 8.0, усредненный из 100 вызовов. Параметры компилятора, данные для nvrtcCompileProgram() является -arch compute_50,

Это ожидается?

Answer 1

Сделал тщательный багсвип в кодовой базе. Оказывается, мое приложение было встроено DEBUG Режим. Это вызывает дополнительные флаги -G а также -lineinfo перешел к nvrtcCompileProgram()

От nvcc справочная страница:

--device-debug(-G)

Генерация отладочной информации для кода устройства. Отключает все оптимизации. Не используйте для профилирования; используйте вместо этого -lineinfo.