\chapter{Parallel Computing Platform}

\section{OpenCL}
OpenCL とは、Multi Core CPU と GPU のようなヘテロジニアスな環境を利用した並列計算を支援するフレームワークである。

OpenCL には主に2つの仕様がある。
\begin{itemize}
  \item OpenCL C
  \item OpenCL Runtime API
\end{itemize}
OpenCL C は演算用プロセッサ(本研究では GPU)上で動作する、C 言語を拡張したプログラミング言語である。
一方で、OpenCL Runtime API は OpenCL C で記述したプログラムを GPU 上で実行させるため、制御用のプロセッサ(本研究では CPU)上で利用する API である。

OpenCL ではオペレーティングシステムなどが処理されるメイン CPU などのことを host、GPGPU を搭載したグラフィックボードなどのことを device と定義している。
OpenCL では device に CPU を割り当てることも可能である。OpenCL Application は host 側のプログラムと device 側のプログラムが一体となって動作する。
この device 上で動作するプログラムを OpenCL では、特別に kernel と呼ぶ。

\subsection{Command Queue}
OpenCL では、デバイスの操作に Command Queue を使用する。Command Queue は device に OpenCL の Operation を送るために仕組みである。
Command Queue は clCreateCommandQueue という OpenCL API に所属するコンテキストと実行対象となる device を指定することで生成される。

Command Queue では kernel の実行、input buffer の読み込み、output buffer への書き込みといった Operation が投入された順序で実行される。
Command Queue を作成するとき CL\_QUEUE\_OUT\_OF\_ORDER\_EXEC\_MODE\_ENABLE のプロパティを指定することで Operation を順序を無視して実行することが可能になる。
Operation を out of order で実行する場合、データの依存関係を記述する必要がある。
各 Operation には event\_wait\_list と event を指定することができ、これらを利用してデータの依存関係を記述することができる。
しかし、この CL\_QUEUE\_OUT\_OF\_ORDER\_EXEC\_MODE\_ENABLE のプロパティをサポートしている device は少なく、Mac OS X では OS レベルでサポートしていない。
パイプライン実行を行うためには kernel の実行やデータ転送を out of order で実行する必要がある。
CL\_QUEUE\_OUT\_OF\_ORDER\_EXEC\_MODE\_ENABLE のプロパティが無効の場合、複数の Command Queue を生成し、Command Queue を複数投入することで Operation を out of order で実行することが可能になる。

\subsection{Memory Access} \label{sec:memory_access}
host 側は主にデータを input/output する Memory の確保を行う。
GPU の Memory 空間(図:\ref{fig:gpuarch})や Cell の Memory 空間(図:\ref{fig:cellarch})は Multi Core CPU(図:\ref{fig:cpuarch})とは異なり、Shared Memory ではないため host と device 間でデータの共有ができない。
アクセスするには Memory 空間ごとコピーしなければならない。

OpenCL では host 側で Memory Buffer を作成して Memory のコピーを行う。
データの読み込みは clEnqueueReadBuffer、書き込みは clEnqueueWriteBuffer という API でそれぞれ行われる。
前節で述べた通り、これらの Operation を Command Queue に Enqueue する。そして、データの依存解消が解消されると実行される。

\begin{figure}[htpd]
  \begin{center}
    \includegraphics[scale=0.4]{./images/gpu_arch.pdf}
  \end{center}
  \caption{Gpu Architecture}
  \label{fig:gpuarch}
\end{figure}

\newpage

\begin{figure}[htpd]
  \begin{minipage}[c]{0.5\hsize}
    \begin{center}
      \includegraphics[scale=0.7]{./images/cell_arch.pdf}
    \end{center}
    \caption{Cell Architecture}
    \label{fig:cellarch}
  \end{minipage}
  \begin{minipage}[t]{0.5\hsize}
    \begin{center}
      \includegraphics[scale=0.7]{./images/cpu_arch.pdf}
    \end{center}
    \caption{Cpu Architecture}
    \label{fig:cpuarch}
  \end{minipage}
\end{figure}
  
    
\subsection{Data Parallel Execution}
3D グラッフィクのような多次元のデータ構造がある場合に高い並列度を保つには、それを分割して並列に実行する機能が必要である。
これを OpenCL ではデータ並列と呼んでいる。
OpenCL では次元数に対応する index があり、OpenCL は一つの記述から index の異なる複数の kernel を自動生成する。
その添字を global\_id と呼ぶ。このとき入力されたデータは WorkItem という処理単位に分割される。

OpenCL は WorkItem に対して、それぞれを識別する ID(global\_id) を割り当てる。
kernel は get\_global\_id という API によって ID を取得し、取得した ID に対応するデータに対して処理を行うことでデータ並列を実現する。
また、WorkItem は3次元までデータを渡すことができる。

データ並列による kernel 実行の場合、clEnqueueNDRangeKernel API を使用する。この関数の引数として WorkItem の数と次元数を指定することでデータ並列で実行できる。

\begin{figure}[htpd]
  \begin{center}
    \includegraphics[scale=0.7]{./images/workitem.pdf}
  \end{center}
  \caption{WorkItem ID}
  \label{fig:workitem}
\end{figure}

\newpage
\subsection{Work Group}
前節で WorkItem という処理単位について述べたが、さらに複数個の WorkItem を WorkGroup という単位にまとめることができる。
WorkGroup 内では同期やローカルメモリの共有が可能になる。

グローバルワークアイテム(ワークアイテム全体)の個数と、ローカルワークアイテム(グループ一つ辺りのアイテム)の個数を指定することでワークアイテムを分割する。
なお、このときグローバルワークアイテム数はローカルアイテム数の整数倍でなければ clEnqueueNDRangeKernel API の呼び出しは失敗する。

ローカルアイテムに0を指定することで、コンパイル時に最適化させることができる。
したがって、ローカルアイテムのサイズは0を指定するのが一般的である。

なお、WorkGroup を設定した場合、global\_id の他に work\_group\_id, local\_id がそれぞれの kernel に割り当てられる(図:\ref{fig:workitem})。

kernel からそれぞれ ID に対応した API を使用して、各 ID を取得する。
取得した ID から自分が担当する index を計算して導く。
表:\ref{table:kernel_id_api}は kernel で使用できる ID を取得するための API となる。

\begin{table}[htpb]
  \begin{center}
    \small
    \begin{tabular}[htpb]{|c|l|} \hline
      get\_group\_id  & work\_group\_id を取得  \\ \hline
      get\_local\_id  & local\_id を取得 \\ \hline
      get\_global\_id & global\_id を取得 \\ \hline
    \end{tabular}
    \caption{kernel で使用する ID 取得の API}
    \label{table:kernel_id_api}
  \end{center}
\end{table}

なお、local\_id, global\_id を取得する API は引数に0, 1, 2の値を指定することができる。
id は x, y, z 座標があり、それぞれが0, 1, 2に対応している。
例えば get\_global\_id(1) と呼び出した場合は y 座標の、get\_global\_id(2) と呼び出した場合は z 座標の global\_id を取得する。

\section{CUDA}
CUDA とは、半導体メーカー NVIDIA 社が提供する GPU コンピューティング向けの総合開発環境でコンパイラ、ライブラリ、デバッガなどから構成される。
プログラム言語である CUDA C は C 言語ベースに拡張を加えたものである。

CUDA には CUDA Runtime API と CUDA Driver API の2種類がある。
Driver API は Runtime API と比べてプログラマが管理すべきリソースが多い。しかし、Runtime API より柔軟な処理を行うことができる。
今回は Driver API を使用して実装した。

CUDA も OpenCL と同様に、制御を行う CPU 側を host、GPU 側を device と定義している。また、device 上で動作するプログラムも OpenCL と同様に kernel と呼ぶ。

\subsection{Stream} \label{sec:stream}
CUDA には OpenCL の Command Queue と似たような仕組みとして Stream がある。
Stream は host 側の発行された Operation を一連の動作として device で実行する。Stream 内の Operation は発行された順序で実行されることが保証されている。
異なる Stream での Operation の依存関係が解消され実行可能な場合、Operation を同時に実行することができる。
例として、ある Stream に kernel を実行する Operation があり、それとは異なる Stream に依存関係がないデータを転送する Operation があった場合、kernel の実行中にデータ転送を行うことが可能になる。

Stream は cuStreamCreate という Driver API で生成される。
OpenCL と異なり、コンテキストと実行対象となる device を指定する必要はないが、コンテキストを作成した Thread と同一の Thread でないと Stream が生成できないという制約がある。
引数に Stream を指定しない API はすべて host 側をブロックする同期的な処理となる。複数の Stream を同時に走らせ Operation を並列に実行するためには非同期処理を行う API を利用する必要がある。

Stream 内の Operation を同期する方法はいくつかある。
一つ目は cuStreamSynchronize API を利用した同期方法である。
cuStraemSynchronize API の引数に Stream を指定すると、指定した Stream に発行されたすべての Operation が終了するまで host をブロックする。
host をブロックすることなく、Stream に発行された Operation が終了したかどうかを調べるには cuStreamQuery API を利用する。
二つ目は event を利用した同期方法である。
各 Operation に event\_wait\_list や event を指定して同期を取る OpenCL とは異なり、CUDA では cuStreamWaitEvent で Stream と event を指定して同期を取る。
cuStreamWaitEvent は指定した event が cuEventRecord されるまで指定された Stream をブロックする。
event は cuEventCreate で作成され、Operation の直後で cuEventRecord に event を指定することで直前に行う Operation の終了を待つことが可能になる。
この event は別の Stream で cuEventRecord されるものでも待つことができる。

\subsection{Memory Access}
CUDA も OpenCL と同様に Shared Memory ではないため host と device 間でデータの共有ができない。アクセスするには Memory 空間ごとコピーする必要がある。
CUDA でのデータの読み込みは cuMemcpyHtoD、書き込みは cuMemcpyDtoH という API でそれぞれ行われる。しかし、これらの API は同期的に実行されてしまう。
非同期処理にしたい場合、読み込みで cuMemcpyHtoDAsync、書き込みで cuMemcpyDtoHAsync という API をそれぞれ利用することで非同期に行うことができる。

\subsection{Data Parallel Execution}
CUDA では OpenCL の WorkItem に相当する単位を Thread と定義してる。この Thread をまとめたものを block と呼ぶ。
CUDA でデータ並列による kernel 実行を行う場合、cuLaunchKernel API を使用する。
この関数の引数として各座標の block 数と各座標の block 一つ当たりの thread 数を指定することでデータ並列で実行できる。
cuLaunchKernel で kernel を起動すると各 Thread に対して block ID と Thread ID が付与される。
CUDA には OpenCL とは異なり、ID を取得する API は存在しない。
代わりとして、kernel には組み込み変数が準備されており、それを参照することで対応するデータに対し処理を行うことでデータ並列を実現する。
組み込み変数は以下の通りである。
\begin{itemize}
  \item uint3 blockDim
  \item uint3 blockIdx
  \item uint3 threadIdx
\end{itemize}
各組み込み変数はベクター型で、blockDim.x とすると x 座標の Thread 数を参照することができる。
blockIdx.x とすると x 座標の block ID が参照でき、threadIdx.x とすると x 座標の thread ID を参照することができる。
blockDim.x * blockIdx.x + threadIdx.x を計算すると OpenCL の get\_global\_id(0) で取得できる ID に相当する ID を得ることができる。
例えば、ある kernel で get\_global\_id(0) の返り値が13の場合、 CUDA では図:\ref{fig:calculate_index}のようにすることで ID を算出することができる。

\begin{figure}[!h]
  \begin{center}
    \includegraphics[scale=0.4]{./images/culculate_index.pdf}
  \end{center}
  \caption{Calculate Index}
  \label{fig:calculate_index}
\end{figure}

\newpage

\section{Porting to OpenCL to CUDA}
本項では OpenCL で記述された Application を CUDA に移植する方法について説明する。
以下の表は OpenCL と CUDA の用語および修飾子、ID の参照、Object、API の対応表である。
\begin{table}[!h]
  \begin{center}
    \small
    \begin{tabular}[htpb]{|l|l|} \hline
      OpenCL & CUDA \\ \hline \hline
      WorkItem  & Thread  \\ \hline
      WorkGroup  & Block \\ \hline
      Global Memory & Global Memory \\ \hline
      Local Memory & Shared Memory \\ \hline
      Private Memory & Local Memory \\ \hline
    \end{tabular}
    \caption{用語}
    \label{table:terminology_comp}
  \end{center}
\end{table}

\begin{table}[!h]
  \begin{center}
    \small
    \begin{tabular}[htpb]{|l|l|} \hline
      OpenCL & CUDA \\ \hline \hline
      \_\_kernel function & \_\_global\_\_ function \\ \hline
      No necessary & \_\_device\_\_ function (not callable from host) \\ \hline
      \_\_constant variable & \_\_constant\_\_ variable \\ \hline
      \_\_global variable & \_\_device\_\_ variable \\ \hline
      \_\_local variable & \_\_shared\_\_ variable \\ \hline
    \end{tabular}
    \caption{修飾子}
    \label{table:Qualifiers_comp}
  \end{center}
\end{table}

\begin{table}[!h]
  \begin{center}
    \small
    \begin{tabular}[htpb]{|l|l|} \hline
      OpenCL & CUDA \\ \hline \hline
      get\_num\_groups() & gridDim \\ \hline
      get\_local\_size() & blockDim \\ \hline
      get\_group\_id() & blockIdx \\ \hline
      get\_local\_id() & threadIdx \\ \hline
      get\_global\_id() & blockDim * blockIdx + threadIdx \\ \hline
      get\_global\_size() & gridDim * blockDim \\ \hline
    \end{tabular}
    \caption{kernel Indexing}
    \label{table:kernel_api_comp}
  \end{center}
\end{table}

\newpage

\begin{table}[!h]
  \begin{center}
    \small
    \begin{tabular}[htpb]{|l|l|} \hline
      OpenCL & CUDA \\ \hline \hline
      cl\_device\_id & CUdevice \\ \hline
      cl\_context & CUcontext \\ \hline
      cl\_program & CUmodule \\ \hline
      cl\_kernel & CUfunction \\ \hline
      cl\_mem & CUdeviceptr \\ \hline
      cl\_command\_queue & CUstream(but imperfection) \\ \hline
    \end{tabular}
    \caption{Objects}
    \label{table:object_comp}
  \end{center}
\end{table}

\begin{table}[!h]
  \begin{center}
    \small
    \begin{tabular}[htpb]{|l|l|} \hline
      OpenCL & CUDA \\ \hline \hline
      No required & cuInit() \\ \hline
      clGetContextInfo() & cuDeviceGet() \\ \hline
      clCreateContext() & cuCtxCreate() \\ \hline
      clCreateCommandQueue() & cuStreamCreate() \\ \hline
      clCreateProgramWithSource() & cuModuleLoad() \\ \hline
      clBuildProgram() & CUDA programs are compiled offline \\ \hline
      clCreateKernel() & cuModuleGetFunction() \\ \hline
      clCreateBuffer() & cuMemAlloc() \\ \hline
      clEnqueueWriteBuffer() & cuMemcpyHtoD() \\ \hline
      clEnqueueReadBuffer() & cuMemcpyDtoH() \\ \hline
      clEnqueueNDRangeKernel() & cuLaunchKernel() \\ \hline
      clSetKernelArg() & Functonality in cuLaunchKernel() \\ \hline
      clReleaseMemObj() & cuMemFree() \\ \hline
    \end{tabular}
    \caption{APIs}
    \label{table:api_comp}
  \end{center}
\end{table}

\newpage

\subsection{Sequential Execution} \label{sec:seq}
OpenCL および CUDA で逐次実行するプログラムを例として変換方法を説明する。与えられた二つの input data を乗算し、指定された領域に output する kernel を複数回起動している。

\subsubsection{Initialize}
\begin{verbatim}
  // initialize(OpenCL)
  cl_platform_id platform_id;
  cl_uint num_platforms;
  cl_device_id device_id;
  cl_uint num_devices;
  cl_uint ret;
  cl_command_queue command_queue;

  clGetPlatformIDs(1, &platform_id, &num_platforms);
  clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU,
                 1, &device_id, &num_devices);
  cl_context context = clCreateContext(NULL, 1, &device_id,
                                       NULL, NULL, &ret);
  command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
\end{verbatim}

OpenCL での初期化は上記のようになる。

OpenCL は様々なメーカー(NVIDIA, AMD など) GPU の対応しているため platform\_id を取得し、それをもとに Context を生成する。

\begin{verbatim}
  // initialize(CUDA)
  CUdevice device;
  CUcontext context;
  
  cuInit(0);
  cuDeviceGet(&device, 0);
  cuCtxCreate(&context, CU_CTX_SCHED_SPIN, device);
\end{verbatim}

CUDA での初期化は上記のようになる。

CUDA は NVIDIA が提供する GPU にのみ対応しているので個別の id を取得する必要がない。
また、CUDA には Null Stream と呼ばれるデフォルトの Stream がある。
OpenCL の Command Queue のように必ず生成する必要はない。
\subsubsection{Load Kernel}
\begin{verbatim}
  // load kernel(OpenCL)
  const char* filename = "multiply.cl";
  const char* functionname = "multiply";
  
  int fp = open(filename, O_RDONLY);
  
  struct stat stats;
  fstat(fp,&stats);
  off_t size = stats.st_size;
  
  char *kernel_src_str = (char*)alloca(size+1);
  size_t kernel_code_size = read(fp, kernel_src_str, size);
  close(fp);
  kernel_src_str[size] = 0;


  cl_program program;
  program = clCreateProgramWithSource(context, 1,
                                      (const char **)&kernel_src_str, 0, &ret);
  clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
  cl_kernel kernel = clCreateKernel(program,functionname, &ret);
\end{verbatim}

OpenCL での kernel の読み込みは上記のようになる。

OpenCL ではプログラム内で kernel をコンパイルすることが可能である。
専用のコンパイラを必要としないため、様々なプラットフォームへの展開が可能になる。
また、Linux 環境では Fixstars 社が提供する foxc、Mac OS X では openclc を利用することで事前にコンパイルすることもできる。

\begin{verbatim}
  // load kernel(CUDA)
  CUmodule module;
  CUfunction function;

  cuModuleLoad(&module, "multiply.ptx");
  cuModuleGetFunction(&function, module, "multiply");
\end{verbatim}

CUDA での kernel の読み込みは上記のようになる。

CUDA では CUDA に付属されている専用コンパイラ nvcc を使って事前に kernel をコンパイルする必要がある。

\subsubsection{Memory Allocate}
\begin{verbatim}
  // memory allcate(OpenCL)
  cl_mem memA = clCreateBuffer(context, CL_MEM_READ_ONLY,
                               WORKS*sizeof(float), NULL, &ret);
  cl_mem memB[num_exec];
  cl_mem memOut[num_exec];
  for (int i=0;i<num_exec;i++) {
    memB[i] = clCreateBuffer(context, CL_MEM_READ_ONLY,
                             sizeof(float), NULL, &ret);
    memOut[i] = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
                               WORKS*sizeof(float), NULL, &ret);
  }
\end{verbatim}

OpenCL での device 上の memory 領域の確保は上記のようになる。

今回のプログラムでは一つ目の input data は共通の値なので、一つ目の input data を格納する領域は一つのみ確保している。
二つ目の input data および output data は kernel ごとに異なる値なので、実行する kernel の数だけ領域を確保している。

\begin{verbatim}
  // memory allocate(CUDA)
  CUdeviceptr devA;
  CUdeviceptr devB[num_exec];
  CUdeviceptr devOut[num_exec];
  
  cuMemAlloc(&devA, LENGTH*THREAD*sizeof(float));
  for (int i=0;i<num_exec;i++) {
    cuMemAlloc(&devB[i], sizeof(float));
    cuMemAlloc(&devOut[i], LENGTH*THREAD*sizeof(float));
  }
\end{verbatim}

CUDA での device 上の memory 領域の確保は上記のようになる。

CUDA では memory 領域を確保するとき、context を生成した Thread と同一の Thread である必要があるため引数に context を指定する必要はない。
\newpage
\subsubsection{Data Transfer(host to device)}
\begin{verbatim}
  // data transfer(OpenCL)
  clEnqueueWriteBuffer(command_queue, memA, CL_TRUE, 0,
                       WORKS*sizeof(float), A, 0, NULL, NULL);

  for (int i = 0;i<num_exec;i++){
    B[i] = (float)(i+1);
    clEnqueueWriteBuffer(command_queue, memB[i], CL_TRUE, 0,
    sizeof(float), &B[i], 0, NULL, NULL);
  }
\end{verbatim}

OpenCL での host から device へのデータ転送は上記のようになる。

clEnqueueWriteBuffer の第三引数に CL\_TRUE を指定することで、データ転送が完了するまで host をブロックする同期的な Operation となる。
\begin{verbatim}
  // data transfer(CUDA)
  cuMemcpyHtoD(devA, A, LENGTH*THREAD*sizeof(float));
  
  for (int i=0;i<num_exec;i++) {
    B[i] = (float)(i+1);
    cuMemcpyHtoD(devB[i], &B[i], sizeof(float));
  }
\end{verbatim}

CUDA での host から device へのデータ転送は上記のようになる。

cuMemcpyHtoD はデータ転送が完了するまで host をブロックする同期的な Operation である。
\subsubsection{Launch Kernel}
\begin{verbatim}
  // launch kernel(OpenCL)
  clSetKernelArg(kernel, 0, sizeof(cl_mem), &memA);
  for (int i = 0;i<num_exec;i++){
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &memB[i]);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &memOut[i]);
    clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
                           &num_work, NULL, 0, NULL, NULL);
    }
\end{verbatim}

OpenCL での kernel の実行は上記のようになる。

clSetKernelArg で kernel の引数に対して値を渡すことができる。

clEnqueueNDRangeKernel の第三引数に WorkItem の次元数を指定し、第五引数に WorkItem の数を指定することでデータ並列で実行される。
\begin{verbatim}
  // launch kernel(CUDA)
  for (int i=0;i<num_exec;i++) {
    void* args[] = {&devA, &devB[i], &devOut[i]};
    cuLaunchKernel(function,
                   LENGTH, 1, 1,
                   THREAD, 1, 1,
                   0, 0, args, NULL);
  }
\end{verbatim}

CUDA での kernel の実行は上記のようになる。

clLaunchKernel の第二、第三、第四引数に各次元のブロック数を指定し、第五、第六、第七引数に各次元の Thread 数を指定することでデータ並列で実行される。
今回は、LENGTH 個のブロックがあり、各ブロックに THREAD 個の Thread が起動する。
また、CUDA には clSetKernelArg に相当する API はない。clLaunchKernel の第十引数に指定することで kernel の引数に対して値を渡すことができる。
\subsubsection{Data Transfer(device to host)}
\begin{verbatim}
  // data transfer(OpenCL)
  for (int i=0;i<num_exec;i++)
    clEnqueueReadBuffer(command_queue, memOut[i], CL_TRUE, 0,
                        WORKS*sizeof(float), result[i], 0, NULL, NULL);
\end{verbatim}

OpenCL での device から host へのデータ転送は上記のようになる。

clEnqueueWriteBuffer と同様に第三引数に CL\_TRUE を指定することでデータ転送が完了するまで host をブロックする同期的な Operation となる。
\begin{verbatim}
  // data transfer(CUDA)
  for (int i=0;i<num_exec;i++)
    cuMemcpyDtoH(result[i], devOut[i], LENGTH*THREAD*sizeof(float));
\end{verbatim}

CUDA での device から host へのデータ転送は上記のようになる。

cuMemcpyDtoH は cuMemcpyHtoD と同様にデータ転送が完了するまで host をブロックする同期的な Operation である。
\newpage
\subsubsection{Memory Release}
\begin{verbatim}
  // memory release(OpenCL)
  clReleaseMemObject(memA);
  for (int i=0;i<num_exec;i++) {
    clReleaseMemObject(memB[i]);
    clReleaseMemObject(memOut[i]);
  }
  clReleaseCommandQueue(command_queue);
  clReleaseKernel(kernel);
  clReleaseProgram(program);
  clReleaseContext(context);
\end{verbatim}

OpenCL での memory 領域の開放やオブジェクトの開放は上記のようになる。

\begin{verbatim}
  // memory release(CUDA)
  cuMemFree(devA);
  for (int i=0;i<num_exec;i++) {
    cuMemFree(devB[i]);
    cuMemFree(devOut[i]);
  }
  cuModuleUnload(module);
  cuCtxDestroy(context);
\end{verbatim}

CUDA での memory 領域の開放やオブジェクトの開放は上記のようになる。

\subsubsection{Kenrel}
\begin{verbatim}
  // kernel(OpenCL)
  __kernel void
  multiply(__global float* A, __global float* B, __global float* C) {
    long index = (long)get_global_id(0);
    C[index] = A[index] * B[0];
  }
\end{verbatim}

OpenCL での kernel の記述は上記のようになる。
get\_global\_id で index を取得し、担当する部分の演算を行っている。

\begin{verbatim}
  // kernel(CUDA)
  extern "C" {
    __global__ void multiply(float* A, float* B, float* C) {
      int index = blockIdx.x * blockDim.x + threadIdx.x;
      C[index] = A[index] * B[0];
    }
  }
\end{verbatim}

CUDA での kernel の記述は上記のようになる。
組み込み変数である blockIdx, blockDim, threadIdx から index を算出し、担当する部分の演算を行なっている。

\subsection{Parallel Execution}
項:\ref{sec:seq}で例として用いたプログラムは一つ目の input data 以外に依存関係が存在しない。
つまり、OpenCL では Command Queue、CUDA では Stream を複数生成し、Operation を並列に実行して良い。
Operation を並列に実行することでデータ転送をオーバーラップすることができる。
このプログラムを kernel を実行する数だけ Command Queue および Stream を生成し、並列実行されるように変更する。
変更がある部分のみ記載する。それ以外の部分は項:\ref{sec:seq}と同じである。
実行の様子は図:\ref{fig:parallel_exec}のようになる。
\begin{figure}[!htpd]
  \begin{center}
    \includegraphics[scale=0.5]{./images/paralle_exec.pdf}
  \end{center}
  \caption{Parallel Execution}
  \label{fig:parallel_exec}
\end{figure}

\subsubsection{Initialize}
\begin{verbatim}
  // initialize(OpenCL)
  cl_command_queue command_queue[num_exec];
  
  for (int i=0;i<num_exec;i++)
    command_queue[i] = clCreateCommandQueue(context, device_id, 0, &ret);
\end{verbatim}

OpenCL では上記のように変更する。

Command Queue を配列として宣言し、要素の数だけ Command Queue を生成する。
\begin{verbatim}
  // initialize(CUDA)
  CUstream stream[num_exec];
  
  for (int i=0;i<num_exec;i++)
    cuStreamCreate(&stream[i], 0);
\end{verbatim}

CUDA では上記のように変更する。

Stream を生成するとき、OpenCL と同じように context と device\_id を渡す必要はない。
これは項:\ref{sec:stream}で述べたとおり、コンテキストを作成した Thread と同一の Thread でないと Stream を生成できないからである。

OpenCL と同様に Stream を配列として宣言し、要素の数だけ Stream を生成する。

\subsubsection{Data Transfer(host to device)}
\begin{verbatim}
  // data transfer(OpenCL)

  for (int i = 0;i<num_exec;i++){
    B[i] = (float)(i+1);
    clEnqueueWriteBuffer(command_queue[i], memB[i], CL_FALSE, 0,
    sizeof(float), &B[i], 0, NULL, NULL);
  }
\end{verbatim}

OpenCL では上記のように変更する。

clEnqueueWriteBuffer の第三引数に CL\_FALSE を指定することで、非同期な Operation として実行できる。つまり、Command Queue に Operation を投入し、投入した Operation の完了を待たずに host に制御を返す。
\begin{verbatim}
  // data transfer(CUDA)
  
  for (int i=0;i<num_exec;i++) {
    B[i] = (float)(i+1);
    cuMemcpyHtoDAsync(devB[i], &B[i], sizeof(float), stream[i]);
  }
\end{verbatim}

CUDA では上記のように変更する。

cuMemcpyHtoDAsync は非同期な Operation として実行される。第四引数に関連する Stream を指定することができる。

\subsubsection{Launch Kernel}
\begin{verbatim}
  // launch kernel(OpenCL)
  for (int i = 0;i<num_exec;i++){
    clEnqueueNDRangeKernel(command_queue[i], kernel, 1, NULL,
                           &num_work, NULL, 0, NULL, NULL);
    }
\end{verbatim}

OpenCL では上記のように変更する。

関連する Command Queue を指定するだけで良い。
\begin{verbatim}
  // launch kernel(CUDA)
  for (int i=0;i<num_exec;i++) {
    cuLaunchKernel(function,
                   LENGTH, 1, 1,
                   THREAD, 1, 1,
                   0, stream[i], args, NULL);
  }
\end{verbatim}

CUDA では上記のように変更する。

第九引数に関連する Stream を指定するだけで良い。
\subsubsection{Data Transfer(device to host)}
\begin{verbatim}
  // data transfer(OpenCL)
  for (int i=0;i<num_exec;i++)
    clEnqueueReadBuffer(command_queue[i], memOut[i], CL_FALSE, 0,
                        WORKS*sizeof(float), result[i], 0, NULL, NULL);
\end{verbatim}

OpenCL では上記のように変更する。

clEnqueueWriteBuffer と同様に第三引数に CL\_FALSE を指定することで非同期な Operation となる。
\begin{verbatim}
  // data transfer(CUDA)
  for (int i=0;i<num_exec;i++)
    cuMemcpyDtoHAsync(result[i], devOut[i], LENGTH*THREAD*sizeof(float), stream[i]);
\end{verbatim}

CUDA では上記のように変更する。

cuMemcpyDtoHAscyn は cuMemcpyHtoDAsync と同様に非同期な Operation として実行される。

\subsubsection{Wait for Exec}
host をブロックしない非同期な Operation はどのタイミングで実行されるかわからない。
全ての Operation が完了していないのに memory やオブジェクトを開放してしまう場合がある。
その場合、処理が正しく行われない。
この問題を回避するために、全ての Operation を待つという処理を host 側に記述する必要がある。
\begin{verbatim}
  // wait for exec(OpenCL)
  for (int i=0;i<num_exec;i++)
    clFinish(command_queue[i]);
\end{verbatim}

OpenCL での投入された Operation の完了を待つのは上記のようになる。

clFinish に完了を待ちたい Command Queue を指定することで、投入された Operation が全て完了するのを待つことができる。
\newpage
\begin{verbatim}
  // wait for exec(CUDA)
  for (int i=0;i<num_exec;i++)
      cuStreamSynchronize(stream[i]);
\end{verbatim}

CUDA での投入された Operation の完了を待つのは上記のようになる。

OpenCL と同様に cuStreamSynchronize に完了を待ちたい Steram を指定することで、Operation の終了を待つことができる。
cuStreamSynchronize は指定された Stream が完了するまで host をブロックする。
完了したかどうかだけを知りたい場合は、cuStreamQuery を利用する。

\section{StarPU}
ヘテロジニアス環境下でのタスクベースなプログラミングをサポートするものとして INRIA から提供されている StarPU というものがある。
StarPU は CPU, CUDA device, OpenCL device など複数のアーキテクチャに対応している。
StarPU には Codelet と呼ばれる device に Task を渡す仕組みがある。Task は非同期に実行され、host をブロックしない。
しかし、Task の依存関係やスケジューリングをプログラマが記述する必要がある。
この記述は非常に煩雑でプログラマに大きな負担をかけることになる。
Cerium ではそれらを自動化することでプログラマの負担を軽減したい。