Mercurial > hg > Papers > 2013 > yuhi-sigos

changeset 2:771aaa69c616

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
commit
author Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
date Mon, 01 Apr 2013 21:16:36 +0900
parents 764783c403bd
children bb436935f877
files paper/cerium.tex paper/cerium_gpu.tex paper/conclusion.tex paper/data_parallel.tex paper/images/createTask.graffle paper/images/createTask.pdf paper/opencl.tex paper/sigos.pdf paper/sigos.tex
diffstat 9 files changed, 165 insertions(+), 70 deletions(-) [+]
line wrap: on
line diff
--- a/paper/cerium.tex	Mon Apr 01 19:09:39 2013 +0900
+++ b/paper/cerium.tex	Mon Apr 01 21:16:36 2013 +0900
@@ -16,8 +16,8 @@
 Task は ActiveTaskList に移される。さらに、Scheduler に転送しやすい TaskList に変換してから各 Scheduler に
 転送される。
 \subsection{taskの生成}
-以下にTaskを生成する部分を示す。
-例題としてinput dataを二つ用意し、input同士を乗算し、outputに格納するmultiplyというプログラムを使用している。
+以下にTaskを生成する例題を示す。
+input dataを二つ用意し、input同士を乗算し、outputに格納するmultiplyという例題となる。
 \begin{verbatim}
 void
 multiply_init
@@ -94,16 +94,3 @@
     \end{center}
   \end{table}
 \end{tiny}
-
-\label{fig:cpuarch}
-\begin{figure}[ht]
-  \begin{center}
-    \includegraphics[scale=0.5]{./images/cpu_arch.pdf}
-  \end{center}
-  \caption{Cpu Archtecure}
-  \label{fig:cpuarch}
-\end{figure}
-
-\subsection{メモリアクセス}
-CPUは、CellやGPUと違い、メインメモリに直接アクセスできるという大きな違いがある。(図:\ref{fig:cpuarch})
-環境にもよるが、複数のcore(汎用コア)を持つ。
--- a/paper/cerium_gpu.tex	Mon Apr 01 19:09:39 2013 +0900
+++ b/paper/cerium_gpu.tex	Mon Apr 01 21:16:36 2013 +0900
@@ -4,7 +4,8 @@
 TaskList からメモリバッファを作成し、clEnqueuWriteBuffer, clEnqueueTask, clEnqueuReadBuffer の順に 
 CommandQueueu に enqueue する。
 Task の投入は CommandQueue を2つ用意しパイプライン的に実行を行う。Task の終了は、
-clWaitForEvent によって検出し、TaskManger 間の通信を担当する同期キューである mail を使って通知する。
+clWaitForEvent によって検出し、TaskManger 間の通信を担当する
+同期キューである mail を使って通知する(図:\ref{fig:createTask})。
 
 GpuScheduler内でplatformやdeviceのIDの取得、contextの生成、
 kernel の build と load 、等も行っているため並列計算のみに集中できる。
--- a/paper/conclusion.tex	Mon Apr 01 19:09:39 2013 +0900
+++ b/paper/conclusion.tex	Mon Apr 01 21:16:36 2013 +0900
@@ -6,6 +6,7 @@
 
 さらにCeriumにデータ並列の機構を実装する事でデータやtaskのコピーを減らし、
 メモリへの負荷とオーバーヘッドを減らすことに成功した。
+これはmany coreでも有効だと考えられる。
 
 Cerium では Task 自体に依存関係を明示的に記述しているが、OpenCL ではメモリバッファの依存関係で
 暗黙的に指定する方法がある。Cerium 側にもデータ依存関係を導入するのが望ましいと考えられる。
--- a/paper/data_parallel.tex	Mon Apr 01 19:09:39 2013 +0900
+++ b/paper/data_parallel.tex	Mon Apr 01 21:16:36 2013 +0900
@@ -1,16 +1,18 @@
 \section{Ceriumにおけるデータ並列}
 OpenCLで充分な並列度を得るには、データ並列による実行をサポートした方が良い。
+ceriumでopenclのデータ並列を使うために、iteratorというAPIを用意した。
+
 ベンチマークをとるために、まずはCPU(many core)上でデータ並列の機構を実装した。
 OpenCLでデータ並列を行う際は、NDRangeの引数でワークアイテムのサイズを設定し、以下のようにkernelを書けばよい。
 
 \begin{verbatim}
 __kernel void
-multi(__global const float *A, 
-      __global const float*B,
-      __global float *C)
+multi(__global const float *i_data1, 
+      __global const float *i_data2,
+      __global float *o_data)
 {
     int i = get_global_id(0);
-    C[i] = A[i]*B[i];
+    o_data[i] = i_data1[i]*i_data2[i];
 }
 
 \end{verbatim}
@@ -39,6 +41,8 @@
     return 0;
 }
 \end{verbatim}
+
+\subsection{Ceriumでのデータ並列におけるindex割り当ての実装}
 taskを生成するとき、dimensionとワークアイテムのサイズをもとに各taskが担当するindexを計算し、set\_paramする。
 kernelはget\_paramでそのindexを取得してデータ並列で実行する。
 get\_param APIがopenCLのget\_global\_id APIに相当する。
@@ -71,6 +75,7 @@
       \end{tabular}
     \end{center}
   \end{table}
+
 \end{tiny}
 この実装により、Ceriumでデータ並列の実行が可能になった。
 並列プログラミングだと、並列化するtaskが全部同一であるという事は少なくない。
--- a/paper/images/createTask.graffle	Mon Apr 01 19:09:39 2013 +0900
+++ b/paper/images/createTask.graffle	Mon Apr 01 21:16:36 2013 +0900
@@ -6,7 +6,7 @@
 	<integer>0</integer>
 	<key>ApplicationVersion</key>
 	<array>
-		<string>com.omnigroup.OmniGrafflePro</string>
+		<string>com.omnigroup.OmniGraffle</string>
 		<string>139.16.0.171715</string>
 	</array>
 	<key>AutoAdjust</key>
@@ -14,7 +14,7 @@
 	<key>BackgroundGraphic</key>
 	<dict>
 		<key>Bounds</key>
-		<string>{{0, 0}, {559, 783}}</string>
+		<string>{{0, 0}, {558.99997329711914, 783}}</string>
 		<key>Class</key>
 		<string>SolidGraphic</string>
 		<key>ID</key>
@@ -57,6 +57,80 @@
 			<key>Head</key>
 			<dict>
 				<key>ID</key>
+				<integer>56</integer>
+			</dict>
+			<key>ID</key>
+			<integer>57</integer>
+			<key>Points</key>
+			<array>
+				<string>{279.5, 80}</string>
+				<string>{279.49998664855957, 114.99999999999994}</string>
+			</array>
+			<key>Style</key>
+			<dict>
+				<key>stroke</key>
+				<dict>
+					<key>HeadArrow</key>
+					<string>0</string>
+					<key>Legacy</key>
+					<true/>
+					<key>TailArrow</key>
+					<string>0</string>
+				</dict>
+			</dict>
+			<key>Tail</key>
+			<dict>
+				<key>ID</key>
+				<integer>3</integer>
+				<key>Info</key>
+				<integer>1</integer>
+			</dict>
+		</dict>
+		<dict>
+			<key>Bounds</key>
+			<string>{{227.49998664855957, 114.99999999999994}, {104, 27}}</string>
+			<key>Class</key>
+			<string>ShapedGraphic</string>
+			<key>ID</key>
+			<integer>56</integer>
+			<key>Magnets</key>
+			<array>
+				<string>{0, 1}</string>
+				<string>{0, -1}</string>
+				<string>{1, 0}</string>
+				<string>{-1, 0}</string>
+				<string>{1, 1}</string>
+				<string>{1, -1}</string>
+				<string>{-1, 1}</string>
+				<string>{-1, -1}</string>
+			</array>
+			<key>Shape</key>
+			<string>Rectangle</string>
+			<key>Style</key>
+			<dict>
+				<key>shadow</key>
+				<dict>
+					<key>Draws</key>
+					<string>NO</string>
+				</dict>
+			</dict>
+			<key>Text</key>
+			<dict>
+				<key>Text</key>
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
+\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
+{\colortbl;\red255\green255\blue255;}
+\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
+
+\f0\fs24 \cf0 mail}</string>
+			</dict>
+		</dict>
+		<dict>
+			<key>Class</key>
+			<string>LineGraphic</string>
+			<key>Head</key>
+			<dict>
+				<key>ID</key>
 				<integer>53</integer>
 			</dict>
 			<key>ID</key>
@@ -100,7 +174,7 @@
 			<integer>54</integer>
 			<key>Points</key>
 			<array>
-				<string>{279, 80}</string>
+				<string>{279.49998664855957, 141.99999999999994}</string>
 				<string>{94, 212}</string>
 			</array>
 			<key>Style</key>
@@ -115,6 +189,13 @@
 					<string>0</string>
 				</dict>
 			</dict>
+			<key>Tail</key>
+			<dict>
+				<key>ID</key>
+				<integer>56</integer>
+				<key>Info</key>
+				<integer>1</integer>
+			</dict>
 		</dict>
 		<dict>
 			<key>Bounds</key>
@@ -147,7 +228,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -186,7 +267,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -225,7 +306,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -264,7 +345,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -303,7 +384,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -347,7 +428,7 @@
 				<key>Pad</key>
 				<integer>0</integer>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -395,7 +476,7 @@
 				<key>Pad</key>
 				<integer>0</integer>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -535,7 +616,7 @@
 			<integer>43</integer>
 			<key>Points</key>
 			<array>
-				<string>{279, 80}</string>
+				<string>{279.49998664855957, 141.99999999999994}</string>
 				<string>{362, 212}</string>
 			</array>
 			<key>Style</key>
@@ -570,7 +651,7 @@
 			<integer>42</integer>
 			<key>Points</key>
 			<array>
-				<string>{279.5, 80}</string>
+				<string>{279.49998664855957, 141.99999999999994}</string>
 				<string>{210, 212}</string>
 			</array>
 			<key>Style</key>
@@ -588,7 +669,9 @@
 			<key>Tail</key>
 			<dict>
 				<key>ID</key>
-				<integer>3</integer>
+				<integer>56</integer>
+				<key>Info</key>
+				<integer>1</integer>
 			</dict>
 		</dict>
 		<dict>
@@ -733,7 +816,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -772,7 +855,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -811,7 +894,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -850,7 +933,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -889,7 +972,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -928,7 +1011,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -967,7 +1050,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -1006,7 +1089,7 @@
 			<key>Text</key>
 			<dict>
 				<key>Text</key>
-				<string>{\rtf1\ansi\ansicpg65001\cocoartf1187\cocoasubrtf340
+				<string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf370
 \cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;}
 {\colortbl;\red255\green255\blue255;}
 \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc
@@ -1062,9 +1145,9 @@
 	<key>MasterSheets</key>
 	<array/>
 	<key>ModificationDate</key>
-	<string>2013-02-18 08:52:29 +0000</string>
+	<string>2013-04-01 12:12:25 +0000</string>
 	<key>Modifier</key>
-	<string>Shinji KONO</string>
+	<string>yuhi</string>
 	<key>NotesVisible</key>
 	<string>NO</string>
 	<key>Orientation</key>
@@ -1093,7 +1176,7 @@
 		<key>NSPaperSize</key>
 		<array>
 			<string>size</string>
-			<string>{595, 842}</string>
+			<string>{594.99997329711914, 842}</string>
 		</array>
 		<key>NSPrintReverseOrientation</key>
 		<array>
@@ -1143,7 +1226,7 @@
 			</dict>
 		</array>
 		<key>Frame</key>
-		<string>{{1269, 383}, {693, 922}}</string>
+		<string>{{32, 116}, {693, 922}}</string>
 		<key>ListView</key>
 		<true/>
 		<key>OutlineWidth</key>
Binary file paper/images/createTask.pdf has changed
--- a/paper/opencl.tex	Mon Apr 01 19:09:39 2013 +0900
+++ b/paper/opencl.tex	Mon Apr 01 21:16:36 2013 +0900
@@ -25,33 +25,51 @@
 
 \subsection{メモリアクセス}
 host側は主にdataをinput/outputするメモリ資源の確保を行う。
-GPU上のメモリはCPUのメモリ空間と異なるので共有することができない(図\ref{fig:gpu_memory})。
-よってhost側で memory bufferを作成してメモリのコピーを行う。
+GPUのメモリ空間(図\ref{fig:gpuarch})やCellのメモリ空間(図\ref{fig:cellarch})
+はマルチコアCPU(図\ref{fig:cpuarch})と違い、共有メモリでないためhostとkernel(task)間でdataの共有ができない。
+アクセスするにはメモリ空間間でコピーしなければならない。
+
+GPGPUではhost側で memory bufferを作成してメモリのコピーを行う。
 これらの処理やTaskはCommand Queueにenqueueすることで実行される。
 \begin{figure}[htb]
   \begin{center}
-    \includegraphics[scale=0.40]{./images/gpu_arch.pdf}
+    \includegraphics[scale=0.3]{./images/gpu_arch.pdf}
   \end{center}
-  \caption{Gpu Memory Archtecture}
-  \label{fig:gpu_memory}
+  \caption{Gpu Archtecture}
+  \label{fig:gpuarch}
+\end{figure}
+
+\begin{figure}[ht]
+  \begin{center}
+    \includegraphics[scale=0.5]{./images/cell_arch.pdf}
+  \end{center}
+  \caption{Cell Archtecture}
+  \label{fig:cellarch}
 \end{figure}
 
-\newpage
+\begin{figure}[ht]
+  \begin{center}
+    \includegraphics[scale=0.5]{./images/cpu_arch.pdf}
+  \end{center}
+  \caption{Cpu Archtecture}
+  \label{fig:cpuarch}
+\end{figure}
 
 \subsection{データ並列}
-データを順に処理するだけでは並列度はあまり上がらない。
-そこで、データを2から3次元に分割し、分割した部分に対して並列に実行する。これはデータ並列と呼ばれている。
-OpenCLでデータ並列を行う場合、複数のKernelを作成しそれぞれのデータに割り当てる事で実現する。
+多次元のデータ構造がある場合に高い並列度を保つには、それを分割して並列に実行する機能が必要である。
+これをOpen CLではデータ並列と読んでいる。
+OpenCLは次元数に対応するindexがあり、openclは一つの記述から異なるindexを持つ複数のkernelを自動生成する。
+その添字をglobal\_idとよぶ
 この時入力されたデータはワークアイテムという処理単位に分割される。
 
-OpenCLはワークアイテムに対してそれぞれを識別するID(グローバルID)割り当てる。
+OpenCLはワークアイテムに対してそれぞれを識別するID(global\_id)を割り当てる。
 kernelはget\_global\_id APIによってIDを取得し、取得したIDに対応するデータに対して処理を行い、
 データ並列を実現する。
 このIDによって取得してきたワークアイテムをグローバルワークアイテムという。
-また、ワークアイテムは3次元までのデータを渡すことができ、渡したデータの次元数をdimensionという。
+また、ワークアイテムは3次元までのデータを渡すことができる。
 
 データ並列によるkernel実行の場合はclEnqueueNDRangeKernel APIを使用するが、
-この関数の引数としてワークアイテムのサイズ、dimensionを指定することでデータ並列で実行できる。
+この関数の引数としてワークアイテムのサイズと次元数を指定することでデータ並列で実行できる。
 
 \begin{figure}[htb]
   \begin{center}
@@ -63,8 +81,8 @@
 
 \subsection{ワークグループ}
 前節でワークアイテムという処理単位について述べたが、
-さらに複数個のグローバルワークアイテムをワークグループという単位にまとめることができる。
-ワークグループ内では同期やローカルメモリの共有が可能となる。
+さらに複数個のグローバルワークアイテムをwork\_groupという単位にまとめることができる。
+work\_group内では同期やローカルメモリの共有が可能となる。
 
 グローバルワークアイテム(ワークアイテム全体)の個数と、
 ローカルワークアイテム(グループ一つ辺りのアイテム)の個数を指定することでワークアイテムを分割する。
@@ -74,8 +92,7 @@
 ローカルアイテム数は0を指定することで、コンパイル時に最適化させることができる。
 したがってローカルアイテムのサイズは0を指定するのが一般的である。
 
-なお、ワークグループを設定した場合はグローバルIDの他にワーク
-グループID、ローカルIDが
+なお、work\_groupを設定した場合はglobal\_idの他にwork\_group\_id、local\_idが
 それぞれのkernelに割り当てられる(図:\ref{fig:workitem_id})。
 
 kernel側からそれぞれIDに対応したAPIを使用して、各IDを取得する。
@@ -89,17 +106,17 @@
       \small
       \begin{tabular}[t]{c|l}
         \hline
-        get\_group\_id & ワーグループIDを取得  \\
+        get\_group\_id & work\_group\_idを取得  \\
         \hline
-        get\_local\_id & ローカルIDを取得 \\
+        get\_local\_id & local\_idを取得 \\
         \hline
-        get\_global\_id &  グローバルIDを取得 \\
+        get\_global\_id &  global\_idを取得 \\
         \hline
       \end{tabular}
     \end{center}
   \end{table}
 \end{tiny}
-なお、ローカルid、グローバルidを取得するAPIは引数に0、1、2の値をsetすることができる。
+なお、local\_id、global\_idを取得するAPIは引数に0、1、2の値をsetすることができる。
 idはx,y,z座標があり、それぞれが0,1,2に対応している。
 例えばget\_global\_id(1)と呼び出した場合はy座標の、
-get\_global\_id(1)と呼び出した場合はz座標のグローバルidを取得する。
+get\_global\_id(1)と呼び出した場合はz座標のglobal\_idを取得する。
Binary file paper/sigos.pdf has changed
--- a/paper/sigos.tex	Mon Apr 01 19:09:39 2013 +0900
+++ b/paper/sigos.tex	Mon Apr 01 21:16:36 2013 +0900
@@ -47,11 +47,12 @@
 \end{abstract}
 % 英文概要
 \begin{eabstract}
-Cerium Task Manager is parallel programming frame work.
-Generally, in order to demonstrate sufficient performance in GPGPU using Open CL, various tuning is needed.
+Cerium Task Manager is a parallel programming framework.
+To achieve good performance in GPGPU using Open CL, various tuning is needed.
 In particular, it is necessary to implement the dependency of task in Cerium by the function of Open CL.
-However, the flexibility of  framework will be spoiled if it specializes in Open CL. 
-Flexibility and the balance of performance is considered based on Sort, Word count, and FFT example.
+But, to match specialization for OpenCL spoiles of flexibility of framework. 
+Balance of flexibility and the performance is considered.
+We evaluate example  Sort, Word count, and FFT.
 
 \end{eabstract}