Описание тега ptx

Графические процессоры nVIDIA имеют разные микроархитектуры, изменения между которыми не всегда являются инкрементными (например, добавление инструкций к x86-64 с последовательными расширениями avx). Однако все они имеют общий промежуточный (виртуальный) набор инструкций, чем-то похожий на промежуточное представление компилятора. В частности, это своего рода параллель с представлением, связанным со стандартом OpenCL, spir-v. Продолжая инструментальную цепочку компиляции, PTX компилируется в один из нескольких языков сборки ( sass), специфичных для графической микроархитектуры, для фактического выполнения.

Вот пример простого ядра CUDA и PTX, полученного в результате его компиляции:

__global__ void square(int *array, int length) {
    int pos = threadIdx.x + blockIdx.x * blockDim.x;
    if (pos < length)
        array[pos] = array[pos] * array[pos];
}

Полученный PTX (после удаления названия):

.visible .entry square(int*, int)(
        .param .u64 square(int*, int)_param_0,
        .param .u32 square(int*, int)_param_1
)
{
        ld.param.u64        %rd1, [square(int*, int)_param_0];
        ld.param.u32        %r2, [square(int*, int)_param_1];
        mov.u32             %r3, %tid.x;
        mov.u32             %r4, %ntid.x;
        mov.u32             %r5, %ctaid.x;
        mad.lo.s32          %r1, %r4, %r5, %r3;
        setp.ge.s32         %p1, %r1, %r2;
        @%p1 bra            BB0_2;

        cvta.to.global.u64  %rd2, %rd1;
        mul.wide.s32        %rd3, %r1, 4;
        add.s64             %rd4, %rd2, %rd3;
        ld.global.u32       %r6, [%rd4];
        mul.lo.s32          %r7, %r6, %r6;
        st.global.u32       [%rd4], %r7;

        ret;
}

Для получения дополнительной информации о PTX в целом, а также о конкретных инструкциях и синтаксисе доступа к данным в приведенном выше примере обратитесь к nVIDIA PTX Referene.