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

In this post I will use the same vector addition kernel that was used before and show some peculiarities for each PTX file, let's start by looking at the PTX for the OpenCL case:

<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];
    // inline asm
    mov.u32     %r5, %envreg3;
    // inline asm
    // inline asm
    mov.u32     %r6, %ntid.x;
    // inline asm
    // inline asm
    mov.u32     %r7, %ctaid.x;
    // inline asm
    // inline asm
    mov.u32     %r8, %tid.x;
    // inline asm
    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>

It 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.

Immediately after the header, the entry point of the kernel is found, starting with the <code>.entry</code> directive followed by the kernel name. The parameters for these kernels are three pointers to 32-bit integers to global memory (notice that there is no distiction 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. 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.

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 jump instruction, for example.

Following the register definitions, we see the <code>ld.param</code> directive, which has the ability to copy a function parameter into a register. This copy to a register is necessary as most of PTX instructions cannot handle function parameters directly. Note that all data related directives end with the data type that is being processed.

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 in x dimension, Cooperative Thread Array, 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, equilalent to get_local_id(0))</li>
</ul>

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 this is not directly part of the code but a driver-specific computation, I won't go into more details about it (note that this does not exist in PTX of the CUDA code). With the <code>mad</code> instruction (short for multiply-add), the global thread ID is calculate (get_global_id(0)), which is done by multiplying second and third arguments, with the lower bits (identified by the <code>.lo</code> suffix of <code>mad</code> instruction) being later added to the fourth argument and storing back to the first argument.
