CUDA はカーネル実行時の動的共有メモリ割り当てをサポートしていますが、メカニズムは OpenCL とは少し異なります。CUDA ランタイム API では、動的に割り当てられた/サイズ変更された共有メモリを使用するカーネルと、メモリのサイズを変更するための起動は、次の構文を使用します。
__global__ void kernel(...)
{
extern __shared__ typename buffer[];
....
}
....
kernel <<< griddim, blockdim, sharedmem, streamID >>> (...)
はsharedmem
、バッファに割り当てられるブロックごとの合計バイト数です。
PyCUDA では、同じメカニズムが次のように機能します。
mod = SourceModule("""
__global__ void kernel(...)
{
extern __shared__ typename buffer[];
....
}
""")
func = mod.get_function("kernel")
func.prepare(..., shared=sharedmem)
func.prepared_call(griddim,blockdim,...)
メソッドに渡された共有メモリの割り当てサイズを使用しprepare
ます。