メニーコアと格闘ブログ:PGIコンパイラの新機能の性能検証~その1:実装熱伝導解析への適用~

出典: トータル・ディスクロージャ・サイト(事実をありのままに)

目次

PGIコンパイラ

CUDAを使わなくてもPGI コンパイラ(ver 9.03)を用いてfortran、Cで書かれたプログラムで並列可能な部分をGPU上で、それ以外の部分をCPU上で実行することができます。
使用方法は簡単で、GPU上で並列に計算したいプログラム領域(たとえばforループ)にディレクティブ(Directive)と呼ばれる指示文を挿入し、-ta=nvidiaというオプションでコンパイルします。
下記のサイトが役に立ちます。

http://www.pgroup.com/resources/accel.htm
http://www.pgroup.com/lit/whitepapers/pgi_accel_prog_model_1.0.pdf

このPGIコンパイラの特徴、性能などを把握するため、それを実装熱伝導解析プログラムに適用することを試みました。途中までの結果をここに掲載します。

使用環境


  • CPU: Intel(R) Xeon(R) CPU E5540 @ 2.53GHz
  • GPU: Tesla C1060

結果1(CPU上での実行)

熱伝導解析プログラム名(sfdm2d.txt)を使用しました。
メディア:sfdm2d.txt
CPU上での実行すると34秒かかります。

[kim@localhost ~]$ pgcc sfdm2d.c
[kim@localhost ~]$ time ./a.out
gamma = 1.000244e-02
t = 0.000000
q_in = 3.200000e+06, q_out = 0.000000e+00, dQ = 3.200000e+06
FILENAME=0.000000.bmp
t = 10.000000
q_in = 6.786950e+04, q_out = 7.252066e-01, dQ = 6.786877e+04
FILENAME=10.000000.bmp
t = 10.000000
q_in = 6.786950e+04, q_out = 7.252066e-01, dQ = 6.786877e+04
FILENAME=10.000000.bmp
real    0m34.222s
user    0m34.206s
sys     0m0.008s

結果2

次にディレクティブを挿入することを試みます。
まず、271行目付近のforループの上に#pragma acc regionというディレクティブを入れてコンパイルします。
メディア:sgfdm2d01.txt
すると下記のような結果が返ってきます。

[kim@localhost ~]$ pgcc -ta=nvidia,time -Minfo sgfdm2d01.c
main:
   271, Accelerator region ignored
   272, Accelerator restriction: loop has multiple exits
        Accelerator restriction: loop contains scalar assignment live after the loop
   273, Accelerator restriction: scalar variable live-out from loop: t
   284, Accelerator restriction: scalar variable live-out from loop: dq
   292, Accelerator restriction: scalar variable live-out from loop: nprint
   299, Memory copy idiom, loop replaced by call to __c_mcopy8

変数t、q_in、q_outなどがforループの後で使われるため並列化ができません。
これらの変数が現れるところとそのほか、とりあえず、重要ではないところをコメントアウトします(行の先頭に//の挿入)。(sgfdm2d02.c)

[kim@localhost ~]$ pgcc -ta=nvidia,time -Minfo sgfdm2d.c
main:
   271, Accelerator region ignored
   272, Accelerator restriction: function/procedure calls are not supported
   275, Accelerator restriction: function/procedure calls are not supported
   299, Memory copy idiom, loop replaced by call to __c_mcopy8

並列化される部分(アクセラレータ領域という)では関数の呼び出しができないため、オプション-Minlineを追加してコンパイルします。
メディア:sgfdm2d02.txt
するとコンパイルできますが、これを実行すると8m55sかかります。
213行目と298行目のforループの計算にほとんどの時間をとられています。

[kim@localhost ~]$ pgcc -ta=nvidia,time -Minfo -Minline sgfdm2d02.c
print_temp:
   123, write_bmp_mono inlined, size=95, file sgfdm2d02.c (38)
main:
   263, calc_q inlined, size=13, file sgfdm2d02.c (127)
   268, print_temp inlined, size=15, file sgfdm2d02.c (109)
   271, Generating copy(T[0:255][0:255])
        Generating local(T_new[0:255][0:255])
   272, Loop carried dependence due to exposed use of T[255][0:128] prevents parallelization
        Loop carried dependence due to exposed use of T[127:255][0] prevents parallelization
        Loop carried dependence due to exposed use of T[0:254][0:255] prevents parallelization
        Loop carried dependence due to exposed use of T[0:255][1:255] prevents parallelization
        Loop carried dependence due to exposed use of T_new[0:255][0:255] prevents parallelization
        Sequential loop scheduled on host
        Generating copyin(T[0:255][1:255])
   275, solve inlined, size=121, file sgfdm2d02.c (139)
        145, Loop is parallelizable
             Accelerator kernel generated
            145, #pragma for parallel, vector(254)
                 Cached references to size [2x256] block of 'T'
             Generating copyout(T_new[0][1:254])
        154, Loop is parallelizable
             Accelerator kernel generated
            154, #pragma for parallel, vector(127)
                 Cached references to size [2x129] block of 'T'
             Generating copyout(T_new[0:255][1:254])
        163, Loop is parallelizable
             Accelerator kernel generated
            163, #pragma for parallel, vector(32)
                 Non-stride-1 accesses for array 'T'
                 Non-stride-1 accesses for array 'T_new'
             Generating copyout(T_new[0:255][1:255])
        172, Loop is parallelizable
             Accelerator kernel generated
            172, #pragma for parallel, vector(32)
                 Non-stride-1 accesses for array 'T'
                 Non-stride-1 accesses for array 'T_new'
             Generating copyout(T_new[0:255][0:255])
        213, Loop is parallelizable
             Accelerator kernel generated
            213, #pragma for parallel, vector(254)
             Generating copyout(T_new[0:255][0:255])
        214, Loop is parallelizable
   278, Loop is parallelizable
        Accelerator kernel generated
       278, #pragma for parallel, vector(32)
            Non-stride-1 accesses for array 'T_new'
            Non-stride-1 accesses for array 'T'
        Generating copyout(T[0:127][0])
        Generating copyout(T_new[0:255][0:255])
   283, calc_q inlined, size=13, file sgfdm2d02.c (127)
        129, Accelerator restriction: scalar variable live-out from loop: q_in
             Scalar last value needed after loop for q_in
        130, Scalar last value needed after loop for q_in
             Loop carried scalar dependence for q_in
             Inner sequential loop scheduled on host
             Generating copyin(T_new[0:127][0:1])
        131, Accelerator restriction: scalar variable live-out from loop: q_in
        132, Accelerator restriction: scalar variable live-out from loop: q_out
        133, Scalar last value needed after loop for q_out
        134, Scalar last value needed after loop for q_out
             Loop carried scalar dependence for q_out
             Inner sequential loop scheduled on host
             Generating copyin(T_new[254:255][128:255])
             Generating copyout(T[0:127][0])
        135, Accelerator restriction: scalar variable live-out from loop: q_out
   298, Loop is parallelizable
        Accelerator kernel generated
       298, #pragma for parallel, vector(256)
        Generating copyin(T_new[0:255][0:255])
        Generating copyout(T_new[0:255][0:255])
   299, Loop is parallelizable
   308, print_temp inlined, size=15, file sgfdm2d02.c.c (109)
[kim@localhost ~]$ time ./a.out
gamma = 1.000244e-02
t = 0.000000
q_in = 3.200000e+06, q_out = 0.000000e+00, dQ = 3.200000e+06
FILENAME=0.000000.bmp
t = 0.000000
q_in = 3.200000e+06, q_out = 1.920000e+06, dQ = 3.200000e+06
FILENAME=0.000000.bmp
Accelerator Kernel Timing data
sgfdm2d02.c.c
 main
   271: region entered 1 time
       time(us): total=535441063 init=1463245 region=533977818
                 kernels=232007561 data=301970257
       w/o init: total=533977818 max=533977818 min=533977818 avg=533977818
       145: kernel launched 100000 times
           grid: [1]  block: [254]
           time(us): total=2325919 max=18895 min=19 avg=23
       154: kernel launched 100000 times
           grid: [1]  block: [127]
           time(us): total=2276789 max=18150 min=18 avg=22
       163: kernel launched 100000 times
           grid: [8]  block: [32]
           time(us): total=2371997 max=22890 min=20 avg=23
       172: kernel launched 100000 times
           grid: [4]  block: [32]
           time(us): total=2445192 max=40027 min=18 avg=24
       213: kernel launched 100000 times
           grid: [1]  block: [254]
           time(us): total=158174076 max=41245 min=1352 avg=1581
       278: kernel launched 100000 times
           grid: [4]  block: [32]
           time(us): total=2177553 max=24029 min=17 avg=21
       298: kernel launched 100000 times
           grid: [1]  block: [256]
           time(us): total=62236035 max=40492 min=487 avg=622
real    8m55.539s
user    3m11.322s
sys     4m31.183s

この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:メニーコアと格闘ブログ:PGIコンパイラの新機能の性能検証~その1:実装熱伝導解析への適用~

個人用ツール