機械学習基礎理論独習

誤りがあればご指摘いただけると幸いです。数式が整うまで少し時間かかります。リンクフリーです。

勉強ログです。リンクフリーです
目次へ戻る

【VS2022】OpenCL のメモ【実験】

Visual Studio

if文で遅くなる？

GPU のコードは if 文を避けるべき、ということはなんとなく知っていたので実験した。

__kernel void test_if(__global float* out, int mode) {
  int id = get_global_id(0);
  float x = (float)id;

  if (mode == 0) {
    // 発散なし（全員同じ経路）
    for (int i = 0; i < 2000000; i++) {
      x = x * 1.000001f + 0.000001f;
    }
  }
  else {
    // 発散あり（8 分岐）
    int branch = id & 7; // 0〜7 の 8 パターン

    if (branch == 0) {
      for (int i = 0; i < 2000000; i++) x = x * 1.000001f + 0.000001f;
    }
    else if (branch == 1) {
      for (int i = 0; i < 2000000; i++) x = x * 1.000002f + 0.000002f;
    }
    else if (branch == 2) {
      for (int i = 0; i < 2000000; i++) x = x * 1.000003f + 0.000003f;
    }
    else if (branch == 3) {
      for (int i = 0; i < 2000000; i++) x = x * 1.000004f + 0.000004f;
    }
    else if (branch == 4) {
      for (int i = 0; i < 2000000; i++) x = x * 1.000005f + 0.000005f;
    }
    else if (branch == 5) {
      for (int i = 0; i < 2000000; i++) x = x * 1.000006f + 0.000006f;
    }
    else if (branch == 6) {
      for (int i = 0; i < 2000000; i++) x = x * 1.000007f + 0.000007f;
    }
    else { // branch == 7
      for (int i = 0; i < 2000000; i++) x = x * 1.000008f + 0.000008f;
    }
  }

  out[id] = x;
}

Nをいくつにしたが忘れたが大きい数字にして、実行すると以下のようになった。
mode0: 0.334 秒 → mode1: 2.59 秒（7.754 倍）
やはり分岐は避けたほうが良さそうだ。

float2 の受け渡し

OpenCL C には float2 が存在する。
ポイントだけ残しておく。

const int N = 24;
std::vector<cl_float2> a(N), b(N), out(N);
for (int i = 0; i < N; i++) {
	a[i].s[0] = i * 1.0f;
	a[i].s[1] = i * 2.0f;
	b[i].s[0] = i * 10.0f;
	b[i].s[1] = i * 20.0f;
}

cl_mem bufA = clCreateBuffer(
	ctx.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
	sizeof(float) * N, a.data(), &err);
cl_mem bufB = clCreateBuffer(
	ctx.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
	sizeof(float) * N, b.data(), &err);
cl_mem bufOut = clCreateBuffer(
	ctx.context, CL_MEM_WRITE_ONLY,
	sizeof(cl_float2) * N, nullptr,	&err);
clSetKernelArg(kb.kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kb.kernel, 1, sizeof(cl_mem), &bufB);
clSetKernelArg(kb.kernel, 2, sizeof(cl_mem), &bufOut);

構造体の受け渡し

特に気を付ける点は無いように思う。
OpenCL C

typedef struct {
  float v0;
  float v1;
  int   i0;
} MyStruct;

__kernel void add_struct(__global const MyStruct* a,
                         __global const MyStruct* b,
                         __global MyStruct* out) {
  int id = get_global_id(0);
  out[id].v0 = a[id].v0 + b[id].v0;
  out[id].v1 = a[id].v1 + b[id].v1;
  out[id].i0 = a[id].i0 + b[id].i0;
}

CPP

typedef struct MyStruct {
	float v0;
	float v1;
	int   i0;
} MyStruct;

const int N = 24;
std::vector<MyStruct> a(N), b(N), out(N);
for (int i = 0; i < N; i++) {
	a[i].v0 = i * 1.0f;
	a[i].v1 = i * 2.0f;
	a[i].i0 = i * 2;
	b[i].v0 = i * 10.0f;
	b[i].v1 = i * 20.0f;
	b[i].i0 = i * 20;
}

cl_mem bufA = clCreateBuffer(
	ctx.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
	sizeof(MyStruct) * N, a.data(), &err);
cl_mem bufB = clCreateBuffer(
	ctx.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
	sizeof(MyStruct) * N, b.data(), &err);
cl_mem bufOut = clCreateBuffer(
	ctx.context, CL_MEM_WRITE_ONLY,
	sizeof(MyStruct) * N, nullptr,	&err);
clSetKernelArg(kb.kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kb.kernel, 1, sizeof(cl_mem), &bufB);
clSetKernelArg(kb.kernel, 2, sizeof(cl_mem), &bufOut);

値渡しの場合は clCreateBuffer() は不要

int の例

int mode = 3;
clSetKernelArg(kernel, 1, sizeof(int), &mode);

cl_float2 の例

cl_float2 offset;
offset.s[0] = 10.0f;
offset.s[1] = 20.0f;
clSetKernelArg(kb.kernel, 1, sizeof(cl_float2), &offset);

OpenCL C には #include がある

GLSL にはありませんが OpenCL C には #include があります。
ただし include path を設定する必要があります。

err = clBuildProgram(out.program, 1, &ctx.device,
                                   "-I kernels",   // ← include パスを渡す
                                   nullptr, nullptr);

kernel1で計算したBufferをkernel2で使う

kernel1実行後にclEnqueueReadBufferする必要は無いってことです。
kernel1,2をまとめて実行できます。

const int N = 8;
std::vector<int> A(N), C(N);

// A を CPU 側でセット
for (int i = 0; i < N; i++) {
	A[i] = i;
}	

// バッファ作成
cl_mem bufA = clCreateBuffer(ctx.context,
														 CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
														 sizeof(int) * N, A.data(), &err);

cl_mem bufB = clCreateBuffer(ctx.context,
														 CL_MEM_READ_WRITE,
														 sizeof(int) * N, nullptr, &err);

cl_mem bufC = clCreateBuffer(ctx.context,
														 CL_MEM_WRITE_ONLY,
														 sizeof(int) * N, nullptr, &err);

clu::CLKernelBundle kb1;
clu::BuildCLKernel(ctx, "kernels/kernel1.cl",
									 "kernel1", "-I kernels", kb1);
clSetKernelArg(kb1.kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kb1.kernel, 1, sizeof(cl_mem), &bufB);

clu::CLKernelBundle kb2;
clu::BuildCLKernel(ctx, "kernels/kernel2.cl",
									 "kernel2", "-I kernels", kb2);
clSetKernelArg(kb2.kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kb2.kernel, 1, sizeof(cl_mem), &bufB);
clSetKernelArg(kb2.kernel, 2, sizeof(cl_mem), &bufC);

// 実行
size_t global = N;

// kernel1 → kernel2 の順に enqueue（GPU 内でパイプラインがつながる）
clEnqueueNDRangeKernel(ctx.queue, kb1.kernel, 1, nullptr, &global, nullptr, 0, nullptr, nullptr);
clEnqueueNDRangeKernel(ctx.queue, kb2.kernel, 1, nullptr, &global, nullptr, 0, nullptr, nullptr);

// 最後にまとめて同期
clFinish(ctx.queue);

// 結果を CPU に読み戻すのは c だけ
clEnqueueReadBuffer(ctx.queue, bufC, CL_TRUE, 0, sizeof(int) * N, C.data(), 0, nullptr, nullptr);

// 結果表示
for (int i = 0; i < N; i++) {
	char buf[256];
	sprintf_s(buf, "c[%d] = %d\n", i, C[i]);
	OutputDebugStringA(buf);	
}

Bufferの書き換え

Bufferの中身の変更の仕方です。
再作成より中身を変更する方が速いようです。

const int N = 8;
std::vector<int> A(N), B(N), C(N);

// --- CPU: A[i] = i ---
for (int i = 0; i < N; i++) A[i] = i;

// GPU バッファ作成
cl_mem bufA = clCreateBuffer(ctx.context,
														 CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
														 sizeof(int) * N, A.data(), &err);

cl_mem bufB = clCreateBuffer(ctx.context,
														 CL_MEM_READ_WRITE,
														 sizeof(int) * N, nullptr, &err);

cl_mem bufC = clCreateBuffer(ctx.context,
														 CL_MEM_READ_WRITE,
														 sizeof(int) * N, nullptr, &err);

// --- kernel1: b[i] = a[i] + 1 ---
clu::CLKernelBundle kb1;
clu::BuildCLKernel(ctx, "kernels/kernel_a1.cl",
									 "kernel_a1", "-I kernels", kb1);
clSetKernelArg(kb1.kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kb1.kernel, 1, sizeof(cl_mem), &bufB);

size_t global = N;
clEnqueueNDRangeKernel(ctx.queue, kb1.kernel, 1, nullptr, &global, nullptr, 0, nullptr, nullptr);

// ★ CPU が b を読む前に同期
clFinish(ctx.queue);

// --- CPU: b を読み出す ---
clEnqueueReadBuffer(ctx.queue, bufB, CL_TRUE, 0,
										sizeof(int) * N, B.data(), 0, nullptr, nullptr);

// ★ debugoutputa で B を出力
OutputDebugStringA("=== B after kernel1 ===\n");
for (int i = 0; i < N; i++) {
	char buf[256];
	sprintf_s(buf, "B[%d] = %d\n", i, B[i]);
	OutputDebugStringA(buf);
}

// --- CPU: A[i] = i * 10 ---
for (int i = 0; i < N; i++) A[i] = i * 10;

// ★ GPU の bufA に書き戻す（これが必須）
clEnqueueWriteBuffer(ctx.queue, bufA, CL_TRUE, 0,
										 sizeof(int) * N, A.data(), 0, nullptr, nullptr);

// --- kernel2: c[i] = a[i] + 1 ---
clu::CLKernelBundle kb2;
clu::BuildCLKernel(ctx, "kernels/kernel_a2.cl",
									 "kernel_a2", "-I kernels", kb2);
clSetKernelArg(kb2.kernel, 0, sizeof(cl_mem), &bufA);
clSetKernelArg(kb2.kernel, 1, sizeof(cl_mem), &bufC);

clEnqueueNDRangeKernel(ctx.queue, kb2.kernel, 1, nullptr, &global, nullptr, 0, nullptr, nullptr);

// ★ CPU が c を読む前に同期
clFinish(ctx.queue);

// --- CPU: c を読み出す ---
clEnqueueReadBuffer(ctx.queue, bufC, CL_TRUE, 0,
										sizeof(int) * N, C.data(), 0, nullptr, nullptr);

// ★ debugoutputa で C を出力
OutputDebugStringA("=== C after kernel2 ===\n");
for (int i = 0; i < N; i++) {
	char buf[256];
	sprintf_s(buf, "C[%d] = %d\n", i, C[i]);
	OutputDebugStringA(buf);
}

Bank Conflict 対応

GPUの Local Memory は高速化のためは一般的に32個の独立した Bank で構成されています。
Bank Conflictとは、複数のスレッドが同バンクの異なるアドレスにアクセスした場合、アクセスの数だけ処理が繰り返され、速度が低下しすることです。
なので float tile[16][16]; の場合、以下のように Bank が割り振られます。

Address Bank
bank[0][0] 0%32=0
bank[1][0] 16%32=16
bank[2][0] 32%32=0
bank[3][0] 48%32=16

以上のように同じ Bank を参照しまいがちで、行列の積計算の時には Bank Conflict が発生します。

これを避けるために float tile[16][17]; とします。

Address Bank
bank[0][0] 0%32=0
bank[1][0] 17%32=17
bank[2][0] 34%32=2
bank[3][0] 51%32=19

以上のように衝突しにくくなります。

目次へ戻る