numba内でcudaを使用する方法を実験しています。しかし、私は自分の予想とは異なる何かに遭遇しました。これが私のコードです
from numba import cuda
@cuda.jit
def matmul(A, B, C):
"""Perform square matrix multiplication of C = A * B
"""
d=cuda.local.array((3,3),dtype=numba.float64)
i, j = cuda.grid(2)
if i < C.shape[0] and j < C.shape[1]:
tmp = 0.
for k in range(A.shape[1]):
tmp += A[i, k] * B[k, j]
C[i, j] = tmp
これは、numba.cudaを使用してテストするために私が自己定義した行列関数です。テストを実行する前に、次のコードで配列もロードしました。
import numpy as np
a=np.random.rand(2000,2000)
b=np.random.rand(2000,2000)
c=np.empty((2000,2000))
a1=cuda.to_device(a)
b1=cuda.to_device(b)
c1=cuda.to_device(c)
次に、実験に次のコードを使用しました。
from time import time
count =0
start=time()
for i in range(2000):
matmul[(256,256),(16,16)](a1,b1,c1)
count +=1
print(count)
forループは、最初の1028回の実行で非常に高速に実行されました。ただし、1028日以降は実行速度が非常に遅くなりました。正確な原因と修正方法を教えてください。ちなみに私はwin10で走っています。
これがnumba.cudaから呼び出された私のcuda情報です
from numba import cuda
gpu = cuda.get_current_device()
print("name = %s" % gpu.name)
print("maxThreadsPerBlock = %s" % str(gpu.MAX_THREADS_PER_BLOCK))
print("maxBlockDimX = %s" % str(gpu.MAX_BLOCK_DIM_X))
print("maxBlockDimY = %s" % str(gpu.MAX_BLOCK_DIM_Y))
print("maxBlockDimZ = %s" % str(gpu.MAX_BLOCK_DIM_Z))
print("maxGridDimX = %s" % str(gpu.MAX_GRID_DIM_X))
print("maxGridDimY = %s" % str(gpu.MAX_GRID_DIM_Y))
print("maxGridDimZ = %s" % str(gpu.MAX_GRID_DIM_Z))
print("maxSharedMemoryPerBlock = %s" %
str(gpu.MAX_SHARED_MEMORY_PER_BLOCK))
print("asyncEngineCount = %s" % str(gpu.ASYNC_ENGINE_COUNT))
print("canMapHostMemory = %s" % str(gpu.CAN_MAP_HOST_MEMORY))
print("multiProcessorCount = %s" % str(gpu.MULTIPROCESSOR_COUNT))
print("warpSize = %s" % str(gpu.WARP_SIZE))
print("unifiedAddressing = %s" % str(gpu.UNIFIED_ADDRESSING))
print("pciBusID = %s" % str(gpu.PCI_BUS_ID))
print("pciDeviceID = %s" % str(gpu.PCI_DEVICE_ID))
出力は次のとおりです。
名前= b'GeForce GTX 1050 Ti '
maxThreadsPerBlock = 1024
maxBlockDimX = 1024
maxBlockDimY = 1024
maxBlockDimZ = 64
maxGridDimX = 2147483647
maxGridDimY = 65535
maxGridDimZ = 65535
maxSharedMemoryPerBlock = 49152
asyncEngineCount = 2
canMapHostMemory = 1
multiProcessorCount = 6
ワープサイズ= 32
UnifiedAddressing = 1
pciBusID = 3
pciDeviceID = 0
これは、GPUカーネルの起動に関連付けられた非同期起動キューが原因で発生します。
numbaにGPUカーネルを送信するように指示する場合：
matmul[(256,256),(16,16)](a1,b1,c1)
このリクエストはキューに入り、GPUカーネルがまだ完了していないか、まだ開始されていなくても、そのカーネル呼び出しを発行したCPUスレッド（つまりPython）は続行できます。
CUDAランタイムはこれらのリクエストをキューに入れ、GPUがさらなる作業の準備ができるとそれらを発行します。
forループの非常に高速なインクリメント中に最初に目撃しているのは、キューが作業要求でいっぱいになっていることです。これは、GPUが作業を実行するために必要な実際の時間を表すものではありません。
最終的にキューがいっぱいになり、CUDAランタイムはカーネルの起動時にキュースロットが開くまでCPUスレッド（つまりPython）を停止します。その時点で、forループはもう1回繰り返されます。この時点（おそらく約1028回の反復）で「スローダウン」が見られ始めます。その後、forループは、GPUカーネルが実行され、処理キューから削除される速度とほぼ同じ速度で進行します。
ここで修正するものはありません。これは予想される動作です。
GPUカーネルが実際に実行される速度でのみ、forループを続行する場合は、forループに同期関数を挿入する必要があります。
たとえば、numbaはnumba.cuda.synchronize（）を提供します。したがって、forループを次のように変更すると次のようになります。
for i in range(2000):
matmul[(256,256),(16,16)](a1,b1,c1)
cuda.synchronize()
count +=1
print(count)
forループは、「キュー充填」速度ではなく、GPU作業完了の実際の速度で進行することがわかります。
この記事はインターネットから収集されたものであり、転載の際にはソースを示してください。
侵害の場合は、連絡してください[email protected]
コメントを追加