In my <a href="http://sandbox.arrayfire.com/generating-ptx-files-from-opencl-code/">previous post</a>, I have shown how to generate PTX files from both CUDA and OpenCL kernels. But how does a PTX file looks like, and more importantly, how to understand all those complicated instructions on PTX files?

In this post I will use the same vector addition kernel that was used in that <a href="http://sandbox.arrayfire.com/generating-ptx-files-from-opencl-code/">previous post</a> (the complete code can be found <a href="https://github.com/pentschev/ptxtract">here</a>), showing what a PTX file looks like and explaining it in detail.

For this post, I will use the OpenCL PTX file, in my next post I will show the differences between PTX files of OpenCL and CUDA code. Let's start by looking at the complete PTX code:

<pre>
//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sun May 18 04:44:51 2014 (1400399091)
// Driver 331.79
//

.version 3.0
.target sm_21, texmode_independent
.address_size 32


.entry add_vectors(
    .param .u32 .ptr .global .align 4 add_vectors_param_0,
    .param .u32 .ptr .global .align 4 add_vectors_param_1,
    .param .u32 .ptr .global .align 4 add_vectors_param_2,
    .param .u32 add_vectors_param_3
)
{
    .reg .pred      %p<2>;
    .reg .s32       %r<21>;


    ld.param.u32    %r9, [add_vectors_param_3];
    mov.u32         %r5, %envreg3;
    mov.u32         %r6, %ntid.x;
    mov.u32         %r7, %ctaid.x;
    mov.u32         %r8, %tid.x;
    add.s32         %r10, %r8, %r5;
    mad.lo.s32      %r4, %r7, %r6, %r10;
    setp.lt.s32     %p1, %r4, %r9;
    @%p1 bra        BB0_2;

    ret;

BB0_2:
    shl.b32         %r11, %r4, 2;
    ld.param.u32    %r18, [add_vectors_param_0];
    add.s32         %r12, %r18, %r11;
    ld.param.u32    %r19, [add_vectors_param_1];
    add.s32         %r13, %r19, %r11;
    ld.global.u32   %r14, [%r13];
    ld.global.u32   %r15, [%r12];
    add.s32         %r16, %r14, %r15;
    ld.param.u32    %r20, [add_vectors_param_2];
    add.s32         %r17, %r20, %r11;
    st.global.u32   [%r17], %r16;
    ret;
}
</pre>

The file starts with a header showing some compiler information in comments, followed by three lines containing the PTX ISA version, target architecture (compute capability) and addressing that will used.

<pre>
//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sun May 18 04:44:51 2014 (1400399091)
// Driver 331.79
//

.version 3.0
.target sm_21, texmode_independent
.address_size 32
</pre>

Immediately after the header, the entry point of the kernel is found, starting with the <code>.entry</code> directive followed by the kernel name.

<pre>
.entry add_vectors(
...
)
</pre>

The parameters for this kernel are three pointers to 32-bit integers global memory arrays (notice that there is no distinction from constant to non-constant pointers and memory!), and one 32-bit integer argument (the length of input and output vectors).

Each parameter of a function starts with the <code>.param</code> directive, followed by its data type, <code>.u32</code> in our case, which in the example are all compiled into 32-bit unsigned integers (for extended information on data types, refer to <a href="http://docs.nvidia.com/cuda/parallel-thread-execution/">PTX ISA documentation</a>). The pointer arguments are followed by a <code>.ptr</code> directive, the address space where they are located (<code>.global</code> for global space) and finally the data alignment, which is <code>.align 4</code> for 4 bytes alignment.

<pre>
    .param .u32 .ptr .global .align 4 add_vectors_param_0,
    .param .u32 .ptr .global .align 4 add_vectors_param_1,
    .param .u32 .ptr .global .align 4 add_vectors_param_2,
    .param .u32 add_vectors_param_3
</pre>

Immediately after the function prototype, the code block begins. First thing in the block are register definitions, they always start with <code>.reg</code> directive, followed by the data type and its name (note that references to registers are always preceded by the <code>%</code> sign). Registers with the <code>.pred</code> data type are used for conditional assignments, to control a branch instruction, for example. The PTX format also supports virtual registers, and a set of registers can be defined by including the number N of virtual registers inside <code>< ></code>, with them being addressable as <code>r0, r1, ..., rN-1</code>, when <code>r</code> is the name given to the set of registers.

<pre>
    .reg .pred      %p<2>;
    .reg .s32       %r<21>;
</pre>

Following the register definitions, we see the <code>ld.param</code> directive, which has the ability to copy a function parameter into a register. Because the second argument is an address, it must be inside brackets, as in <code>[add_vectors_param_3]</code>. This copy to a register is necessary as most of PTX instructions can not handle function parameters directly. See that all data related directives end with the data type being processed.

<pre>
    ld.param.u32    %r9, [add_vectors_param_3];
</pre>

Later, the values of some special registers are copied to GPU registers, they are:
<ul>
<li><code>envreg3</code>: read-only driver-defined special register;</li>
<li><code>ntid.x</code>: number of threads per CTA, Cooperative Thread Array, in x dimension, equivalent to get_local_size(0);</li>
<li><code>ctaid.x</code>: identifier of CTA within a grid, equivalent to get_group_id(0);</li>
<li><code>tid.x</code>: thread identifier in the x dimension of CTA, equivalent to get_local_id(0).</li>
</ul>

<pre>
    mov.u32         %r5, %envreg3;
    mov.u32         %r6, %ntid.x;
    mov.u32         %r7, %ctaid.x;
    mov.u32         %r8, %tid.x;
</pre>

After copying these registers into locally defined registers, the real fun starts. The values of <code>envreg3</code> and <code>tid.x</code> are added and stored in register <code>r10</code>, since <code>envreg3</code> is not directly part of the code, but instead a driver-specific computation, I will not go into more details about it (note that <code>envreg</code> does not appear in PTX of the equivalent CUDA kernel).

<pre>
    add.s32         %r10, %r8, %r5;
</pre>

With the <code>mad</code> instruction (short for multiply-add), the global thread ID is calculated (get_global_id(0)), which is done by multiplying second and third arguments (<code>r7</code> and <code>r6</code>, holding values of <code>ntid.x</code> and <code>ctaid.x</code>), with the lower 16 bits (identified by the <code>.lo</code> suffix of <code>mad</code> instruction) of the result being later added to the fourth argument (<code>r10</code>, that holds <code>tid.x</code>) and stored back to the first argument (<code>r4</code>).

<pre>
    mad.lo.s32      %r4, %r7, %r6, %r10;
</pre>

The first conditional case is then reached (the only one in this example), which is composed by the <code>setp</code> instruction, followed by <code>@</code> and <code>bra</code> instructions. The <code>setp</code> instruction sets a predicate (first argument) by comparing second argument (in our case, the global thread ID) to the third argument (number of elements of input/output vectors) using the condition stated after <code>setp</code>, in this case <code>lt</code> (lower than). The <code>@</code> instruction then executes the predicate, if true, branches (by executing <code>bra</code>) to target <code>BB0_2</code> (target must be a label or register pointing to the address of a label) and continues execution from there, if false, just continues execution without branching. In the case of the condition being false, our kernel would simply return by calling the <code>ret</code> instruction that comes immediately after.

<pre>
    setp.lt.s32     %p1, %r4, %r9;
    @%p1 bra        BB0_2;

    ret;
</pre>

The real data processing is done by the set of instructions under label <code>BB0_2</code>, which calculates data pointers, loads data from global memory, performs calculations and stores results back to global memory. Since we are dealing with 32-bit data pointers, the first instruction <code>shl</code> takes 32-bit untyped binary registers and shift bits left by 2 positions (equivalent to an integer multiplication by 4), this is in fact calculating the zero-based position of data according to the global thread ID. The result in <code>r11</code> is then added to the addresses pointed by the global memory pointers (after loading them to registers by instruction <code>ld.param</code>), resulting in the addresses where the data that needs to be processed by the thread lies in memory.

<pre>
BB0_2:
    shl.b32         %r11, %r4, 2;
    ld.param.u32    %r18, [add_vectors_param_0];
    add.s32         %r12, %r18, %r11;
    ld.param.u32    %r19, [add_vectors_param_1];
    add.s32         %r13, %r19, %r11;
</pre>

Data is loaded from global memory with instruction <code>ld.global</code>, the second parameter is a register that holds the pointer to a memory address and must be within brackets, just as we have seen before with <code>ld.param</code>. The actual data read from memory is then stored into the register passed as the first argument. After the resulting value is calculated by adding both values read from memory, and memory address where result will be stored is calculated, it is stored in global memory with complimentary instruction to <code>ld.global</code>, that is <code>st.global</code>. In the storage to memory case, the first argument is the address, that is the argument inside brackets. Finally, the return instruction <code>ret</code> is called and the processing is finished.

<pre>
    ld.global.u32   %r14, [%r13];
    ld.global.u32   %r15, [%r12];
    add.s32         %r16, %r14, %r15;
    ld.param.u32    %r20, [add_vectors_param_2];
    add.s32         %r17, %r20, %r11;
    st.global.u32   [%r17], %r16;
    ret;
</pre>

As the PTX might not be as straightforward to understand as the C/C++ code, in my opinion is still easier to understand than pure assembly for most architectures, as we do not have to handle lots of interrupts, for example. Understanding PTX code can significantly help you understand what the compiler is actually doing with you C/C++ code, and how the program will flow through hardware. In some cases, this can be an important tool to put your GPU to full power.

