C++を使ってRubyのメソッドを書く(その3) openACCを使ってGPUで計算 #C++

Cを使ってRubyのメソッドを書く(その1)
Cを使ってRubyのメソッドを書く(その2) Numo::NArray
C++を使ってRubyのメソッドを書く(その1)
C++を使ってRubyのメソッドを書く(その2) ベンチマーク
C++を使ってRubyのメソッドを書く(その3) openACCを使ってGPUで計算 <- ここ

参考

C/C++で書いたループを高速化する(その1) ベクトル化
C/C++で書いたループを高速化する(その2) openACC

はじめに

これまでSWIGを使ってRubyの拡張ライブラリを作る方法を書いてきましたが、今回はopenACCを使ってGPUで計算する方法を紹介します。
まず、Nvidiaから無料で利用できるコンパイラが公開さえていますのでをインストールしてください。参考にある「C/C++で書いたループを高速化する(その1)」にリンクがあります。

GPUでの計算はそこそこの計算量がないと有り難みがわからないので、XYZの3次元座標の相互距離の平均値を計算するプログラムを作ってみました。プログラムは下に書いてあります。
普通のC++プログラムですが、openACCを使うためにループの直前に#pragma acc 〜の1行が追加されています。これだけで二重のforループがGPUで実行されるようになります。

average_distance1.cpp

#include <cmath>

inline double distance(double x1, double y1, double z1, double x2, double y2, double z2){
    return(sqrt((x1-x2)*(x1-x2)+(y1-y2)*(y1-y2)+(z1-z2)*(z1-z2)));
}

double average_distance1(int n, const double xyz[][3]){
    double sum=0.0;
    #pragma acc kernels
    for(int i=0; i<n; i++)
        for(int j=i+1; j<n; j++)
            sum+=distance(xyz[i][0], xyz[i][1], xyz[i][2], xyz[j][0], xyz[j][1], xyz[j][2]);
    return(sum/(n*(n-1.0)/2.0));
}

次にSWIG用のインターフェースファイルを作ります。
二次元配列を関数に渡すのにNArrayを使っています。この辺りの書き方は「Cを使ってRubyのメソッドを書く(その2) Numo::NArray」も参考にしてください。

test.i

%module testF
%{
#include "numo/narray.h"
extern double average_distance1(int n, const double xyz[][3]);
%}

%typemap(in) (int LENGTH, const double xyz[][3]){
  narray_t *nary;
  if (rb_obj_class($input)!=numo_cDFloat){
      $input = rb_funcall(numo_cDFloat, rb_intern("cast"), 1, $input);
  }
  GetNArray($input, nary);
  if (NA_TYPE(nary)==NARRAY_VIEW_T){
    $input = rb_funcall($input, rb_intern("dup"), 0);
    GetNArray($input, nary);
  }
  $2 = ($2_ltype)na_get_pointer_for_read($input);
  $1 = NA_SHAPE(nary)[1];
}
extern double average_distance1(int LENGTH, const double xyz[][3]);

続いてMakefile作成用のプログラムです。numo/narray.hを使っていますので、その場所を指定できるようにしています。g++を前提にMakefileが作られますが、今回はnvc++を使います。3行目、4行目はnvc++が解釈できないオプションを回避するためです。
実際にnvc++でコンパルしてみるとSWIGで作られたtest_wrap.cxxでエラーが出てしまいます。しかたがないので、これだけはg++でコンパイルするようにしました。

extconf.rb

require 'mkmf'
dir_config("numo/narray")
have_header("numo/narray.h")
$warnflags="-Wall"
$LDFLAGS="-L. -rdynamic -Wl,-export-dynamic"
create_makefile("testF")

全部のファイルを1つのフォルダに入れて、以下のコマンドを実行します。

swig -c++ -ruby test.i
ruby extconf.rb  -- --with-numo/narray-include=/opt/lib/ruby/gems/2.6.0/gems/numo-narray-0.9.1.8/lib/numo/
nvc++ -acc -fast -fPIC -Minfo=all -c average_distance1.cpp
make

２行目でnumo/narray.hの場所を教えていますので、環境に合わせて書き換えてください。narray.hではなくこれが入っているnumoフォルダに位置ですので間違えないように。

average_distance1.cppはnvc++でGPUが使えるようにコンパイルしています。makeで残りのファイルtest_wrap.cxxをg++でコンパイルし、g++でリンクします。
しかし、g++でリンクした場合には、nvc++で必要なライブラリがリンクされていませんので、nvc++でリンクをやり直します。

rm testF.so
make --dry-run

とすると、リンクコマンドが出力されます。
先頭のg++部分をnvc++ -acc -fastに書き換えて実行します。
こちらの環境では次のようになりました。

nvc++ -acc -fast -shared -o testF.so average_distance1.o test_wrap.o -L. -L/opt/lib -Wl,-rpath,/opt/lib -L. -rdynamic -Wl,-export-dynamic -Wl,--compress-debug-sections=zlib    -Wl,-rpath,/opt/lib -L/opt/lib -lruby  -lm   -lc

以上で拡張ライブラリファイル testF.so ができます。
まとめると、test_wrap.cxxだけはg++でコンパイルし、それ以外はnvc++でコンパイル、リンクをしています。

テスト用のプログラムです。"./testF" より先に "numo/narray"を読みこまないとエラーになります。

test.rb

require "numo/narray"
require "./testF"

data=Numo::DFloat.new(1024,3).rand*100
p TestF.average_distance1(data)

GPUが動いているかは、ターミナルで次のコマンドを実行しておくと簡易プロファイラからの出力が追加されます。
export PGI_ACC_TIME=1

GPUを使わずにCPUだけを使う場合

nvc++のコマンドで-accを抜くとCPUだけで実行できるようになります。

CPU Intel Core i7-2600 CPU 3.40GHz GPU GTX 750
CentOS 8.2, SWIG 3.0.12, nvc++ 20.7

C++を使ってRubyのメソッドを書く(その3) openACCを使ってGPUで計算

目次

参考

はじめに

GPUを使わずにCPUだけを使う場合