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

　行列積の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)