GPGPUで爆速でオセロを解く 【KMC Advent Calendar 2016 19日目】

はじめに

この記事はKMC Advent Calendar 2016の19日目の記事です。

www.adventar.org

GPGPUとは

GPGPU（General-purpose computing on graphics processing units; GPUによる汎用計算）とは、GPUの演算資源を画像処理以外の目的に応用する技術のことである。

(出典: GPGPU - Wikipedia )

GPUは並列演算に特化しており、大量のスレッドをハードウェアで処理でき、汎用CPUに比べて演算処理性能が高いという特徴があるので、これを用いて演算を高速に行おうというのがGPGPUのアイデアです。

ただ、良いことばかりではなく、汎用CPUに比べてプログラミング上の制約が大きいなど、性能を出し切るのはかなり難しいです。

GPGPUでオセロを解く

今回はGPGPUを用いてオセロの終盤の完全読み*1を高速化します。

オセロを解く際に基本となるのがアルファベータ法です。

詳しい解説は他に譲るとして、このアルゴリズムの肝はすでに計算した部分の結果を利用して枝刈りを行うことです。そのため、並列化が難しいアルゴリズムになっています。 今回はアルゴリズムレベルでの並列化は後回しにして、たくさんの局面を並列で完全読みすることを考えます。

各スレッドに一つの局面を割り当てて完全読みを行うのが素直な実装ですね。

アルファベータ法は再帰を用いると自然に書けるので再帰で実装してみます。 今回はCUDAで実装します(したがって、この記事では以後GPUはNVIDIAのGPUと仮定します)。CUDAで再帰するときは必要なスタック数が静的に見積もれないので必要に応じて動的にスタックサイズを指定してやります（今回は4096バイト/スレッドにしました）。*2 高速化のためビット演算を用いてひっくり返る石の位置を計算しています。詳しくはReversi - bitboard and GPU computingを参照してください。

CUDAカーネル部分だけ示します（CPU側のコードはスタックサイズの指定を除いて後で示す改良バージョンと同じです）。

constexpr int threadsPerBlock = 128;

using ull = unsigned long long;

struct Board {
  ull player;
  ull opponent;
  __host__ __device__ Board() = default;
  __host__ __device__ Board(ull p, ull o) : player(p), opponent(o) {}
  __host__ __device__ Board(const Board &) = default;
  __host__ __device__ Board(Board &&) = default;
  __host__ __device__ Board &operator=(const Board &) = default;
  __host__ __device__ Board &operator=(Board &&) = default;
};

__shared__ int count[threadsPerBlock];

__device__ char score(const Board &bd) {
  char pnum = __popcll(bd.player);
  char onum = __popcll(bd.opponent);
  if (pnum == onum) return 0;
  if (pnum > onum) return 64 - 2*onum;
  else return 2*pnum - 64;
}
__constant__ ull mask1[4] = {
  0x0080808080808080ULL,
  0x7f00000000000000ULL,
  0x0102040810204000ULL,
  0x0040201008040201ULL
};
__constant__ ull mask2[4] = {
  0x0101010101010100ULL,
  0x00000000000000feULL,
  0x0002040810204080ULL,
  0x8040201008040200ULL
};

__device__ ull flip(const Board &bd, int pos, int index) {
  ull om = bd.opponent;
  if (index) om &= 0x7E7E7E7E7E7E7E7EULL;
  ull mask = mask1[index] >> (63 - pos);
  ull outflank = (0x8000000000000000ULL >> __clzll(~om & mask)) & bd.player;
  ull flipped = (-outflank * 2) & mask;
  mask = mask2[index] << pos;
  outflank = mask & ((om | ~mask) + 1) & bd.player;
  flipped |= (outflank - (outflank != 0)) & mask;
  return flipped;
}

__device__ ull flip_all(const Board &bd, int pos) {
  return flip(bd, pos, 0) | flip(bd, pos, 1) | flip(bd, pos, 2) | flip(bd, pos, 3);
}

__device__ Board move_pass(const Board &bd) {
  return Board(bd.opponent, bd.player);
}

__device__ char alpha_beta(const Board &bd, char alpha, char beta, bool passed) {
  ++count[threadIdx.x];
  ull puttable = ~(bd.player | bd.opponent);
  if (puttable == 0) return score(bd);
  bool pass = true;
  for (; puttable;) {
    ull bit = puttable & -puttable;
    puttable ^= bit;
    int pos = __popcll(bit-1);
    ull flipped = flip_all(bd, pos);
    if (flipped) {
      pass = false;
      Board next(bd.opponent ^ flipped, (bd.player ^ flipped) | bit);
      alpha = max(alpha, -alpha_beta(next, -beta, -alpha, false));
      if (alpha >= beta) return alpha;
    }
  }
  if (pass) {
    if (passed) {
      return score(bd);
    } else {
      return -alpha_beta(move_pass(bd), -beta, -alpha, true);
    }
  }
  return alpha;
}

__global__ void search_noordering(const Board *bd_ary, int *res_ary, int *nodes_count, const size_t size) {
  size_t index = threadIdx.x + blockIdx.x * blockDim.x;
  while(index < size) {
    count[threadIdx.x] = 0; 
    res_ary[index] = alpha_beta(bd_ary[index], -64, 64, false);
    nodes_count[index] = count[threadIdx.x];
    index += blockDim.x * gridDim.x;
  }
}

GTX1080(ZOTAC Geforce GTX 1080 AMP Edition, Base 1683MHz/Boost 1822MHz)でGGSの棋譜から生成した10石空きの問題1091780局面の完全読みを実行してみると5分48秒かかりました。

一方、CPUで各種高速化を施したソルバー*3ではCore i7-6700K上で同じ処理を行うのに1分13秒しかかかりませんでした。これではGPGPUにした意味がほとんどありません。

高速化する

2018/03/31追記

このプログラムが遅い最も大きな要因は、以下で説明するレイテンシの隠蔽失敗によるものではなく、再帰的な分岐により演算器が遊んでしまうことによるものです。 実際、Shared Memoryを少し節約しつつ、1スレッドで1局面読むようにすると下記の結果の9秒より速く計算することが出来ます。

追記終わり

GPUはSM(Streaming Multiprocessor)という単位でワープを管理していて、GTX 1080（や他のCompute Capability 6.1のGPU)では同時に4つのワープを実行できます。しかし、実際にはSMは4つよりも多くのワープを管理でき、実行可能なワープを選んで実行することで、命令やメモリの読み書きのレイテンシを隠蔽しています。逆に言うと、同時にたくさんのワープを管理できなければレイテンシを隠蔽できず、多くの時間演算器は遊んでしまうことになります。

今回の場合、再帰でレジスタをかなりたくさん消費してしまっているので、一つのSMで持てるワープは数個しかありません。

そこで、Shared Memoryに再帰のスタックの内容を移し、再帰をループに直すことにします。 こうした場合今度はShared Memoryの使用量がSMの持てるワープ数を制限してしまうので、4スレッドで1局面を読むことで使用量を減らし、その代わり4スレッドをSIMD的に使うことでひっくり返る石の場所の計算を高速化します。

また、一つの局面を読み終わったらすぐ次の局面を読めるように変更しました。これにより、ワープの中で最も遅いスレッドが律速になるのを軽減できます。

その他の細かい改良を含めた結果、実行時間は最終的に9秒になりました。CPUの8倍高速です。

ソースコード及びベンチデータ

ソースコードは以下の通り。 問題ファイルの読み込みはBase81という1文字で4マス分表すフォーマットです。詳しくはGitHub - primenumber/issen: オセロAIの高速な実装を参照。

CUDA Othello Solver · GitHub

ベンチマークに使用した入力ファイルは次のリンクからダウンロードできます

Google Drive: Sign-in

今後の課題

• 枝刈りの成功率を上げるため、次の手を読む順番を良さそうな順番になるように並べ替える。 並べ替えのためにメモリをかなり消費するので、グローバルメモリ上に置かないと実行効率が上がらないと予想される。

• 多数の局面を並列に読むのではなく、アルゴリズムを工夫して一つの局面を並列に読めるようにする。 あわよくばCPUで読んだときより速くならないだろうか…

参考文献

Reversi - bitboard and GPU computing 五石空きをminimax法で完全読みするのをGPUにやらせてみたがあまり速くなかったという結果。　ビット演算部分はここを参考にした。

GPU Acceleration for Board Games | NVIDIA Developer この記事と似ていて、並列にオセロのAIを走らせているらしい。Windowsを持ってないので動作を確かめられない。

*1:その局面から両者が最善を尽くしたときに、勝敗及び石差がどうなるかを判定すること。

*2:http://primenumber.hatenadiary.jp/entry/2016/11/27/031429

*3:GitHub - primenumber/issen: オセロAIの高速な実装