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 instruction 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> and show what a PTX file looks like and also explain 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 size being 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 one 32-bit integer argument (the length of our vectors). Each parameter start 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 arguments which are pointers are then 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 (note that references to registers are always preceded by the <code>%</code> sign). Registers with the data type defined by <code>.pred</code> directive are used for conditional assignments, to control a branch instruction, for example.

<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, but the second argument is an address, so 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 cannot handle function parameters directly. See that all data related directives end with the data type that is being processed.

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

Later, the values of some special registers are copied to the registers previously defined, 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 a driver-specific computation, I won't go into more details about it (note that <code>envreg</code> does not appear in PTX of the CUDA code for this 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, 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 vector) using the condition the condition stated after <code>setp</code>, in this case being <code>lt</code> (lower than). The <code>@</code> instruction then executes the predicate, if true, branches (by executing <code>bra</code>) to target and continue execution (target must be a label, or register pointing to address of a label), if false, just continues execution. In the case of the condition being false, our kernel would simply return by calling the <code>ret</code> instruction.

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

    ret;
</pre>

The real data processing is done by instructions under label BB0_2, which calculates data pointers, loads data from global memory, perform calculations and stores it 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, which is then added to the address pointed by the global memory pointers (after loading them to registers by instruction <code>ld.param</code> to find the where the data that will 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 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 the stored to the first argument, a register. 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 instruction <code>st.global</code>, in this case the first argument is the address, so that is the one 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, it is still not that complicated, and it is very useful to understand it when you want to extract the most out of a GPU.