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 <a href="http://docs.nvidia.com/cuda/parallel-thread-execution" title="Parallel Thread eXecution">Parallel Thread eXecution</a>, 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:

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

    if (idx >= n) return;

    c[idx] = a[idx] + b[idx];
}
</pre>

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

<pre>
nvcc --ptx add_vectors.cu
</pre>

The output is a simple text file named <code>add_vectors.ptx</code>.

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.

The OpenCL kernel will be very similar to the CUDA kernel, but not exactly equal, below you can see the adapted version of the previous vector addition kernel to OpenCL:

<pre>
__kernel void add_vectors(
    __global const int* a,
    __global const int* b,
    __global       int* c,
    int n)
{
    int idx = get_global_id(0);

    if (idx >= n) return;

    c[idx] = a[idx] + b[idx];
}
</pre>

Finally, we need the host-side code. A minimalist version (without actually executing the OpenCL kernel):

<pre>
#include <stdio.h>
#include <stdlib.h>

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

int main(int argc, char* argv[])
{
    // Loads add_vectors.cl
    FILE* fp;
    fp = fopen("add_vectors.cl", "r");
    if (!fp) {
        fprintf(stderr, "Error loading kernel.\n");
        exit(1);
    }

    fseek(fp, 0, SEEK_END);
    size_t kernel_sz = ftell(fp);
    rewind(fp);

    char* kernel_str = (char*)malloc(kernel_sz);
    fread(kernel_str, 1, kernel_sz, fp);
    fclose(fp);

    cl_platform_id platform;
    cl_device_id device;
    cl_uint num_devices, num_platforms;
    cl_int err = clGetPlatformIDs(1, &platform, &num_platforms);
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1,
                         &device, &num_devices);

    cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);

    cl_command_queue command_queue = clCreateCommandQueue(context, device, 0, &err);

    cl_program program = clCreateProgramWithSource(context, 1,
            (const char **)&kernel_str, (const size_t *)&kernel_sz, &err);

    err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);

    cl_kernel kernel = clCreateKernel(program, "add_vectors", &err);

    size_t binsz;
    err = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binsz, NULL);

    unsigned char *bin = (unsigned char *)malloc(binsz);
    err = clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(unsigned char *), &bin, NULL);

    fp = fopen("add_vectors_ocl.ptx", "wb");
    fwrite(bn, sizeof(char), bin, fp); // Save the binary, but my file stay empty  
    fclose(fp);

    free(bn);

    err = clFlush(command_queue);
    err = clFinish(command_queue);
    err = clReleaseKernel(kernel);
    err = clReleaseProgram(program);
    err = clReleaseCommandQueue(command_queue);
    err = clReleaseContext(context);
    return 0;
}
</pre>