Outline

• Simple Spawn Sync Example Java
• Java Spawn Sync Inner Product
• In-Class Exercise
• Java OpenCL bindings
• PRAMs

Introduction

• We have been talking about multithreaded algorithms.
• These were algorithms which run on a Dynamic Multithreading concurrency platform, a concurrency platform with a load balancing scheduler and supported nested parallelism (spawning) and parallel loops (parallel-for).
• We have shown several such algorithms such as parallel Fibonacci computation, matrix vector multiplication, and matrix multiplication.
• We haven't actually done any programming in a language that could be used to implement these algorithms such as Cilk.
• We start today by seeing how we can fake the spawn sync mechanism using Java. Recall, if we can do spawn-sync, we can do parallel for.

A Simple Multithreaded Program

Let A =[0, 0] be a global array of length 2

A = [0,0]
SpawnSync(true, 0)

SpawnSync(parent, location):
1 if(parent):
      spawn SpawnSync(false, location + 1)
2 A[location] = location + 1
3 sync 
4 if(parent):
     for i =0 to A.length - 1: print A[i]

The above code should output:

1
2

Java Threads

• To simulate this in Java we make use of Threads.
• To use a Java thread, you extend the class Thread, say with a class MyThread. You then override in MyThread, Thread's run() method with the code you want to run independent of the main thread.
• The Java runtime typically maps threads to operating systems threads and the OS determines whose run method should execute on which processor and when.
• To use MyThread, one could make a new object, myThread, of type MyThread, then call its start() method to make its run() method available for scheduling.
• Threads have two main states: runnable and blocked. If it is runnable, then the scheduler can schedule it or has scheduled its run method to execute. If it is blocked, then run method won't execute until the block condition is removed.
• Threads can be blocked if they are waiting to be notified by another object or if their sleep() method is called.
• To get a thread to wait on another object we get the thread to call a synchronized method/block of that other object and then call wait() in that method or block.
• Only when that other object calls notify() or its thread's run method completes does the original thread become runnable again.
• Let's use all this now to make a Java implement of our SpawnSync code.

A Simple Spawn Sync Example Java

import java.lang.Thread;

public class SpawnSyncDemo extends Thread
{
    public SpawnSyncDemo(SpawnSyncDemo spawner, int location)
    {
        this.location = location;
        this.spawner = spawner;
        done = false;
    }
    public void run()
    {
        SpawnSyncDemo child = null;
        done = false;
        if(spawner == null)
        {
            child = new SpawnSyncDemo(this, location + 1);
            child.start();
        }
        a[location] = location + 1;
        done = true;
        if(child != null)
        {
            child.sync();
        }
        if(spawner == null)
        {
            for(int i = 0; i < a.length; i++)
            {
                System.out.println(a[i]);
            }
        }
    }
    public synchronized void sync()
    {
        if(!done)
        {
            try
            {
                wait(); /*parent executes this code
                          and waits until child thread
                          completes
                         */
            }
            catch(InterruptedException ie)
            {
                ie.printStackTrace();
            }
        }
    }
    public static void main(String args[])
    {
        SpawnSyncDemo parent = new SpawnSyncDemo(null, 0);
        parent.start();
    }
    int location;
    SpawnSyncDemo spawner;
    boolean done;
    static int a[] = {0, 0};
}

Thread Sum Demo

• The example below shows how we can use this basic spawn sync mechanism with recursion to compute the sum of an array.
• This technique would allow us to simulate parallel for loops.

import java.lang.Thread;

public class SumDemo extends Thread
{
    public SumDemo(SumDemo spawner, int low, int high)
    {
        this.high = high;
        this.low = low;
        this.spawner = spawner;
        done = false;
        c = 0;
    }
    public void run()
    {
        SumDemo firstHalf = null;
        SumDemo secondHalf = null;
        done = false;
        if(low == high)
        {
            c = a[low];
        }
        else
        {
            int mid = (low + high) / 2;
            firstHalf = new SumDemo(this, low, mid);
            firstHalf.start();
            secondHalf = new SumDemo(this, mid + 1, high);
            secondHalf.start();
        }
        if(low != high ){
            firstHalf.sync();
            secondHalf.sync();
            c = firstHalf.c + secondHalf.c;
        }
        done = true;
        if(spawner == null)
        {
            System.out.println(c);
        }
    }
    public synchronized void sync()
    {
        if(!done)
        {
            try
            {
                wait(); 
            }
            catch(InterruptedException ie)
            {
                ie.printStackTrace();
            }
        }
    }
    public static void main(String args[])
    {
        SumDemo parent = new SumDemo(null, 0, a.length - 1);
        parent.start();
    }
    int low;
    int high;
    int c = 0;
    SumDemo spawner;
    boolean done;
    static int a[] = {4, 3, 9, -1, 1};

}

In-Class Exercise

• Modify the code of the last slide so instead of computing a sum, it computes the minimum of the values in a[].
• Post your solutions to the Feb 20 In-Class Exercise Thread.

OpenCL

• At some point in the mid 2000s it was realized that the processor units that make up the graphics processing units in modern computers were capable enough to do serious computing, not just shaders.
• GPUs typically have several hundred or thousands of processing units rather than the 1,2,4, or 8 of a typical CPU.
• Apple in 2006 developed OpenCL as a common library to interface from code running on your CPU to code that runs on your GPU.
• Since then the Khronos Group with input from AMD, IBM, Qualcomm, Intel, and Nvidia has been developing this library.
• In 2018, Apple deprecated OpenCL in favor of its Metal technology.
• Other vendors are moving to the cross-platform Vulkan technology, however, as yet, there is a dearth of good simple examples for this, so I am still doing things with OpenCL.
• Various languages have bindings to OpenCL. Java's binding is called JOCL.
• It can be downloaded from JOCL at Maven Central and it consists of a jar file you can use in your classpath.
• For example, you might use:

  javac -classpath "./jocl-2.0.1.jar" MyJOCL.java
  

  to compile such a program, and use:

  java -classpath ".:./jocl-2.0.1.jar" MyJOCL
  

  to run it.
• The next slide shows a simple example of code using this library.

Java Open CL Example

• In the code below, the part that runs on the GPU is stored in the Java String programSource.
• OpenCL programs are essentially C programs with a few restrictions and some additional keywords and special functions.
• For example, the new keyword __global says that the variable is global to all processors
• The function get_global_id(0) gets the id of the current processor.
• To use the GPU code we need to compile it and specify bindings between variables in GPU land and those in CPU land.
• Then we can execute it on inputs from the CPU and read the output compute back from the GPU to the CPU.

/*
 * JOCL - Java bindings for OpenCL
 * 
 * Copyright 2009 Marco Hutter - http://www.jocl.org/
 */


import static org.jocl.CL.*;

import org.jocl.*;

/**
 * A small JOCL sample.
 */
public class JOCLSample
{
    /**
     * The source code of the OpenCL program to execute
     */
    private static String programSource =
        "__kernel void "+
        "sampleKernel(__global const float *a,"+
        "             __global const float *b,"+
        "             __global float *c)"+
        "{"+
        "    int gid = get_global_id(0);"+
        "    c[gid] = a[gid] * b[gid];"+
        "}";
    

    /**
     * The entry point of this sample
     * 
     * @param args Not used
     */
    public static void main(String args[])
    {
        // Create input- and output data 
        int n = 10;
        float srcArrayA[] = new float[n];
        float srcArrayB[] = new float[n];
        float dstArray[] = new float[n];
        for (int i=0; i<n; i++)
        {
            srcArrayA[i] = i;
            srcArrayB[i] = i;
        }
        Pointer srcA = Pointer.to(srcArrayA);
        Pointer srcB = Pointer.to(srcArrayB);
        Pointer dst = Pointer.to(dstArray);

        // The platform, device type and device number
        // that will be used
        final int platformIndex = 0;
        final long deviceType = CL_DEVICE_TYPE_ALL;
        final int deviceIndex = 0;

        // Enable exceptions and subsequently omit error checks in this sample
        CL.setExceptionsEnabled(true);

        // Obtain the number of platforms
        int numPlatformsArray[] = new int[1];
        clGetPlatformIDs(0, null, numPlatformsArray);
        int numPlatforms = numPlatformsArray[0];

        // Obtain a platform ID
        cl_platform_id platforms[] = new cl_platform_id[numPlatforms];
        clGetPlatformIDs(platforms.length, platforms, null);
        cl_platform_id platform = platforms[platformIndex];

        // Initialize the context properties
        cl_context_properties contextProperties = new cl_context_properties();
        contextProperties.addProperty(CL_CONTEXT_PLATFORM, platform);
        
        // Obtain the number of devices for the platform
        int numDevicesArray[] = new int[1];
        clGetDeviceIDs(platform, deviceType, 0, null, numDevicesArray);
        int numDevices = numDevicesArray[0];
        
        // Obtain a device ID 
        cl_device_id devices[] = new cl_device_id[numDevices];
        clGetDeviceIDs(platform, deviceType, numDevices, devices, null);
        cl_device_id device = devices[deviceIndex];

        // Create a context for the selected device
        cl_context context = clCreateContext(
            contextProperties, 1, new cl_device_id[]{device}, 
            null, null, null);
        
        /* Create a command-queue for the selected device
           This syntax is not deprecated in OpenCL 2 in favor of
           cl_queue_properties cmd_props = new cl_queue_properties();
           int[] errors = new int[1];
           cl_command_queue commandQueue =
               clCreateCommandQueueWithProperties(context, device, cmd_props,
               errors);
            However, the new syntax requires native compilation I have not done
         */
        cl_command_queue commandQueue = 
            clCreateCommandQueue(context, device, 0, null);

        // Allocate the memory objects for the input- and output data
        cl_mem memObjects[] = new cl_mem[3];
        memObjects[0] = clCreateBuffer(context, 
            CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
            Sizeof.cl_float * n, srcA, null);
        memObjects[1] = clCreateBuffer(context, 
            CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
            Sizeof.cl_float * n, srcB, null);
        memObjects[2] = clCreateBuffer(context, 
            CL_MEM_READ_WRITE, 
            Sizeof.cl_float * n, null, null);
        
        // Create the program from the source code
        cl_program program = clCreateProgramWithSource(context,
            1, new String[]{ programSource }, null, null);
        
        // Build the program
        clBuildProgram(program, 0, null, null, null, null);
        
        // Create the kernel
        cl_kernel kernel = clCreateKernel(program, "sampleKernel", null);
        
        // Set the arguments for the kernel
        clSetKernelArg(kernel, 0, 
            Sizeof.cl_mem, Pointer.to(memObjects[0]));
        clSetKernelArg(kernel, 1, 
            Sizeof.cl_mem, Pointer.to(memObjects[1]));
        clSetKernelArg(kernel, 2, 
            Sizeof.cl_mem, Pointer.to(memObjects[2]));
        
        // Set the work-item dimensions
        long global_work_size[] = new long[]{n};
        long local_work_size[] = new long[]{1};
        
        // Execute the kernel
        clEnqueueNDRangeKernel(commandQueue, kernel, 1, null,
            global_work_size, local_work_size, 0, null, null);
        
        // Read the output data
        clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE, 0,
            n * Sizeof.cl_float, dst, 0, null, null);
        
        // Release kernel, program, and memory objects
        clReleaseMemObject(memObjects[0]);
        clReleaseMemObject(memObjects[1]);
        clReleaseMemObject(memObjects[2]);
        clReleaseKernel(kernel);
        clReleaseProgram(program);
        clReleaseCommandQueue(commandQueue);
        clReleaseContext(context);
        
        // Verify the result
        boolean passed = true;
        final float epsilon = 1e-7f;
        for (int i=0; i<n; i++)
        {
            float x = dstArray[i];
            float y = srcArrayA[i] * srcArrayB[i];
            boolean epsilonEqual = Math.abs(x - y) <= epsilon * Math.abs(x);
            if (!epsilonEqual)
            {
                passed = false;
                break;
            }
        }
        System.out.println("Test "+(passed?"PASSED":"FAILED"));
        if (n <= 10)
        {
            System.out.println("Result: "+java.util.Arrays.toString(dstArray));
        }
    }
}

OpenCL Compute Units and Processing Elements

• In OpenCL, a function that executes on the GPU is called a kernel.
• In OpenCL you can determine which compute devices like CPUs and GPUs are available.
• Such devices in turn can have many compute units (can process different instructions in parallel) which are in turn made up of processing elements (which process the same instruction, potentially on different data, in parallel like a SIMD).
• When we enqueue our kernel to the GPU using a command like clEnqueueNDRangeKernel we specify the global work size and local work size.
• These can be thought of as the total number of virtual processing elements (work items) and how many of these items should be fit into a virtual compute unit (work group).
• The OpenCL scheduler does the scheduling and managing of GPU/CPU resources from the selected devices for a job to then "fake" running our kernels on the requested virtual compute devices.

OpenCL Kinds of Memory and Barriers

• In addition to the __global keyword of the last slide, we can declare a kernel variable to be __local, if it is to be read/writable among a compute units worth of processing elements.
• If we want a variable only visible to a processing element, we can declare it __private.
• Finally, GPU processing element read-only, CPU writable memory can be declared using the __constant keyword.
• Although it is up to the scheduler to determine how a work item is processed, you can tell the scheduler when memory within a work group needs to be synchronized by issuing a barrier(CLK_LOCAL_MEM_FENCE) command in your kernel.
• Upon seeing such a command the scheduler won't continue processing that work item until all other work items in the work group have also reached that command.
• If one wants similar kind of synchronization but among all processing elements, one can call barrier(CLK_GLOBAL_MEM_FENCE).