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

　行列積のプログラムを、CUDAぽく、GridとBlockを使ってWVさん(知らない人はこちらへ)に書いてもらったのが以下のプログラムです。どうも私にはまだ理解できません。GridとBlockが出てくると、拒絶反応を起こしてしまいます。逆に言えば、GridとBlockの使い方が理解できれば、GPGPUを半分くらい使いこなせるような気がします。後の半分は、Shared Memoryを使いこなす事だと思います。こちらも私にはまだできていません。WVさんにすべてお任せですね。

　ともあれこのプログラムをコンパイルして実行してみると、実行時間は5.3秒と結構速くなりました。CPU1Coreと比べると約8.5倍速くなっています。ここまで来ると、苦労してGPGPUを使う価値もあるねえ、というところでしょうか。

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

PROGRAM MATMUL3
  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 MATMUL3
$ pgf95 -fastsse matmul3.cuf -o matmul3
$ ./matmul3
    5.321247     SEC
 CHECKSUM =    16017338556.46359     
$

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