Here at ArrayFire, we develop code to work efficiently on both CUDA and OpenCL devices. Therefore, it is not uncommon that CUDA code on NVIDIA GPUs run faster than OpenCL. A very good way to understand what is behind the curtains is to generate the PTX file for both cases and compare them.

PTX stands for Parallel Thread eXecution [1], which is a low-level virtual machine and instruction set architecture (ISA). For those familiar with assembly language, the PTX instruction set is nothing much more complicated as a single processor assembly code, except that now we are thinking in massively parellel execution.

Retrieving the PTX file from a CUDA kernel is a pretty straightforward process. First, we need a CUDA kernel, we will use a simple vector addition kernel for this article, as shown below:

&lt;pre&gt;__global__ void addVectors(
const float* a,
const float* b,
float *c,
const int n)
{
const int idx = blockIdx.x * blockDim.x + threadIdx.x;

c[idx] = a[idx] + b[idx];
}

After the code is properly saved to a file with the `.cu` extension, here the file name is `add_vectors.cu`, we simply call the NVCC compiler with the `--ptx` argument.

nvcc --ptx add_vectors.cu

The output is a simple text file named `add_vectors.ptx`.

For the OpenCL case, as most people familiar with both technologies would expect, the process to retrieve the PTX file is a bit more complicated, it requires the host-side code and the file will only be extracted at runtime, which means that the host-side code must be compiled.



[1] http://docs.nvidia.com/cuda/parallel-thread-execution