Outline

• Some remarks on JOCL
• PRAMs
• Parallel Complexity Classes
• In-Class Exercise
• Sorting on PRAMs

Introduction

• Last day, we introduced and talked about parallel processing using Java Threads and gave an example of JOCL.
• We'll begin today with a few remarks on common terms used in JOCL programming.
• We then start a new topic: time and space complexity of parallel algorithms.
• To measure this, we will make use of an abstract machine called a PRAM.

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).
• If you want to execute a sequence of consecutive kernels jobs you can always enqueue additional kernels.
• 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 our example program, 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).

The PRAM Model

• We now consider a second model of parallel computation based on synchronous parallel random access machines (PRAMs) (Fortune Wyllie 1978).
• We can think of this as a model for running parallel machine code.
• There are several equivalent ways we could formulate PRAMs, I am fixing below one way that we can do this for this CS255 class.
• The original paper linked above has a FORK instruction and is not as concrete on its instruction list.
• Goldshlager 1982, A Unified Approach to Models of Synchronous Parallel Machines, has a model closer to ours with a slightly different instruction set that also includes multiply, this can be simulated (with some slowdown using the instruction below).
• A parallel processor will consist of p processors, each with its own processor id, and each of which can be a viewed as supporting the RAM model of computation:
  • Each processor has a fixed finite set of local registers 1, ..., m reserved to itself, we'll call accumulators.
  • The `p`-processors share a global memory consisting of `M` registers, that each processor may read from and write to.
  • Each processor gets the same program with each line numbered 1 to n, a program counter (PC) is used to say the current line.
  • Each processor has a finite instruction set supporting direct and indirect addressing, simple branching, and arithmetic instructions. Below are some examples, of the kinds of instructions our RAMs support:
    • Add j, k -- add the value of register j to the accumulator k's value
    • AddAcc j, k -- add the value of accumulator j to the accumulator k's value
    • Half k -- divide accumulator k's value by 2 round down (aka shift left)
    • Halt -- halts execution of current processor.
    • Jump j -- set program counter to j
    • JNeg j, k -- set program counter to j if accumulator k is negative
    • JPos j, k -- set program counter to j if accumulator k is positive
    • JZero j, k -- set program counter to j if accumulator k is zero
    • LoadProcid k -- load the current processor id into accumulator k.
    • Load x,k -- load integer x into accumulator k
    • Mult j,k -- multiply the value of register j from the accumulator k's value
    • MultAcc j,k -- multiply the value of accumulator j from the accumulator k's value
    • Read j, k -- read register j into accumulator k
    • ReadAcc j, k -- read accumulator j into accumulator k
    • Read (j),k -- look up value v in register j, then read register v into accumulator k
    • ReadAcc (k1),k2 -- look up value v in accumulator k1, then read register v into accumulator k2
    • Store j, k -- store value of accumulator k into register j
    • StoreAcc j, k -- store value of accumulator k into raccumulator j
    • Store(j), k -- look up value v in register j, then store in register v value of accumulator k
    • StoreAcc (k1), k2 -- look up value v in accumulator k1, then store in register v value of accumulator k2
    • Sub j,k -- subtract the value of register j from the accumulator k's value
    • SubAcc j,k -- subtract the value of accumulator j from the accumulator k's value
  • Each register and accumulator can hold one integer. I.e., a number like -20, 57, 0, etc.
• At the start of a computation the program counter of each processor is 0, and all accumulators are 0. The input is written as a sequence of integers in the global memory registers.
• A computation ends when all processors halt. The content of the global memory is the output. We say our program accepts an input if the 0th global memory location has a positive value when then computation halts.

More on the PRAM Model

A computation in the PRAM model proceeds as a series of parallel steps:

• In each such step, each processor executes the single instruction given by its program counter. This can involve:
  • Choosing a global memory location to read from
  • Executing the instruction on the operand fetched together with any of the operands in local registers.
  • Finally, writing a value to a single global memory location.
  • Updating the program counter either by adding 1 (default) or according to what the instruction executed says.
• We insist on synchrony: that is, each processor finishes the parallel step i before the parallel step i+1 begins.
• Conflict resolution can be handled in different ways: Exclusive Read/Exclusive Write (EREW), Concurrent Read/Exclusive Write(CREW), and Concurrent Read/Concurrent Write (CRCW).
• Exclusive means only one processor can do that action on a given memory location at a given time. Concurrent is the opposite of this.
• We will mainly be interested in EREW and CREW PRAMs.

Example PRAM Program

0. LoadProcid 0
1. ReadAcc (0),1
2. JNeg 5, 1
3. Load 1, 1
4. Jump 6
5. Load 0, 1
6. StoreAcc (0), 1
7. Halt

• Assume we have `M`-processors and our global memory is also `M`-integers.
• The idea is this code replaces each nonnegative integer in the input with 1 and each negative integer with 0.
• So if the input was (-1,0, 10, 5, -50) the output will be (0, 1, 1, 1, 0).
• Since it can be quite painful to fully spec out machine code for our algorithms we tend to describe our algorithms at a slightly higher level, while remaining careful that any of the things we are saying a PRAM could do, could in principle be written in code as above.

In-Class Exercise

• Suppose we have a PRAM with M processors and 3M global registers.
• Suppose initially M is stored in register 0.
• Write a PRAM program that takes the values in locations i+1 and i+ M +1 and writes their sum in location i+2M+1 for all `0 <= i < M`.
• Please post your solution to the Feb 23 In-Class Exercise Thread.