CUDAを使ったプログラミング その2 総当たり計算

前回は、データ検索プログラムにCUDAを使おうとして失敗しました

その原因は、検索結果の出力位置が決まっていないことと、データをGPUにコピーする時間がデータ比較より遅かったことでした

今回は、正しい方法でCUDAを使い、プログラムの高速化に挑戦してみたいと思います

各点から最も近い点を求める

サンプルプログラムはこちら

開発環境は Visal Studio 2022 C++(MFC)とCUDA 12.0です

プログラムの内容は、画面内にランダムな点を配置し、その点を一定の方向、一定の速度で直進させ、最も近い点同士をリアルタイムに直線で結ぶという簡単なものです

内容は簡単ですが、画面内の各点から最も近い点を求める処理は、点の個数×点の個数という大きな計算負荷が掛かります。

CPUを使ったプログラム(最も近い点を求める部分)の抜粋

// 全オブジェクト、一番距離の近い点を計算
	int calc_near_id() {
		int start = clock();
		int n = chars.size();
		float dist, d;
		for (int i = 0; i < n; i++) {
			dist = 10000;
			for (int j = 0; j < n; j++) {
				if (i == j)
					continue;
				d = sqrtf(
					powf(chars[i].x - chars[j].x, 2) +
					powf(chars[i].y - chars[j].y, 2) );
				if (dist > d) {
					dist = d;
					chars[i].near_id = j;
				}
			}
		}
		ms = clock() - start;
		return true;

内容は三平方の定理を使い、距離が最も近いIDを求めるという簡単なものです

2重ループで総当たりの計算処理をしています

実行の結果です(CPU処理)

点を2,000個配置すると4,000,000回の計算になります。一度の総当たり計算に65ミリ秒掛かってます

次に、点を10,000個置いたときの結果です

描画がカクカクになりました。左上を見ると一度の計算に約1.6秒掛かってることが分かります

10,000個の点では100,000,000回の計算回数、データ量が5倍に増えると計算量は25倍になるということですね

CPUを使った計算ではこの辺が限界です

GPUには数千個の計算コアがあり、このようなものを同時並列的に計算をするのが得意です

次が、GPUを使ったプログラム(の抜粋)です

	// 全オブジェクト、一番距離の近い点を計算(CUDA使用
	int calc_near_id_cuda() {
		int start = clock();

		for (int i = 0; i < count; i++) {
			pBufX[i] = chars[i].x;
			pBufY[i] = chars[i].y;
		}
		cuda_near_search(pBufX, pBufY, pBufResult, count);

		for (int i = 0; i < count; i++) 
			chars[i].near_id = pBufResult[i];		

		ms = clock() - start;
		return 0;
	}

CUDAのカーネルを呼び出すときは、構造体配列を渡すのではなく、それぞれの要素を配列変数にして渡します

[X座標の配列]、[Y座標の配列]、[最も近いIDを格納する配列]といった感じです

ループのたびにメモリ確保(malloc)をしなくて良いように、あらかじめ配列を確保しています

CUDA側の処理は

#include <stdio.h>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cudaMain.cuh"

__global__ void gpu_near_search(float* dev_x, float* dev_y, int* results, int cnt)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;	// インデックスを計算
	int i;
	float dist = 10000, d;
	for (i = 0; i < cnt; i++) {
		if (idx == i)
			continue;
		d = sqrtf(
			powf(dev_x[idx] - dev_x[i], 2) +
			powf(dev_y[idx] - dev_y[i], 2));
		if (dist > d) {
			dist = d;
			results[idx] = i;
		}
	}
}

int cuda_near_search(float* host_x, float* host_y, int* results, int cnt) {
	memset(results, -1, cnt);

	float* dev_x, * dev_y;
	int* dev_results;
	cudaMalloc(&dev_x, cnt * sizeof(float));
	cudaMalloc(&dev_y, cnt * sizeof(float));
	cudaMalloc(&dev_results, cnt * sizeof(int));

	cudaMemcpy(dev_x, host_x, cnt * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_y, host_y, cnt * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_results, results, cnt * sizeof(int), cudaMemcpyHostToDevice);

	gpu_near_search <<< (cnt + 255) / 256, 256 >>> (dev_x, dev_y, dev_results, cnt);

	cudaMemcpy(results, dev_results, cnt * sizeof(int), cudaMemcpyDeviceToHost);
	return 0;
}

GPUメモリに配列データをコピーして、共通の関数(カーネル)を呼び出すという手順です

cuda_near_searchが外部から呼び出すインターフェースで、gpu_near_searchがGPU用のカーネルになります

実行の結果です(点を2,000個置いた場合)

65ms→10msと高速化してますね

次が点を10,000個置いたとき

圧倒的な速度向上です

このように2重ループで総当たりするような計算はCUDA(GPU)を使うことで高速化することが出来ます

どんな事に使える？

CUDAは主に科学計算やAIなどに用いられると聞きますが、今回のように、2重ループで総当たりするような計算ならば何にでも使える思います

例えば、多人数が参加するオンラインゲームなどでパケットを伝える相手を絞り込みたいとき、シミュレーションゲームで近くにいるキャラ同士でイベントを起こしたいとき、シューティングゲームの当たり判定にも良いですね。(ゲームばかり)

他には迷路探索などでも力を発揮しそうです。CUDAは順次的な処理はできませんが、1歩進むごとに関数を1回呼び出す、つまりループでカーネルを呼び出せば、ゴールまでの最短回数を呼び出すだけで迷路が解けそうです。ゴールからも逆算させれば、半分の呼び出し回数で済みそうですね。

個人的な用途では、自作AIの情報類似度を計算する部分にCUDAを組み込んでみようと思います。以前作ったボロノイ図アプリもCUDAを使えば高速化できそうですね。

描画の観察

適当に考えて作ったプログラムですが、眺めていると面白い法則に気付きます

• 線が他の線を跨がない
• 三角形(閉じた図形)は作られない
• 近いもの同士でペアにならず複数の点が連結される
• 放っておくと中心に空洞が出来る

全ての点がランダムな方向に直進しているだけなのに、調和の取れた結果が出てきます

ここから、世界の仕組みについて気付けることがないでしょうか

例えば宇宙が膨張する理由。すべての物体がランダムな方向に直進すると、中心に空洞ができ、速度ごとに層ができ、ある程度時間が経つと、全ての点から見て、全ての点が遠ざかっているように見えます。

人間関係にも例えられそうです。最も近いもの同士を線で結んでも、ペアにはならない。AにとってBが最も大事な人であっても、BにとってはCが最も大事な人かもしれない。ここに組織形成の原理みたいなものが感じ取れます。片思いによって社会が大きくなる説を提唱します(おい

応用として、知識を空間上にランダムに配置し、知識ベクトル方向にランダムな速度で直進させる。最も近いもの同士を結び付け、新しいアイデアを生み出させる。というのはどうでしょうか