Путаница с кодом CUDA PTX и память регистра

Question

:) Пока я пытался управлять ресурсами ядра, я решил посмотреть в PTX, но есть несколько вещей, которые я не понимаю. Вот очень простое ядро, которое я написал:

__global__ 
void foo(float* out, float* in, uint32_t n) 
{ 
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; 
    uint32_t one = 5; 
    out[idx] = in[idx]+one; 
} 

Тогда я скомпилировал его помощью: nvcc --ptxas-options=-v -keep main.cu и я получил этот вывод на консоль:

ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z3fooPfS_j' for 'sm_10' 
ptxas info : Used 2 registers, 36 bytes smem 

И в результате PTX является следующее:

.entry _Z3fooPfS_j (
      .param .u64 __cudaparm__Z3fooPfS_j_out, 
      .param .u64 __cudaparm__Z3fooPfS_j_in, 
      .param .u32 __cudaparm__Z3fooPfS_j_n) 
    { 
    .reg .u16 %rh<4>; 
    .reg .u32 %r<5>; 
    .reg .u64 %rd<8>; 
    .reg .f32 %f<5>; 
    .loc 15 17 0 
$LDWbegin__Z3fooPfS_j: 
    .loc 15 21 0 
    mov.u16  %rh1, %ctaid.x; 
    mov.u16  %rh2, %ntid.x; 
    mul.wide.u16 %r1, %rh1, %rh2; 
    cvt.u32.u16  %r2, %tid.x; 
    add.u32  %r3, %r2, %r1; 
    cvt.u64.u32  %rd1, %r3; 
    mul.wide.u32 %rd2, %r3, 4; 
    ld.param.u64 %rd3, [__cudaparm__Z3fooPfS_j_in]; 
    add.u64  %rd4, %rd3, %rd2; 
    ld.global.f32 %f1, [%rd4+0]; 
    mov.f32  %f2, 0f40a00000;  // 5 
    add.f32  %f3, %f1, %f2; 
    ld.param.u64 %rd5, [__cudaparm__Z3fooPfS_j_out]; 
    add.u64  %rd6, %rd5, %rd2; 
    st.global.f32 [%rd6+0], %f3; 
    .loc 15 22 0 
    exit; 
$LDWend__Z3fooPfS_j: 
    } // _Z3fooPfS_j 

Теперь есть некоторые вещи, которые я не понимаю:

• В соответствии с сборкой ptx используются 4 + 5 + 8 + 5 = 22 регистра. Тогда почему он пишет used 2 registers во время компиляции?
• Глядя на сборку, я понял, что тип данных threadId, blockId и т. Д. - u16. Это определено в спецификации CUDA? Или это может варьироваться в зависимости от разных версий драйвера CUDA?
• Может кто-нибудь объяснить мне эту строку: mul.wide.u16 %r1, %rh1, %rh2;? %r1 is u32, почему wide вместо u32 используется?
• Как выбираются имена регистров? В моей вазе я понимаю часть %r, но я не понимаю h, (null), d часть. Выбирается ли она на основе длины данных? т.е.: h для 16 бит, null для 32 бит, d для 64 бит?
• Если я заменю последние 2 строки моего ядра на этот out[idx] = in[idx];, тогда, когда я скомпилирую программу, он говорит, что используются 3 регистра! Как теперь можно использовать больше регистров?

Пожалуйста, проигнорируйте тот факт, что мое тестовое ядро ​​не проверяет, является ли индекс массива вне пределов.

спасибо.

Answer 1

PTX - это промежуточный язык, который предназначен для переносимости нескольких архитектур графического процессора. Он компилируется компонентом компилятора PTXAS в конечный машинный код, также называемый SASS для конкретной архитектуры. Опция nvcc -Xptxas -v заставляет PTXAS сообщать о различных статистических данных о сгенерированном машинный код, включая количество физических регистров, используемых в машинных кодах. Вы можете проверить код машины, разобрав его cuobjdump --dump-sass.

Таким образом, количество регистров, которые, как видно, используется в коде PTX, не имеет значения, поскольку это виртуальные регистры. Компилятор CUDA генерирует код PTX в так называемой форме SSA (статическое одиночное назначение, см. http://en.wikipedia.org/wiki/Static_single_assignment_form). Это в основном означает, что каждому новому результату присваивается новый регистр.

Инструкция mul.wide описана в спецификации PTX, текущая версия которой (3.1) вы можете найти здесь: http://docs.nvidia.com/cuda/parallel-thread-execution/index.html. В вашем примере кода суффикс .u16 означает, что он умножает две неподписанные 16-битные величины и возвращает 32-битный результат без знака, то есть он вычисляет полное произведение двойной ширины исходных операндов.

Набираются виртуальные регистры в PTX, но их имена могут быть выбраны свободно, независимо от типа. Компилятор CUDA, похоже, придерживается определенных соглашений, которые (насколько мне известно) не документированы, поскольку они являются внутренними артефактами реализации. Глядя на кучу кода PTX, ясно, что имя регистра в настоящее время генерирует информацию типа кодирования, это может быть сделано для облегчения отладки: p<num> используется для предикатов, r<num> для 32-разрядных целых чисел, rd<num> для 64-битных целых чисел, f<num> для 32-битных поплавков и fd<num> для 64-битных удвоений. Вы можете легко это увидеть сами, посмотрев директивы .reg в коде PTX, которые создают эти виртуальные регистры.