Почему компилятор OpenCL Nvidia (nvcc) не использует регистры дважды?

Я выполняю небольшой тест OpenCL с использованием драйверов Nvidia, мое ядро ​​​​выполняет 1024 фьюза, умножая и добавляя, и сохраняю результат в массиве:

#define FLOPS_MACRO_1(x)    { (x) = (x) * 0.99f + 10.f; } // Multiply-add
#define FLOPS_MACRO_2(x)    { FLOPS_MACRO_1(x) FLOPS_MACRO_1(x) }
#define FLOPS_MACRO_4(x)    { FLOPS_MACRO_2(x) FLOPS_MACRO_2(x) }
#define FLOPS_MACRO_8(x)    { FLOPS_MACRO_4(x) FLOPS_MACRO_4(x) }
// more recursive macros ...
#define FLOPS_MACRO_1024(x) { FLOPS_MACRO_512(x) FLOPS_MACRO_512(x) }

__kernel void ocl_Kernel_FLOPS(int iNbElts, __global float *pf)
{
   for (unsigned i = get_global_id(0); i < iNbElts; i += get_global_size(0))  
   {
      float f = (float) i;
      FLOPS_MACRO_1024(f)
      pf[i] = f;
    }   
}

Но когда я смотрю на сгенерированный PTX, я вижу это:

    .entry ocl_Kernel_FLOPS(
    .param .u32 ocl_Kernel_FLOPS_param_0,
    .param .u32 .ptr .global .align 4 ocl_Kernel_FLOPS_param_1
)
{
    .reg .f32   %f<1026>; // 1026 float registers !
    .reg .pred  %p<3>;
    .reg .s32   %r<19>;    

    ld.param.u32    %r1, [ocl_Kernel_FLOPS_param_0];
    // some more code unrelated to the problem
    // ...

BB1_1:
    and.b32     %r13, %r18, 65535;
    cvt.rn.f32.u32  %f1, %r13;
    fma.rn.f32  %f2, %f1, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f3, %f2, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f4, %f3, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f5, %f4, 0f3F7D70A4, 0f41200000;
    // etc
    // ...

Если я прав, PTX использует 1026 регистров с плавающей запятой для выполнения 1024 операций и никогда не использует регистр повторно дважды, даже если он может выполнять все операции умножения-сложения, используя только 2 регистра. 1026 намного превышает максимальное количество регистров, которое может иметь поток (согласно specs), так что я предполагаю, что это приводит к утечке памяти.

Это ошибка компилятора или я чего-то не понимаю?

Я использую nvcc версии 6.5 на графическом процессоре Quadro K2000.

ИЗМЕНИТЬ

На самом деле я что-то упустил в спецификациях:

«Поскольку PTX поддерживает виртуальные регистры, внешний интерфейс компилятора довольно часто генерирует большое количество имен регистров. Вместо того, чтобы требовать явного объявления каждого имени, PTX поддерживает синтаксис для создания набора переменных, имеющих общую строку префикса, добавленную с целочисленные суффиксы. Например, предположим, что программа использует большое количество, скажем, сто, переменных .b32 с именами %r0, %r1, ..., %r99"


nvidia opencl nvcc ptx
person GaTTaCa    schedule 27.02.2015    source источник

Ответы (1)


arrow_upward
5
arrow_downward

формат файла PTX предназначен для описания архитектуры виртуальной машины и набора инструкций:

PTX определяет виртуальную машину и ISA для выполнения параллельных потоков общего назначения. Программы PTX транслируются во время установки в набор инструкций целевого оборудования. Транслятор PTX-to-GPU и драйвер позволяют использовать графические процессоры NVIDIA в качестве программируемых параллельных компьютеров.

Таким образом, вывод PTX, который вы получаете, не является формой "ассемблера графического процессора". Это всего лишь промежуточное представление, предназначенное для описания практически любой формы параллельных вычислений.

Затем представление PTX компилируется в фактические двоичные файлы для соответствующего целевого графического процессора. Это важно для того, чтобы можно было абстрагироваться от фактической архитектуры — в частности, в отношении вашего примера: должна быть возможность использовать одно и то же представление программы PTX независимо от того, количества регистров, доступных на конкретной целевой машине. 1026 «регистров», которые вы видите, являются «виртуальными» регистрами и, в конце концов, могут быть сопоставлены с (немногими) реальными аппаратными регистрами, которые действительно доступны. Вы можете добавить аргумент --ptxas-options=-v в NVCC во время компиляции, чтобы получить дополнительную информацию об использовании регистра.

(Это примерно та же идея, что и в LLVM, а именно, иметь представление, которое можно оптимизировать и аргументировать о том, как абстрагируясь от исходного исходного кода и от фактической целевой архитектуры).

person Marco13    schedule 27.02.2015
comment
Четкий ответ. Спасибо. - person GaTTaCa; 27.02.2015