参考文献：Interface 2022年8月号 第５部第２章 OpenCL を試す
学習記録です。
環境
ノート PC
- HP Pavilion Laptop 15-eh3xxx
- AMD Ryzen 5 7530U with Radeon Graphics
OpenCL の導入
NVIDIA 製の GPU でしか動かない CUDA とは違い、OpenCl はメーカーや設計が異なるハードウェアを同じコードで効率よく動かすことができます。
OpenCL の SDK（Software Development Kit：ソフトウェア開発キット）をインストールします。
OpenCL は共通規格ですが、実際に計算を行うのは NVIDIA や AMD、Intelの チップです。そのため、各メーカーが自社のチップを効率よく動かせるように、専用の SDK を配布しています。また、最近では、Khronos OpenCL SDK という、メーカーを問わず使える公式の標準SDKがGitHubなどで公開されています。
AMD が独自の OpenCL SDK の単体配布を終了した？という記事が出てきたので、Khronos 公式の SDK をインストールしました。
Visual Studio が Khronos 公式のヘッダーとライブラリを認識するようにします。
git clone https://github.com/microsoft/vcpkg
.\vcpkg\bootstrap-vcpkg.bat
.\vcpkg\vcpkg install opencl:x64-windows
.\vcpkg\vcpkg integrate install
認識するようにしたのですが、なぜかパスが通っていなかったので、手動で設定しました。プロジェクトのプロパティ ＞ 追加のインクルードディレクトリ にパスを設定します。
"（vcpkg を置いた場所）\vcpkg\installed\x64-windows\include"
ライブラリもパスを通します。（リンカー ＞ 全般）
"（vcpkg を置いた場所）\vcpkg\installed\x64-windows\lib"
ノート PC の内蔵 GPU
インテルや AMD の CPU を搭載しているノート PC には内蔵 GPU も搭載されており、内蔵 GPU は iGPU（Integrated GPU）と呼ばれます。
CPU と iGPU の情報を取得する
テキストで紹介されていたプログラムにメモリ解放部分がない？ので追加しました。
#include <vector>
#include <stdio.h> // printfのために追加
// OpenCLのヘッダ・ファイル
#include <CL/cl.h>
#include <CL/cl_platform.h>
// OpenCLライブラリ
#pragma comment(lib, "OpenCL.lib")
int main()
{
cl_uint num_of_platforms = 0;
// プラットフォーム数を取得
cl_int err = clGetPlatformIDs(0, 0, &num_of_platforms);
std::vector<cl_platform_id> platforms(num_of_platforms);
// プラットフォーム識別子を取得
err = clGetPlatformIDs(num_of_platforms, &platforms[0], 0);
// 各プラットフォームごとに(通常は1つ)
for (cl_uint i = 0; i < num_of_platforms; i++)
{
cl_uint num_of_devices = 0;
// OpenCL対応デバイス数を取得
err = clGetDeviceIDs(
platforms[i],
CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_CPU, // CPUとGPUを取得
0,
0,
&num_of_devices
);
cl_device_id* id = new cl_device_id[num_of_devices];
// デバイス識別子を取得
err = clGetDeviceIDs(
platforms[i],
CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_CPU, // CPUとGPUを取得
num_of_devices,
id,
0
);
// 各デバイスについて
for (cl_uint j = 0; j < num_of_devices; j++)
{
size_t length;
cl_uint freq, units;
char deviceName[128];
cl_device_svm_capabilities svmCapability;
// デバイス情報を取得
err = clGetDeviceInfo(id[j],
CL_DEVICE_MAX_CLOCK_FREQUENCY,
sizeof(cl_uint), &freq, &length);
err = clGetDeviceInfo(id[j],
CL_DEVICE_MAX_COMPUTE_UNITS,
sizeof(cl_uint), &units, &length);
err = clGetDeviceInfo(id[j], CL_DEVICE_NAME,
sizeof(deviceName), deviceName, NULL);
// 変数名を svmCapability に統一
err = clGetDeviceInfo(id[j],
CL_DEVICE_SVM_CAPABILITIES,
sizeof(svmCapability), &svmCapability, NULL);
if (svmCapability & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER)
{
printf("%s %d Compute units %dMHz SVM supported\n",
deviceName, units, freq);
}
else
{
printf("%s %d Compute units %dMHz\n",
deviceName, units, freq);
}
}
// メモリの解放を追加
delete[] id;
}
return 0;
}
実行結果 ↓
gfx90c 7 Compute units 2000MHz SVM supported
GPU は認識されましたが、CPU が認識されていません。CPU が OpenCL から認識されるためには、OpenCL CPU Runtime という専用のソフトが必要だそうです。
OpenCLを動かすための最も標準的なソフトが Intel CPU Runtime で、AMD の CPU でも問題なく動作するとのことで、ダウンロードページより、ダウンロードしました。
Starting in the 2020 February release ("igfx_win10_100.7870.exe"), the CPU runtime for OpenCL is no longer included with the OpenCL Driver package. とのことです。
再度実行 ↓
gfx90c 7 Compute units 2000MHz SVM supported
AMD Ryzen 5 7530U with Radeon Graphics 12 Compute units 0MHz SVM supported
GPU, CPU 共に情報を取得できました。
「CPU」「OpenMP（CPU並列）」「iGPU」の3パターンで計測・比較する
まず、Visual Studio の OpenMP を有効にします。
プロジェクトのプロパティ ＞ C/C++ ＞ 言語 ＞ OpenMP サポート を 「はい (/openmp)」 に設定しました。
演算用のコードについて、
テキストでは命令キュー作成時に clCreateCommandQueue が使われていますが、この関数は OpenCL 2.0 から「古い形式」とされ現在は clCreateCommandQueueWithProperties が使われているそうです
一度実行した際に GPU の結果しか表示されなかったので、GPUから戻ってきた後、CPUが再び触る前にもう一度 Map を入れました、CPU,
OpenMP の結果も表示されるようになりました。
#include <stdio.h>
#include <vector>
#include <windows.h>
#include <CL/cl.h>
#pragma comment(lib, "OpenCL.lib")
// --- リスト2: 配列同士の足し算をするカーネル ---
const char* kernelSource =
"__kernel void vectorAdd(__global double* a, "
" __global double* b, "
" __global double* c, "
" const unsigned int n) "
"{ "
" int id = get_global_id(0); "
" if (id < n) "
" c[id] = a[id] + b[id]; "
"} ";
#define VECTOR_SIZE 50000000
int main()
{
cl_int err;
cl_uint num_platforms;
clGetPlatformIDs(0, NULL, &num_platforms);
std::vector<cl_platform_id> platforms(num_platforms);
clGetPlatformIDs(num_platforms, &platforms[0], NULL);
cl_device_id deviceId;
// 最初のプラットフォームからGPUを取得
clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, 1, &deviceId, NULL);
// --- リスト3: カーネルを実行する流れ ---
// コンテキスト作成
cl_context context = clCreateContext(NULL, 1, &deviceId, NULL, NULL, &err);
// 命令キュー作成
cl_command_queue queue = clCreateCommandQueueWithProperties(context, deviceId, NULL, &err);
// カーネルコードからプログラムを生成
cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, NULL, &err);
// ビルド（SVM使用のためバージョン2.0以上を指定）
err = clBuildProgram(program, 0, NULL, "-cl-std=CL2.0", NULL, NULL);
// カーネル生成
cl_kernel vectorAdd = clCreateKernel(program, "vectorAdd", &err);
// --- リスト4: 演算対象のデータをデバイス側に渡す (SVM編) ---
// Shared Virtual Memory を確保
double* A = (double*)clSVMAlloc(context, CL_MEM_READ_ONLY, sizeof(double) * VECTOR_SIZE, 0);
double* B = (double*)clSVMAlloc(context, CL_MEM_READ_ONLY, sizeof(double) * VECTOR_SIZE, 0);
double* C = (double*)clSVMAlloc(context, CL_MEM_WRITE_ONLY, sizeof(double) * VECTOR_SIZE, 0);
// この領域へホスト側でアクセスする (Map)
clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE, A, sizeof(double) * VECTOR_SIZE, 0, NULL, NULL);
clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE, B, sizeof(double) * VECTOR_SIZE, 0, NULL, NULL);
// 入力データを設定
for (int i = 0; i < VECTOR_SIZE; i++) {
A[i] = i;
B[i] = (double)VECTOR_SIZE - i;
}
// ホスト側のアクセス終了 (Unmap)
clEnqueueSVMUnmap(queue, A, 0, NULL, NULL);
clEnqueueSVMUnmap(queue, B, 0, NULL, NULL);
unsigned int vectorSize = VECTOR_SIZE;
size_t globalSize = VECTOR_SIZE;
size_t localSize = 8;
// カーネルのパラメータを設定 (SVMポインタを使用)
err = clSetKernelArgSVMPointer(vectorAdd, 0, A);
err |= clSetKernelArgSVMPointer(vectorAdd, 1, B);
err |= clSetKernelArgSVMPointer(vectorAdd, 2, C);
err |= clSetKernelArg(vectorAdd, 3, sizeof(unsigned int), &vectorSize);
LARGE_INTEGER freq, start, stop;
QueryPerformanceFrequency(&freq);
// --- リスト5: GPUでの実行時間を計測 ---
QueryPerformanceCounter(&start);
// カーネルを命令キューに追加して実行
err = clEnqueueNDRangeKernel(queue, vectorAdd, 1, NULL, &globalSize, &localSize, 0, NULL, NULL);
// 命令キューの終了を待つ
err = clFinish(queue);
QueryPerformanceCounter(&stop);
printf("GPU %f msec\n", 1000.0 * (double)(stop.QuadPart - start.QuadPart) / freq.QuadPart);
// --- GPUの終了を待った後、CPUがメモリを触るための許可を再取得する ---
clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, A, sizeof(double) * VECTOR_SIZE, 0, NULL, NULL);
clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, B, sizeof(double) * VECTOR_SIZE, 0, NULL, NULL);
clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, C, sizeof(double) * VECTOR_SIZE, 0, NULL, NULL);
// --- リスト5: CPUで実行時間を計測 ---
QueryPerformanceCounter(&start);
for (int i = 0; i < VECTOR_SIZE; i++) {
C[i] = A[i] + B[i];
}
QueryPerformanceCounter(&stop);
printf("CPU %f msec\n", 1000.0 * (double)(stop.QuadPart - start.QuadPart) / freq.QuadPart);
// --- リスト5: OpenMPで実行時間を計測 ---
QueryPerformanceCounter(&start);
#pragma omp parallel for
for (int i = 0; i < VECTOR_SIZE; i++) {
C[i] = A[i] + B[i];
}
QueryPerformanceCounter(&stop);
printf("OpenMP %f msec\n", 1000.0 * (double)(stop.QuadPart - start.QuadPart) / freq.QuadPart);
// 解放
clSVMFree(context, A);
clSVMFree(context, B);
clSVMFree(context, C);
clReleaseKernel(vectorAdd);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}
実行結果 ↓
GPU 334.404900 msec
CPU 66.563400 msec
OpenMP 65.212900 msec
計算が簡単だったので、メモリの読み書きの速度で CPU が勝ったのかなと考察しました。
なので以下のように sin と cos の計算を加えて、重い処理に変えて実行してみました。
// --- リスト2: 配列同士の足し算をするカーネル ---
const char* kernelSource =
"__kernel void vectorAdd(__global double* a, "
" __global double* b, "
" __global double* c, "
" const unsigned int n) "
"{ "
" int id = get_global_id(0); "
" if (id < n) "
" c[id] = a[id] + b[id]; "
" for (int k = 0; k < 100; k++) { "
" c[id] = sin(c[id]) + cos(c[id]); "
" } "
"} ";
// --- リスト5: CPUで実行時間を計測 ---
QueryPerformanceCounter(&start);
for (int i = 0; i < VECTOR_SIZE; i++) {
C[i] = A[i] + B[i];
for (int k = 0; k < 100; k++) {
C[i] = sin(C[i]) + cos(C[i]);
}
}
QueryPerformanceCounter(&stop);
printf("CPU %f msec\n", 1000.0 * (double)(stop.QuadPart - start.QuadPart) / freq.QuadPart);
// --- リスト5: OpenMPで実行時間を計測 ---
QueryPerformanceCounter(&start);
#pragma omp parallel for
for (int i = 0; i < VECTOR_SIZE; i++) {
C[i] = A[i] + B[i];
for (int k = 0; k < 100; k++) {
C[i] = sin(C[i]) + cos(C[i]);
}
}
QueryPerformanceCounter(&stop);
printf("OpenMP %f msec\n", 1000.0 * (double)(stop.QuadPart - start.QuadPart) / freq.QuadPart);
実行結果 ↓
GPU 7035.128600 msec
CPU 82467.545500 msec
OpenMP 11607.859300 msec
実行速度は予想通り、iGPU ＞ OpenMP（CPU並列）＞ CPU の順になりました。
カメラからの画像を iGPU でリアルタイム処理する
３つのフェーズでリアルタイム処理を実現します。
- OpenCV（カメラ担当）: カメラから今この瞬間の「1枚の写真」を撮る
- OpenCL（加工担当）: GPUのパワーを使って、その写真を猛スピードで加工する
- OpenGL（テレビ担当）: 加工が終わった写真を画面に映し出す
OpenCV をインストールする
.\vcpkg\vcpkg install opencv:x64-windows
.\vcpkg\vcpkg integrate install
OpenGL 用のライブラリ freeglut をインストールする
.\vcpkg\vcpkg install freeglut:x64-windows
.\vcpkg\vcpkg integrate install
コード
OpenCV 系で警告は出ましたが、以下のコードでエラーなく実行できました。
#define CL_TARGET_OPENCL_VERSION 200
#include <opencv2/opencv.hpp>
#include <CL/cl.h>
#include <GL/freeglut.h>
#include <iostream>
#include <vector>
#include <windows.h>
#pragma comment(lib, "OpenCL.lib")
// --- リスト6 & リスト8：ガウシアンおよびメディアン・フィルタのカーネル ---
const char* kernelSource =
"__constant sampler_t NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;"
// リスト6: 3x3のガウシアン・フィルタ
"kernel void gaussian3x3(read_only image2d_t src, write_only image2d_t dst) {"
" int x = get_global_id(0); int y = get_global_id(1);"
" uint4 p[3][3];"
" p[0][0] = read_imageui(src, NEAREST, (int2)(x-1, y-1)); p[0][1] = read_imageui(src, NEAREST, (int2)(x, y-1)); p[0][2] = read_imageui(src, NEAREST, (int2)(x+1, y-1));"
" p[1][0] = read_imageui(src, NEAREST, (int2)(x-1, y)); p[1][1] = read_imageui(src, NEAREST, (int2)(x, y)); p[1][2] = read_imageui(src, NEAREST, (int2)(x+1, y));"
" p[2][0] = read_imageui(src, NEAREST, (int2)(x-1, y+1)); p[2][1] = read_imageui(src, NEAREST, (int2)(x, y+1)); p[2][2] = read_imageui(src, NEAREST, (int2)(x+1, y+1));"
" uint4 pixel;"
" pixel.x = ((p[0][0].x + p[0][2].x + p[2][0].x + p[2][2].x) + (p[0][1].x + p[1][0].x + p[1][2].x + p[2][1].x)*2 + p[1][1].x*4) / 16;"
" pixel.y = ((p[0][0].y + p[0][2].y + p[2][0].y + p[2][2].y) + (p[0][1].y + p[1][0].y + p[1][2].y + p[2][1].y)*2 + p[1][1].y*4) / 16;"
" pixel.z = ((p[0][0].z + p[0][2].z + p[2][0].z + p[2][2].z) + (p[0][1].z + p[1][0].z + p[1][2].z + p[2][1].z)*2 + p[1][1].z*4) / 16;"
" pixel.w = 255; write_imageui(dst, (int2)(x, y), pixel);"
"}"
// リスト8: メディアン・フィルタ
"kernel void median3x3(read_only image2d_t src, write_only image2d_t dst) {"
" int x = get_global_id(0); int y = get_global_id(1);"
" uint i, j, k; uint4 p[9], q[5], min;"
" p[0] = read_imageui(src, NEAREST, (int2)(x-1, y-1)); p[1] = read_imageui(src, NEAREST, (int2)(x, y-1)); p[2] = read_imageui(src, NEAREST, (int2)(x+1, y-1));"
" p[3] = read_imageui(src, NEAREST, (int2)(x-1, y)); p[4] = read_imageui(src, NEAREST, (int2)(x, y)); p[5] = read_imageui(src, NEAREST, (int2)(x+1, y));"
" p[6] = read_imageui(src, NEAREST, (int2)(x-1, y+1)); p[7] = read_imageui(src, NEAREST, (int2)(x, y+1)); p[8] = read_imageui(src, NEAREST, (int2)(x+1, y+1));"
" for(i=0; i<5; i++) { min=(uint4)(255,0,0,0); k=0; for(j=0; j<9; j++) { if(p[j].x < min.x) { min=p[j]; k=j; } } q[i]=min; p[k].x=255; }"
" write_imageui(dst, (int2)(x, y), q[4]);"
"}"
// 5x5 メディアン（拡張）
"kernel void median5x5(read_only image2d_t src, write_only image2d_t dst) {"
" int x = get_global_id(0); int y = get_global_id(1);"
" uint i, j, k; uint4 p[25], q[13], min;"
" k=0; for(int jj=-2; jj<=2; jj++) for(int ii=-2; ii<=2; ii++) p[k++] = read_imageui(src, NEAREST, (int2)(x+ii, y+jj));"
" for(i=0; i<13; i++) { min=(uint4)(255,0,0,0); k=0; for(j=0; j<25; j++) { if(p[j].x < min.x) { min=p[j]; k=j; } } q[i]=min; p[k].x=255; }"
" write_imageui(dst, (int2)(x, y), q[12]);"
"}";
// リスト9準拠の列挙型 (Windowsとの衝突回避のため接頭辞付与)
enum filter_type { FILTER_NONE, FILTER_GAUSSIAN3x3, FILTER_MEDIAN3x3, FILTER_MEDIAN5x5 };
filter_type current_filter = FILTER_NONE;
double msec = 0;
// グローバル
cv::VideoCapture camera(0);
cl_context context; cl_command_queue queue;
cl_kernel kGaus, kMed3, kMed5, activeK = NULL;
cl_mem input, output;
int width, height;
// --- リスト7：画像処理の演算フロー ---
void display() {
cv::Mat image, rgba;
camera.read(image);
if (image.empty()) return;
cv::cvtColor(image, rgba, cv::COLOR_BGR2BGRA); // 32ビットBGRA
if (current_filter != FILTER_NONE && activeK != NULL) {
LARGE_INTEGER freq, start, stop;
QueryPerformanceFrequency(&freq);
QueryPerformanceCounter(&start);
size_t origin[3] = { 0, 0, 0 }, region[3] = { (size_t)width, (size_t)height, 1 };
cl_event wait, finish;
// clEnqueueWriteImage
clEnqueueWriteImage(queue, input, CL_FALSE, origin, region, 0, 0, rgba.data, 0, NULL, &wait);
// clEnqueueNDRangeKernel
size_t global[2] = { (size_t)width, (size_t)height };
clEnqueueNDRangeKernel(queue, activeK, 2, NULL, global, NULL, 1, &wait, &finish);
// clEnqueueReadImage
clEnqueueReadImage(queue, output, CL_TRUE, origin, region, 0, 0, rgba.data, 1, &finish, NULL);
QueryPerformanceCounter(&stop);
msec = 1000.0 * (double)(stop.QuadPart - start.QuadPart) / freq.QuadPart;
}
// 表示処理
glClear(GL_COLOR_BUFFER_BIT);
glPixelZoom(1.0, -1.0); glRasterPos2f(-1.0, 1.0);
glDrawPixels(width, height, GL_BGRA_EXT, GL_UNSIGNED_BYTE, rgba.data);
glutSwapBuffers();
glutPostRedisplay();
}
// --- リスト9：キー入力処理 ---
void keyboard(unsigned char key, int x, int y) {
switch (key) {
case 'Q': case 'q': exit(0); break;
case 'n': case 'N': current_filter = FILTER_NONE; puts("Filter:NONE"); break;
case 'g': case 'G': current_filter = FILTER_GAUSSIAN3x3; activeK = kGaus; puts("Filter:Gaussian 3x3"); break;
case '3': current_filter = FILTER_MEDIAN3x3; activeK = kMed3; puts("Filter:Median 3x3"); break;
case '5': current_filter = FILTER_MEDIAN5x5; activeK = kMed5; puts("Filter:Median 5x5"); break;
case 'f': printf("%f msec\n", msec); break;
}
}
int main(int argc, char** argv) {
if (!camera.isOpened()) return -1;
width = (int)camera.get(cv::CAP_PROP_FRAME_WIDTH);
height = (int)camera.get(cv::CAP_PROP_FRAME_HEIGHT);
cl_platform_id p; clGetPlatformIDs(1, &p, NULL);
cl_device_id d; clGetDeviceIDs(p, CL_DEVICE_TYPE_GPU, 1, &d, NULL);
context = clCreateContext(NULL, 1, &d, NULL, NULL, NULL);
queue = clCreateCommandQueueWithProperties(context, d, NULL, NULL);
cl_image_format fmt = { CL_BGRA, CL_UNSIGNED_INT8 };
cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, (size_t)width, (size_t)height, 0, 0, 0, 0, 0, 0 };
input = clCreateImage(context, CL_MEM_READ_ONLY, &fmt, &desc, NULL, NULL);
output = clCreateImage(context, CL_MEM_WRITE_ONLY, &fmt, &desc, NULL, NULL);
cl_program prog = clCreateProgramWithSource(context, 1, &kernelSource, NULL, NULL);
clBuildProgram(prog, 1, &d, NULL, NULL, NULL);
kGaus = clCreateKernel(prog, "gaussian3x3", NULL);
kMed3 = clCreateKernel(prog, "median3x3", NULL);
kMed5 = clCreateKernel(prog, "median5x5", NULL);
clSetKernelArg(kGaus, 0, sizeof(cl_mem), &input); clSetKernelArg(kGaus, 1, sizeof(cl_mem), &output);
clSetKernelArg(kMed3, 0, sizeof(cl_mem), &input); clSetKernelArg(kMed3, 1, sizeof(cl_mem), &output);
clSetKernelArg(kMed5, 0, sizeof(cl_mem), &input); clSetKernelArg(kMed5, 1, sizeof(cl_mem), &output);
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
glutInitWindowSize(width, height);
glutCreateWindow("OpenCL Realtime Filter (Interface)");
glutDisplayFunc(display);
glutKeyboardFunc(keyboard);
glutMainLoop();
return 0;
}