As <a title="ArrayFire function list" href="http://www.sandbox.arrayfire.com/docs/group__arrayfire__func.htm" target="_blank">extensive</a> as ArrayFire is, there are a few cases where you are still working with custom CUDA or OpenCL kernels. You may want to integrate ArrayFire into an existing code base for productivity or you may want to keep around the old implementation for testing purposes. In this post we are going to talk about how you can integrate your custom kernels into ArrayFire in a seamless fashion.
<h3>In and Out of ArrayFire</h3>
First let us look at the code and then break it down bit by bit.
<pre>int main()
{
    af::array x = af::randu(num, 1);

    af::array y = af::array(num, 1);

    float *d_x = x.device&lt;float&gt;();
    float *d_y = y.device&lt;float&gt;();

    af::sync();

    launch_simple_kernel(d_y, d_x, num);
    x.unlock();
    y.unlock();

    float err = af::sum(y - 1);
    printf("Error: %f\n", err);
    return 0;
}
</pre>
Firstly you need an input to the kernel.
<pre>// Generate input data
af::array x = af::randu(num, 1);
</pre>
In this case, it is the variable <strong> x </strong> that has been created as a uniformly random array. If you have data coming from other computations, you can use that instead.

The necessary memory required for the output needs to be allocated before we launch the kernel.
<pre>// Create empty array with required dimensions
af::array y = af::array(num, 1);
</pre>
In this sample, the output y has the same number of elements as <strong> x </strong>. Note that we are not specifying the output data type. In such cases ArrayFire assumes the data type is single precision floating point (<a title="ArrayFire data types" href="http://www.sandbox.arrayfire.com/docs/namespaceaf.htm#acfe99d230e216901bd782cc580e4e815"><strong> af::f32 </strong></a>). You can add the necessary data type at the end of the <a title="array constructors" href="http://www.sandbox.arrayfire.com/docs/classaf_1_1array.htm"><strong> array(..)</strong></a> constructor.

Once you have the input and output arrays, you will need to extract the device pointers / objects using <a title="Getting device pointers from ArrayFire" href="http://www.sandbox.arrayfire.com/docs/group__array__func__cpp__device.htm">array::device&lt;float&gt;()</a> method in the following manner.
<pre> // Get device pointers
 float *d_x = x.device&lt;float&gt;();
 float *d_y = y.device&lt;float&gt;();
</pre>
Before you launch your custom kernel, it would be best to make sure all the computations performed by ArrayFire have finished. This can be called by using <strong><a href="http://www.sandbox.arrayfire.com/docs/group__array__func__cpp__sync.htm"> af::sync()</a>; </strong>
<pre>// Finish the tasks arrayfire was doing
af::sync();

// Launch kernel to do the following operations
// y = sin(x)^2 + con(x)^2
launch_simple_kernel(d_y, d_x, num);
</pre>
The function <strong> launch_simple_kernel </strong> will handle the launching of your custom kernel. We will have a look at how to do this in CUDA and OpenCL later in the post.

Once you have finished your computations, you have to tell ArrayFire to take control of the memory objects.
<pre>// Tell arrayfire it controls the pointers again
x.unlock();
y.unlock();
</pre>
This is a very crucial step because if ArrayFire believes the user still controls the pointer, it does not perform garbage collection on those arrays and will potentially lead to memory leaks.

You can now proceed with the rest of the program. In our particular example, we are just performing an error check and exiting.
<pre>// Check for errors
// sin(x)^ + cos(x)^2 == 1
// The following should print 0
float err = af::sum(y - 1);
printf("Error: %f\n", err);
</pre>
One of the benefits of ArrayFire is that it enables you to use the same code base for both CUDA and OpenCL devices. In the same vein, this particular example can be used to extend using a CUDA kernel or an OpenCL kernel. The externally defined <strong>launch_simple_kernel</strong> will be the only part of your code base that will change when moving from CUDA to OpenCL (or vice versa!).
<h3>Launching a CUDA kernel</h3>
If you are integrating a CUDA kernel into your ArrayFire code base, it is a fairly straight forward process. You need to set the launch configuration parameters, launch the kernel and wait for the computations to finish. You can see how to do that using a simple kernel below.
<pre>__global__
static void simple_kernel(float *d_y,
                          const float *d_x,
                          const int num)
{
    const int id = blockIdx.x * blockDim.x + threadIdx.x;

    if (id &lt; num) {
        float x = d_x[id];
        float sin_x = sin(x);
        float cos_x = cos(x);
        d_y[id] = (sin_x * sin_x) + (cos_x * cos_x);
    }
}

void inline launch_simple_kernel(float *d_y,
                                 const float *d_x,
                                 const int num)
{
    // Set launch configuration
    const int threads = 256;
    const int blocks = (num / threads) + ((num % threads) ? 1 : 0);
    simple_kernel&lt;&lt;&lt;blocks, threads&gt;&gt;&gt;(d_y, d_x, num);
    // Synchronize and check for error
    cudaDeviceSynchronize();
}
</pre>
<h3>Launching an OpenCL kernel</h3>
If you are integrating an OpenCL kernel into your ArrayFire code base, launching a kernel bit more complicated.
<pre>void inline launch_simple_kernel(float *d_y,
                                 const float *d_x,
                                 const int num)
{
    std::string simple_kernel = get_kernel_string("simple.cl");

    cl_context context = get_context((cl_mem)d_x);
    cl_program program = build_program(context, simple_kernel);
    cl_kernel   kernel = create_kernel(program, "simple_kernel");
    cl_command_queue queue = create_queue(context);

    cl_int err = CL_SUCCESS;
    int arg = 0;

    err |= clSetKernelArg(kernel, arg++, sizeof(cl_mem), &amp;d_y);
    err |= clSetKernelArg(kernel, arg++, sizeof(cl_mem), &amp;d_x);
    err |= clSetKernelArg(kernel, arg++, sizeof(int   ), &amp;num);

    size_t local  = 256;
    size_t global = local * (num / local + ((num % local) ? 1 : 0));
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &amp;global, &amp;local, 0, NULL, NULL);
    clFinish(queue);
    return;
}
</pre>
Since ArrayFire uses its own context internally, you need to get the context from the memory object. Once you have access to the same context ArrayFire is using, the rest of the process is exactly the same as launching a stand alone OpenCL context. The code snippet above uses helper functions. You can find the helper functions <a title="OpenCL helper functions." href="https://github.com/arrayfire/arrayfire_examples/blob/master/opencl/helper.h" target="_blank">over here</a>

The OpenCL kernel itself is fairly simple as seen below.
<pre>__kernel
void simple_kernel(__global float *d_y,
                   __global const float *d_x,
                   const int num)
{
    const int id = get_global_id(0);

    if (id &lt; num) {
        float x = d_x[id];
        float sin_x = sin(x);
        float cos_x = cos(x);
        d_y[id] = (sin_x * sin_x) + (cos_x * cos_x);
    }
}
</pre>
You can find the code samples at the following links:
<ul>
	<li><a title="ArrayFire CUDA kernel example" href="https://github.com/arrayfire/arrayfire_examples/tree/master/cuda" target="_blank">CUDA</a></li>
	<li><a title="ArrayFire OpenCL kernel example" href="https://github.com/arrayfire/arrayfire_examples/tree/master/opencl" target="_blank">OpenCL</a></li>
</ul>