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. Below we provide the vector addition kernel adapted 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);

    // Query platforms and devices
    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);

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

    // Create OpenCL command queue
    cl_command_queue command_queue = clCreateCommandQueue(context, device, 0, &err);

    // Create OpenCL program for add_vectors.cl
    cl_program program = clCreateProgramWithSource(context, 1,
            (const char **)&kernel_str, (const size_t *)&kernel_sz, &err);

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

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

    // Queries binary (PTX file) size
    size_t binsz;
    err = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binsz, NULL);

    // Reads binary (PTX file) to memory buffer
    unsigned char *bin = (unsigned char *)malloc(binsz);
    err = clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(unsigned char *), &bin, NULL);

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

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