tkokamoの日記

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

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_structblockedreal_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)-1 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()によりnewfoopにセットされていればその値を読み出し出力する、という関数ですが、
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 -> loadload -> storestore -> loadstore -> storeの4種類のどのアクセス順序を変更して良いか、というポリシーが各CPUで決められています。
x86は比較的メモリアクセス順序を守るほうなのですが、store -> loadに関してはその順序を入れ替えても良い(コア間で見え方が違っても良い)、というポリシーになっています。(そうでないx86のCPUもあります)
上記の例では、このリオーダリングが見事に現れました。
つまり、スレッド1とスレッド2が別コアで動いた時、それぞれのarg->a = 1arg->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->afoop->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っぽい話ができるかと思います。

RCU(Read Copy Update)をちゃんと知る(1)-2 実装よりの概要

はじめに

前回(http://tkokamo.hateblo.jp/entry/2017/05/31/213427)は、かなりざっくりとした理解だったのでもう少し実装に近づいた概要です。

(1) 概要(概要)← 前回と今回
(2) Linuxでの実装
(3) rcu-walk(RCUの応用先)

今回までで概要は終わらして、次回からは実際の実装を見ていこうと思っています。

RCUの種類

現在のLinuxには大きく以下のRCUの実装があるようです。(https://lwn.net/Articles/541037/)

参照クリティカルセクション内でブロック不可 ブロック可
SMP非対応 Tiny RCU Tiny Sleepable RCU
SMP対応 Tree RCU (これを理解したい) Tree Sleepable RCU

Sleepable RCU(SRCU)について


クリティカルセクション内で眠れるなんて、すばらしいじゃないか!と思いますが、実際はそんなに使われていません。 (4.13-rc3では、rcu_read_lock:875件、srcu_read_lock:57件)

SRCUの実装はよくわかっていませんが、参照クリティカルセクションでブロックが発生すると参照期間が当然伸びます。そして、これは古いデータが解放されるまでの時間(grace period)が伸びるということを意味します。
grace periodが長くなると、その間に複数世代のデータが作られメモリを圧迫してしまう危険があるので、好ましくありません。
SRCUについてはこれ以上深くはつっこみませんが、興味がある方はhttp://www.rdrop.com/~paulmck/RCU/srcu.2007.01.14a.pdfを見ると良いかと思います。(少し古いです)

Classic RCUとTree RCU


v2.6で当初実装されたRCUはClassic RCUと呼ばれていますが、CPU数が増えると更新側の性能がスケールしない、無駄に電力を使ってしまう、などといった問題がありました。(※CPU数が増えると、といっても数百とかいったレベルらしい)
これらの問題を解決したRCU実装がTree RCUで、現在はTree RCUが用いられています。
参考:http://www.atmarkit.co.jp/flinux/rensai/watch2009/watch04a.html RCUの全面書き直しも! 2.6.29は何が変わった?(1/2) − @IT

Tree RCU(とrcu walk)を理解するのが目的ですが、SMP環境におけるRCUの仕組みを知るには、まずはClassic RCUから見ると良いかと思うので、Classic RCUについて見ていってからTree RCUに手を伸ばそうと思います。(今回ではないです)

RCUのAPIと処理概要

以下に参照側と更新側が用いるRCUのAPIとgrace periodの関係などを表してみました。

f:id:tkokamo:20170812174518p:plain

参照側が呼び出す関数は以下の通りです。

関数 説明
rcu_read_lock() 参照側クリティカルセクションを開始する。実際はただプリエンプションを無効にするだけ
rcu_read_unlock() 参照側クリティカルセクションを終了する。実際はただプリエンプションを有効にするだけ
rcu_dereference() RCUで管理されているデータのポインタを参照する。

更新側が呼び出す関数は以下の通りです。

関数 説明
rcu_assign_pointer() RCUで管理されているデータのポインタを更新する。
call_rcu() 更新前の古いデータへの参照がなくなった時に呼ばれる回収用のコールバック関数を登録する。
synchronize_rcu() grace periodがすぎるまで待ち合わせる。古いLinuxではsynchronize_kernel()

上の図では、CPU0の更新者(update)がcall_rcu()を呼んだ時、参照側クリティカルセクション内にいるCPU2、CPU3の参照者(read)が参照側クリティカルセクションを抜けるまで、古いデータを保持しておく必要があります。
CPU1の2つ目の参照者はcall_rcu()後に参照を始めているため、最新の値を参照しています。そのため、このタスクが参照をやめるまで待つ必要はありません。

実際には、call_rcu()呼び出し後、全てのCPUについてcontext switchが起きたらcall_rcu()で登録されたコールバック関数を呼び出し、古いデータの回収を行うことができます。
最初は「?」となりましたが、以下のように少し整理すると、正しいことがわかります。
参照側クリティカルセクションではプリエンプションが禁止されており、この間にcontext switchが起きることはありません。
逆にいうと、context switchが起きたCPUは参照側クリティカルセクションにいないことが保証されます。
そのため、call_rcu()後、すべてのCPUでcontext switchが起きた時、古いデータの参照は終わっていると判断できるのです。

おわりに

前回と今回でRCUの大雑把な理解はできると思います。
間違いなどありましたら、指摘いただけると助かります。

次回からは、実際にRCUのAPIのコードを以下のような流れで見ていこうと思います。

  • rcu_dereference(), rcu_assign_pointer()
  • Classic RCUの回収処理など
  • Tree RCUの回収処理など