How to access/use the E16/64 from the R prompt - Part 2

Moderator: censix

How to access/use the E16/64 from the R prompt - Part 2

Postby censix » Thu Jun 13, 2013 8:19 pm

.... continuing from part 1. We are still running R as root !!

We can finally, finally run a bit of proper OpenCL code on the E16/64. The following example was kindly provided by Willem Ligtenberg, (http://openanalytics.eu) who is the author of the "ROpenCL" package !!! We could not have achievd all of this progress in such a short time without it. The package was actually designed with other types of GPU in mind, but we were able to use it for our purposes with very few changes. ;)

The example is a toy vector addition that will use all 16/64 Epiphany cores. Feel free to experiment with longer vectors ...

Paste the following to the R prompt, to create the kernel:

Code: Select all
    require(ROpenCL)   
    p <- getPlatformIDs()
    d <- getDeviceIDs(p[[1]])
    context <- createContext(d[[2]])  # d[[2]] is Epiphany!
    queue <- createCommandQueue(context, d[[2]])  # d[[2]] is Epiphany!
       
    a <- seq(256)/10
    b <- seq(256)
    out <- rep(0.0, length(a))
     
    localWorkSize = 16
    globalWorkSize = ceiling(length(a)/localWorkSize)*localWorkSize

    inputBuf1 <- createBuffer(context, "CL_MEM_READ_ONLY", globalWorkSize, a)
    inputBuf2 <- createBuffer(context, "CL_MEM_READ_ONLY", globalWorkSize, b)
    outputBuf1 <- createBufferFloatVector(context, "CL_MEM_WRITE_ONLY", globalWorkSize)

    kernel <- "__kernel void VectorAdd(__global const float* a, __global const int* b, __global float* c, int iNumElements)
    {
        // get index into global data array
        int iGID = get_global_id(0);

        // bound check (equivalent to the limit on a 'for' loop for standard/serial C code
        if (iGID >= iNumElements)
        {   
            return;
        }
        // add the vector elements
        c[iGID] = a[iGID] + b[iGID];
    }"

    kernel <- createProgram(context, kernel, "VectorAdd", inputBuf1, inputBuf2, outputBuf1, length(out))


Paste this to write the two input vectors to the device==E16/64 and to launch the compute kernel.
Code: Select all
enqueueWriteBuffer(queue, inputBuf1, globalWorkSize, a)
enqueueWriteBuffer(queue, inputBuf2, globalWorkSize, b)
enqueueNDRangeKernel(queue, kernel, globalWorkSize, localWorkSize)


That will produce the following barrage of messages...

[2019] clmesg info: e32pth_engine_needham.c(100): e32_engine_startup: engine is path-through
(0,0) 4 4 4 4 4 reg: config=0x0 status=0x0 pc=0x0 imask=0x0 ipend=0x0
(0,1) 4 4 4 4 4 reg: config=0x0 status=0x0 pc=0x0 imask=0x0 ipend=0x0
(0,2) 4 4 4 4 4 reg: config=0x0 status=0x0 pc=0x0 imask=0x0 ipend=0x0
(0,3) 4 4 4 4 4 reg: config=0x0 status=0x0 pc=0x0 imask=0x0 ipend=0x0
(1,0) 4 4 4 4 4 reg: config=0x0 status=0x0 pc=0x0 imask=0x0 ipend=0x0
(1,1) 4 4 4 4 4 reg: config=0x0 status=0x0 pc=0x0 imask=0x0 ipend=0x0
(1,2) 4 4 4 4 4 reg: config=0x0 status=0x0 pc=0x0 imask=0x0 ipend=0x0
....


Now we collect the result of the vector addition and print it.
Code: Select all
result <- enqueueReadBuffer(queue, outputBuf1, globalWorkSize, out)
result

To verify correctness, the result of the previous section should be the same as a+b:
Code: Select all
a+b


To get an idea about the size of the error introduced by using float instead of double
Code: Select all
sum(result-(a+b))


Worked fine for me!!

Soren

http://censix.com
censix
 
Posts: 49
Joined: Sun Dec 16, 2012 7:54 pm
Location: europe

Re: How to access/use the E16/64 from the R prompt - Part 2

Postby censix » Thu Jun 12, 2014 6:46 pm

Important note:

please note that the Epiphany always shows up as the second device in the list of opencl devices! Device no.1 is always the ARM host.
censix
 
Posts: 49
Joined: Sun Dec 16, 2012 7:54 pm
Location: europe


Return to R

Who is online

Users browsing this forum: No registered users and 1 guest