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

前回のコードを見るとカーネルコードの実装以前に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

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

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

tkokamoの日記

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

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

今回のコード

関数の中身

main()

clPrintDevInfo