■Pinned Memory
Fortran で Pinned Memory というものを使うと転送が早くなるというので、それを用いることを考えます。第1巻: CUDAプログラミング入門 (日本語版) http://www.nvidia.co.jp/docs/IO/59373/VolumeI.pdf の 88 ページに記述があります。
まずは NVIDIA 提供のアプリで適正な値を調べておきます。bandwidthtest.exe でオプション -memory=pinned をつけることで調べることが可能です。
■実行結果
CPU 側から GPU 側へメモリー内容を転送し、GPU 内部でコピーを行い、そのコピーしたものを GPU 側から CPU 側に転送して戻してやるということを、3セット*500回行います。これは NVIDIA の入門文書に出てくる典型的な例です。
ここで cudaMalloc で GPU 上に確保できる配列の最大値は 2**24 =64Mbyte でした。またこの大きさの配列は合計5個までしか確保できませんでした。 64Mbyte*5=320Mbyte この数値が何を意味するのか分かりません。
CPU 側での配列の確保は ALLOCATE ではなく、cudaMallocHost で行います。またポインターの結びつけは、NVIDIA の文書にしたがって Fortran2003 の ISO_C_BINDING の機能で行います。Cray 型のポインターより少し厄介です。
転送=コピー=転送のサイクルで 64Mbyte * 3 * 500 = 96Gbyte(片道) のデータを約 40 秒で処理しています。これは約4.8Gbyte/s に相当します。個別の転送過程のみを調べると、CPU−GPU 間の転送は 5.5 Gbyte/sec 前後で行われていることがわかりました。これは、bandwidthtest の結果にほぼ一致します。また CPU→GPU 型の転送の方が GPU→CPU 型よりやや早いなどという傾向も一致します。
その一方 GPU-GPU 間の転送は、CPU側のクロックでは計れなかったので、全過程から CPU→GPU、GPU→CPU にかかった時間を差し引くことで見積もりました。その結果は 18Gbyte/sec 程度で bandwidthtest の結果の丁度半分になりました。理由はいまのところ分かりませんが、送・受信でファクター2が出るのかもしれません。
■Fortran ソース
必要なライブラリは cudart.lib のみです。すべて Fortran 側から CUDA API を呼び出すので、cu ファイルは必要ありません。
MODULE m_cuda USE, INTRINSIC :: ISO_C_BINDING IMPLICIT NONE ! ENUM, BIND(C) !cudaMemcpyKind ENUMERATOR :: cudaMemcpyHostToHost = 0 ENUMERATOR :: cudaMemcpyHostToDevice ENUMERATOR :: cudaMemcpyDeviceToHost ENUMERATOR :: cudaMemcpyDeviceToDevice END ENUM ! ENUM, BIND(C) ! cudaError_enum { ENUMERATOR :: CUDA_SUCCESS = 0 ENUMERATOR :: CUDA_ERROR_INVALID_VALUE = 1 ENUMERATOR :: CUDA_ERROR_OUT_OF_MEMORY = 2 ENUMERATOR :: CUDA_ERROR_NOT_INITIALIZED = 3 ENUMERATOR :: CUDA_ERROR_DEINITIALIZED = 4 ENUMERATOR :: CUDA_ERROR_NO_DEVICE = 100 ENUMERATOR :: CUDA_ERROR_INVALID_DEVICE = 101 ENUMERATOR :: CUDA_ERROR_INVALID_IMAGE = 200 ENUMERATOR :: CUDA_ERROR_INVALID_CONTEXT = 201 ENUMERATOR :: CUDA_ERROR_CONTEXT_ALREADY_CURRENT = 202 ENUMERATOR :: CUDA_ERROR_MAP_FAILED = 205 ENUMERATOR :: CUDA_ERROR_UNMAP_FAILED = 206 ENUMERATOR :: CUDA_ERROR_ARRAY_IS_MAPPED = 207 ENUMERATOR :: CUDA_ERROR_ALREADY_MAPPED = 208 ENUMERATOR :: CUDA_ERROR_NO_BINARY_FOR_GPU = 209 ENUMERATOR :: CUDA_ERROR_ALREADY_ACQUIRED = 210 ENUMERATOR :: CUDA_ERROR_NOT_MAPPED = 211 ENUMERATOR :: CUDA_ERROR_INVALID_SOURCE = 300 ENUMERATOR :: CUDA_ERROR_FILE_NOT_FOUND = 301 ENUMERATOR :: CUDA_ERROR_INVALID_HANDLE = 400 ENUMERATOR :: CUDA_ERROR_NOT_FOUND = 500 ENUMERATOR :: CUDA_ERROR_NOT_READY = 600 ENUMERATOR :: CUDA_ERROR_LAUNCH_FAILED = 700 ENUMERATOR :: CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701 ENUMERATOR :: CUDA_ERROR_LAUNCH_TIMEOUT = 702 ENUMERATOR :: CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING = 703 ENUMERATOR :: CUDA_ERROR_UNKNOWN = 999 END ENUM ! TYPE, BIND(C) :: t_cudaDeviceProp CHARACTER (256) :: name INTEGER(C_SIZE_T) :: totalGlobalMem ! t_size = C_SIZE_T = 8bytes INTEGER(C_SIZE_T) :: sharedMemPerBlock INTEGER(C_INT) :: regsPerBlock INTEGER(C_INT) :: warpSize INTEGER(C_SIZE_T) :: memPitch INTEGER(C_INT) :: maxThreadsPerBlock INTEGER(C_INT) :: maxThreadsDim(3) INTEGER(C_INT) :: maxGridSize(3) INTEGER(C_INT) :: clockRate INTEGER(C_SIZE_T) :: totalConstMem INTEGER(C_INT) :: major INTEGER(C_INT) :: minor INTEGER(C_SIZE_T) :: textureAlignment INTEGER(C_INT) :: deviceOverlap INTEGER(C_INT) :: multiProcessorCount INTEGER(C_INT) :: kernelExecTimeoutEnabled INTEGER(C_INT) :: cudaReserved(39) END TYPE ! INTERFACE SUBROUTINE cudaGetDeviceCount(n) BIND(C, name='cudaGetDeviceCount') INTEGER, INTENT(OUT) :: n END SUBROUTINE cudaGetDeviceCount ! INTEGER FUNCTION cudaGetDeviceProperties(cudaDeviceProp, idev) BIND(C, name='cudaGetDeviceProperties') USE, INTRINSIC :: ISO_C_BINDING IMPORT TYPE (t_cudaDeviceProp), INTENT(OUT) :: cudaDeviceProp INTEGER, VALUE, INTENT(IN) :: idev END FUNCTION cudaGetDeviceProperties ! INTEGER FUNCTION cudaSetDevice(idev) BIND(C, name='cudaSetDevice') INTEGER, VALUE, INTENT(IN) :: idev END FUNCTION cudaSetDevice ! INTEGER FUNCTION cudaMalloc(ip, nn) BIND(C, name='cudaMalloc') USE, INTRINSIC :: ISO_C_BINDING TYPE (C_PTR), INTENT(OUT) :: ip INTEGER (C_SIZE_T), VALUE, INTENT(IN) :: nn END FUNCTION cudaMalloc ! INTEGER FUNCTION cudaFree(ip) BIND(C, name='cudaFree') USE, INTRINSIC :: ISO_C_BINDING TYPE (C_PTR), VALUE, INTENT(IN) :: ip END FUNCTION cudaFree ! INTEGER FUNCTION cudaMallocHost(ip, nn) BIND(C, name='cudaMallocHost') USE, INTRINSIC :: ISO_C_BINDING TYPE (C_PTR), INTENT(OUT) :: ip INTEGER (C_SIZE_T), VALUE, INTENT(IN) :: nn END FUNCTION cudaMallocHost ! INTEGER FUNCTION cudaFreeHost(ip) BIND(C, name='cudaFreeHost') USE, INTRINSIC :: ISO_C_BINDING TYPE (C_PTR), VALUE, INTENT(IN) :: ip END FUNCTION cudaFreeHost ! INTEGER FUNCTION cudaMemcpy(p_out, p_in, nsize, itype) BIND(C, name='cudaMemcpy') USE, INTRINSIC :: ISO_C_BINDING TYPE (C_PTR), VALUE, INTENT(IN) :: p_out TYPE (C_PTR), VALUE, INTENT(IN) :: p_in INTEGER (C_SIZE_T), VALUE, INTENT(IN) :: nsize INTEGER (C_INT), VALUE, INTENT(IN) :: itype END FUNCTION cudaMemcpy END INTERFACE CONTAINS !--------------------------------------------------------------------- SUBROUTINE printCudaLastError() USE, INTRINSIC :: ISO_C_BINDING IMPLICIT NONE ! INTERFACE INTEGER FUNCTION cudaGetLastError() BIND(C, name='cudaGetLastError') END FUNCTION cudaGetLastError ! FUNCTION cudaGetErrorString(n) RESULT(res) BIND(C, name='cudaGetErrorString') USE, INTRINSIC :: ISO_C_BINDING INTEGER, VALUE, INTENT(IN) :: n TYPE (C_PTR) :: res END FUNCTION cudaGetErrorString END INTERFACE ! CHARACTER (512), POINTER :: buff TYPE (C_PTR) :: ibuff INTEGER :: nerr, nlen nerr = cudaGetLastError() ibuff = cudaGetErrorString( nerr ) CALL C_F_POINTER(ibuff, buff) nlen = INDEX(buff, ACHAR(0)) - 1 PRINT '(a)', buff(1:nlen) RETURN END SUBROUTINE printCudaLastError !--------------------------------------------------------------------- SUBROUTINE CheckCudaRet(iret, text) USE, INTRINSIC :: ISO_C_BINDING IMPLICIT NONE INTEGER, INTENT(IN) :: iret CHARACTER (*), INTENT(IN), OPTIONAL :: text IF (iret /= CUDA_SUCCESS) THEN IF (PRESENT(text)) PRINT '(a)', text CALL printCudaLastError() STOP END IF RETURN END SUBROUTINE CheckCudaRet !--------------------------------------------------------------------- SUBROUTINE gpu_info() USE, INTRINSIC :: ISO_C_BINDING IMPLICIT NONE TYPE (t_cudaDeviceProp) :: prop INTEGER :: idev, ndev, iret ! CALL cudaGetDeviceCount(ndev) PRINT *, 'Number of Device(s) =', ndev DO idev = 0, ndev - 1 iret = cudaGetDeviceProperties(prop, idev) CALL CheckCudaRet(iret, 'cudaGetDeviceProperties') PRINT '(1x, a, i4)', 'Device No.:', idev PRINT '(1x, a, a )', 'name: ', prop%name( 1:INDEX(prop%name, ACHAR(0)) - 1 ) PRINT '(1x, a25, i10, a, i5, a)', 'totalGlobalMem: ', prop%totalGlobalMem, ' bytes =', & prop%totalGlobalMem / 2**20, ' Mbytes' PRINT '(1x, a25, i10, a, i5, a)', 'sharedMemPerBlock: ', prop%sharedMemPerBlock, ' bytes =', & prop%sharedMemPerBlock / 2**10, ' Kbytes' PRINT '(1x, a25, i10, a, i5, a)', 'regsPerBlock: ', prop%regsPerBlock, ' bytes =', & prop%regsPerBlock / 2**10, ' Kbytes' PRINT '(1x, a25, i10)', 'warpSize: ', prop%warpSize PRINT '(1x, a25, i10)', 'memPitch: ', prop%memPitch PRINT '(1x, a25, i10)', 'maxThreadsPerBlock: ', prop%maxThreadsPerBlock PRINT '(1x, a25, 3i10)', 'maxThreadsDim: ', prop%maxThreadsDim PRINT '(1x, a25, 3i10)', 'maxGridSize: ', prop%maxGridSize PRINT '(1x, a25, i10, a, i5, a)', 'totalConstMem: ', prop%totalConstMem, ' bytes =', & prop%totalConstMem / 2**10, ' Kbytes' PRINT '(1x, a25, 2i10)', 'major, minor: ', prop%major, prop%minor PRINT '(1x, a25, i10, a, f8.2, a)', 'clockRate: ', prop%clockRate, ' Hz =', & prop%clockRate / 1.0e6, ' GHz' PRINT '(1x, a25, i10)', 'textureAlignment: ', prop%textureAlignment PRINT '(1x, a25, i10)', 'deviceOverlap: ', prop%deviceOverlap PRINT '(1x, a25, i10)', 'multiProcessorCount: ', prop%multiProcessorCount PRINT '(1x, a25, i10)', 'kernelExecTimeoutEnabled:', prop%kernelExecTimeoutEnabled PRINT * END DO END SUBROUTINE gpu_info !--------------------------------------------------------------------- END MODULE m_cuda !========================================================================= PROGRAM CudaTest USE, INTRINSIC :: ISO_C_BINDING USE :: m_cuda IMPLICIT NONE ! !define length of the array INTEGER, PARAMETER :: N = 2**24 REAL (8) :: t0, t1, tt0, tt1 CHARACTER(10) :: time0, time1 INTEGER :: i, iret, ic0, ic1, icrate TYPE (C_PTR) :: h_c, h_b, d_p, d_q, d_r(10) REAL (C_FLOAT), POINTER:: c(:), b(:) INTEGER (C_SIZE_T) :: nn ! CALL gpu_info() iret = cudaSetDevice( 0 ) CALL CheckCudaRet(iret, 'cudaSetDevice') ! nn = sizeof(C_FLOAT) * n PRINT *, nn, 'bytes ', nn / 2**20, 'Mbytes' iret = cudaMallocHost(h_c, nn) CALL CheckCudaRet(iret, 'cudaMallocHost1') CALL C_F_POINTER( h_c, c, [n] ) iret = cudaMallocHost(h_b, nn) CALL CheckCudaRet(iret, 'cudaMallocHost2') CALL C_F_POINTER( h_b, b, [n] ) ! c = 1.0 ! iret = cudaMalloc(d_p, nn) CALL CheckCudaRet(iret, 'cudaMalloc1') iret = cudaMalloc(d_q, nn) CALL CheckCudaRet(iret, 'cudaMalloc2') DO i = 1, 3 iret = cudaMalloc(d_r(i), nn) PRINT *, 'Alloc ', i CALL CheckCudaRet(iret, 'cudaMalloc3') END DO ! CALL CPU_TIME(t0) CALL SYSTEM_CLOCK(ic0) CALL date_and_time(time = time0) DO i = 1, 500 iret = cudaMemcpy( d_p, h_c, nn, cudaMemcpyHostToDevice ) CALL CheckCudaRet(iret, 'cudaMemcpy1') iret = cudaMemcpy( d_q, d_p, nn, cudaMemcpyDeviceToDevice ) CALL CheckCudaRet(iret, 'cudaMemcpy2') iret = cudaMemcpy( h_b, d_q, nn, cudaMemcpyDeviceToHost ) CALL CheckCudaRet(iret, 'cudaMemcpy3') ! iret = cudaMemcpy( d_r(1), h_c, nn, cudaMemcpyHostToDevice ) CALL CheckCudaRet(iret, 'cudaMemcpy1') iret = cudaMemcpy( d_r(2), d_r(1), nn, cudaMemcpyDeviceToDevice ) CALL CheckCudaRet(iret, 'cudaMemcpy2') iret = cudaMemcpy( h_b, d_r(2), nn, cudaMemcpyDeviceToHost ) CALL CheckCudaRet(iret, 'cudaMemcpy3') ! iret = cudaMemcpy( d_r(3), h_c, nn, cudaMemcpyHostToDevice ) CALL CheckCudaRet(iret, 'cudaMemcpy1') iret = cudaMemcpy( d_r(1), d_r(3), nn, cudaMemcpyDeviceToDevice ) CALL CheckCudaRet(iret, 'cudaMemcpy2') iret = cudaMemcpy( h_b, d_r(1), nn, cudaMemcpyDeviceToHost ) CALL CheckCudaRet(iret, 'cudaMemcpy3') END DO CALL date_and_time(time = time1) CALL CPU_TIME(t1) CALL SYSTEM_CLOCK(ic1, icrate) ! READ(time0, '(F10.3)') tt0 READ(time1, '(F10.3)') tt1 PRINT *, 'Copy Host->GPU->GPU->HOST: cpu_time = ', t1 - t0, 'sec' PRINT *, ' DATE_AND_TIME = ', tt1 - tt0, 'sec' !, time0, time1 ! IF ( ANY(c /= b) ) THEN PRINT *, 'error: mismatch!' ELSE PRINT *, 'Copy Check OK', SUM(b) * sizeof(C_FLOAT), KIND(1.0_C_FLOAT) END IF ! iret = cudaFree(d_p) CALL CheckCudaRet(iret, 'cudaFree') iret = cudaFree(d_q) CALL CheckCudaRet(iret, 'cudaFree') DO i = 1, 3 iret = cudaFree(d_r(i)) PRINT *, 'Free ', i CALL CheckCudaRet(iret, 'cudaFree') END DO ! iret = cudaFreeHost(h_c) CALL CheckCudaRet(iret, 'cudaFreeHost') iret = cudaFreeHost(h_b) CALL CheckCudaRet(iret, 'cudaFreeHost') ! STOP END PROGRAM CudaTest