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 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 host-side code and the file will only be extracted at runtime, which means that the host-side code must be compiled first.

The OpenCL kernel will be very similar to the CUDA kernel. Below the vector addition kernel adapted to OpenCL is provided.

<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) follows.

<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(bin, sizeof(char), binsz, 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>

To compile the code above with GCC, it is necessary only to pass the OpenCL include files to the compiler and the OpenCL library to the linker. The code should also compile under Visual Studio with little or no changes.

<pre>
gcc -Wall -g add_vectors.c -o add_vectors -I/usr/local/cuda/include -lOpenCL
</pre>

Finally, the code can be executed and the PTX will be saved to the file <code>add_vectors_ocl.ptx</code>.

If you think the host-side code is too long and complicated, don't worry, check the <a href="http://sandbox.arrayfire.com/quest-for-the-smallest-opencl-program/">Quest for the Smallest OpenCL Program</a>. One of our engineers at ArrayFire wrote that article to explain how to write the possible shortest functional code ever for OpenCL, with few changes, it is possible to use that code to extract the PTX file from OpenCL kernels.