GPGPUのShared Memoryを使う

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

 行列積のFortranプログラムを、CPUで実行(約45秒)、ディレクティブを入れるだけでGPGPUで実行(約9.3秒)、CUDAぽいプログラムにして実行(約5.3秒)、といろいろテストしてきました。やはりGPGPUは速いねと思うか、たいしたことないねと思うかは、人それぞれでしょう。

 GPGPUは、以下のプログラムで使っているShared Memoryを有効活用しなければ、宝の持ち腐れになります。これまではまだ使っていませんでした。Shared MemoryとはGPGPUの中にある非常に高速なメモリの事です。しかし、同一Block内からしかアクセスできない、容量が16KBと非常に小さい、という2つの大きな制限があります。この制限を乗り越えて使いこなせれば、得るものは大きいというわけです。

 わかっていない私が能書きを言っていても仕方がないので、早速コンパイルして実行してみると、0.72秒と非常に高速です。CPUの45秒に比べると、なんと約76倍高速になりました。ここまで来ると、GPGPUを使わずにはいられないという気になります。如何でしょうか?

$ cat matmul2.cuf
MODULE CUDAMOD
  USE CUDAFOR
  INTEGER, PARAMETER :: N = 4000
  CONTAINS
    ATTRIBUTES(GLOBAL) SUBROUTINE MATMUL(A, B)
      REAL, DIMENSION(N, *) :: A, B
      REAL, SHARED :: AS1(16, 16), AS2(16, 16)
      IDX = THREADIDX%X
      IDY = THREADIDX%Y
      II = (BLOCKIDX%X - 1) * 16 + IDX
      JJ = (BLOCKIDX%Y - 1) * 16 + IDY
      BI = 0.0
      DO I = 1, N, 16
         AS1(IDX, IDY) = A(II, I + IDY - 1)
         AS2(IDX, IDY) = A(I + IDX - 1, JJ)
         CALL SYNCTHREADS()
         DO J = 1, 16
            BI = BI + AS1(IDX, J) * AS2(J, IDY)
         END DO
         CALL SYNCTHREADS()
      END DO
      B(II, JJ) = BI
    END SUBROUTINE MATMUL
END MODULE CUDAMOD

PROGRAM MATMUL2
  USE CUDAFOR
  USE CUDAMOD
  REAL :: A(N, N), B(N, N)
  REAL*8 :: SUM
  REAL, DEVICE :: AD(N, N), BD(N, N)
  TYPE(DIM3) :: DIMGRID, DIMBLOCK
  DO J = 1, N
     DO I = 1, N
        A(I, J) = REAL(I * N + J) / REAL(N * N)
     END DO
  END DO
  CALL CPU_TIME(T1)
  AD = A
  DIMGRID = DIM3(N / 16, N / 16, 1)
  DIMBLOCK = DIM3(16, 16, 1)
  CALL MATMUL<<<DIMGRID, DIMBLOCK>>>(AD, BD)
  B = BD
  CALL CPU_TIME(T2)
  WRITE(*, *) T2 - T1, 'SEC'
  SUM = 0.0
  DO J = 1, N
     DO I = 1, N
        SUM = SUM + B(I, J)
     END DO
  END DO
  WRITE(*, *) 'CHECKSUM =', SUM
END PROGRAM MATMUL2
$ pgf95 -fastsse matmul2.cuf -o matmul2
$ ./matmul2
   0.7210441     SEC
 CHECKSUM =    16017338556.46359     
$

--ケンちゃん 2010年3月6日 (土) 14:45 (JST)


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

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

<comments hideform="false" />


Comments

ノート:GPGPUのShared Memoryを使う

個人用ツール