福野泰介の一日一創 - create every day

GPUはじめのいっぽでOpenCLのスゴさが分かったので、次は実運用に向けサーバーで動かしてみます。
GPUといえば、さくらの高火力
1時間357円でレンタルできる、NVIDIA Tesla P100モデルを使ってみました。(T100モデルは月額プランのみ

契約してから使えるようになるまで何分か時間がかかります。利用開始時間として表示される前から実際使えるおまけがありました。(おまけ時間中は終了できないので注意!)


浮動小数点演算(float)9.3TFLOPSという性能。OpenCLではこのように見えました。 OpenCLバージョン1.2、コンピューティングユニット56、1.3GHz、最大メモリ確保量4GB、ローカルメモリ48KB、ワークグループサイズは1024。(MacBookProのGPUでは、それぞれ1.2、48、1Hz、384MB、64KB、256 と極端な差は見えず)

高火力サーバー、ユーザー名ubuntu、管理画面のパスワードとIPで、sshを使ってログイン。
NVIDIAのサイトからドライバをダウンロード用スクリプトを取得します(Linux/x86_64/Ubuntu/14.04/Type...dev(network))

sudo apt update sudo apt upgrade sudo apt-get install wget sudo apt-get install software-properties-common sudo apt install ocl-icd-libopencl1

などで、必要なものと合わせてセットアップ!(参考、CUDA Toolkit/GPUカードドライバー導入手順 – さくらのサポート情報
OpenCLのヘッダファイルを探してみます。

$ find / -name cl.h 2> /dev/null /usr/local/cuda-10.1/targets/x86_64-linux/include/CL/cl.h /usr/include/nvidia-418/CL/cl.h

MacとLinuxでヘッダファイルが違うので、編集

#ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/cl.h> #endif

あとは、コンパイルして実行するだけ!

$ g++ -I/usr/include/nvidia-418/ test.c -lOpenCL $ ./a.out

結果、1000万個の配列の足し算の結果。

time[sec]=0.375 cpu[sec]=0.375 cpu usage=99.94% - Tesla P100 GPU time[sec]=5.382 cpu[sec]=0.091 cpu usage=1.69% - MacOS GPU

14倍速!

もっと差がでると思いましたが、例が悪かったのでしょう。
手順は押さえたので、また何か作った時に試してみようと思います。

もうちょっと手頃なGPU付きクラウド、登場してほしいところです。

いろんなプログラミング言語はありますが、IchigoJamを作っている言語でもある、C言語を押さえておくと楽しさぐっと広がります。

コンピューターのネイティブ言語はマシン語ですが、C言語はそれを手軽に生成してくれます。 今回は、パスワード解読で活躍したGPUプログラミングにC言語を使って挑戦してみましょう!(Mac向け)

マルチコアなCPUや、いろんなメーカーのGPUを統一的に使えるOpenCLを使うと、C言語で簡単に並列プログラミングができます。
MacBookProが対応しているOpenCLのバージョンは1.2。仕様書を読んで、まずは対応デバイスを調べます。

cl_platform_id platform_id[MAX_PLATFORMS]; cl_uint num_platforms; cl_int ret = clGetPlatformIDs(MAX_PLATFORMS, platform_id, &num_platforms) printf("num of platforms: %d\n", num_platforms); // ... つづく詳細は末尾

プラットフォームを取得、デバイスを取得し、デバイスの特性を表示してみると、CPUとGPUが1つずつ得られました。

num of platforms: 1 platform num: 0 profile: FULL_PROFILE version: OpenCL 1.2 (Jun 23 2019 21:50:55) name: Apple vendor: Apple extensions: cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event n devices: 2 device num: 0 (platform 0) device name: Intel(R) Core(TM) i5-7360U CPU @ 2.30GHz device vendor: Intel driver vesion: 1.1 device_type: CPU compute units: 4 max clock frequency: 2300MHz max alloc size: 4096MB local mem size: 32kB work group size: 1024 work item dimensions: 3 max work item sizes: 1024 1 1 = 1024 device num: 1 (platform 0) device name: Intel(R) Iris(TM) Plus Graphics 640 device vendor: Intel Inc. driver vesion: 1.2(Aug 20 2019 17:28:33) device_type: GPU compute units: 48 max clock frequency: 1000MHz max alloc size: 384MB local mem size: 64kB work group size: 256 work item dimensions: 3 max work item sizes: 256 256 256 = 16777216

演算処理を行う(compute units)が、CPUに4コ。GPUに48コ。MacBookProには52コの演算器があるんですね!

100万個の配列2つの足し算を計算してみましょう!

for (int i = 0; i < 1000000; i++) { c[i] = a[i] + b[i]; }

ひとまず1000回ループさせて、純粋にCPUのみで計算させて計測して、1.98秒(useopencl=0)
同じ条件で、CPUをデバイスとして使うと、0.57秒(CPU時間は2.6倍とマルチコアが活きてます!)
そしてGPUデバイスを使うと、0.35秒と更に速い!

__kernel void vadd(global const float *a, global const float *b, global float *c, int n) { int i = get_global_id(0); if (i < n) { c[i] = a[i] + b[i]; } }

こちらがOpenCLで使うC言語(vadd.cl として保存)は、再起呼び出しが使えない、可変長引数非対応など一部制限ありますが、ほぼ普通のC言語! 同じパソコン上であっても、独立して動く演算器同士、ネットワーク越しで使うようにつなぐところがおもしろいですね。


意外と簡単に使えることが分かったので、何か実用的なコードづくりにチャレンジしてみたいと思います!

下記、今回のプログラム全文(参考、9.1 OpenCLとは - www.e-em.co.jp/tutorial/chap9.htm

#include <stdio.h> #include <stdlib.h> #include <time.h> #include <sys/time.h> #include <OpenCL/opencl.h> #define MAX_SOURCE_SIZE 100000 #define MAX_PLATFORMS 10 #define MAX_DEVICES 20 cl_device_id getDeviceCL(cl_device_type type) { cl_device_id res = NULL; cl_platform_id platform_id[MAX_PLATFORMS]; cl_uint num_platforms; cl_int ret = clGetPlatformIDs(MAX_PLATFORMS, platform_id, &num_platforms); printf("num of platforms: %d\n", num_platforms); printf("\n"); size_t len; cl_device_id device_id[MAX_DEVICES]; cl_uint num_devices; cl_uint idx_devices = 0; for (int i = 0; i < num_platforms; i++) { cl_platform_id pid = platform_id[i]; #define MAX_BUFFER 1024 char name[MAX_BUFFER]; len = MAX_BUFFER; printf("platform num: %d\n", i); ret = clGetPlatformInfo(pid, CL_PLATFORM_PROFILE, len, name, &len); // FULL_PROFILE or EMBEDDED_PROFILE printf("profile: %s\n", name); len = MAX_BUFFER; ret = clGetPlatformInfo(pid, CL_PLATFORM_VERSION, len, name, &len); // FULL_PROFILE or EMBEDDED_PROFILE printf("version: %s\n", name); len = MAX_BUFFER; ret = clGetPlatformInfo(pid, CL_PLATFORM_NAME, len, name, &len); // FULL_PROFILE or EMBEDDED_PROFILE printf("name: %s\n", name); len = MAX_BUFFER; ret = clGetPlatformInfo(pid, CL_PLATFORM_VENDOR, len, name, &len); // FULL_PROFILE or EMBEDDED_PROFILE printf("vendor: %s\n", name); len = MAX_BUFFER; ret = clGetPlatformInfo(pid, CL_PLATFORM_EXTENSIONS, len, name, &len); // FULL_PROFILE or EMBEDDED_PROFILE printf("extensions: %s\n", name); ret = clGetDeviceIDs(pid, CL_DEVICE_TYPE_ALL, MAX_DEVICES - idx_devices, device_id, &num_devices); // CL_DEVICE_TYPE_ALL -> 2 (Apple) // ret = clGetDeviceIDs(pid, CL_DEVICE_TYPE_DEFAULT, MAX_DEVICES - idx_devices, device_id, &num_devices); // CL_DEVICE_TYPE_DEFAULT -> 1 printf("n devices: %d\n", num_devices); printf("\n"); for (int j = 0; j < num_devices; j++) { cl_device_id did = device_id[j]; printf("device num: %d (platform %d)\n", j, i); len = MAX_BUFFER; ret = clGetDeviceInfo(did, CL_DEVICE_NAME, len, name, &len); printf("device name: %s\n", name); len = MAX_BUFFER; ret = clGetDeviceInfo(did, CL_DEVICE_VENDOR, len, name, &len); printf("device vendor: %s\n", name); len = MAX_BUFFER; ret = clGetDeviceInfo(did, CL_DRIVER_VERSION, len, name, &len); printf("driver vesion: %s\n", name); cl_device_type device_type; ret = clGetDeviceInfo(did, CL_DEVICE_TYPE, sizeof(device_type), &device_type, &len); if (device_type == CL_DEVICE_TYPE_CPU) printf("device_type: CPU\n"); else if (device_type == CL_DEVICE_TYPE_GPU) printf("device_type: GPU\n"); else if (device_type == CL_DEVICE_TYPE_ACCELERATOR) printf("device_type: ACCELERATOR\n"); else printf("device_type: other %llu\n", device_type); if (device_type == type) { res = did; } cl_uint compute_units; ret = clGetDeviceInfo(did, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(compute_units), &compute_units, &len); printf("compute units: %u\n", compute_units); // CPU: 4, GPU: 48 cl_uint clock_freq; ret = clGetDeviceInfo(did, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(clock_freq), &clock_freq, &len); printf("max clock frequency: %uMHz\n", clock_freq); // CPU: 2300MHz, GPU: 1000MHz cl_ulong max_alloc_size; ret = clGetDeviceInfo(did, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(max_alloc_size), &max_alloc_size, &len); printf("max alloc size: %lluMB\n", max_alloc_size / (1024 * 1024)); // CPU: 4096MB, GPU: 384MB cl_ulong local_mem_size; ret = clGetDeviceInfo(did, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(local_mem_size), &local_mem_size, &len); printf("local mem size: %llukB\n", local_mem_size / 1024); // CPU: 32kB, GPU: 64kB size_t work_group_size; ret = clGetDeviceInfo(did, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(work_group_size), &work_group_size, &len); printf("work group size: %zu\n", work_group_size); // CPU: 1024, GPU: 256 cl_uint work_item_dimensions; ret = clGetDeviceInfo(did, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(work_item_dimensions), &work_item_dimensions, &len); printf("work item dimensions: %u\n", work_item_dimensions); // CPU: 3, GPU: 3 (minimum 3) size_t work_item_sizes[work_item_dimensions]; ret = clGetDeviceInfo(did, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(work_item_sizes), work_item_sizes, &len); printf("max work item sizes: "); // CPU: 1024 1 1 = 1024, GPU: 256 256 256 = 16777216 size_t maxitems = 1; for (int k = 0; k < work_item_dimensions; k++) { printf("%zu ", work_item_sizes[k]); maxitems *= work_item_sizes[k]; } printf("= %zu\n", maxitems); printf("\n"); } } return res; } int main(int argc, char** argv) { //getDeviceCL(NULL); //return 0; const int n = 1 * 1000 * 1000; // 1M float* a = (float*)malloc(n * sizeof(float)); float* b = (float*)malloc(n * sizeof(float)); float* c = (float*)malloc(n * sizeof(float)); for (int i = 0; i < n; i++) { a[i] = i + 1; b[i] = i + 1; c[i] = 0; // a[i] + b[i] } int useopencl = 1; cl_device_type type = CL_DEVICE_TYPE_GPU; //cl_device_type type = CL_DEVICE_TYPE_CPU; cl_device_id device_id = getDeviceCL(type); const int nloop = 1000; // sec -> msec cl_int ret; cl_context context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret); cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret); // load program for GPU FILE* fp = fopen("vadd.cl", "r"); if (!fp) { fprintf(stderr, "kernel source open error\n"); exit(1); } char* source_str = (char*)malloc(MAX_SOURCE_SIZE * sizeof(char)); size_t source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp); fclose(fp); cl_program program = clCreateProgramWithSource(context, 1, (const char**)&source_str, (const size_t*)&source_size, &ret); if (ret != CL_SUCCESS) { fprintf(stderr, "clCreateProgramWithSource() error\n"); exit(1); } free(source_str); if (clBuildProgram(program, 1, &device_id, NULL, NULL, NULL) != CL_SUCCESS) { fprintf(stderr, "clBuildProgram() error\n"); exit(1); } cl_kernel kernel = clCreateKernel(program, "vadd", &ret); if (ret != CL_SUCCESS) { fprintf(stderr, "clCreateKernel() error\n"); exit(1); } // memory cl_mem d_a = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret); cl_mem d_b = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret); cl_mem d_c = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret); // host to device clEnqueueWriteBuffer(command_queue, d_a, CL_TRUE, 0, n * sizeof(float), a, 0, NULL, NULL); clEnqueueWriteBuffer(command_queue, d_b, CL_TRUE, 0, n * sizeof(float), b, 0, NULL, NULL); // args for program clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&d_a); clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&d_b); clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&d_c); clSetKernelArg(kernel, 3, sizeof(int), (void*)&n); // work item size_t local_item_size = 128; // limit on MacBookPro: <=128(CPU) CL_INVALID_WORK_GROUP_SIZE(-54), <=256(GPU) CL_INVALID_WORK_ITEM_SIZE(-55) size_t global_item_size = ((n + local_item_size - 1) / local_item_size) * local_item_size; printf("item size: global %zu, local %zu\n", global_item_size, local_item_size); /// timer clock_t startClock = clock(); struct timeval startTime; gettimeofday(&startTime, NULL); // run if (useopencl) { for (int loop = 0; loop < nloop; loop++) { ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_item_size, &local_item_size, 0, NULL, NULL); if (ret != CL_SUCCESS) { fprintf(stderr, "clEnqueueNDRangeKernel() error %d\n", ret); exit(1); } } // device to host clEnqueueReadBuffer(command_queue, d_c, CL_TRUE, 0, n * sizeof(float), c, 0, NULL, NULL); } else { for (int loop = 0; loop < nloop; loop++) { for (int i = 0; i < n; i++) { c[i] = a[i] + b[i]; } } } // timer clock_t endClock = clock(); double cpu = (double)(endClock - startClock) / CLOCKS_PER_SEC; struct timeval endTime; gettimeofday(&endTime, NULL); time_t diffsec = difftime(endTime.tv_sec, startTime.tv_sec); suseconds_t diffsub = endTime.tv_usec - startTime.tv_usec; double time = diffsec + diffsub * 1e-6; // output double sum = 0; for (int i = 0; i < n; i++) { //printf("%f ", c[i]); sum += c[i]; } double exact = n * (n + 1.0); printf("n=%d nloop=%d %e(%.6e)\n", n, nloop, sum, exact); printf("time[sec]=%.3f cpu[sec]=%.3f cpu usage=%.2f%%\n", time, cpu, cpu / time * 100); // release clFlush(command_queue); clFinish(command_queue); clReleaseMemObject(d_a); clReleaseMemObject(d_b); clReleaseMemObject(d_c); clReleaseKernel(kernel); clReleaseProgram(program); clReleaseCommandQueue(command_queue); clReleaseContext(context); // free free(a); free(b); free(c); return 0; }

上記 test.c として保存して、下記コマンドでコンパイル&実行!

$ gcc test.c -framework OpenCL $ ./a.out

links
- Khronos OpenCL Registry - The Khronos Group Inc
- 9.1 OpenCLとは - www.e-em.co.jp/tutorial/chap9.htm
- MacのGPUでも700倍速! パスワード20文字時代の盾と矛 / さくらクラウドで分散探査する方法
- パスワード付きZIPファイルの有効性検証! Go言語を使ったマルチコア対応パスワードチャレンジプログラム

Tweet
クリエイティブ・コモンズ・ライセンス
この作品は「Creative Commons — CC BY 4.0」の下に提供されています。
CC BY 福野泰介 - Taisuke Fukuno / @taisukef / アイコン画像 / プロフィール画像