\chapter{GPGPU への対応}
Cerium の新たな演算資源として GPU の使用を可能にした。\cite{yuhi:2013a}
現在、GPU のような異なる種類のアーキテクチャを搭載した CPU 、つまりヘテロジニアスな CPU が増えている。
特定の計算に特化した Task の生成やスケジューリングを行い、 GPGPU により高い並列度を出す研究は様々な分野で行われている。
本研究では Cerium を特定の計算に限らず、 GPU を用いて汎用計算を行えるフレームワークに改良する。


\section{OpenCL および CUDA による実装}
OpenCL 、CUDA による GPGPU 対応を行った。
Scheduler と CpuThreads に対応させる形で
OpenCL を用いた GpuScheduler と GpuThreads、
CUDA を用いた CudaScheduler と CudaThreads を実装した。
それぞれの Scheduler 内で各フレームワークの API を用いて GPU の制御を行っている。

 TaskManager から受け取った TaskList をもとに Device 上のメモリバッファを作成する。
その後 CommandQueue、 Stream といったそれぞれの Queue に Device 制御用の Command を Queueing していく。

Command は Queueing した順に実行されるので、以下のように Command を Queueing する。
\begin{enumerate}
\item Host から Device へのデータ転送
\item kernel の実行
\item Device から Host へのデータ転送
\end{enumerate}

データの転送や kernel の実行は非同期 API を用いることで並列に行うことができる。

通常、フレームワークが依存関係を解決して実行するが、
非同期 API を用いる場合はユーザが依存関係を考慮する必要がある。
しかし Task の依存関係は TaskManager が既に解決した状態で送ってくるので、
 Scheduler は依存関係を考慮せずに実行して問題ない。

GPGPU 用の Scheduler は CommandQueue を2つ持っており、Task をパイプライン的に実行する。
GpuScheduler のパイプライン処理部分をソースコード:\ref{src:pipeline_gpu}に示す。
\newpage
\begin{lstlisting}[frame=lrbt,label=src:pipeline_gpu,caption=GpuSchedulerにおけるパイプライン処理,numbers=left]
void
GpuScheduler::run() {
  for (;;) {
    memaddr params_addr = connector->task_list_mail_read();
    // read task list mail from DmaManager
    
    while (params_addr) {
      // since we are on the same memory space, we don't has to use dma_load here
      tasklist = (TaskListPtr)connector->dma_load(this, params_addr,sizeof(TaskList),
                                                  DMA_READ_TASKLIST);
                                                  
    for (TaskPtr nextTask = tasklist->tasks; nextTask < tasklist->last();
         nextTask = nextTask->next()) {
                   
            kernel[cur] = clCreateKernel(program, function, &ret);

            int param = 1;
            for(int i=0;i<nextTask->inData_count;i++) {
                ListElement *input_buf = nextTask->inData(i);
                if (input_buf->size==0) break;
                createBuffer(&memin[cur], param, context, mem_flag, input_buf->size, &ret);
                ret = clEnqueueWriteBuffer(command_queue[cur], memin[cur].buf[param],
                                           CL_FALSE, 0, input_buf->size,
                                           input_buf->addr, 0, NULL, NULL);
                ret = clSetKernelArg(kernel[cur],  param, sizeof(memaddr),
                                     (void *)&memin[cur].buf[param]);
                param++;
            }
            memin[cur].size  = param; // +1 means param
            
            for(int i = 0; i<nextTask->outData_count;i++) { // set output data
                ListElement *output_buf = nextTask->outData(i);
                if (output_buf->size==0) break;
                createBuffer(&memout[cur], i, context, CL_MEM_WRITE_ONLY, output_buf->size, &ret);
                ret = clSetKernelArg(kernel[cur],  param,
                                     sizeof(memaddr), (void *)&memout[cur].buf[i]);
                param++;
            }
            memout[cur].size = param - memin[cur].size;

            ret = clEnqueueTask(command_queue[cur], kernel[cur], 0, NULL, NULL);
                
            for(int i=0;i<nextTask->outData_count;i++) { // read output data
                ListElement *output_buf = nextTask->outData(i);
                if (output_buf->size==0) break;
                GpuBufferPtr mem = memout ;
                ret = clEnqueueReadBuffer(command_queue[cur], mem[cur].buf[i0], CL_FALSE, 0,
                                          output_buf->size, output_buf->addr, 0,
                                          NULL,&memout[cur].event[i]);
                }
            cur++;
            if (STAGE <= cur) cur = 0;
            wait_for_event(kernel_event, memout, tasklist, cur);
         }
         reply = (memaddr)tasklist->waiter;
         params_addr = (memaddr)tasklist->next;
      }

      wait_for_event(kernel_event, memout, tasklist, cur);
    
      unsigned long long wait = 0;
      (*connector->end_dmawait_profile)(&wait, &(connector->start_time), &(connector->stop_time));
      connector->mail_write((memaddr)MY_SPE_STATUS_READY);
    }
    /* NOT REACHED */
}

\end{lstlisting}

\begin{itemize}
\item 4行目 : DMAManager から tasklist mail を取得
\item 9行目 : 取得した mail から TaskList を read をする。 
  TaskList に格納されている Task 全てを実行し終わるまでループする
\item 15行目 : 二段のパイプラインを形成するため、kernel を2つ持作る
\item 18行目 : パイプラインの read 部分。input データを全て kernel の引数として MemoryBuffer に書き込み、
  kernel に Buffer をset する
\item 31行目 : Output データを書き込む MemoryBuffer を用意し、kernel にset している。
\item 40行目 : パイプラインの exec 部分。kernel を実行する
\item 42行目 : パイプラインの write 部分。kernel は実行した結果を MemoryBuffer に書き込む。
  MemoryBuffer に書き込まれた値をここで読み出している。
\item 52行目 : パイプラインのステージの切り替えを行っている。
  wait\_for\_event 内で依存関係を解決したら実行の終わったステージの MemoryBuffer を delete し、次のステージへ移行する
\item 55行目 : 次の task を読み出し、このループを終了する
\end{itemize}

DMAManager から転送されてきた Task を読み込み、Input/Output データを取り出す。
データは OpenCL の API を介して GPU の MemoryBuffer に送信され、kernel が実行される。
実行終了後は MemoryBuffer から戻り値を読み取り、Cerium に Output データとして返している。
一連の処理は CommandQueue を介して GPU で実行される。
GpuScheduler は CommandQueue を2つ持っており、二段のパイプラインが形成される。

全ての Task が終了すると、
TaskManager 間の通信を担当する同期キューである mail を通して TaskManager に Task の終了を通知する。
終了が通知されると TaskManager でその TaskList に関する依存関係が解消される。

Scheduler 内で Platform や Device ID の取得、 Context の生成、 Kernel の Build と Load等も行っており、
並列処理したい計算のみに集中できる。

\section{データ並列}
並列プログラミングにおいて、明示的な並列化部分はループ部分である。
GPU は数百個のコアを有しており、ループ部分に対してデータ並列で処理を行うことで CPU より高速に演算を行う事ができる。
プログラムの大部分がループであれば、データ並列による実行だけでプログラムの性能は向上する。

OpenCL 、 CUDA ともにデータ並列をサポートしている。\cite{yuhi:2014a}\cite{kkb:2014a}
OpenCL と CUDA はTask を実行する際にデータをどう分割するか指定し、
kernel にデータ並列用の処理を加えることで可能となる。
\ref{sec:multicore_dataparallel}節で Cerium でマルチコア CPU におけるデータ並列を可能にした。
GPGPU においてもデータ並列実行をサポートする。
GPU 上でのデータ並列実行もマルチコア CPU と変わらず、iterate API によりデータ並列用の Task を生成することができる。
iterate で Task を生成することで Scheduler が OpenCL 及び CUDA の API に適切なパラメタを渡している。
Task の生成部分は マルチコア CPU と GPU で完全に同じ形式で記述できる。

データ並列実行の際、Task は以下のように記述する。
なお、例題は multiply を用いている。

\begin{lstlisting}[frame=lrbt,label=src:multiply_opencl,caption=Multiply(OpenCL),numbers=left]
__kernel void
multiply(__global const long  *params,
         __global const float *input1,
         __global const float *input2,
         __global const float *output) {

    long id = get_global_id(0);

    output[id] = input1[id] * input2[id];
}
\end{lstlisting}

\begin{lstlisting}[frame=lrbt,label=src:multiply_cuda,caption=Multiply(CUDA),numbers=left]
__global__ void 
multiply(__global const long  *params,
         __global const float *input1,
         __global const float *input2,
         __global const float *output) {

    int id = blockIdx.x * blockDim.x + threadIdx.x;

    output[id] = input1[id] * input2[id];
}
\end{lstlisting}

このような Task を分割数分生成する。
分割数は Task それぞれのフレームワークが用意している API を用いて指定する。

\begin{itemize}
\item 自分の計算する範囲を取得(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}の7行目)
\item 取得した範囲を計算(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}の9行目)
\end{itemize}
いずれの Task も上記の手順で処理を行っている。
計算する範囲について、 OpenCL では取得用の API を用い、 CUDA では kernel の持つ組み込み変数から算出する。
マルチコア CPU では引数としてデータを直接渡していたが、OpenCL 、CUDA では上記の方法でメモリバッファから Load し、計算を行う。
値渡しや修飾子等若干の違いはあるが、OpenCL 、CUDA ともにマルチコア CPU(ソースコード:\ref{src:multicore_cpu}) とほぼ同じ形式で kernel を記述することができる。
CPU、 OpenCL、 CUDA いずれか1つの記述から残りのコードも生成できるようにする事が望ましい。

データ並列で実行する場合、 Input と Output を各 Task 間で共有するため、少ないコピーに抑えられる。
CPU ではメモリ領域を節約する事はできるが、 Task と Manager でメモリ領域が同じ(\ref{sec:shared_memory}節)なため、
コピーによるオーバーヘッドは少ない。

しかし GPU は SharedMemory ではなく、データの転送がオーバーヘッドとなるため、コピーを減らす事で並列度の向上が見込める。