\begin{figure}[t]
\begin{lstlisting}[xleftmargin=0.35cm, xrightmargin=0.25cm]
public class ArraySum {
  static { System.loadLibrary("libCalc"); }
  public static native void openCL_Kernel0(...);
  public static void main(String[] args) {
  ...
  isSafe = ...;
  if (isSafe) {
    /* Safe Region */
    openCL_Kernel0(arrays, result); // JNI call
  } else {
    /* Unsafe Region, run in JVM */
    forall (point [i]:[0:N-1]) {
      result[i] = 0;
      for(int j=0; j<M; j++) {
        result[i] += arraysView[i,j];
      }
    }
  }
}
\end{lstlisting}
\caption{Generated code by the HJ compiler}
\label{fig:transformation}
\end{figure}
\begin{figure*}[t]
\begin{center}
\begin{lstlisting}[xleftmargin=0.35cm, xrightmargin=0.25cm]
cl_context      context;
cl_command_queue command_queue;
JNIEXPORT void  JNICALL Java_ArraySum_initGPU(JNIEnv * env, jclass cls) {
   cl_int          ret;
   /* Get Device */
   clGetPlatformIDs(...);
   clGetDeviceIDs(...);
   /* OpenCL create Context and Queue */
   context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
   command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
}
extern char *HjKernel0; // from Kernel.c
JNIEXPORT void JNICALL Java_ArraySum_openCL_1Kernel0(JNIEnv * env, jclass cls, jintArray arrays, jintArray result, ... ) {
    void *aptr = (env)->GetPrimitiveArrayCritical(arrays, 0);
    void *rptr = (env)->GetPrimitiveArrayCritical(result, 0);
    ...
    /* Create Buffer */
    cl_mem Aobj = clCreateBuffer(context, ...);
    cl_mem Robj = clCreateBuffer(context, ...);
    /* Host to Device Communication */
    clEnqueueWriteBuffer(command_queue, Aobj, ..., aptr, ...); // transfer aptr
    clEnqueueWriteBuffer(command_queue, Robj, ..., rptr, ...); // transfer rptr
    /* Kernel Compilation */
    clCreateProgramWithSource(context, 1, (const char **) &HjKernel0, ...);
    clCreateKernel();
    clSetKernelArg(...);
    /* Invoke Kernel */
    clEnqueueNDRangeKernel(command_queue);
    /* Device to Host Communication */    
    clEnqueueReadBuffer(command_queue); // transfer rptr
    /* Clean up */
    clFinish(command_queue);
    ...
    (env)->ReleasePrimitiveArrayCritical(arrays, aptr, 0);
    (env)->ReleasePrimitiveArrayCritical(result, rptr, 0);
}
\end{lstlisting}
\end{center}
\caption{Generated JNI and OpenCL code (\textit{OpenCL\_hjstub.c})}
\label{fig:jnicode}
\end{figure*}

While HJ-OpenCL provides language features which allow the programmer to
identify parallel loop kernels in their Java code, mapping these 
kernels to OpenCL devices requires compile-time and run-time support. This section introduces the code generation 
techniques used at compile-time to take JVM bytecode as input and produce OpenCL kernels and the required glue 
code (i.e. the JNI stub).

The HJ-OpenCL compiler leverages APARAPI~\cite{APARAPI}, a comprehensive, open-source framework for executing computational kernels from Java applications on 
OpenCL devices. For this work we extended the APARAPI component that generates OpenCL 
code from Java bytecode.  
Our major extension to APARAPI has been to enable static translation
from Java bytecode to OpenCL, whereas the publicly available APARAPI
framework works in a dynamic compilation mode.

In addition to OpenCL kernels, glue code must be automatically generated to transfer execution and 
data from the JVM to the OpenCL device and back. This functionality is provided internally by 
the HJ-OpenCL compiler, and includes the generation of JNI functions, OpenCL API calls, and transformed bytecode.

The rest of this Section is organized as follows:
Section~\ref{subsec:flow} introduces the code generation steps. %compilation flow and runtime environment.
Section~\ref{subsec:barrier} discusses the solution to barrier synchronization support.
Section~\ref{subsec:multidimension} discusses the HJ array view construct which enables
efficient support of multi-dimensional arrays in parallel loops translated to OpenCL.

\subsection{Compilation Flow}\label{subsec:flow}

Figure~\ref{fig:flow} illustrates the compilation flow and how the compiler interacts with the HJ-OpenCL runtime.
The HJ-OpenCL compiler takes an HJ program (see Section~\ref{sec:hj}) as input,
and produces the following outputs:
\begin{enumerate}
\item
Java CLASS files intended for execution in the JVM.
\item
JNI glue code between the JVM and OpenCL kernels, contained in the file \textit{OpenCL\_hjstub.c}.
\item
A Java CLASS file named \textit{OpenCLKernel.class} which contains all the bytecode to be translated to 
OpenCL kernels by the APARAPI bytecode translator.
\end{enumerate}

To explain the compilation work flow, let us again consider the ArraySum example from 
Section~\ref{sec:motivation}.
The ArraySum implementation from Figure~\ref{fig:hjopencl-example} is translated to the Java code shown in 
Figure~\ref{fig:transformation}. During code generation, the HJ-OpenCL compiler identifies the input and output data relationships of 
the \hjforall{} loop, and replaces the original \hjforall{} loop with a native method named \textit{openCL\_Kernel0()} 
(line 9). 
This native method invokes the OpenCL API to execute computations on an OpenCL device. Its definition
is contained in a library, \textit{libCalc} (line 2).
%\textit{System.loadLibrary(``libCalc'')} in line 2 is automatically added to define the native method.
These transformations translate the original HJ program with \hjforall{} constructs to JVM bytecode with JNI stubs for OpenCL kernel invocation.
%These transformations are emitted as bytecode in the CLASS file for the Java class containing the parallel 
%\textit{forall} loop.

A powerful feature of Java is runtime exceptions, thrown when some safety constraints are violated.
For example, the array access \textit{result[i] = 0;} in 
Figure~\ref{fig:hjopencl-example} might throw {\tt IndexOutOfBoundsException} or {\tt NullPointerException}.
Running natively in OpenCL removes all safety guarantees from Java exceptions.
HJ-OpenCL code generation maintains Java exception semantics
by generating two versions of the same \hjforall{} loop: one to run when the programmer
guarantees no exceptions will be thrown, and one to run when the exception-safe conditions do not hold.
The code region to generate these two versions from is specified 
by the {\textit safe} language construct, which takes user-specified conditions as input (see Section~\ref{subsec:safe}).
If these conditions hold at runtime, HJ-OpenCL assumes no exception will be thrown inside the {\textit safe} block.
For example, in Figure~\ref{fig:hjopencl-example} the programmer has set the value of \textit{isSafe}. 
This definition is kept unchanged through to the transformed version in Figure~\ref{fig:transformation}. If the 
exception-safe condition is met, the runtime uses \textit{openCL\_Kernel0()} to launch OpenCL execution.
Otherwise the computation is conservatively run in the JVM on the HJ runtime to preserve Java exception semantics.
These changes are only visible in the emitted bytecode.
In the case that programmers do not provide a \safe{} region around a parallel loop but have provided compiler flags indicating that OpenCL execution is
desired, the HJ-OpenCL compiler and runtime assumes that no exception is thrown.

\textit{OpenCL\_hjstub.c} in Figure~\ref{fig:jnicode} is the auto-generated bridge between the JVM and OpenCL devices. 
It uses JNI to invoke the OpenCL API to %transfer the state variables from the JVM to the OpenCL device address space, 
initialize the OpenCL device, transfer data, compile and launch the OpenCL kernel, and return data and control to 
the JVM.
\textit{OpenCL\_hjstub.c} and \textit{Kernel.c}, which contains the OpenCL kernels translated from HJ \hjforall{} loops, %discussed later in this section) 
are generated by the HJ-OpenCL compiler and built into a library 
which the JVM links to at runtime.

The first function in Figure~\ref{fig:jnicode}, \textit{Java\_ArraySum\_initGPU},
checks for OpenCL devices and creates an OpenCL context and command queue. The HJ-OpenCL compiler inserts an asynchronous 
invocation of this function at the very beginning of the input HJ program to reduce the overhead of OpenCL device 
initialization.

The second function, \textit{Java\_ArraySum\_openCL\_1Kernel0}, performs the following steps to run the 
ArraySum example in OpenCL:

\begin{enumerate}
\item
The JNI call \textit{GetPrimitiveArrayCritical} is used to retrieve a bare pointer to the contents of Java arrays 
(line 14, 15).
\item
The OpenCL API is invoked to transfer scalar and vector inputs from the JVM into the OpenCL address space
(line 17 - 22).
\item
The ArraySum kernel is compiled and executed on an OpenCL device (line 23-28).
\item
The OpenCL API is invoked to retrieve results from the device if they are marked as output (line 29-32).
\item
The JNI call \textit{ReleasePrimitiveArrayCritical} is used to release access to the 
Java arrays after copying the results into them and before returning to JVM execution (line 34, 35).
\end{enumerate}

The paragraphs above have covered the generation of all bytecode and JNI stubs intended for execution in the host system.
To generate the OpenCL kernel code, the APARAPI translator is passed a single CLASS file, \textit{OpenCLKernel.class}.
The HJ-OpenCL compiler generates \textit{OpenCLKernel.class} by extracting the body of parallel \hjforall{} loops 
and translating them into a format which the APARAPI translator can process. Figure~\ref{fig:kernelclass} shows the 
equivalent Java of the bytecode generated in \textit{OpenCLKernel.class}.
Line 1 shows that \textit{OpenCLKernel} extends \textit{com.amd.aparapi.Kernel}, as required by the APARAPI translator.
The body of the original \hjforall{} loop is placed in the \textit{run\(\)} method (lines 6-11).
The APARAPI translator automatically identifies the \textit{run\(\)} method of any class extending
\textit{com.amd.aparapi.Kernel} as the entry point of an OpenCL kernel, and can be used to emit an equivalent OpenCL kernel.
The generated OpenCL kernel is expressed as array of characters(\textit{HjKernel0}) 
which is stored in \textit{Kernel.c} and later compiled with \textit{OpenCL\_hjstub.c}(see line 12, 24 in Figure~\ref{fig:jnicode}).

\begin{figure}
\begin{center}
\begin{lstlisting}[xleftmargin=0.35cm, xrightmargin=0.25cm]
public class OpenCLKernel 
        extends com.amd.aparapi.Kernel {
    int[] result, array;
    int OpEnCl_size;
    public OpenCLKernel(); 
    public void run() {
        int gid = this.get_global_id(0);
        if (gid >= OpEnCl_size) { return; }
        for (int j = 0; j < M; j++) {
          result[gid] += array[gid*M + j];
        }
    }
}
\end{lstlisting}
\end{center}
\caption{Example code of OpenCLKernel.class}
\label{fig:kernelclass}
\end{figure}

\subsection{Barrier Synchronization}\label{subsec:barrier}

HJ-OpenCL also supports all-to-all barrier synchronization across parallel iterations of a \hjforall{} loop.
The programmers can specify synchronization points in a loop using the \textit{next} statement.

Consider the HJ-OpenCL barrier example in Fig~\ref{fig:barrier-example}.
Each iteration of the \hjforall{} loop calls \textit{method1} and \textit{method2}. 
Following execution of \textit{method1(i, j)} or \textit{method2(i, j)} each \hjforall{} iteration is blocked until all other 
iterations also complete the corresponding calls in their own execution streams.
Because OpenCL does not support all-to-all barrier synchronization as a kernel language feature,
the HJ-OpenCL compiler partitions the \hjforall{} loop body into blocks separated by synchronization points.
Rather than performing a single kernel invocation for a single \hjforall{} loop in \textit{OpenCL\_hjstub.c},
the HJ-OpenCL compiler must enqueue multiple kernels for execution, where each kernel implements a different block
of the original \hjforall{} loop body. This takes advantage of OpenCL's implicit all-to-all synchronization between
kernel launches in the same command queue.

Fig~\ref{fig:barrier-example-class} shows the Java equivalent of the generated bytecode in \textit{OpenCLKernel.class}
for the example in Figure~\ref{fig:barrier-example}, which will eventually be translated to an OpenCL kernel.
\textit{OpenCLKernel.class} declares \textit{passId}, which identifies a block of the body of the original \hjforall{} loop.
Fig~\ref{fig:barrier-example-stub} shows the sequence of OpenCL API calls for this barrier example.
Note that the value of \textit{passId} passed to the OpenCL kernel is incremented before each kernel launch,
leading to execution of a different block of code for each kernel invocation as the kernel progresses through synchronization points.

\begin{figure}
\begin{small}
\begin{lstlisting}[xleftmargin=0.35cm, xrightmargin=0.25cm]            
forall (point [i]:[0:n-1]) {
  method1(i, j);
  // synchronization point 1
  next;
  method2(i, j);
  // synchronization point 2
  next;
}
\end{lstlisting}
\end{small}
\caption{HJ-OpenCL Barrier Example.}
\label{fig:barrier-example}
\end{figure}

\begin{figure}
\begin{center}
\begin{lstlisting}[xleftmargin=0.35cm, xrightmargin=0.25cm]
public class OpenCLKernel 
		extends com.amd.aparapi.Kernel {
    ...
    int passId;
    public void run() { 
      // For synchronization point 1
      if (passId == 0) {
          method1(gid, j);   
      }
      // For synchronization point 2
      if (passId == 1) { 
          method2(gid, j);
      }
    }
}
\end{lstlisting}
\end{center}
\caption{Barrier Example in OpenCLKernel.class}
\label{fig:barrier-example-class}
\end{figure}

\begin{figure}
\begin{center}
\begin{lstlisting}[xleftmargin=0.35cm, xrightmargin=0.25cm]
cl_int passid;
// For method1
passid = 0;
ret = clSetKernelArg(..., (void *) &passid);
ret = clEnqueueNDRangeKernel(...);
clEnqueueBarrier(command_queue);
// For method2
passid = 1;
ret = clSetKernelArg(..., (void *) &passid);
ret = clEnqueueNDRangeKernel(...);
clEnqueueBarrier(command_queue);
\end{lstlisting}
\end{center}
\caption{Generated kernel invocation code of barrier example by the HJ-Compiler}
\label{fig:barrier-example-stub}
\end{figure}

\begin{table*}
\small{
\begin{center}
\begin{tabularx}{\linewidth}{|c|X|c|c|c|}
\hline
{\bf Benckmark} & {\bf Summary} & {\bf Data Size} & {\bf ArrayView?} & {\bf Next?} \\
\hline
Blackscholes & Data-parallel financial application which calculates the price of European put and call options & \shortstack{16,777,216 virtual options} & \shortstack{No} & \shortstack{No}\\
\hline
Crypt & Cryptographic application from the Java Grande Benchmarks~\cite{jgf} & Size C with {\tt N}= 50,000,000 & No & No \\
\hline
MatMult & A standard dense matrix multiplication: $C = A . B$ & 1024$\times$1024 & Yes & No \\
\hline
Doitgen & Multi-resolution analysis kernel from PolyBench~\cite{polybench}, ported to Java & 128$\times$128$\times$128 & Yes & No \\
\hline
MRIQ & Three-dimensional medical benchmark from Parboil~\cite{parboil}, ported to Java & large size(64$\times$64$\times$64) & No & No \\
\hline
Syrk & Symmetric rank-1 computation: $C = \alpha . A . A^T + \beta . C$ from PolyBench~\cite{polybench}, ported to Java & 2048$\times$2048 & Yes & No \\
\hline
Jacobi & 1-D Jacobi stencil computation from PolyBench~\cite{polybench}, ported to Java & {\tt T} = 50 and {\tt N} = 134,217,728 & No & No \\
\hline
SparseMatmult & Sparse matrix multiplication from the Java Grande Benchmarks~\cite{jgf} & Size C with {\tt N} = 500,000 & No & No \\
\hline
Spectral-norm & Eigenvalue computation using the power method from the Computer Language Benchmarks Game & {\tt N} = 2,000 & No & Yes \\
\hline 
SOR & Successive over relaxation from the Java Grande Benchmarks~\cite{jgf} & Size C with {\tt N} = 2,000 & Yes & Yes \\
\hline
\end{tabularx}
\caption{Details on the benchmarks used to evaluate HJ-OpenCL}
\label{tab:benchmark}
\end{center}
}
\end{table*}

\subsection{Multidimensional Array Issues}\label{subsec:multidimension}
\input{mdarray}
