tkokamoの日記

HPCの研究開発なのでそんなことをかきたい

FUSEでfilesystemを使ってみる。

仕事で分散ファイルシステムを作っているんだけども、カーネルレイヤでのコーディングなので、非常に神経使うし面倒（楽しいけれども）。
Filesystem in userspace(FUSE)というものの存在はずいぶん前から認知していて学生時代にsshfsなどで使っていた。
FUSEは実行パスが長いのでカーネル内でに比べるとソフトオーバーヘッドがでかいとか欠点もあるけれども、
アイデアを形にするときには重宝すると思ったので、まずは Writing a FUSE Filesystem: a Tutorial に沿って動かしてみる。

fuseのinstall

tutorialは2.9.Xを想定しているみたいだったが、手元のCentOS 7.2でyum使ってinstallしたら問題なくinstallされた。

[root@localhost fuse-tutorial-2018-02-04]# uname -r
3.10.0-693.17.1.el7.x86_64
[root@localhost fuse-tutorial-2018-02-04]# yum install fuse fuse-devel
...
Updated:
  fuse.x86_64 0:2.9.2-11.el7                                                 fuse-devel.x86_64 0:2.9.2-11.el7

Dependency Updated:
  fuse-libs.x86_64 0:2.9.2-11.el7

Complete!

触ってみる

Compiling and Runningに書かれていることをやるだけ。
bbfsのコードはhttp://www.cs.nmsu.edu/~pfeiffer/fuse-tutorial.tgz

# rootではやらせてくれなかったので通常ユーザになって続行

takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ ../src/bbfs rootdir/ mountdir/
Fuse library version 2.9
about to call fuse_main
takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ ls -lR
.:
total 8
-rw-r--r-- 1 takuya takuya  185 Jul 11 20:31 Makefile
-rw-r--r-- 1 takuya takuya 2003 Jul 11 20:32 bbfs.log
drwxr-xr-x 2 takuya takuya   22 Jul 11 20:31 mountdir
drwxr-xr-x 2 takuya takuya   22 Jul 11 20:31 rootdir

./mountdir:
total 4
-rw-r--r-- 1 takuya takuya 11 Jul 11 20:31 bogus.txt

./rootdir:
total 4
-rw-r--r-- 1 takuya takuya 11 Jul 11 20:31 bogus.txt
takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ touch rootdir/hoge　★rootdirにhogeを作った。
takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ ls -lR
.:
total 12
-rw-r--r-- 1 takuya takuya  185 Jul 11 20:31 Makefile
-rw-r--r-- 1 takuya takuya 6639 Jul 11 20:32 bbfs.log
drwxr-xr-x 2 takuya takuya   33 Jul 11 20:32 mountdir
drwxr-xr-x 2 takuya takuya   33 Jul 11 20:32 rootdir

./mountdir:
total 4
-rw-r--r-- 1 takuya takuya 11 Jul 11 20:31 bogus.txt
-rw-r--r-- 1 takuya takuya  0 Jul 11 20:32 hoge ★mountdirにもできている！

./rootdir:
total 4
-rw-r--r-- 1 takuya takuya 11 Jul 11 20:31 bogus.txt
-rw-r--r-- 1 takuya takuya  0 Jul 11 20:32 hoge

とりあえず触るのは問題なさそう。

レイテンシはいかに示す通り、FUSEを通すとかなり大きくなっている。
open（create含む）で約9倍、writeで約15倍。操作的にディスクアクセスがないのでソフトオーバーヘッドがもろに聞いている様子。

takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ strace -T touch rootdir/a 2>&1 | grep "rootdir/a"
execve("/usr/bin/touch", ["touch", "rootdir/a"], [/* 34 vars */]) = 0 <0.000195>
open("rootdir/a", O_WRONLY|O_CREAT|O_NOCTTY|O_NONBLOCK, 0666) = 3 <0.000063>
takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ strace -T touch mountdir/b 2>&1 | grep "mountdir/b"
execve("/usr/bin/touch", ["touch", "mountdir/b"], [/* 34 vars */]) = 0 <0.000290>
open("mountdir/b", O_WRONLY|O_CREAT|O_NOCTTY|O_NONBLOCK, 0666) = 3 <0.000533>

takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ strace -T dd if=/dev/zero of=rootdir/file00 bs=1M count=1 2>&1 | grep "write"
write(1, "\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0"..., 1048576) = 1048576 <0.001646>
write(2, "1+0 records in\n1+0 records out\n", 311+0 records in
write(2, "1048576 bytes (1.0 MB) copied", 291048576 bytes (1.0 MB) copied) = 29 <0.000037>
write(2, ", 0.00409826 s, 256 MB/s\n", 25, 0.00409826 s, 256 MB/s
takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ strace -T dd if=/dev/zero of=mountdir/file01 bs=1M count=1 2>&1 | grep "write"
write(1, "\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0"..., 1048576) = 1048576 <0.023698>
write(2, "1+0 records in\n1+0 records out\n", 311+0 records in
write(2, "1048576 bytes (1.0 MB) copied", 291048576 bytes (1.0 MB) copied) = 29 <0.000032>
write(2, ", 0.0261279 s, 40.1 MB/s\n", 25, 0.0261279 s, 40.1 MB/s

ディスクから読む場合も3倍程度かかっている。

# 読み込む前にecho 3 > /prco/sys/vm/drop_cachesでキャッシュを破棄する
takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ strace -T dd of=/dev/null if=rootdir/file00 bs=1M count=1 2>&1 | grep "read.*1048576"
read(0, "\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0"..., 1048576) = 1048576 <0.031173>
takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ strace -T dd of=/dev/null if=mountdir/file01 bs=1M count=1 2>&1 | grep "read.*1048576"
read(0, "\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0\0"..., 1048576) = 1048576 <0.092250>

少し中を見てみる

普通ファイルシステムはmountコマンド（mountシステムコール）で特定のディレクトリにマウントする。
その中でファイルシステムごとの初期化処理が行われるが、fuseの場合はmountのifはなくプログラムのmain関数で初期化を行う。

# 以下でマウント相当の初期化処理が行われる。
takuya@localhost:~/fuse_tutorial/fuse-tutorial-2018-02-04/example$ ../src/bbfs rootdir/ mountdir/

初期化としてやっているのは、以下の二つの認識
* ファイルシステム固有の初期化処理 * operationの登録 大雑把に言ってしまえばカーネルレイヤで作るFSも基本的には同じことをやっている。

ファイルシステム固有の初期化処理

bbfsは単にファイルアクセスをリダイレクトする（別ディレクトリ上で同じものを見せる）だけなので、大層な初期化は行っていない。

871 int main(int argc, char *argv[])
872 {
873     int fuse_stat;
874     struct bb_state *bb_data;
875
... // ここは入力チェックなど
901     bb_data = malloc(sizeof(struct bb_state)); // bbfsの固有データ構造の獲得
902     if (bb_data == NULL) {
903         perror("main calloc");
904         abort();
905     }
906
907     // Pull the rootdir out of the argument list and save it in my
908     // internal data
909     bb_data->rootdir = realpath(argv[argc-2], NULL); ★

ここでやっているのはrootdir（rootdir/の絶対パス）の登録。
ファイルの実態はrootdir/にあるので、mountdir/に対するアクセスがあったらrootdir/にアクセスするために内部でその情報をひかえている。

// 続き
906
907     // Pull the rootdir out of the argument list and save it in my
908     // internal data
909     bb_data->rootdir = realpath(argv[argc-2], NULL);
910     argv[argc-2] = argv[argc-1]; ★
911     argv[argc-1] = NULL;
912     argc--;
913
914     bb_data->logfile = log_open();
915
916     // turn over control to fuse
917     fprintf(stderr, "about to call fuse_main\n");
918     fuse_stat = fuse_main(argc, argv, &bb_oper, bb_data); // operationの登録
919     fprintf(stderr, "fuse_main returned %d\n", fuse_stat);
920
921     return fuse_stat;
922 }

★ではargv、argcを少しいじっている。この処理はいらないのでは？と思っている。
あまりfuseの内部構造については知らないがおそらくargv[argc-1]がfuseでアクセスするマウントポイント（mountdir/）となることがfuseの仕様で決まっているんだと思う。
最後にoperationの登録、その他を行っている。

operationの登録

個人的にはfuseの一つの大きな利点は、実装するべき関数群がカーネルレイヤで作る場合に比べてとても少ないことだと思っている。
bbfsでは34個の関数を実装しているが、カーネル内で作るのに比べてこれでもかなり少ないはず。
そして、それぞれの関数群が割とユーザが使うIFと一致しているので、カーネル内部の仕組みを理解する必要がほとんどない。 bbfsで登録されているのはbb_operでそのリストは以下の通り。

822 struct fuse_operations bb_oper = {
823   .getattr = bb_getattr,
824   .readlink = bb_readlink,
825   // no .getdir -- that's deprecated
826   .getdir = NULL,
827   .mknod = bb_mknod,
828   .mkdir = bb_mkdir,
829   .unlink = bb_unlink,
830   .rmdir = bb_rmdir,
831   .symlink = bb_symlink,
832   .rename = bb_rename,
833   .link = bb_link,
834   .chmod = bb_chmod,
835   .chown = bb_chown,
836   .truncate = bb_truncate,
837   .utime = bb_utime,
838   .open = bb_open,
839   .read = bb_read,
840   .write = bb_write,
841   /** Just a placeholder, don't set */ // huh???
842   .statfs = bb_statfs,
843   .flush = bb_flush,
844   .release = bb_release,
845   .fsync = bb_fsync,
822 struct fuse_operations bb_oper = {
823   .getattr = bb_getattr,
824   .readlink = bb_readlink,
825   // no .getdir -- that's deprecated
826   .getdir = NULL,
827   .mknod = bb_mknod,
828   .mkdir = bb_mkdir,
829   .unlink = bb_unlink,
830   .rmdir = bb_rmdir,
831   .symlink = bb_symlink,
832   .rename = bb_rename,
833   .link = bb_link,
834   .chmod = bb_chmod,
835   .chown = bb_chown,
836   .truncate = bb_truncate,
837   .utime = bb_utime,
838   .open = bb_open,
839   .read = bb_read,
840   .write = bb_write,
841   /** Just a placeholder, don't set */ // huh???
842   .statfs = bb_statfs,
843   .flush = bb_flush,
844   .release = bb_release,
845   .fsync = bb_fsync,

試しにopen処理を見てみる。

281
282 /** File open operation
283  *
284  * No creation, or truncation flags (O_CREAT, O_EXCL, O_TRUNC)
285  * will be passed to open().  Open should check if the operation
286  * is permitted for the given flags.  Optionally open may also
287  * return an arbitrary filehandle in the fuse_file_info structure,
288  * which will be passed to all file operations.
289  *
290  * Changed in version 2.2
291  */
292 int bb_open(const char *path, struct fuse_file_info *fi)
293 {
294     int retstat = 0;
295     int fd;
296     char fpath[PATH_MAX];
297
298     log_msg("\nbb_open(path\"%s\", fi=0x%08x)\n",
299             path, fi);
300     bb_fullpath(fpath, path); （１）
301
302     // if the open call succeeds, my retstat is the file descriptor,
303     // else it's -errno.  I'm making sure that in that case the saved
304     // file descriptor is exactly -1.
305     fd = log_syscall("open", open(fpath, fi->flags), 0); （２）
306     if (fd < 0)
307         retstat = log_error("open");
308
309     fi->fh = fd;
310
311     log_fi(fi);
312
313     return retstat;
314 }

やっていることは、（１）実際に読み込むrootdir/配下のパス名の作成、（２）open(2)の実行だけ。
（bbfsは言ってしまえば何もやっていないので、実際に独自のファイルシステムを作ろうとするとこの程度ではすまない）

fuseを使ってやろうとしていること

ファイルシステムの研究開発をしていているので、既存研究のシステムを使ってみたいことはある。
が、まぁやはり研究用のソースなのでまともに動かなくて萎えてしまうことが多い。
fuseをちょろっと見た感じ非常に簡単にPoCが作れそうなので、動かないなら余暇につくってしまおう、というのが魂胆。

openCLのコードを少し調べてみる（カーネルの実行）

openCLの最初の記事で動かしたコードについて見て行く。

記事：openCLを動かしてみる - tkokamoの日記

動かしたコードのソース：OpenCLで行列の掛け算をしてみた : 試行錯誤な日々

そもそも何をするコードだったのか、、、

並列かしているのは、行列のかけ算、と書かれていたが行列を一次元配列で表現しているらしい。
それを踏まえた上でカーネルコードをまずは見てみる。

~ 113         /* Set source code of kernel*/
~ 114         char *source_str =
~ 115         "__kernel void matrix_dot_matrix( \
~ 116                 __global const float* A, \  // 行列
~ 117                 __global const float* B, \  // 行列
~ 118                 __global float* Result, \  // かけ算結果いれる配列
~ 119                 const int wA, \    // Aのwidth (列数
~ 120                 const int wB) { \  // Bのwidth (列数
~ 121                 const int x = get_global_id(0); \
~ 122                 const int y = get_global_id(1); \
~ 123                 float value = 0; \
~ 124                 for (int i = 0; i < wA; ++i) { \
                                  // 行列を一次元で書いているので少しわかりづらいが
                                  // y*wAはy行目を見ている、ということ
~ 125                         int index_a = y * wA + i; \
                                  // こっちはx列目を見ている、ということ。
~ 126                         int index_b = i * wB + x; \
                                  // 下は各行列の要素取得して掛けているだけ
~ 127                         float elementA = A[index_a]; \
~ 128                         float elementB = B[index_b]; \
~ 129                         value = value + elementA * elementB; \
~ 130                         } \
~ 131                 Result[wB * y + x] = value; \
~ 132         }";

上のコードでopenCL specificな部分は大体以下のものかと思う。

~ 115         "__kernel void matrix_dot_matrix( \
__kernelはカーネルコードにつける。

~ 116                 __global const float* A, \  // 行列
~ 117                 __global const float* B, \  // 行列
~ 118                 __global float* Result, \  // かけ算結果いれる配列
__globalはデバイスのグローバルメモリにつける。このメモリは全てのwork groupから参照可能なデバイスメモリ。

~ 121                 const int x = get_global_id(0); \
~ 122                 const int y = get_global_id(1); \
get_global_idは、work itemのidを取得する函数。引数は次元。
CUDAでいうblockIdxとかthreadIdxを取ってきているイメージかな？

カーネル実行までの道程を読む

例のごとく、マニュアルは手元に、、、 前回（openCLのコードを少し調べてみる（デバイスの情報取得） - tkokamoの日記）の結果からopenCL 1.2であることが分かったので、1.2のmanualをみる。 OpenCL 1.2 Reference Pages

contextの作成

~ 135         /* Create OpenCL context */
~ 136         context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);

" Contexts are used by the OpenCL runtime for managing objects such as command-queues, memory, program and kernel objects and for executing kernels on one or more devices specified in the context."
Contextはよくあるものなので、特に説明は不要。
openClのruntimeがprogramなどを管理するために使う。

command queueの作成

~ 142         command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

コマンドキューは、デバイスに実行させる処理を入れるキュー。
今回のコードでは行列のかけ算をさせるカーネル一つを入れているが、
複数入れるとそれらは逐次dispatchされるんだろうか、、、

カーネル実行の準備

それぞれをいまいちちゃんと理解しないんだけど、実行前にいくつかやることがあるみたい。

~ 145         /* Create Kernel Program from the source */
~ 146         program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *      )&source_size, &ret);

マニュアルを見ると、

  • プログラムオブジェクトを作成してsource_strで指定されたプログラムをプログラムオブジェクトにロードする。
  • プログラムオブジェクトはcontextを作成したデバイスに結びつけられる。
    らしい。なんか必要そうなのはわかるけど、openCLのモデルを理解していないからふわっとしている。
~ 148         /* Build Kernel Program */
~ 149         ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

プログラムをビルドする。まぁ複数デバイスサポートしているのこれは必要でしょう。

~ 159         /* Create OpenCL Kernel */
~ 160         kernel = clCreateKernel(program, "matrix_dot_matrix", &ret);

カーネルオブジェクトの作成。ビルド後のプログラムのオブジェクト、と思えばよいのだろうか。

~ 167         /* Create Memory Buffer */
~ 168         matrixAMemObj = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, matrixAMem      Size, matrixA, &ret);
~ 169         matrixBMemObj = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, matrixBMem      Size, matrixB, &ret);
~ 170         matrixRMemObj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, matrixRMemSize, NULL, &ret);

カーネルの引数となるデバイスメモリの獲得を行っている。
ここで指定しているのは、

  • context
  • 獲得方法を決めるフラグ
    • CL_MEM_READ_ONLYはreadonlyのメモリ
    • CL_MEM_COPY_HOST_PTRは獲得したデバイスメモリに第3引数のホストメモリの値をコピーする
  • ホストメモリのポインタ
  • メモリのサイズ
  • エラーコードを格納する変数のアドレス

戻り値はbuffer object。フラグについては結構種類があるようだが、今後自分でプログラムを書くときに必要であれば調べてみる。

~ 176         /* Set OpenCL Kernel Parameters */
~ 177         ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&matrixAMemObj);
~ 178         ret |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&matrixBMemObj);
~ 179         ret |= clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&matrixRMemObj);
~ 180         ret |= clSetKernelArg(kernel, 3, sizeof(int), (void *)&wA);
~ 181         ret |= clSetKernelArg(kernel, 4, sizeof(int), (void *)&wB);

カーネルに渡す引数の設定

カーネルの実行

~ 256         /* Execute OpenCL Kernel */
~ 257         ret = clEnqueueNDRangeKernel(command_queue, kernel, workDim, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);

command_queueにkernelをenqueueする。
この時指定するのは、（第3引数から）

  • work itemが使う次元数。ここでは2次元
  • global work offset。ここではNULLだが、指定されるとglobalIDの開始offsetが指定したIDからになる。
  • global_work_size。各次元におけるwork sizeを指定する。全体的なwork sizeは各次元のサイズの積になる。この値については後述。
  • local_work_size。work groupのサイズを指定する。この値については後述。openCLではworkは（こういう言い方でよいのか？）複数のwork groupから構成され、work groupは複数のwork itemから構成される。work group内ではデバイスのlocal memoryが使えるので適切にworkを分割すると計算性能の向上が図れるんだろう、、、多分
  • num_events_in_wait_list。カーネルが実行される前に待ち合わせるイベントの数。同期に使われる。
  • event_wait_list。待ち合わせるeventの配列。
  • event。ここではNULLだが、ポインタが指定されるとここにenqueueしたカーネルのevent objectが入れられる。

local_work_sizeを決める為に、
まずはデバイスの最大work groupサイズと最大次元数を取得している。

~ 189         size_t maxWorkGroupSize;
~ 190         {
~ 191                 ret = clGetDeviceInfo(device_id, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGroupSize), &maxWorkGroupSize, NULL);


~ 198         size_t* maxLocalSizes;
~ 199         {
~ 200                 cl_uint maxWorkItemDims;
~ 201                 ret = clGetDeviceInfo( device_id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(cl_uin      t), &maxWorkItemDims, NULL);

次に得られた次元数分size_tの配列を獲得し、各次元の最大work itemサイズを取得している。

~ 206                 maxLocalSizes = (size_t*)malloc(maxWorkItemDims * sizeof(size_t));

~ 208                 ret = clGetDeviceInfo(device_id, CL_DEVICE_MAX_WORK_ITEM_SIZES, maxWorkItemDims* sizeof(size_t), maxLocalSizes, NULL);

そして最終的にlocal work sizeを以下の処理で決めている。

~ 220         size_t localWR = wR, localHR = hR;
~ 221         while (maxLocalSizes[0] < localWR ||
~ 222                         maxLocalSizes[1] < localHR ||
~ 223                         localWR * localHR > maxWorkGroupSize) {
~ 224                 if (maxLocalSizes[0] < localWR) {
~ 225                         if (localWR % 2 == 0) {
~ 226                                 localWR /= 2;
~ 227                         } else {
~ 228                                 localWR = getMaxCommonFactorOf2Pow(localWR);
~ 229                         }
~ 230                 } else if (maxLocalSizes[1] < localHR) {
~ 231                         if (localWR % 2 == 0) {
~ 232                                 localHR /= 2;
~ 233                         } else {
~ 234                                 localHR = getMaxCommonFactorOf2Pow(localHR);
~ 235                         }
~ 236                 } else if (localHR != 1) {
~ 237                         if (localHR % 2 == 0) {
~ 238                                 localHR /= 2;
~ 239                         } else {
~ 240                                 localHR = getMaxCommonFactorOf2Pow(localHR);
~ 241                         }
~ 242                 } else {
~ 243                         if (localWR % 2 == 0) {
~ 244                                 localWR /= 2;
~ 245                         } else {
~ 246                                 localWR = getMaxCommonFactorOf2Pow(localWR);
~ 247                         }
~ 248                 }
~ 249         }

結果のreadback

~ 264         ret = clEnqueueReadBuffer(command_queue, matrixRMemObj, CL_TRUE, 0, matrixRMemSize, matrixR      , 0, NULL, NULL);

終了処理

~ 271         /* Finalization */
~ 272         ret = clFlush(command_queue);
~ 273         ret = clFinish(command_queue);
~ 274         ret = clReleaseKernel(kernel);
~ 275         ret = clReleaseProgram(program);
~ 276         ret = clReleaseMemObject(matrixAMemObj);
~ 277         ret = clReleaseMemObject(matrixBMemObj);
~ 278         ret = clReleaseMemObject(matrixRMemObj);
~ 279         ret = clReleaseCommandQueue(command_queue);
~ 280         ret = clReleaseContext(context);

こんなところか、、、次は自前コードを動かそう。

openCLのコードを少し調べてみる（デバイスの情報取得）

前回はとりあえずコードを動かした。 http://tkokamo.hateblo.jp/

openCLを動かしてみる - tkokamoの日記

前回のコードを見るとカーネルコードの実装以前にclGetPlatformIDsやclGetDeviceIDsという函数が呼ばれていて気になったので、デバイスの情報取得周りについて調べてみた。

基本は前回のコードベースで気になったことを気が済むまで掘り下げる、ということにする。

今回のコード

下のコードがとりあえずまとまっていてよいなと思ったので、とりあえずこいつを動かしてみる。

https://gist.githubusercontent.com/tzutalin/51a821f15a735024f16b/raw/3f7217e967d73fcd522f337d1ffa6410f5d6ec4d/clDeviceQuery.cpp

[takuya@localhost device]$ g++ dev.cpp -lOpenCL
[takuya@localhost device]$ ./a.out 
clDeviceQuery Starting...

1 OpenCL Platforms found

 CL_PLATFORM_NAME:  Intel(R) OpenCL
 CL_PLATFORM_VERSION:   OpenCL 1.2 LINUX
OpenCL Device Info:

 1 devices found supporting OpenCL on: Intel(R) OpenCL

 ----------------------------------
 Device Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz
 ---------------------------------
  CL_DEVICE_NAME:           Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz
  CL_DEVICE_VENDOR:             Intel(R) Corporation
  CL_DRIVER_VERSION:            1.2.0.25
  CL_DEVICE_TYPE:           CL_DEVICE_TYPE_CPU
  CL_DEVICE_MAX_COMPUTE_UNITS:      8
  CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:   3
  CL_DEVICE_MAX_WORK_ITEM_SIZES:    8192 / 8192 / 8192 
  CL_DEVICE_MAX_WORK_GROUP_SIZE:    8192
  CL_DEVICE_MAX_CLOCK_FREQUENCY:    3500 MHz
  CL_DEVICE_ADDRESS_BITS:       64
  CL_DEVICE_MAX_MEM_ALLOC_SIZE:     7971 MByte
  CL_DEVICE_GLOBAL_MEM_SIZE:        31885 MByte
  CL_DEVICE_ERROR_CORRECTION_SUPPORT:   no
  CL_DEVICE_LOCAL_MEM_TYPE:     global
  CL_DEVICE_LOCAL_MEM_SIZE:     32 KByte
  CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:   128 KByte
  CL_DEVICE_QUEUE_PROPERTIES:       CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
  CL_DEVICE_QUEUE_PROPERTIES:       CL_QUEUE_PROFILING_ENABLE
  CL_DEVICE_IMAGE_SUPPORT:      1
  CL_DEVICE_MAX_READ_IMAGE_ARGS:    480
  CL_DEVICE_MAX_WRITE_IMAGE_ARGS:   480

  CL_DEVICE_IMAGE <dim>           2D_MAX_WIDTH     16384
                    2D_MAX_HEIGHT    16384
                    3D_MAX_WIDTH     2048
                    3D_MAX_HEIGHT    2048
                    3D_MAX_DEPTH     2048
  CL_DEVICE_PREFERRED_VECTOR_WIDTH_<t>    CHAR 1, SHORT 1, INT 1, FLOAT 1, DOUBLE 1


clDeviceQuery, Platform Name = Intel(R) OpenCL, Platform Version = OpenCL 1.2 LINUX, NumDevs = 1, Device = Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz

System Info: 

 Local Time/Date =  23:22:49, 04/04/2019
 CPU Name: Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz 
 # of CPU processors: 8
 Linux version 3.10.0-514.16.1.el7.x86_64 (builder@kbuilder.dev.centos.org) (gcc version 4.8.5 20150623 (Red Hat 4.8.5-11) (GCC) ) #1 SMP Wed Apr 12 15:04:24 UTC 2017


TEST PASSED

いくつか？？？となるものがあるので、函数の中身を実際に追っていきたいと思う。

関数の中身

mainから順に見て行く。マニュアルは手元に、、、

OpenCL 2.1 Reference Pages

main()

165         cl_int ciErrNum = clGetPlatformIDs(0, NULL, &num_platforms);

マニュアルには"Obtain the list of platforms available." だけ書かれていてぱっとしないのだが恐らくopenCLの実行環境?のIDを取ることが出来るんだろう。。。
（上では第1引数0で第2引数NULLなので、num_platformsに存在するplatform数を入れるみたい）

182                         ciErrNum = clGetPlatformIDs (num_platforms, clPlatformIDs, NULL);
183                         for(cl_uint i = 0; i < num_platforms; ++i) { // ここからplatform数情報表示する

実際にIDを取ってきている部分。

160         char cBuffer[1024];
...
184                                 ciErrNum = clGetPlatformInfo (clPlatformIDs[i], CL_PLATFORM_NAME, 1024, &cBuffer, NULL);
// CL_PLATFORM_NAME:   Intel(R) OpenCL

この関数は第1引数で指定されたIDのplatformについて第2引数で指定された値を取得する。
ここではplatform名を取得している。

198                                         ciErrNum = clGetPlatformInfo (clSelectedPlatformID, CL_PLATFORM_VERSION, sizeof(cBuffer), cBuffer, NULL);
// CL_PLATFORM_VERSION:    OpenCL 1.2 LINUX

platformのバージョン

210                                         // Get and log OpenCL device info
211                                         cl_uint ciDeviceCount;
212                                         cl_device_id *devices;
213                                         printf("OpenCL Device Info:\n\n");
214                                         ciErrNum = clGetDeviceIDs (clSelectedPlatformID, CL_DEVICE_TYPE_ALL, 0, NULL, &ciDeviceCount);

これはclGetPlatformIDs()のデバイスバージョン。で、ここではデバイス数を取得している。

235                                                 ciErrNum = clGetDeviceIDs (clSelectedPlatformID, CL_DEVICE_TYPE_ALL, ciDeviceCount, devices, &ciDeviceCount);

ここでデバイスIDを取得して、、、

237                                                         for(unsigned int i = 0; i < ciDeviceCount; ++i )  {
238                                                                 printf(" ----------------------------------\n");
239                                                                 clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(cBuffer), &cBuffer, NULL);
240                                                                 printf(" Device %s\n", cBuffer);
//  Device Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz
241                                                                 printf(" ---------------------------------\n");
242                                                                 clPrintDevInfo(devices[i]);
243                                                                 sProfileString += ", Device = ";
244                                                                 sProfileString += cBuffer;
245                                                         }

ここでデバイスの情報を取得／表示している。clPrintDevInfo()はコード内に定義されているので次はこの函数を見てみる。

clPrintDevInfo

ここまできたらなんとなく分かるものが多いので、そういうものははしょる。

 21         clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_string), &device_string, NULL);
//  CL_DEVICE_NAME:            Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz

 25         clGetDeviceInfo(device, CL_DEVICE_VENDOR, sizeof(device_string), &device_string, NULL);
//  CL_DEVICE_VENDOR:          Intel(R) Corporation

 29         clGetDeviceInfo(device, CL_DRIVER_VERSION, sizeof(device_string), &device_string, NULL);
//  CL_DRIVER_VERSION:             1.2.0.25

ドライバーって何のこと？と思ったのでマニュアルを見てみると、 "OpenCL software driver version string in the form major_number.minor_number."と書かれている。
結局OpenCLのバージョン？

 34         clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(type), &type, NULL);
 35         if( type & CL_DEVICE_TYPE_CPU )
 36                 printf("  CL_DEVICE_TYPE:\t\t\t%s\n", "CL_DEVICE_TYPE_CPU");
//  CL_DEVICE_TYPE:            CL_DEVICE_TYPE_CPU

 46         clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(compute_units), &compute_units,     NULL);
 47         printf("  CL_DEVICE_MAX_COMPUTE_UNITS:\t\t%u\n", compute_units);
//  CL_DEVICE_MAX_COMPUTE_UNITS:       8

まぁスレッド数でしょうと思いつつも説明を見てみる。
"The number of parallel compute units on the OpenCL device. A work-group executes on a single compute unit. The minimum value is 1."
GPUとかだとどうなるんだろうか？（ソースコード配布されているgit上でGPUの情報があるが、これは2だった、、、）

 51         clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(workitem_dims), &workitem    _dims, NULL);
//  CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:    3

"Maximum dimensions that specify the global and local work-item IDs used by the data parallel execution model. (Refer to clEnqueueNDRangeKernel). The minimum value is 3 for devices that are not of type CL_DEVICE_TYPE_CUSTOM."
なんだろう、、、CUDAでいうgrid, block, threadみたいなものなんだろうか？

 56         clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(workitem_size), &workitem_size    , NULL);
//  CL_DEVICE_MAX_WORK_ITEM_SIZES: 8192 / 8192 / 8192

各DIMENTIONで指定できる最大のwork-items、、、らしい。

こういうのはカーネルを動かす時にもっとわかると思う。

 61         clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(workgroup_size), &workgroup_si    ze, NULL);
//  CL_DEVICE_MAX_WORK_GROUP_SIZE: 8192

"Maximum number of work-items in a work-group that a device is capable of executing on a single compute unit"
ふむ、これもカーネル動かす時に意識するのかな、、、
clEnqueueNDRangeKernel()が関連しているらしいので、カーネルのコード実行部分のプログラムをみる時に調べてみる。

 66         clGetDeviceInfo(device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(clock_frequency), &clock_frequ    ency, NULL);
//  CL_DEVICE_MAX_CLOCK_FREQUENCY: 3500 MHz

 71         clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS, sizeof(addr_bits), &addr_bits, NULL);
//  CL_DEVICE_ADDRESS_BITS:        64

 76         clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(max_mem_alloc_size), &max_mem_a    lloc_size, NULL);
//  CL_DEVICE_MAX_MEM_ALLOC_SIZE:      7971 MByte

"Max size of memory object allocation in bytes. The minimum value is max(min(102410241024, 1/4th of CL_DEVICE_GLOBAL_MEM_SIZE), 3210241024) for devices that are not of type CL_DEVICE_TYPE_CUSTOM"
なので 1/4th of CL_DEVICE_GLOBAL_MEM_SIZE（メモリ容量）が取られたんでしょう。
これはKBか？

 81         clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(mem_size), &mem_size, NULL);
//  CL_DEVICE_GLOBAL_MEM_SIZE:     31885 MByte

メモリ容量

 86         clGetDeviceInfo(device, CL_DEVICE_ERROR_CORRECTION_SUPPORT, sizeof(error_correction_support)    , &error_correction_support, NULL);
//  CL_DEVICE_ERROR_CORRECTION_SUPPORT:    no

メモリがECCをサポートしているか、これはdmpidecodeでみることができる。

[takuya@localhost ~]$ sudo dmidecode --type memory
[sudo] password for takuya: 
# dmidecode 3.0
# SMBIOS entry point at 0x000f04c0
Found SMBIOS entry point in EFI, reading table from /dev/mem.
SMBIOS 2.7 present.

Handle 0x0042, DMI type 17, 34 bytes
Memory Device
    Array Handle: 0x0043
    Error Information Handle: Not Provided ★
    Total Width: 64 bits ★
    Data Width: 64 bits ★ 
    Size: 8192 MB
    Form Factor: DIMM
    Set: None
    Locator: ChannelA-DIMM0
    Bank Locator: BANK 0
    Type: DDR3
    Type Detail: Synchronous
    Speed: 1333 MHz
    Manufacturer: 0215
    Serial Number: 00000000
    Asset Tag: 9876543210
    Part Number: CMZ32GX3M4A1866C10
    Rank: 2
    Configured Clock Speed: 1333 MHz

 91         clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_TYPE, sizeof(local_mem_type), &local_mem_type, N    ULL);

  CL_DEVICE_LOCAL_MEM_TYPE:       global

デバイス固有にsramのようなメモリ(キャッシュは対応しない？、、、)を持っている場合はlocal、そうでない場合はglobalらしい。
globalやlocalのモデルは下の記事が詳しい。openCLのモデルについては後日調べるので今回はつっこまない。

https://www.mql5.com/ja/articles/407

 95         clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(mem_size), &mem_size, NULL);

  CL_DEVICE_LOCAL_MEM_SIZE:     32 KByte

ローカルメモリサイズ。ローカル、といっても全てのwork itemからアクセス可能なメモリになるらしい。（上記の記事から）
manualには最低32KBと書かれているので、それが出ているのか、
使っているCPUのL1 DCACHEのサイズが32KBだからそれが出ているのか分からない。
（L1はコアごとだけどコヒーレンシあるので、全work itemからアクセス可能とみなせるだろう、、、多分）

 99         clGetDeviceInfo(device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(mem_size), &mem_size, NUL    L);
  CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:   128 KByte

コンスタントバッファとして獲得できる最大サイズ。globalメモリの一部らしい。

  CL_DEVICE_QUEUE_PROPERTIES:        CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
  CL_DEVICE_QUEUE_PROPERTIES:       CL_QUEUE_PROFILING_ENABLE
  CL_DEVICE_IMAGE_SUPPORT:      1
  CL_DEVICE_MAX_READ_IMAGE_ARGS:    480
  CL_DEVICE_MAX_WRITE_IMAGE_ARGS:   480

  CL_DEVICE_IMAGE <dim>           2D_MAX_WIDTH     16384
                    2D_MAX_HEIGHT    16384
                    3D_MAX_WIDTH     2048
                    3D_MAX_HEIGHT    2048
                    3D_MAX_DEPTH     2048
  CL_DEVICE_PREFERRED_VECTOR_WIDTH_<t>    CHAR 1, SHORT 1, INT 1, FLOAT 1, DOUBLE 1

イメージとは、、、なんだろう、、、とりあえず置いとく。

適当だけどひとまずこんなところか

openCLを動かしてみる

HPC関連の仕事してるくせに（システムソフトを作っているので）書いたことないから、作業記録を込めて簡単に動かしてみる。

openCLとは

wikipediaをみれば多分わかる。（自分は専門じゃないのでwikipediaが合っているかどうかはわからない） OpenCL - Wikipedia

GPUだとCUDAのイメージだけど、CUDAで買いたカーネルコードはマルチコアCPUでは動かないよね。

openCLはクロスプラットフォーム対応なので、ハードウェアの違いを吸収してくれるからうれしい、、、らしい。

もちろんCUDAにはCUDAだからGPUを使いこなせる部分があるらしいが、、、こういう話はおいおいで良い。

動かすコード

ぱっと調べた感じ下の記事がよさそうだったので、紹介されているコードを試してみた。

OpenCLで行列の掛け算をしてみた : 試行錯誤な日々

普通に動いた。

[takuya@localhost c_opencl_practice]$ ls
README.md  hello.cl           hello_in_2_dims.cl   matrix_dot_matrix_on_cpu.c
hello.c    hello_in_2_dims.c  matrix_dot_matrix.c  multi_size_test.sh
[takuya@localhost c_opencl_practice]$ gcc matrix_dot_matrix.c -lOpenCL
[takuya@localhost c_opencl_practice]$ ./a.out 
matrix
   0.00  10.00  20.00  30.00  40.00  50.00  60.00  70.00  80.00  90.00
   1.00  11.00  21.00  31.00  41.00  51.00  61.00  71.00  81.00  91.00
   2.00  12.00  22.00  32.00  42.00  52.00  62.00  72.00  82.00  92.00
   3.00  13.00  23.00  33.00  43.00  53.00  63.00  73.00  83.00  93.00
   4.00  14.00  24.00  34.00  44.00  54.00  64.00  74.00  84.00  94.00
   5.00  15.00  25.00  35.00  45.00  55.00  65.00  75.00  85.00  95.00
   6.00  16.00  26.00  36.00  46.00  56.00  66.00  76.00  86.00  96.00
   7.00  17.00  27.00  37.00  47.00  57.00  67.00  77.00  87.00  97.00
   8.00  18.00  28.00  38.00  48.00  58.00  68.00  78.00  88.00  98.00
   9.00  19.00  29.00  39.00  49.00  59.00  69.00  79.00  89.00  99.00
matrix
   1.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00
   0.00   2.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00
   0.00   0.00   1.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00
   0.00   0.00   0.00   1.00   0.00   0.00   0.00   0.00   0.00   0.00
   0.00   0.00   0.00   0.00   1.00   0.00   0.00   0.00   0.00   0.00
   0.00   0.00   0.00   0.00   0.00   1.00   0.00   0.00   0.00   0.00
   0.00   0.00   0.00   0.00   0.00   0.00   1.00   0.00   0.00   0.00
   0.00   0.00   0.00   0.00   0.00   0.00   0.00   1.00   0.00   0.00
   0.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00   1.00   0.00
   0.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00   0.00   1.00
CL_DEVICE_MAX_WORK_GROUP_SIZE: 8192
work item sizes:  8192 8192 8192
localWorkSize: 10, 10
matrix
   0.00  20.00  20.00  30.00  40.00  50.00  60.00  70.00  80.00  90.00
   1.00  22.00  21.00  31.00  41.00  51.00  61.00  71.00  81.00  91.00
   2.00  24.00  22.00  32.00  42.00  52.00  62.00  72.00  82.00  92.00
   3.00  26.00  23.00  33.00  43.00  53.00  63.00  73.00  83.00  93.00
   4.00  28.00  24.00  34.00  44.00  54.00  64.00  74.00  84.00  94.00
   5.00  30.00  25.00  35.00  45.00  55.00  65.00  75.00  85.00  95.00
   6.00  32.00  26.00  36.00  46.00  56.00  66.00  76.00  86.00  96.00
   7.00  34.00  27.00  37.00  47.00  57.00  67.00  77.00  87.00  97.00
   8.00  36.00  28.00  38.00  48.00  58.00  68.00  78.00  88.00  98.00
   9.00  38.00  29.00  39.00  49.00  59.00  69.00  79.00  89.00  99.00
Set matrix time: 0.000000 sec
Load device time: 0.000000 sec
Set kernel time: 0.230000 sec
Set memory time: 0.000000 sec
Set work size time: 0.000000 sec
Calc time: 0.000000 sec
Release resource time: 0.010000 sec
Total time: 0.240000 sec

ちゃんと並列化されているんだろうか、、、と思ったらちゃんと並列化されていない版もあった。 確かに違う。

●並列化した版
[takuya@localhost c_opencl_practice]$ ./a.out 2000
CL_DEVICE_MAX_WORK_GROUP_SIZE: 8192
work item sizes:  8192 8192 8192
localWorkSize: 2000, 1
Set matrix time: 0.030000 sec
Load device time: 0.010000 sec
Set kernel time: 0.170000 sec
Set memory time: 0.000000 sec
Set work size time: 0.000000 sec
Calc time: 5.090000 sec
Release resource time: 0.010000 sec
Total time: 5.310000 sec

●並列化しない版
[takuya@localhost c_opencl_practice]$ gcc matrix_dot_matrix_on_cpu.c 
[takuya@localhost c_opencl_practice]$ ./a.out 2000
Set matrix time: 0.040000 sec
Calc time: 26.190000 sec
Total time: 26.230000 seconds

とりあえず、今日はここまで。 次回はコードを見ていきたい。

「メモ」hwlocでノード構成を表示

NUMAのマシンなどでどのノードにどのデバイスが接続されていることを知るのにhwlocコマンドが役立つ。 以下はqemuでNUMA2ノード構成、nvmeデバイスをエミュレートした場合の構成

f:id:tkokamo:20180505192424p:plain

シグナルのブロック／無視の覚書き

少し調べていたのでメモ。基本的にはカーネル内の話。コードはv4.9.6

シグナルに関しては言葉がややこしいので、以下のようにこの記事では使う。

  • シグナルの登録 シグナルが生成され対象のタスクのpendingリストに繋がった状態。

  • シグナルの受信 シグナルが登録されたことがタスクに通知された状態。（TIF_SIGPENDINGがたっている状態）

  • シグナルの配信 受信したシグナルに応じた処理を行う。

Linuxのシグナルには「ブロックする」と「無視する」の二つがある。 違いは以下の通り

  • ブロック

シグナルは対象のタスクに登録されるが、受信されない。 struct task_structのblocked、real_blockedで指定されたシグナルがブロックされる。

struct task_struct {
#ifdef CONFIG_THREAD_INFO_IN_TASK
        /*   
         * For reasons of header soup (see current_thread_info()), this
         * must be the first element of task_struct.
         */
...
        sigset_t blocked, real_blocked;
...
  • 無視

そもそもシグナルが登録されない。

シグナルを送る__send_signal()では以下のようにブロックと無視が実現されている。

static int __send_signal(int sig, struct siginfo *info, struct task_struct *t,
                        int group, int from_ancestor_ns)
{
        struct sigpending *pending;
...
        result = TRACE_SIGNAL_IGNORED;
        if (!prepare_signal(sig, t, ★
                        from_ancestor_ns || (info == SEND_SIG_FORCED)))
                goto ret; 

        pending = group ? &t->signal->shared_pending : &t->pending;
...
        q = __sigqueue_alloc(sig, t, GFP_ATOMIC | __GFP_NOTRACK_FALSE_POSITIVE,
                override_rlimit);
        if (q) {
                // ここでリストに登録
                list_add_tail(&q->list, &pending->list);
                switch ((unsigned long) info) {
...
out_set:
        signalfd_notify(t, sig);
        sigaddset(&pending->signal, sig);
        complete_signal(sig, t, group); ★
ret:
        trace_signal_generate(sig, info, t, group, result);
        return ret;
}

そもそもシグナルを登録するかどうかがprepare_signal()で判断される。 ブロックするものは登録される。無視されるものは登録すらされない。

static bool prepare_signal(int sig, struct task_struct *p, bool force)
{
        struct signal_struct *signal = p->signal;
        struct task_struct *t;
        sigset_t flush;
...
        return !sig_ignored(p, sig, force); ★
}


static int sig_ignored(struct task_struct *t, int sig, bool force)
{
        /*
         * Blocked signals are never ignored, since the
         * signal handler may change by the time it is
         * unblocked.
         */
// コメントに書かれている通り、blockedに登録されていてもシグナルは無視されない
// ブロック状態のシグナルでもprepare_signal()が真になり、登録される。
        if (sigismember(&t->blocked, sig) || sigismember(&t->real_blocked, sig))
                return 0;

        if (!sig_task_ignored(t, sig, force)) ★
                return 0;

        /*
         * Tracers may want to know about even ignored signals.
         */
// ptraceされていない場合はシグナルは無視される。
        return !t->ptrace;
}


static int sig_task_ignored(struct task_struct *t, int sig, bool force)
{
        void __user *handler;

        handler = sig_handler(t, sig);

        if (unlikely(t->signal->flags & SIGNAL_UNKILLABLE) &&
                        handler == SIG_DFL && !force)
                return 1;

        return sig_handler_ignored(handler, sig); ★
}

static int sig_handler_ignored(void __user *handler, int sig)
{
        /* Is it explicitly or implicitly ignored? */
// handlerがSIG_IGNになっている場合は無視
// handlerがSIG_DFLの場合は、sig_kernel_ignoreが真なら無視（↓参照）
        return handler == SIG_IGN ||
                (handler == SIG_DFL && sig_kernel_ignore(sig));
}

#define sig_kernel_ignore(sig)          siginmask(sig, SIG_KERNEL_IGNORE_MASK)

#define SIG_KERNEL_IGNORE_MASK (\
        rt_sigmask(SIGCONT)   |  rt_sigmask(SIGCHLD)   | \
        rt_sigmask(SIGWINCH)  |  rt_sigmask(SIGURG)    )

parepare_signal()が真の場合、シグナルが生成され登録される。 最後にタスクを起床し、受信させるかどうかの判断はcomplete_signal()で行われる。

static void complete_signal(int sig, struct task_struct *p, int group)
{
        struct signal_struct *signal = p->signal;
        struct task_struct *t;

        /*
         * Now find a thread we can wake up to take the signal off the queue.
         *
         * If the main thread wants the signal, it gets first crack.
         * Probably the least surprising to the average bear.
         */
        if (wants_signal(sig, p))  ★
                t = p;  
        else if (!group || thread_group_empty(p))
                /*
                 * There is just one thread and it does not need to be woken.
                 * It will dequeue unblocked signals before it runs again.
                 */
// wants_signal()が成立せず、プロセスグループ内にスレッドが1個なら何もせず復帰する。
                return;
...
        /*
         * The signal is already in the shared-pending queue.
         * Tell the chosen thread to wake up and dequeue it.
         */
        signal_wake_up(t, sig == SIGKILL); // ここで起床する。
        return;
}

static inline int wants_signal(int sig, struct task_struct *p)
{
// blockedになっていたら0が返る。
        if (sigismember(&p->blocked, sig)) 
                return 0; 
        if (p->flags & PF_EXITING)
                return 0;
        if (sig == SIGKILL)
                return 1;
        if (task_is_stopped_or_traced(p))
                return 0;
        return task_curr(p) || !signal_pending(p);
}

無視されず、ブロックされないシグナルの場合、対象のタスクはsignal_wake_up()で起床させられる。

カーネルスレッドでは親のkthreaddが全シグナルのハンドラをSIG_IGNにしているため、 通常は全シグナルが無視される。

int kthreadd(void *unused)
{
        struct task_struct *tsk = current;

        /* Setup a clean context for our children to inherit. */
        set_task_comm(tsk, "kthreadd");
        ignore_signals(tsk); ★
...

void ignore_signals(struct task_struct *t)
{
        int i; 

        for (i = 0; i < _NSIG; ++i)
                t->sighand->action[i].sa.sa_handler = SIG_IGN;

        flush_signals(t);
}

sigfillset()などの関数を使うと、blockedを一括してセットすることが可能。

static inline void sigfillset(sigset_t *set)
{
        switch (_NSIG_WORDS) {
        default:
                memset(set, -1, sizeof(sigset_t));
                break;
        case 2: set->sig[1] = -1;
        case 1: set->sig[0] = -1;
                break;
        }
}

RCU(Read Copy Update)をちゃんと知る（2）-１ RCUにおけるポインタの更新と参照

はじめに

前回(http://tkokamo.hateblo.jp/entry/2017/08/12/191142)で概要は終わりにし、 今回からはコードを見ていこうと思います。

大体以下のような流れでRCUの実装を見ていこうかと思います。

  • rcu_dereference(), rcu_assign_pointer() ← 今回
  • Classic RCU
  • Tree RCU

rcu_read_lock()、rcu_read_unlock()はプリエンプションの無効化/有効化である、ということは前回述べたので触れないです。

RCUでのポインタの更新（と参照）

概要(http://tkokamo.hateblo.jp/entry/2017/05/31/213427)で述べたように、
RCUにおけるデータの更新は、更新後データの初期化を行った後、古いデータを指していたポインタを新しいデータに置き換えることで行われます。（下図参照）

f:id:tkokamo:20170816144144p:plain

更新後データは、値の初期化後にポインタによって指されるため、参照側に中途半端な状態(更新中のデータ)が見えることは無いはずですが、
現在のシステムでこれを保証するには少しだけ細工をしなければいけません。

データの初期化後にポインタで指す、という例として以下にコードを示します。

struct foo {
   int a;
   int b;
};

struct foo *foop = NULL;
struct foo new;

/*foop!=NULLで値を出力*/
void reader()
{
    int a, b;
    if (foop != NULL) {
        a = foop->a;
        b = foop->b;
        printk("Set. a = %d, b = %d\n", a, b);
    }
    return;
}

/*newを初期化しfoopにセット*/
void writer()
{
    new.a = 1;     // (1)
    new.b = 1;     // (2)
    foop = &new;   // (3)
    return;
}

reader()は、writer()によりnewがfoopにセットされていればその値を読み出し出力する、という関数ですが、
writer()の(1), (2), (3)の順序が変わってしまう場合には、Set. a = 1, b = 1が出力されるとは限りません。
そして、これらは以下の2つのメモリアクセス順序の変更によって実際に起きることがあります。

  • コンパイラによるメモリアクセス順序の変更
    静的なリオーダリング
  • CPUによるメモリアクセス順序の変更（目に見えるのはSMP）
    動的なリオーダリング

RCUにおけるポインタの更新と参照は、上記2種類のリオーダリングを適切に防ぐことで実装されています。 以下ではそれぞれのリオーダリングについて見たあとに、rcu_dereference()、rcu_assign_pointer()の実装を見てみます。

コンパイラによるメモリアクセス順序の変更

人間はしばしばおバカなコードを書いてしまい、メモリアクセスの観点から効率が悪いことがありますが、
コンパイラはその尻拭いをしてくれます。

例を下に示します。

// おばかなコード
#include <stdio.h> 
 
int main() 
{ 
        unsigned long arr[16]; 
        int i; 
 
        arr[0] = 1; 
        arr[8] = 1; 
        arr[1] = 1; 
        arr[9] = 1; 
        arr[2] = 1; 
        arr[10] = 1; 
        arr[3] = 1; 
        arr[11] = 1; 
        arr[4] = 1; 
        arr[12] = 1; 
        arr[5] = 1; 
        arr[13] = 1; 
        arr[6] = 1; 
        arr[14] = 1; 
        arr[7] = 1; 
        arr[15] = 1; 
 
        for (i = 0; i < 16; i++) 
                printf("%lu", arr[i]); 
 
        printf("\n"); 
        return 0; 
}

上のコードは二つのキャッシュラインを交互に見て1を代入しているおばかなコードです。（家の環境のi7-2600は確かキャッシュラインサイズが64B）
これをgccで最適化なしでコンパイルすると、、、

; 最適化なしの代入コード
        movq    $1, -144(%rbp)  ;arr[0] = 1
        movq    $1, -80(%rbp)   ;arr[8] = 1
        movq    $1, -136(%rbp)  ;arr[1] = 1
        movq    $1, -72(%rbp)   ;arr[9] = 1
...

馬鹿な人間に忠実なコードを吐いてくれますが、-O3をつけると

; -O3付き
        movdqa  %xmm0, (%rsp)
        movdqa  %xmm0, 16(%rsp)
        movdqa  %xmm0, 32(%rsp)
...

となります。movdqa命令はいわゆるSIMD命令です。
-O3を付けてコンパイルをすると、配列アクセスがプログラマが書いたようにではなく、ストリーム的にアクセスされるよう最適化されています(ストリームというにはしょぼいか)。
上記のようなコードであれば、特に意味も変わりませんが、代入の順番が大きな意味を持つ場合、最適化されてしまっては困ります。

gccではasm volatile("" ::: "memory");を入れることで、その前後のメモリアクセス順序がコンパイラによって変更されることを防ぐことができます。
今回の例の各代入文の間にasm volatile("" ::: "memory");を入れると、-O3を付けたとしても、メモリアクセス順序が人間の意図したままに維持されます。（興味があれば試してみてください）

CPUによるメモリアクセス順序の変更

こちらの方が少し理解に苦しむかも知れませんが、コンパイラ同様CPUもメモリアクセス順序を変更し、プログラム実行の最適化を図っています。
CPUによるリオーダリングが効いてくる例を見てみます。

// mem.c
#include <stdio.h>
#include <pthread.h>

struct arg {
        int a;
        int b;
        int c;
        int d;
};

void *t1_func(void *p)
{
    struct arg *arg = (struct arg *)p;

    arg->a = 1;
    if (arg->b == 0)
        arg->c = 1;

    return NULL;
}

void *t2_func(void *p)
{
    struct arg *arg = (struct arg *)p;

    arg->b = 1;
    if (arg->a == 0)
        arg->d = 1;
    
    return NULL;
}

int main()
{
    pthread_t t1, t2;
    struct arg arg;       

        /* 全部0で初期化 */
    arg.a = 0;
    arg.b = 0;
    arg.c = 0;
    arg.d = 0;

        /* スレッド生成 */
    pthread_create(&t1, NULL, t1_func, &arg);
    pthread_create(&t2, NULL, t2_func, &arg);

    /* スレッド終了待ち */
    pthread_join(t1, NULL);
    pthread_join(t2, NULL);

        /* a,b,c,dを出力 */
    printf("%d%d%d%d\n", arg.a, arg.b, arg.c, arg.d);
    return 0;
}

上記コードは、a,b,c,dを0で初期化したのちに2つのスレッドを生成し、各スレッドは以下の処理を行います。

// スレッド1
    arg->a = 1;
    if (arg->b == 0)     // (1)
        arg->c = 1;
// スレッド2
    arg->b = 1;
    if (arg->a == 0)     // (2)
        arg->d = 1;

親スレッドはスレッド1、スレッド2の終了後、a,b,c,dの値を出力しますが、
このコードは出力が1110または1101になることを期待しています。
その理由は以下の通りです。

  • (1)が成立するということは、スレッド1のarg->a = 1は実行されているが、スレッド2のarg->b = 1は実行されていない状態である。
    この時、スレッド2の(2)は成り立たず、4値の出力は1110となる。

  • (2)が成立するということは、スレッド2のarg->b = 1は実行されているが、スレッド2のarg->a = 1は実行されていない状態である。
    この時、スレッド1の(1)は成り立たず、4値の出力は1101となる。

しかし、SMP環境では1111が出力されてしまうことがあります。
以下では、上記コードを10000回実行し、出力が1111となってしまう回数を調べています。

答えが先になってしまいますが、1111が出力される理由はCPUによってメモリアクセス順序が変更されているからです。
下では、上記コードに対し、特定箇所にコンパイラ/CPUによるリオーダリングを禁止する命令を追加した場合、
シングルコアで上記コードを実行した場合に1111がどれくらい出力されてしまうかも合わせて示しています。

◯ 上で示したコード（SMP環境）
# for i in `seq 1 10000`; do ./mem >> mem.log; done
# grep 1111 mem.log | wc
    470     470    2350  ★470回1111が出力された

◯ 上で示したコード（シングルコア）
# echo 0 > /sys/devices/system/cpu/cpu1/online
# echo 0 > /sys/devices/system/cpu/cpu2/online
# echo 0 > /sys/devices/system/cpu/cpu3/online
# echo 0 > /sys/devices/system/cpu/cpu4/online
# echo 0 > /sys/devices/system/cpu/cpu5/online
# echo 0 > /sys/devices/system/cpu/cpu6/online
# echo 0 > /sys/devices/system/cpu/cpu7/online
# lscpu | grep On-line
On-line CPU(s) list:   0
# for i in `seq 1 10000`; do ./mem >> mem_uni.log; done
# grep 1111 mem_uni.log | wc
      0       0       0  ★1111は出力されなかった

◯ 特定箇所のコンパイラによるメモリリオーダリングを禁止したコード
# for i in `seq 1 10000`; do ./mem_comp >> mem_comp.log; done
# grep 1111 mem_comp.log | wc
    487     487    2435  ★487回1111が出力された

◯ 特定箇所のCPUによるメモリリオーダリングを禁止したコード
# for i in `seq 1 10000`; do ./mem_cpu >> mem_cpu.log; done
# grep 1111 mem_cpu.log | wc
      0       0       0  ★1111は出力されなかった

出力結果が1111となってしまう、ということはスレッド1の(1)、スレッド2の(2)が成り立ってしまった、ということです。
そして、この原因はCPUによるメモリアクセス順序の変更にあります。

/*再掲*/
// スレッド1
    arg->a = 1;
    if (arg->b == 0)     // (1)
        arg->c = 1;
// スレッド2
    arg->b = 1;
    if (arg->a == 0)     // (2)
        arg->d = 1;

SMP環境において、CPUは依存関係の無い二つのメモリアドレスに対するアクセス順序を動的に変更し、プログラム実行を高速化しています。
メモリアクセスはloadかstoreですから、依存関係の無い二つのメモリアドレスに関してload -> load、load -> store、store -> load、store -> storeの4種類のどのアクセス順序を変更して良いか、というポリシーが各CPUで決められています。
x86は比較的メモリアクセス順序を守るほうなのですが、store -> loadに関してはその順序を入れ替えても良い（コア間で見え方が違っても良い）、というポリシーになっています。(そうでないx86のCPUもあります)
上記の例では、このリオーダリングが見事に現れました。
つまり、スレッド1とスレッド2が別コアで動いた時、それぞれのarg->a = 1、arg->b = 1は互いに反映される前に、loadが先行され(1)、(2)が成立してしまったわけです。
コンパイラ同様、CPUによるメモリアクセスのリオーダリングを防ぐ方法もあり、asm volatile("mfence" ::: "memory");がこれに該当します。
この1行を以下の箇所に入れることで、期待通り1111の出力を防ぐことができます。

// mem.cの改良
void *t1_func(void *p)
{
...
    arg->a = 1;
    asm volatile("mfence" ::: "memory");
    if (arg->b == 0)
...
}

void *t2_func(void *p)
{
...
    arg->b = 1;
    asm volatile("mfence" ::: "memory");
    if (arg->a == 0)
...
}

CPUのメモリアクセスのリオーダリングなんて無い（プログラムからは順番どおりに見える）ほうが良いじゃないか！という気持ちはわかりますが、
異なるコア間、ましてやNUMAなどの異なるチップ感でメモリの見え方を完全に一致させるには、同期を行うコストが大きくなり、性能に与える影響も小さくありません。
中国の「神威・太湖之光」やpezy computingのプロセッサでは、CPUによるリオーダリングを自由にし、プログラマが必要な箇所のみで同期を行うことで、データ並列性の高いアプリケーションでの実行速度を上げているようです。

でRCUではなにやってんの？

かなり寄り道してしまった感もありますが、再び最初に見たコードを思い出してみましょう。

struct foo {
   int a;
   int b;
};

struct foo *foop = NULL;
struct foo new;

/*foop!=NULLで値を出力*/
void reader()
{
    int a, b;
    if (foop != NULL) {
        a = foop->a;
        b = foop->b;
        printk("Set. a = %d, b = %d\n", a, b);
    }
    return;
}

/*newを初期化しfoopにセット*/
void writer()
{
    new.a = 1;
    new.b = 1;
    foop = &new;
    return;
}

上記コードでは、参照時にfoop->a、foop->bが1になっていない可能性がある、ということを冒頭で述べました。
x86のCPUを前提とした場合、どうすればこれを防げるか、というと、、、

struct foo {
   int a;
   int b;
};

struct foo *foop = NULL;
struct foo new;

/*foop!=NULLで値を出力*/
void reader()
{
    int a, b;
    if (foop != NULL) {
        a = foop->a;
        b = foop->b;
        printk("Set. a = %d, b = %d\n", a, b);
    }
    return;
}

/*newを初期化しfoopにセット*/
void writer()
{
    new.a = 1;   // (1)
    new.b = 1;   // (2)
    asm volatile("" ::: "memory");  //★
    foop = &new; // (3)
    return;
}

asm volatile("" ::: "memory");がfoop = &new;の前に入っただけです。
「え〜っ」って感じですが、x86に関してはこれだけで、上記コードは思った通りに動作します。
以下その理由です。
x86のCPUは、store -> storeのメモリアクセス順序は変更しません。 つまり、writer()の(1), (2), (3)はコンパイラによってメモリ順序の変更が起きない限り、意図通りに動くわけです。
そのため、asm volatile("" ::: "memory");を(1)、(2)と(3)の間に入れ、この間のコンパイラによるメモリ順序の変更を防いでいます。

そして、何を隠そうこれがRCUにおけるアトミックなポインタの更新と参照になるのです。
※ アーキテクチャが異なると実装も異なります。

/*注意！！！以下は正しい使い方ではありません！！！*/

void reader()
{
    struct foo *p;
    int a, b;

    p = rcu_dereference(foop);  //★
    if (p != NULL) {
        a = p->a;
        b = p->b;
        printk("Set. a = %d, b = %d\n", a, b);
    }
    return;
}

void writer()
{
    new.a = 1;
    new.b = 1;
    rcu_assign_pointer(foop, &new); //★
    return;
}

では、念の為嘘をついていないことを確認しましょう。
以下はv4.3のコードです。

rcu_assign_pointer

#define rcu_assign_pointer(p, v) smp_store_release(&p, RCU_INITIALIZER(v))

#define smp_store_release(p, v)                                         \
do {                                                                    \
        compiletime_assert_atomic_type(*p);                             \
        // コンパイラによるリオーダリングを禁止
        // #define barrier() __asm__ __volatile__("": : :"memory")
        barrier();                                                      \
        // 代入
        WRITE_ONCE(*p, v);                                              \
} while (0)

rcu_dereference

#define rcu_dereference(p) rcu_dereference_check(p, 0)

#define rcu_dereference_check(p, c) \
        __rcu_dereference_check((p), (c) || rcu_read_lock_held(), __rcu)

#define __rcu_dereference_check(p, c, space) \
({ \
        /* Dependency order vs. p above. */ \
        typeof(*p) *________p1 = (typeof(*p) *__force)lockless_dereference(p); \ //★
        RCU_LOCKDEP_WARN(!(c), "suspicious rcu_dereference_check() usage"); \
        rcu_dereference_sparse(p, space); \
        ((typeof(*p) __force __kernel *)(________p1)); \
})

#define lockless_dereference(p) \
({ \
        // 代入しているだけ
        typeof(p) _________p1 = READ_ONCE(p); \
        // x86の場合からっぽ
        smp_read_barrier_depends(); /* Dependency order vs. p above. */ \
        (_________p1); \
})

（嘘を見つけた方はコメント頂けると大変助かります）

おわりに

もともとここはさらっと行こうと思ったのですが、考えてみたら、いつもコード書くときにメモリオーダリングとか全く意識しないなぁ、という感じだったので、自分の勉強がてら調べてみました。
今回はRCUというよりはプリミティブな部分の話でしたが、次回からはRCUっぽい話ができるかと思います。