Radeon7xxxシリーズのレイテンシを見るの回
この記事はGPGPU Advent Calendarの14日目です
いろいろボロボロなのは多分僕の所為なので許してほしい
さて、RadeonとかBulldozerとかいってるAMD党な私です。
ただそういえば、Radeonの詳細なデータとか見た覚えないなー、速い、速い、実際速いとか言いつつ本当に速いの？という疑問についてちゃんとした答えを出そうというのが今回の趣旨です
とはいえ、優秀なMemory Bandwidth計測ツールがSDKに付属してるので基本的にそれ使います
そして最終的には、どの程度使い物になるんだろうか、というところを模索したいところ
さて、前に関東GPGPU勉強会で、「Radeonが遅いんです、キックが遅いっぽいんです」といって嘆いていたのを覚えている方はいるんでしょうか？
体感的に、Radeonはレイテンシが高いように思っていて今の今まで来ています
ではちゃんと計測してみましょう
対象は手元に存在するRadeon HD7770。縁起物っぽいナンバリング。
キックが遅いんです、という話をどうにか考えてみたいと思います。
何もしないカーネルを起動して、終了するまでの時間を計測してみると、簡易的ではあるものの、キック(と終了も含んでしまうが)の時間を調べることができます
#include <iostream> #ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/cl.h> #endif #include <vector> #include <cassert> #include <sstream> #include <fstream> #include <streambuf> const char source[] = "__kernel void nop(void)\n{\n}\n"; int main(int argc, char** argv) { cl_int ret; cl_platform_id platform_id; cl_device_id device; cl_context context; cl_command_queue command_queue; size_t loop_num = 20; cl_ulong kernel_execute_time = 0; if(argc >= 2) { loop_num = atoi(argv[1]); } // get platform id. // not use Intel platform. { std::vector<cl_platform_id> platform_ids; cl_uint ret_num_platforms; ret = clGetPlatformIDs(0, NULL, &ret_num_platforms); assert(ret == CL_SUCCESS); platform_ids.resize(ret_num_platforms); ret = clGetPlatformIDs(ret_num_platforms, &platform_ids[0], NULL); assert(ret == CL_SUCCESS); std::cout << "Platform num : " << ret_num_platforms << std::endl; // use first platform platform_id = platform_ids[0]; } // get Device id. // only to use gpu. { std::vector<cl_device_id> device_ids(1); cl_uint ret_num_devices; ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_ids[0], &ret_num_devices); assert(ret == CL_SUCCESS); if(ret_num_devices > 0) { device = device_ids[0]; } } context = clCreateContext(NULL, 1, &device, NULL, NULL, &ret); assert(ret == CL_SUCCESS); command_queue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &ret); assert(ret == CL_SUCCESS); // get device name. { char msg[1024]; ret = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(msg), msg, NULL); assert(ret == CL_SUCCESS); std::cout << "Device name : " << msg << std::endl; } // create memory buffer cl_mem buffer1; buffer1 = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int) * 1024, NULL, &ret); assert(ret == CL_SUCCESS); // load CL source and create program. cl_program program; { const char* src[] = {source}; size_t size[] = {sizeof(source) - 1}; program = clCreateProgramWithSource(context, 1, src, size, &ret); assert(ret == CL_SUCCESS); } { char option[512]; ret = clBuildProgram(program, 1, &device, option, NULL, NULL); if(ret != CL_SUCCESS) { std::cout << "Error code : " << ret << std::endl; char msg[4096]; int l_ret; l_ret = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, sizeof(msg), msg, NULL); assert(l_ret == CL_SUCCESS); std::cout << msg << std::endl; } } cl_kernel kernel = clCreateKernel(program, "nop", &ret); assert(ret == CL_SUCCESS); cl_event kernel_event; // run first time ret = clEnqueueTask(command_queue, kernel, 0, NULL, &kernel_event); clWaitForEvents(1, &kernel_event); assert(ret == CL_SUCCESS); for(size_t i = 0; i < loop_num; ++i) { ret = clEnqueueTask(command_queue, kernel, 0, NULL, &kernel_event); clWaitForEvents(1, &kernel_event); assert(ret == CL_SUCCESS); cl_ulong start, end; clGetEventProfilingInfo(kernel_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL); clGetEventProfilingInfo(kernel_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL); kernel_execute_time += (end - start); } std::cout << "Kernel execute done" << std::endl; std::cout << "Kernel execute time : " << (static_cast<double>(kernel_execute_time) / loop_num) / 100000.0 << " [ms] " <<std::endl; ret = clFlush(command_queue); ret = clFinish(command_queue); ret = clReleaseKernel(kernel); ret = clReleaseProgram(program); ret = clReleaseMemObject(buffer1); ret = clReleaseCommandQueue(command_queue); ret = clReleaseContext(context); return 0; }
これを実行すると、
$ ./a.out Platform num : 1 Device name : Capeverde Kernel execute done Kernel execute time : 0.0974075 [ms]
という結果になります
デフォルトだと20回ループを回してその平均実行時間をとるものですが、ループ回数をあげていくと同時に数字も変わっていくというなかなかの現象
$ ./a.out 100 Platform num : 1 Device name : Capeverde Kernel execute done Kernel execute time : 0.0963552 [ms] $ ./a.out 1000 Platform num : 1 Device name : Capeverde Kernel execute done Kernel execute time : 0.0643082 [ms] $ ./a.out 10000 Platform num : 1 Device name : Capeverde Kernel execute done Kernel execute time : 0.0583559 [ms]
# ただし、なにもしないカーネルを複数回起動することによって変なことになっている可能性もなきにしもあらず
ちなみに比較対象として、GTX580の場合はというと
$ ./a.out Platform num : 2 Device name : GeForce GTX 580 Kernel execute done Kernel execute time : 0.035984 [ms]
3倍速かった。シャアか。
ここから先はw_o大先生の
http://d.hatena.ne.jp/w_o/20120610#1339330759
この記事とか読んでからのほうがいいと思うんです
それぞれ指定するオプションの意味は
- if/-of 3 : CL_MEM_USE_HOST_PTR
- if/-of 5 : CL_MEM_ALLOC_HOST_PTR
- if/-of 6 : CL_MEM_USE_PERSISTENT_MEM_AMD
です。大先生ありがとうございます
$ ./BufferBandWidth Device 0 Capeverde (略) inputBuffer: CL_MEM_READ_ONLY outputBuffer: CL_MEM_WRITE_ONLY Host baseline (naive): Timer resolution 257.14 ns Page fault 1201.03 ns CPU read 3.90 GB/s memcpy() 4.10 GB/s memset(,1,) 8.96 GB/s memset(,0,) 8.92 GB/s AVERAGES (over loops 2 - 19, use -l for complete log)
-
-
-
-
-
-
- -
-
-
-
-
-
-
-
-
-
-
-
- -
-
-
-
-
-
-
-
-
-
-
-
- -
-
-
-
-
-
-
-
-
-
-
-
- -
-
-
-
-
-
いくらホストからの書き込みが高速とはいえ、カーネルからのアクセスが10倍も違うとなると、USE_HOST_PTRはかなり実用的じゃねーなーと思います
でもそんなにアクセスしないけどかなり巨大なデータを一気に突っ込むとかの用法ならまだましなのか…？
けつろん。
Radeonのバス回りにやや不安が残る
特にカーネルキックのところ。3倍違うとなるといろいろ響いてくる
ただ、俺のプログラムに組み込んだときはもっと差があったように思ったんだけど…気のせいかな…
